4.2 SASS האסמבלי של החומרה פתרון
פתרון - SASS - האסמבלי של החומרה¶
הפתרונות כאן הם דוגמה מייצגת. פלט ה-SASS - בחירת האוגרים, ה-offsets, מספרי הפקודות, ואפילו אילו פקודות בדיוק נוצרות - משתנה בין גרסאות nvcc, גרסאות driver, ודורות חומרה. כל הדאמפים למטה נלקחו על H100 SXM עם sm_90a כסדר גודל וכצורה טיפוסית. מה שחשוב הוא הדפוסים - איזו משפחת פקודות מזהים, איפה הפיצול, מה השתנה בין הדורות - לא ה-offset המדויק. השמות המקושקשים (mangled) עשויים להיראות מעט שונה אצלכם.
פתרון תרגיל 1 - פירוק וזיהוי משפחות פקודות¶
הפירוק:
הקטע של saxpy (מקוצר, עם הערות שלנו):
Function : _Z5saxpyfPKfS0_Pfi
.headerflags ... SM90 ...
/*0000*/ IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
/*0010*/ S2R R2, SR_CTAID.X ; // blockIdx.x
/*0020*/ S2R R3, SR_TID.X ; // 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 ;
/*0060*/ IMAD.WIDE R4, R2, 0x4, c[0x0][0x210] ; // &x[i]
/*0070*/ IMAD.WIDE R6, R2, 0x4, c[0x0][0x208] ; // &y[i]
/*0080*/ LDG.E R4, desc[UR4][R4.64] ; // x[i]
/*0090*/ LDG.E R7, desc[UR4][R6.64] ; // y[i]
/*00a0*/ IMAD.WIDE R8, R2, 0x4, c[0x0][0x220] ; // &out[i]
/*00b0*/ FFMA R4, R4, c[0x0][0x200], R7 ; // a*x[i] + y[i]
/*00c0*/ STG.E desc[UR4][R8.64], R4 ; // out[i] = ...
/*00d0*/ EXIT ;
/*00e0*/ BRA 0xe0 ;
הזיהוי: S2R בשורות 0x10-0x20 (blockIdx.x, threadIdx.x), IMAD ב-0x30 (חישוב i), ISETP.GE.AND + @P0 EXIT ב-0x40-0x50 (בדיקת הגבולות), שתי LDG.E ב-0x80-0x90, FFMA ב-0xb0, STG.E ב-0xc0, ו-EXIT ב-0xd0. יש שתי LDG ואחת STG - בדיוק כמו בקוד: קוראים x[i] ו-y[i], כותבים out[i].
למה זה עבד: ה-SASS משקף ישירות את מבנה ה-kernel. כל גישה למערך global בקוד המקור הפכה לפקודת LDG/STG אחת (הטעינות אינן מאוחדות כי הן ממערכים שונים). blockDim.x מגיע ממאגר הקבועים c[0x0][0x0] ולא מ-Special Register משום שהוא ערך אחיד לכל ה-warp - הcompiler יודע זאת ומניח אותו בקבועים, בעוד threadIdx.x שונה בכל thread ולכן חייב להיקרא דרך S2R SR_TID.X.
איך להכליל: קריאת השלד של כל kernel מתחילה תמיד באותו דפוס: S2R לאינדקסים -> IMAD לחישוב האינדקס הגלובלי -> ISETP+predication לבדיקת הגבולות -> LDG/חישוב/STG. ברגע שאתם מזהים את הדפוס הזה, אתם יכולים לדלג מיד אל הלב החישובי של כל kernel שתפרקו.
פתרון תרגיל 2 - לזהות MUFU ו-HMMA¶
הקטע הרלוונטי של activate:
/*0070*/ LDG.E R0, desc[UR4][R2.64] ; // x[i]
/*0080*/ FMUL R0, R0, -1.4426950216 ; // -x * log2(e) (0xBFB8AA3B)
/*0090*/ MUFU.EX2 R0, R0 ; // 2^(-x*log2(e)) = e^(-x)
/*00a0*/ FADD R0, R0, 1 ; // 1 + e^(-x)
/*00b0*/ MUFU.RCP R0, R0 ; // 1 / (1 + e^(-x))
/*00c0*/ STG.E desc[UR4][R4.64], R0 ;
יש כאן שתי פקודות MUFU: MUFU.EX2 (אקספוננט בבסיס 2) ו-MUFU.RCP (הופכי). ה-expf(-x) הפך ל-MUFU.EX2 כי ה-hardware יודע לחשב רק חזקות של 2; הcompiler משתמש בזהות e^y = 2^(y*log2(e)), ולכן מקדים FMUL ב-log2(e) = 1.4427. גם החלוקה 1.0f / (...) הפכה ל-MUFU.RCP (הופכי מהיר) ולא לפקודת חלוקה מלאה.
הקטע הרלוונטי של wgemm:
הספירה המהירה:
$ cuobjdump --dump-sass kernels.o | grep -c -E 'FFMA'
1
$ cuobjdump --dump-sass kernels.o | grep -c -E 'MUFU'
2
$ cuobjdump --dump-sass kernels.o | grep -c -E 'HMMA'
2
$ cuobjdump --dump-sass kernels.o | grep -c -E 'LDG'
5
$ cuobjdump --dump-sass kernels.o | grep -c -E 'STG'
2
כפל מטריצות 16x16x16 הוא 16*16*16 = 4096 פעולות כפל-חיבור. במקום אלפי FFMA, הcompiler הוציא רק שתי פקודות HMMA.16816 (המספר 16816 = צורת ה-MMA של החומרה, M=16, N=8, K=16; ולכן צריך שתי פקודות, 16/8=2, כדי לכסות את tile ה-WMMA של 16x16x16), שכל אחת מפעילה את ה-Tensor Core על tile שלם. זה החיסכון העצום שראינו ב-1.5: פקודה אחת שמבצעת מטריצה שלמה במקום לולאה ארוכה של FFMAs.
למה זה עבד: כל פונקציה מתמטית "יקרה" בקוד מתגלה ב-SASS כפקודת חומרה ספציפית. expf -> MUFU.EX2, חלוקה -> MUFU.RCP, כפל מטריצות דרך WMMA -> HMMA. אלו רמזי ביצועים ישירים: המון MUFU בלולאה חמה מסמן עומס על ה-SFU (יחידה נדירה, ראו 1.3), ו-HMMA מסמן שאתם באמת על ה-Tensor Cores ולא נופלים בטעות ל-CUDA Cores.
איך להכליל: אחרי profiling, grep על ה-SASS למשפחות המעניינות הוא בדיקת-שפיות מהירה. רוצים לוודא שה-GEMM שלכם רץ על Tensor Cores? חפשו HMMA/IMMA. חושדים שספריית מתמטיקה מוסיפה עלות סמויה? ספרו MUFU. השוואת ספירות בין גרסאות קוד חושפת מיד מה השתנה ברמת החומרה.
פתרון תרגיל 3 - מיפוי SASS למקור עם -lineinfo¶
הרצת nvdisasm -g kernels.cubin, הקטע של saxpy:
//## File "kernels.cu", line 7
/*0010*/ S2R R2, SR_CTAID.X ;
/*0020*/ S2R R3, SR_TID.X ;
/*0030*/ IMAD R2, R2, c[0x0][0x0], R3 ;
//## File "kernels.cu", line 8
/*0040*/ ISETP.GE.AND P0, PT, R2, c[0x0][0x218], PT ;
/*0050*/ @P0 EXIT ;
//## File "kernels.cu", line 9
/*0060*/ IMAD.WIDE R4, R2, 0x4, c[0x0][0x210] ;
/*0080*/ LDG.E R4, desc[UR4][R4.64] ;
/*0090*/ LDG.E R7, desc[UR4][R6.64] ;
/*00b0*/ FFMA R4, R4, c[0x0][0x200], R7 ;
/*00c0*/ STG.E desc[UR4][R8.64], R4 ;
השורה out[i] = a * x[i] + y[i]; (שורה 9) התמפתה לבלוק שכולל את חישוב הכתובות (IMAD.WIDE), שתי הטעינות (LDG.E), ה-FFMA, וה-STG.E. עם cuobjdump --dump-sass -lineinfo kernels.o מקבלים מיפוי דומה, בפורמט הערות מעט שונה. חלק מפקודות ה-IMAD.WIDE עשויות להתמפות לשורת ההכרזה או ללולאה בגלל ש-ptxas הזיז אותן (instruction scheduling).
למה זה עבד: הדגל -lineinfo משבץ ב-cubin טבלה שממפה כל טווח פקודות SASS לקובץ ולשורת המקור. nvdisasm -g קורא את הטבלה הזו ומדפיס את שורת המקור כהערה מעל כל בלוק. בלי -lineinfo, הטבלה חסרה ואין מיפוי. הפיזור ה"לא צפוי" של פקודות בין שורות נובע מכך שהcompiler מסדר מחדש (reorders) פקודות כדי להסתיר latency - למשל מזיז IMAD.WIDE מעלה כדי שחישוב הכתובת יסתיים לפני שצריך אותה.
איך להכליל: בכל profiling אמיתי (Nsight Compute, פרק 8) קומפלו תמיד עם -lineinfo. כך, כשהמפרופיילר מצביע על "השורה החמה", תוכלו לקפוץ מיד מהפקודה ב-SASS אל השורה ב-CUDA C++ שיצרה אותה. -lineinfo זול (לא מכבה אופטימיזציות), בניגוד ל--G שמיועד לניפוי ומאט הכל.
פתרון תרגיל 4 - גרף זרימת-בקרה עם nvdisasm -cfg¶
מבנה ה-CFG של saxpy (תיאור טקסטואלי של הצמתים):
┌──────────────────────────────────────┐
│ BB0: IMAD.MOV R1, ... │
│ S2R R2, SR_CTAID.X │
│ S2R R3, SR_TID.X │
│ IMAD R2, R2, blockDim, R3 │
│ ISETP.GE.AND P0, PT, R2, n │
│ @P0 EXIT │
└───────────┬──────────────────┬────────┘
P0 false │ (i < n) │ P0 true (i >= n)
▼ ▼
┌───────────────────────┐ ┌──────────┐
│ BB1: IMAD.WIDE ... │ │ EXIT │
│ LDG.E x[i] │ └──────────┘
│ LDG.E y[i] │
│ FFMA a*x+y │
│ STG.E out[i] │
│ EXIT │
└───────────────────────┘
הבלוק הבסיסי הראשון (BB0) מסתיים ב-@P0 EXIT ויש לו שתי קשתות יוצאות: כש-P0 אמת (כלומר i >= n) הזרימה קופצת ל-EXIT והthread מסיים; כש-P0 שקר (כלומר i < n) היא נופלת (falls through) לבלוק החישוב. יש כאן שני-שלושה בלוקים בסיסיים, והפיצול היחיד הוא בדיקת הגבולות.
למה זה עבד: nvdisasm -cfg מזהה את גבולות הבלוקים הבסיסיים (רצף פקודות בלי כניסה/יציאה באמצע) ואת הקפיצות ביניהם, ובונה גרף DOT. בדיקת ה-if (i < n) היא מקור הפיצול היחיד ב-saxpy, ולכן ה-CFG פשוט וללא מעגלים - אין לולאות. ב-activate וב-wgemm (או בכל kernel עם לולאה) תראו back-edges - קשתות חוזרות שסוגרות מעגל, סימן ההיכר של לולאה.
איך להכליל: ה-CFG הוא הכלי לנתח warp divergence (פרק 8): כל פיצול הוא נקודה שבה threads שונים באותו warp עלולים ללכת למסלולים שונים, וה-hardware מריץ את שני המסלולים בטור. קריאת ה-CFG מזהה מראש היכן ה-divergence עלול לפגוע, ומיפוי מספר הבלוקים ללולאות עוזר להבין כמה עמוקה הזרימה.
פתרון תרגיל 5 - השוואת SASS בין שתי compute capabilities¶
saxpy על sm_70 (Volta), מקוצר:
/*0020*/ S2R R0, SR_CTAID.X ;
/*0028*/ S2R R2, SR_TID.X ;
/*0030*/ XMAD R0, R0, c[0x0][0x0], R2 ; // Volta: XMAD instead of IMAD
/*0040*/ ISETP.GE.AND P0, PT, R0, c[0x0][0x164], PT ;
/*0048*/ @P0 EXIT ;
/*0050*/ SHL R6, R0, 0x2 ;
/*0060*/ IADD ... ; // more manual address computation
/*0080*/ LDG.E R4, [R4] ; // Volta: without desc[UR..]
/*0090*/ LDG.E R7, [R6] ;
/*00b0*/ FFMA R4, R4, c[0x0][0x160], R7 ;
/*00c0*/ STG.E [R2], R4 ;
saxpy על sm_90a (Hopper) - כפי שראינו בתרגיל 1:
/*0080*/ LDG.E R4, desc[UR4][R4.64] ; // Hopper: descriptor in UR
/*00b0*/ FFMA R4, R4, c[0x0][0x200], R7 ;
/*00c0*/ STG.E desc[UR4][R8.64], R4 ;
שלושה הבדלים בולטים: (1) צורת הטעינה - LDG.E [R4] ב-Volta מול LDG.E desc[UR4][R4.64] ב-Hopper, שמעביר את הכתובת דרך מתאר זיכרון ב-Uniform Register. (2) שימוש ב-UR - Hopper מנצל את המסלול האחיד (UR4), Volta כמעט לא. (3) פקודת החשבון - Volta משתמש ב-XMAD (הכפל בן-16-הסיביות שלו) ולעתים ב-SHL+IADD לחישוב כתובת, בעוד Hopper מאחד לפקודת IMAD.WIDE אחת. גם ה-offsets במאגר הקבועים שונים (ה-ABI השתנה בין הדורות).
למה זה עבד: קוד ה-CUDA C++ זהה, וגם ה-PTX כמעט זהה (cuobjdump --dump-ptx על שתיהן יראה זאת) - אבל ה-SASS שונה כי הוא צמוד לארכיטקטורת ה-SM. ptxas מהדר את אותו PTX ל-SASS שונה לכל דור, מנצל את יכולות החומרה החדשות (המסלול האחיד, מתארי הזיכרון, IMAD.WIDE). זו בדיוק ההבחנה מ-4.1: PTX הוא היעד היציב והנייד, SASS הוא הריאליזציה הצמודה-לdevice ולא-היציבה.
איך להכליל: כשמנתחים ביצועים או מדווחים באג-compiler, תמיד ציינו את גרסת ה-SM המדויקת - sm_70 ו-sm_90a הם עולמות שונים. אותו עיקרון מסביר למה אי אפשר "לשמור" SASS מכוונן-יד ולהריץ אותו על דור הבא: הוא פשוט לא יתאים. לכן מכוונים ב-CUDA C++ או ב-in-line PTX, ונותנים ל-ptxas לייצר את ה-SASS המתאים לכל יעד.
פתרון תרגיל 6 (בונוס) - לזהות register spilling¶
ה-kernel הרעב לאוגרים וcompilation עם דוח ptxas:
$ nvcc -arch=sm_90a -Xptxas -v -cubin -o spill.cubin spill.cu
ptxas info : Compiling entry function '_Z9hungry...' for 'sm_90a'
ptxas info : Used 168 registers, used 0 barriers, 32 bytes spill stores, 32 bytes spill loads
מספר האוגרים הגיע לתקרה (168 קרוב למקסימום של 255 ל-thread), ודווחו spill stores/loads. בפירוק ה-SASS מופיעות פקודות ה-spill:
/*0210*/ STL [R1+0x4], R40 ; // Store Local - spilling a register to local memory
/*0220*/ ...
/*0340*/ LDL R40, [R1+0x4] ; // Load Local - reading it back
עם -maxrregcount=32:
$ nvcc -arch=sm_90a -maxrregcount=32 -Xptxas -v -cubin -o spill32.cubin spill.cu
ptxas info : Used 32 registers, 480 bytes spill stores, 480 bytes spill loads
מספר האוגרים ירד ל-32, אבל ה-spilling זינק מ-32 ל-480 בתים, ובהתאם מספר פקודות ה-STL/LDL ב-SASS גדל משמעותית.
למה זה עבד: ptxas אחראי על הקצאת האוגרים. כשה-kernel דורש יותר ערכים חיים בו-זמנית ממה שיש אוגרים, הוא "שופך" (spills) חלק מהם ל-local memory (שהוא פיזית global memory, איטי מאוד) - ואלו בדיוק פקודות ה-STL/LDL. -maxrregcount=32 כופה תקרה נמוכה, וכדי לעמוד בה ptxas נאלץ לשפוך עוד, ולכן ה-spilling גדל. זהו הטרייד-אוף המרכזי: פחות אוגרים לכל thread מאפשר יותר warps תושבים על ה-SM (occupancy גבוה יותר), אבל spilling מוגזם הורג את הרווח.
איך להכליל: STL/LDL ב-SASS, או bytes spill בדוח ptxas -v, הם הדגל האדום ל-register pressure. הפתרון הוא לצמצם את מספר המשתנים החיים (לפרק את ה-kernel, לצמצם unrolling), לא סתם להנמיך את -maxrregcount. את האיזון המדויק בין אוגרים ל-occupancy נחשב שיטתית בפרק 8.1; כאן ראינו איך הוא נראה מלמטה, בשפת החומרה.