4.2 SASS האסמבלי של החומרה הרצאה
בשיעור 4.1 הכרנו את PTX - ה-Parallel Thread eXecution, ה-ISA הווירטואלי היציב שמשמש "מותן צרה" (narrow waist) בין העולם התוכנתי שמעל לחומרה שמתחת. ראינו ש-PTX הוא ייצוג ביניים (IR) בסגנון LLVM-IR, שהוא נשמר קדימה-תואם (forward compatible), ושה-CUDA Drivers מהדרים אותו בזמן ריצה (JIT) לקוד המכונה האמיתי. השיעור הזה יורד את המדרגה האחרונה: אל SASS - ה-Streaming ASSembler, שפת האסמבלי הילידית של ה-GPU, הרמה הנמוכה ביותר שבה בני אדם עדיין יכולים לקרוא קוד. בניגוד ל-PTX, SASS הוא לא ווירטואלי ולא נייד - הוא צמוד לארכיטקטורת ה-SM המסוימת, ה-SMs מריצים אותו ישירות, וגרסה שנוצרה ל-Hopper לא תרוץ על Blackwell. נלמד לזהות את סוגי האוגרים ב-SASS, לפענח שורת פקודה בודדת, למפות משפחות פקודות שנפגוש שוב ושוב לאורך הקורס, ובעיקר - לתפעל את הכלים (cuobjdump, nvdisasm) שמאפשרים לפרק kernel מהודר ולקרוא בדיוק מה שהחומרה עומדת לבצע. זו מיומנות שנשען עליה בכבדות בפרק 8, כשננתח warp divergence (התפצלות ה-warp) ו-scoreboards (לוחות-תוצאות) של המתזמן.
מה זה SASS ומאיפה השם - Streaming ASSembler¶
השם SASS הוא ראשי תיבות של Streaming ASSembler. המילה "Streaming" מפנה ישירות אל ה-Streaming Multiprocessors (ה-SMs, המעבדים הרב-זרמיים) שאותם שפת האסמבלי הזו מתכנתת - SASS הוא זרם הפקודות הילידי שה-SMs טוענים ומריצים. זו הרמה הנמוכה ביותר שבה עדיין אפשר לכתוב ולקרוא קוד קריא לבני אדם עבור ה-GPU; מתחתיה יש רק ה-microcode הבינארי הספציפי לdevice, שאליו SASS מומר בזמן ההרצה.
העובדה הקריטית, וההבדל המהותי מ-PTX, היא ש-SASS צמוד לגרסת ארכיטקטורת SM מסוימת ואינו נייד בין דורות. לכל ארכיטקטורת SM יש SASS משלה. לדוגמה, ה-glossary משתמש ב-SM90a, ארכיטקטורת ה-Hopper (ה-H100 שלנו). ה-SASS של SM90a שונה מזה של Ampere (SM80) או של Blackwell (SM100), ולא רק בקידוד - גם באוסף הפקודות עצמו: פקודות כמו wgmma (warpgroup matrix-multiply-accumulate) קיימות רק החל מ-Hopper. זו התמונה המשלימה ל-PTX: PTX נכתב פעם אחת ורץ על כל דור בעל compute capability תואם או גבוה יותר, ואילו SASS נוצר מחדש לכל דור.
The compilation chain - from CUDA C++ to the hardware
────────────────────────────────────────────────────────────
CUDA C++ ──(nvcc/cicc)──► PTX ──(ptxas)──► SASS ──► microcode
readable, portable virtual IR assembly binary
stable, portable SM-tied device-tied
(compute_XY) (sm_XY)
Point we'll return to: PTX is embedded in the binary also for forward compatibility,
and then the driver performs JIT from PTX to the SASS of the actual device.
שימו לב לזוג המונחים compute_90 מול sm_90a שהכרנו בשיעור על nvcc: compute_XY מגדיר את גרסת ה-PTX (compute capability), ו-sm_XY מגדיר את גרסת ה-SASS (SM architecture version). ה-a הסופי ב-sm_90a מסמן את התכונות המואצות הצמודות-לארכיטקטורה של Hopper (כמו wgmma/tma) - נחזור לזה כשנקמפל.
אנטומיה של שורת SASS - anatomy of an instruction¶
נתחיל משתי הפקודות המדויקות שה-glossary מביא מ-SM90a, ונפרק כל אחת:
הפקודה הראשונה, FFMA, היא Fused Floating-point Multiply-Add - כפל-חיבור מאוחד של נקודה-צפה. משמעותה: R0 = R7 * R0 + 1.5. שלושת האופרנדים הראשונים הם אוגרים או קבועים; ה-"fused" אומר שהכפל והחיבור מבוצעים בפעולה אחת עם עיגול יחיד בסוף (זה בדיוק ה-fma.rn.f32 שראינו ב-PTX ב-4.1, רק שכאן האוגרים פיזיים ולא ווירטואליים). הפקודה השנייה, S2UR UR4, SR_CTAID.X, מעתיקה את רכיב ה-X של אינדקס ה-Cooperative Thread Array (ה-CTAID, כלומר blockIdx.x) מ-Special Register אל Uniform Register מספר 4.
מכאן נגזרים שלושת סוגי האוגרים שתפגשו ב-SASS:
| סימון | שם מלא | היקף | דוגמה |
|---|---|---|---|
R |
General (per-thread) Register | פרטי לכל thread - ערך שונה בכל אחד מ-32 הthreads של ה-warp | R0, R7, RZ (אוגר האפס) |
UR |
Uniform Register | אחיד (uniform) על פני כל ה-warp - ערך זהה ל-32 הthreads | UR4, URZ |
SR |
Special Register | אוגר מובנה לקריאה בלבד עם ערכי מערכת | SR_CTAID.X, SR_TID.X |
ההבחנה בין R ל-UR היא לב העניין ותחזור בפרק 8. כשערך זהה לכל ה-warp - למשל blockIdx.x, שווה לכל 32 הthreads - חבל להחזיק אותו ב-32 עותקים ולחשב אותו 32 פעמים במסלול ה-vector. החל מ-Volta, ל-GPU יש מסלול נתונים אחיד (uniform datapath) נפרד: יחידה סקלרית קטנה שמחזיקה ערכים אחידים ב-UR ומחשבת אותם פעם אחת ל-warp, ובכך חוסכת אוגרים ואנרגיה במסלול ה-vector. לכן S2UR (Special-to-Uniform-Register) שולפת את CTAID.X ישר אל UR, ולא אל R.
עוד סימונים שיופיעו כמעט בכל dump:
- מאגר הקבועים -
c[0x0][0x28]: הבנק הקבוע (constant bank). כאן ה-hardware מניח את פרמטרי ה-kernel, אתblockDim/gridDim, ומצביע המחסנית. הצורה היאc[bank][offset], שני המספרים בהקסה. - פרדיקטים -
@P0,@!P0,PT: אוגרי פרדיקט (predicate registers) בני סיבית אחת. קידומת@P0לפני פקודה אומרת "בצע רק אםP0אמת";@!P0הוא השלילה;PTהוא הפרדיקט הקבוע "תמיד אמת". זה מנגנון ה-predication שבו ה-GPU מבצע ענפים קצרים בלי branch אמיתי. - מספרי כתובת -
/*0090*/: ה-offset של הפקודה בבתים בתוך ה-kernel, שכלים מדפיסים בהערה בתחילת השורה. - סיומות -
.E,.64,.GE.AND: מגדירות מצב לפקודה, למשלLDG.E= extended (כתובת 64 סיביות),R2.64= זוג אוגרים המחזיק כתובת 64 סיביות,ISETP.GE.AND= השווה עם תנאי גדול-שווה וצרף (AND) לפרדיקט.
מפת משפחות הפקודות - instruction families¶
אין טעם לשנן את מאות פקודות ה-SASS. במקום זאת, למדו לזהות משפחות - ובכל dump תזהו מיד מה ה-kernel עושה. הטבלה הזו היא מפת הדרכים; נפגוש את כל המשפחות שוב לאורך הקורס.
| משפחה (מנמוניקים) | מה היא עושה | היכן תפגשו אותה |
|---|---|---|
FFMA, FADD, FMUL |
אריתמטיקת FP32 (נקודה-צפה 32 סיביות) | כל kernel חישובי; ה-FFMA הוא סוס העבודה |
HFMA2, HADD2, HMUL2 |
אריתמטיקת FP16 מזווגת (שני חצאים במקביל) | kernels של דיוק-מעורב, deep learning |
IMAD, IADD3, SHF, LOP3 |
חשבון שלמים, בעיקר לחישוב אינדקסים וכתובות | בכל kernel, לחישוב blockIdx*blockDim+threadIdx |
HMMA, IMMA |
כפל-מטריצות על Tensor Cores (ליבות טנזור) | GEMM, קונבולוציה, attention; פרק 1.5 והפרויקטים |
LDG, STG |
טעינה/אחסון מglobal memory (זיכרון גלובלי) | כל גישה למערך שהוקצה ב-cudaMalloc |
LDS, STS |
טעינה/אחסון מshared memory (זיכרון משותף) | kernels מרוצפים (tiled), ראו 3.3 |
LDGSTS |
העתקה א-סינכרונית global -> shared (Ampere+) | pipelining של טעינות, ראו 3.3 והמשך |
LDC |
טעינה ממאגר הקבועים (constant) | קריאת פרמטרים, __constant__ |
MUFU |
Multi-Function Unit - פונקציות מיוחדות (טרנסצנדנטיות) | sinf, expf, rsqrtf, __expf; ראו SFU ב-1.3 |
ISETP, FSETP |
קביעת פרדיקט לפי השוואה (Set Predicate) | כל if, כל בדיקת גבולות |
@P0 + BRA, BRX, JMP |
predication וקפיצות (branch) | תנאים, לולאות; לב ה-divergence בפרק 8 |
BAR.SYNC |
מחסום ברמת ה-block - זהו __syncthreads() |
kernels מרוצפים, כל סנכרון block |
S2R, S2UR, CS2R |
העתקת Special Register ל-R/UR |
קריאת threadIdx, blockIdx בתחילת כמעט כל kernel |
MOV, IMAD.MOV, SEL |
העברת ערכים ובחירה מותנית | בכל מקום |
EXIT, BRA (self) |
סיום ה-kernel והלולאה האינסופית הסופית | סוף כל kernel |
שימו לב ל-MUFU: ה-glossary וה-SFU בפרק 1.3 לימדו אותנו שהפונקציות הטרנסצנדנטיות (סינוס, אקספוננט, שורש-הופכי) לא רצות על ה-CUDA Cores הרגילות אלא על יחידה נפרדת ואיטית יותר, ה-Special Function Unit. ב-SASS זה מתגלה בדיוק כפקודת MUFU.SIN, MUFU.EX2, או MUFU.RSQ. אם אתם רואים המון MUFU בלולאה חמה - סימן שאתם שורפים throughput על יחידה נדירה, ואולי כדאי לשקול קירוב מהיר יותר.
איך קוראים SASS בפועל - cuobjdump ו-nvdisasm¶
כמעט אף פעם לא כותבים SASS ביד. תזרים העבודה האמיתי הוא: קוראים את ה-SASS שהcompiler יצר, בזמן profiling, כדי להבין למה ה-kernel איטי - ואז חוזרים ומתקנים את ה-CUDA C++ או, לכל היותר, מזריקים in-line PTX. שני הכלים המרכזיים לקריאה מגיעים עם ה-CUDA Binary Utilities.
הכלי הראשון - cuobjdump. זהו הכלי הרחב: הוא פותח בינארי-host שלם (fat binary) או קובץ cubin בודד ומדפיס את מה שיש בפנים. הדגל המרכזי הוא --dump-sass (בקיצור -sass):
# compile a regular kernel to a host binary (fat binary with SASS embedded)
nvcc -O2 -arch=sm_90a -o app app.cu
# print all the SASS embedded in the binary
cuobjdump --dump-sass app
הפלט מראה, לכל kernel, את שם הפונקציה (מקושקש - mangled), את גרסת ה-SM, ואת רצף פקודות ה-SASS עם ה-offsets. דגלים שימושיים נוספים: --dump-ptx להצגת ה-PTX המשורשר, --list-elf (-lelf) לרשימת ה-cubins שבפנים, ו---extract-elf all (-xelf all) לחילוץ ה-cubins לקבצים נפרדים על הדיסק.
הכלי השני - nvdisasm. זהו ה-disassembler הממוקד, והוא עובד על קובץ cubin (לא על בינארי-host). היכולות שלו עשירות יותר מ-cuobjdump: הוא בונה גרפי זרימת-בקרה (control flow graphs) וממפה פקודות SASS בחזרה לשורות המקור. הדרך הנוחה להשיג cubin היא לקמפל ישירות אליו עם -cubin:
# compile directly to cubin, with line info for mapping to source
nvcc -O2 -arch=sm_90a -lineinfo -cubin -o app.cubin app.cu
# disassemble SASS with basic-block boundaries
nvdisasm -c app.cubin
הדגל -lineinfo הוא קריטי. בלעדיו, ה-cubin לא נושא את הקישור בין פקודות SASS לשורות ה-.cu המקוריות. עם -lineinfo (מיפוי קל, מתאים ל-build אופטימלי; בניגוד ל--G שמכבה אופטימיזציות ומיועד לניפוי) אפשר לבקש מ-nvdisasm להשחיל את מספרי השורות לתוך הפלט:
וכדי לייצר גרף זרימת-בקרה בפורמט DOT ולעבד אותו לתמונה:
לסיכום חלוקת העבודה: cuobjdump הוא "מה יש בבינארי הזה" (רחב, פותח fat binaries ו-cubins), ו-nvdisasm הוא "מה הקוד הזה בעצם עושה" (עמוק, CFG ומיפוי למקור). הם מתחברים: cuobjdump -xelf all כדי לחלץ cubin מבינארי קיים, ואז nvdisasm לניתוח מעמיק. כלי משלים חביב הוא Godbolt (compiler explorer), שמציג CUDA C++, PTX ו-SASS זה לצד זה בדפדפן.
דוגמה מלאה - קריאת SASS של saxpy¶
בואו נראה זאת מקצה לקצה על kernel פשוט אך מייצג - SAXPY, כלומר out = a*x + y:
__global__ void saxpy(float a, const float* x, const float* y,
float* out, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
out[i] = a * x[i] + y[i];
}
נקמפל ונפרק:
פלט SASS מייצג ל-SM90a (הערכים, ה-offsets ובחירת האוגרים משתנים בין גרסאות nvcc ו-driver - קחו אותו כמייצג, לא מדויק לספרה):
/*0000*/ IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ; // initialize stack pointer
/*0010*/ S2R R2, SR_CTAID.X ; // R2 = blockIdx.x
/*0020*/ S2R R3, SR_TID.X ; // R3 = threadIdx.x
/*0030*/ IMAD R2, R2, c[0x0][0x0], R3 ; // i = blockIdx.x*blockDim.x + threadIdx.x
/*0040*/ ISETP.GE.AND P0, PT, R2, c[0x0][0x218], PT ; // P0 = (i >= n)
/*0050*/ @P0 EXIT ; // if i>=n, end the thread
/*0060*/ IMAD.WIDE R4, R2, 0x4, c[0x0][0x210] ; // address of x[i] (64 bits)
/*0070*/ IMAD.WIDE R6, R2, 0x4, c[0x0][0x208] ; // address of y[i]
/*0080*/ LDG.E R4, desc[UR4][R4.64] ; // R4 = x[i] (global load)
/*0090*/ LDG.E R7, desc[UR4][R6.64] ; // R7 = y[i]
/*00a0*/ IMAD.WIDE R8, R2, 0x4, c[0x0][0x220] ; // address of out[i]
/*00b0*/ FFMA R4, R4, c[0x0][0x200], R7 ; // R4 = x[i]*a + y[i]
/*00c0*/ STG.E desc[UR4][R8.64], R4 ; // out[i] = R4 (global store)
/*00d0*/ EXIT ;
/*00e0*/ BRA 0xe0 ; // trap loop at the end of the kernel
נקרא זאת שורה-שורה, כי כל התבניות כאן חוזרות בכל kernel:
-
חישוב האינדקס (offsets 0x10-0x30). שתי פקודות
S2Rשולפות אתblockIdx.xואתthreadIdx.xמה-Special Registers אלR2ו-R3. אז פקודתIMADיחידה מבצעתi = R2 * blockDim.x + R3- הכפל-חיבור השלם. שימו לב ש-blockDim.xנשלף ממאגר הקבועיםc[0x0][0x0], לא מ-Special Register: הcompiler יודע שהוא אחיד ל-warp ומניח אותו בקבועים. -
בדיקת הגבולות וה-predication (0x40-0x50).
ISETP.GE.AND P0, PT, R2, c[0x0][0x218], PTמחשב את הפרדיקטP0 = (i >= n), כאשרnנשלף ממאגר הקבועים. אז@P0 EXITמסיים את ה-thread רק אםP0אמת. זהו התרגום ל-SASS שלif (i < n)- הcompiler הפך אותו לתנאי הפוך ("אם i>=n צא") ומימש אותו כ-EXIT מותנה, בלי branch מלא. -
חישוב הכתובות (0x60-0x70, 0xa0).
IMAD.WIDE R4, R2, 0x4, c[0x0][0x210]מחשב כתובת 64 סיביות:base + i*4(הכפל ב-4 כיfloatהוא 4 בתים). ה-base של המערךxנשלף מהקבועים. הסיומת.WIDEפירושה שהתוצאה היא זוג אוגרים (R4:R5) המחזיק כתובת 64 סיביות. -
הטעינות והחישוב (0x80-0xb0).
LDG.Eהן שתי טעינות ה-global שלx[i]ו-y[i]. שימו לב לתחבירdesc[UR4][R4.64]:UR4הוא Uniform Register המחזיק מתאר זיכרון (memory descriptor) - אחיד ל-warp - ו-R4.64הוא הכתובת הפרטית ל-thread. זו בדיוק ההבחנה R מול UR שדיברנו עליה. לבסוףFFMA R4, R4, c[0x0][0x200], R7מבצעתx[i]*a + y[i]בפעולה מאוחדת אחת, כאשר הסקלרaנשלף מהקבועים. -
האחסון והסיום (0xc0-0xe0).
STG.Eכותבת את התוצאה ל-out[i],EXITמסיים, וה-BRA 0xe0הוא לולאה עצמית - "מלכודת" סטנדרטית שהcompiler שם בסוף כל kernel כדי לתפוס זרימת-בקרה שסטתה.
הנה כל משפחות הפקודות מהטבלה, בפעולה, ב-13 שורות בלבד: S2R, IMAD, ISETP, predication (@P0), LDG/STG, FFMA, EXIT, BRA. כשתפגשו kernel מסובך יותר, אתם כבר יודעים לקרוא את השלד שלו.
תיעוד והנדסה לאחור - documentation and reverse engineering¶
נקודה מפוכחת וחשובה: SASS מתועד בצורה דלה מאוד. פקודות ה-SASS מופיעות ברשימה בתיעוד של ה-CUDA Binary Utilities, אבל הסמנטיקה שלהן אינה מוגדרת שם - אין הסבר פורמלי מה בדיוק כל פקודה עושה עם הדגלים והסיומות. גרוע מכך, המיפוי מהאסמבלי ה-ASCII אל האופקודים והאופרנדים הבינאריים אינו מתועד כלל. הוא כן פוענח בהנדסה-לאחור עבור חלק מהדורות (Maxwell, Lovelace) בידי חוקרים חיצוניים, אבל אין אחריות ואין יציבות.
מכאן שתי מסקנות מעשיות:
- כתיבת SASS ביד היא נדירה מאוד. אין assembler רשמי נתמך שממיר SASS-כטקסט חזרה לבינארי. לכן המסלול הוא חד-כיווני: קוראים את מה שהcompiler יצר, אבל לא כותבים חדש. אם צריך שליטה עדינה יותר ממה ש-
nvccנותן, יורדים ל-in-line PTX (שכן מתועד ויציב), לא ל-SASS. - כל כלי או כוונון ברמת ה-SASS שביר בין דורות. מכיוון שהקידוד הבינארי לא מתועד ומשתנה בכל ארכיטקטורה, כל דבר שנשען עליו נשבר במעבר מ-Hopper ל-Blackwell. זו סיבה נוספת ש-PTX הוא היעד העמיד: אותו קוד ימשיך לרוץ על חומרה עתידית שעדיין לא קיימת.
התזרים הנכון, אם כן, הוא: פרופיילינג -> קריאת SASS -> תיקון ב-CUDA C++ (או, לכל היותר, in-line PTX). ה-SASS הוא כלי אבחון, לא שפת כתיבה. קוראים אותו כדי לענות על שאלות כמו: "האם הcompiler באמת יצר את ה-FFMA שקיוויתי לו?", "האם הלולאה נפתחה (unrolled)?", "האם יש כאן register spill ל-global memory?", "כמה LDG הוא הוציא, ולמה?".
הצצה לבקרה וללוח-התוצאות - control and scoreboard preview¶
יש שכבה נוספת ב-SASS שלא הצגנו עדיין, ושנקדיש לה את פרק 8.3: מידע הבקרה (control information). החל מ-Volta, כל פקודת SASS נושאת, לצד האופקוד והאופרנדים, מילת-בקרה זעירה שמקודדת כיצד המתזמן מריץ אותה: כמה מחזורים לעצור (stall) לפני הפקודה הבאה, האם לוותר (yield) על ה-warp לטובת אחר, ובאילו scoreboards (לוחות-תוצאות) להשתמש כדי לעקוב אחר תלויות.
הרעיון: פקודות בעלות latency קבועה, כמו FFMA, הcompiler יודע מראש כמה זמן הן לוקחות ומקודד stall מדויק. אבל פקודות בעלות latency משתנה - LDG (שתלוי אם הנתון ב-cache או ב-HBM) ו-MUFU - לא ניתן לתזמן סטטית. עבורן, ה-SASS מקצה מחסום scoreboard: הפקודה שמנפיקה את הטעינה מסמנת scoreboard מסוים, וכל פקודה שצורכת את התוצאה "ממתינה" עליו עד שהנתון הגיע. כך הcompiler שוזר תלויות בלי חומרת scoreboard מלאה לכל פקודה.
The scoreboard idea (shown conceptually; the exact encoding is in chapter 8.3)
─────────────────────────────────────────────────────────
LDG.E R4, [R6.64] ; issues a load, marks scoreboard #0
FFMA R9, R9, R1, R2 ; independent work - runs while the load is in flight (latency hiding)
... wait on scoreboard #0 ...
FADD R10, R4, R10 ; consumes R4 - must wait until the load completed
זו התשתית שמאחורי הסתרת ה-latency (latency hiding) שראינו ב-1.4: ה-warp scheduler מריץ עבודה עצמאית בזמן שטעינה בדרך, וה-scoreboards הם המנגנון שמבטיח שלא נצרוך נתון לפני שהגיע. ברוב הכלים (למשל cuobjdump) מידע הבקרה אינו מוצג כברירת מחדל; Nsight Compute חושף אותו, ובו נשתמש בפרק 8.3 כדי לנתח stalls. כאן די להכיר את הרעיון: מתחת לכל פקודת SASS יש שכבת תזמון סטטי שהcompiler קידד, וקריאה שלה היא המפתח לניתוח ביצועים ברמה העמוקה ביותר.
סיכום¶
- ה-SASS (Streaming ASSembler) הוא שפת האסמבלי הילידית של ה-GPU, הרמה הנמוכה ביותר הקריאה לבני אדם, וה-SMs מריצים אותו ישירות; השם "Streaming" מפנה אל ה-Streaming Multiprocessors.
- בניגוד ל-PTX הנייד והיציב, SASS צמוד לגרסת ארכיטקטורת SM מסוימת (למשל SM90a של Hopper) ואינו נייד בין דורות -
sm_XYמגדיר SASS ואילוcompute_XYמגדיר PTX. - ב-SASS שלושה סוגי אוגרים:
R(פרטי לכל thread),UR(אחיד על פני ה-warp, במסלול הנתונים האחיד), ו-SR(Special Register מובנה כמוSR_CTAID.X,SR_TID.X). - כדאי לזהות משפחות פקודות ולא לשנן פקודות בודדות:
FFMA/HFMA/IMAD(אריתמטיקה),HMMA(Tensor Core),LDG/STG(global),LDS/STS(shared),MUFU(פונקציות מיוחדות),ISETP/predication@P0, ו-BRA/EXIT. - הכלים לקריאה הם
cuobjdump --dump-sass(רחב, פותח בינארי-host ו-cubin) ו-nvdisasm(ממוקד על cubin, בונה CFG עם-cfgוממפה למקור עם-g); לקמפל עם-lineinfoהוא תנאי למיפוי SASS חזרה לשורות המקור. - ב-kernel טיפוסי (saxpy) קוראים בבירור את השלד:
S2Rלאינדקסים,IMADלחישובi,ISETP+@P0 EXITלבדיקת גבולות,IMAD.WIDEלכתובות,LDG/FFMA/STGלחישוב, ו-EXIT. - SASS מתועד בדלות - הפקודות מופיעות ברשימה אך הסמנטיקה והקידוד הבינארי אינם מתועדים (פוענחו בהנדסה-לאחור רק לחלק מהדורות); לכן כותבים אותו ביד רק לעתים נדירות, וקוראים אותו בזמן profiling ואז מתקנים ב-CUDA C++ או ב-in-line PTX.
- מתחת לכל פקודת SASS יש מידע-בקרה (stall, yield, scoreboards) שהcompiler קידד סטטית; ה-scoreboards מנהלים תלויות של פקודות latency-משתנה (
LDG,MUFU) ומאפשרים הסתרת latency - נעמיק בזה בפרק 8.3.