לדלג לתוכן

4.2 SASS האסמבלי של החומרה תרגול

תרגול - SASS - האסמבלי של החומרה

בתרגול הזה תלמדו לפרק kernels מהודרים ולקרוא את ה-SASS שהחומרה באמת מריצה. תפרקו kernel עם cuobjdump --dump-sass ותזהו את משפחות הפקודות שהכרנו בהרצאה, תאלצו את הcompiler להוציא MUFU ו-HMMA, תקמפלו עם -lineinfo ותמפו שורת SASS חמה חזרה לשורת המקור עם nvdisasm, תייצרו גרף זרימת-בקרה (CFG) ותקראו את הענפים, ולבסוף תשוו SASS בין שתי compute capabilities שונות. עבדו לפי הסדר - כל תרגיל נשען על הקודם. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 (sm_90a) הוא כרטיס הייחוס שלנו; אם אתם על T4 החליפו sm_90a ב-sm_75 (וזכרו ש-sm_75 לא תומך ב-wgmma; ה-HMMA של תרגיל 2 עדיין יעבוד).

הכנה

צרו קובץ kernels.cu עם שלושה kernels: saxpy (ייצר FFMA ו-LDG/STG), activate (ייצר MUFU דרך expf), ו-wgemm (ייצר HMMA דרך WMMA על Tensor Cores). הקובץ אינו חייב לרוץ נכון מבחינה אלגוריתמית - אנחנו רק מפרקים את ה-SASS שלו - אבל הוא קומפילבילי לחלוטין.

#include <cuda_runtime.h>
#include <mma.h>
using namespace nvcuda;

// a. FFMA + LDG/STG
__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];
}

// b. MUFU (special function) - expf runs on the SFU
__global__ void activate(const float* x, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        out[i] = 1.0f / (1.0f + expf(-x[i]));   // sigmoid
}

// c. HMMA - 16x16x16 matrix multiply on Tensor Cores
__global__ void wgemm(const half* A, const half* B, float* C) {
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c;
    wmma::fill_fragment(c, 0.0f);
    wmma::load_matrix_sync(a, A, 16);
    wmma::load_matrix_sync(b, B, 16);
    wmma::mma_sync(c, a, b, c);
    wmma::store_matrix_sync(C, c, 16, wmma::mem_row_major);
}

שתי צורות הcompilation שנשתמש בהן לאורך התרגול - בינארי-host (ל-cuobjdump) ו-cubin ישיר (ל-nvdisasm):

# full host binary, with line info
nvcc -O2 -arch=sm_90a -lineinfo -c -o kernels.o kernels.cu

# direct cubin (the input to nvdisasm), with line info
nvcc -O2 -arch=sm_90a -lineinfo -cubin -o kernels.cubin kernels.cu

תרגיל 1 - פירוק וזיהוי משפחות פקודות

  1. פרקו את ה-SASS של הבינארי: cuobjdump --dump-sass kernels.o.
  2. אתרו את הקטע של הפונקציה saxpy (חפשו את השם המקושקש, למשל _Z5saxpyfPKfS0_Pfi). בתוכו, סמנו: את שתי פקודות ה-S2R שקוראות את blockIdx.x ו-threadIdx.x, את ה-IMAD שמחשב את האינדקס i, את ה-ISETP + @P0 של בדיקת הגבולות, את שתי טעינות ה-LDG.E, את ה-FFMA, את ה-STG.E, ואת ה-EXIT.
  3. ספרו: כמה LDG וכמה STG יש? האם המספר תואם למספר הגישות ל-global memory בקוד המקור?
  4. הסבירו לעצמכם למה blockDim.x מגיע ממאגר הקבועים c[0x0][...] ולא מ-Special Register כמו threadIdx.x.

רמז: כדי לסנן רק את ה-kernel שמעניין אתכם, הוסיפו -fun עם שם הפונקציה, או פשוט cuobjdump --dump-sass kernels.o | less וחפשו. אם אתם לא בטוחים מהו השם המקושקש, cuobjdump --dump-elf kernels.o | grep -i saxpy או cu++filt <symbol> יפענחו אותו. blockDim אחיד לכל ה-warp, ולכן הcompiler מניח אותו בקבועים; threadIdx שונה בכל thread ולכן נקרא מ-Special Register.

תרגיל 2 - לזהות MUFU ו-HMMA

  1. פרקו שוב את ה-SASS, הפעם של הפונקציה activate. אתרו את פקודת ה-MUFU - איזו וריאנטה היא (MUFU.EX2? MUFU.RCP?). הסבירו למה expf הפך ל-MUFU.EX2 ולא ל-MUFU.EXP (רמז: e^x = 2^(x*log2(e))).
  2. פרקו את הפונקציה wgemm ואתרו את פקודות ה-HMMA. שימו לב לצורת הפקודה, למשל HMMA.16816.F32 - המספרים הם ממדי הtile (M, N, K) שה-Tensor Core מעבד בפקודה אחת.
  3. השוו את מספר פקודות ה-HMMA שנוצרו למספר פקודות ה-FFMA שהיו נדרשות לאותו כפל מטריצות 16x16x16 על CUDA Cores רגילות. כמה FFMAs חוסך HMMA אחד?
  4. הריצו cuobjdump --dump-sass kernels.o | grep -c -E 'FFMA|MUFU|HMMA|LDG|STG' וקבלו ספירה מהירה של כל אחת.

רמז: אם אתם על T4 (sm_75), ה-wgemm עדיין מקמפל ומייצר HMMA (Tensor Cores קיימים מ-Volta), אבל ה-.16816 עשוי להיות .1688 בגלל ממדי tile שונים לפי דור. כפל מטריצות 16x16x16 הוא 16*16*16 = 4096 פעולות כפל-חיבור; HMMA יחיד מבצע tile שלם, בעוד על CUDA Cores היו נדרשות אלפי FFMA. זה בדיוק היתרון של Tensor Cores שראינו ב-1.5.

תרגיל 3 - מיפוי SASS למקור עם -lineinfo

  1. ודאו שקימפלתם עם -lineinfo (ראו ההכנה). הריצו nvdisasm -g kernels.cubin והתמקדו ב-saxpy.
  2. מצאו את השורה out[i] = a * x[i] + y[i]; בקוד המקור, ואתרו איזה בלוק פקודות SASS ממופה אליה. אילו פקודות SASS "נולדו" מהשורה הזו?
  3. עשו את אותו הדבר עם cuobjdump --dump-sass -lineinfo kernels.o (הדגל -lineinfo ל-cuobjdump מוסיף הערות מקור). השוו את שתי צורות המיפוי.
  4. מדוע פקודות בודדות עשויות להתמפות לשורות מקור "לא צפויות" (למשל פקודת IMAD שממופה לשורת ההכרזה של הלולאה)? מה עשה הcompiler?

רמז: nvdisasm -g דורש cubin שקומפל עם -lineinfo (או -G). אם לא רואים מספרי שורות, שכחתם את הדגל בcompilation. אופטימיזציות כמו reordering ו-instruction scheduling מפזרות פקודות שמקורן בשורה אחת על פני הפלט, ולכן המיפוי לא תמיד רציף - זו הסיבה שקוראים SASS לצד המקור ולא במקומו.

תרגיל 4 - גרף זרימת-בקרה עם nvdisasm -cfg

  1. ייצרו CFG של ה-cubin: nvdisasm -cfg kernels.cubin > kernels.dot.
  2. אם מותקן אצלכם graphviz, המירו לתמונה: dot -Tpng kernels.dot -o kernels.png (או -Tsvg). אם לא, פתחו את קובץ ה-.dot בעורך טקסט וקראו את הצמתים והקשתות ידנית.
  3. ב-saxpy, אתרו את הבלוק הבסיסי (basic block) שמסתיים ב-@P0 EXIT. לאן מובילות שתי הקשתות היוצאות ממנו - מה קורה כש-P0 אמת ומה כשהוא שקר?
  4. תארו במילים את מבנה ה-CFG של saxpy: כמה בלוקים בסיסיים יש, והיכן נמצא הפיצול (branch) היחיד.

רמז: ב-CFG כל צומת הוא בלוק בסיסי (רצף פקודות בלי קפיצה פנימה או החוצה באמצע), וכל קשת היא קפיצה אפשרית. בדיקת הגבולות if (i < n) יוצרת פיצול: מסלול אחד ל-EXIT (כש-i>=n), מסלול שני להמשך החישוב. ל-saxpy אין לולאות, ולכן ה-CFG שלו פשוט - בלי מעגלים (back-edges). את ה-wgemm וה-activate השוו כדי לראות CFG עשירים יותר.

תרגיל 5 - השוואת SASS בין שתי compute capabilities

  1. קמפלו את saxpy פעמיים - פעם ל-Volta (sm_70) ופעם ל-Hopper (sm_90a):
    nvcc -O2 -arch=sm_70  -cubin -o saxpy_sm70.cubin kernels.cu
nvcc -O2 -arch=sm_90a -cubin -o saxpy_sm90.cubin kernels.cu
  2. פרקו את שתיהן (cuobjdump --dump-sass saxpy_sm70.cubin וכן ל-sm90) והעמידו את ה-saxpy זו מול זו.
  3. אתרו לפחות שלושה הבדלים. שימו לב במיוחד ל: צורת טעינת ה-global (LDG.E [R.64] מול LDG.E desc[UR..][R.64]), לשימוש באוגרי UR (uniform), ולמספר הפקודות הכולל.
  4. הסבירו למה ה-SASS שונה בכלל, בעוד קוד ה-CUDA C++ זהה לחלוטין - וקשרו זאת להבחנה PTX מול SASS מ-4.1.

רמז: מ-Hopper ואילך, טעינות global עוברות דרך מתאר זיכרון (memory descriptor) שיושב ב-Uniform Register, ולכן תראו desc[UR4][...] ב-sm_90a אבל לא ב-sm_70. המסלול האחיד (uniform datapath) בשל יותר בדורות החדשים. זה בדיוק העניין ש-SASS צמוד-ארכיטקטורה: אותו PTX (ואותו מקור) עובר JIT/compilation ל-SASS שונה לכל דור. אם רוצים לראות שה-PTX זהה בין הדורות, הריצו cuobjdump --dump-ptx על שתיהן והשוו.

תרגיל 6 (בונוס) - לזהות register spilling

  1. כתבו kernel "רעב לאוגרים": לולאה שמחזיקה הרבה משתנים חיים בו-זמנית, למשל 40 מצברים (float acc0..acc39) שכולם מתעדכנים בכל איטרציה. קמפלו עם nvcc -arch=sm_90a -Xptxas -v -cubin -o spill.cubin spill.cu וקראו את שורות ה-ptxas info - כמה registers ל-thread? האם דווח bytes spill stores/loads?
  2. פרקו את ה-SASS וחפשו פקודות STL/LDL (Store/Load Local) - אלו הן פקודות ה-spill: אחסון/טעינה של אוגרים שנשפכו ל-local memory (שהוא פיזית global memory).
  3. הוסיפו -maxrregcount=32 לcompilation, פרקו שוב, וראו כיצד מספר ה-STL/LDL גדל (יותר spilling) בזמן שמספר האוגרים לכל thread קטן. הסבירו את הטרייד-אוף מול ה-occupancy.

רמז: ptxas מקצה אוגרים ומנסה לאזן בין מעט אוגרים לכל thread (יותר occupancy) לבין spilling יקר. פקודות STL/LDL ב-SASS הן הדגל האדום ל-spill - ראינו ב-4.1 שהאוגרים יכולים "להישפך" ל-global memory במחיר ביצועים כבד. -maxrregcount כופה תקרת אוגרים ולכן מכריח יותר spilling כשהתקרה נמוכה מדי; זהו בדיוק חישוב ה-occupancy שנעמיק בו בפרק 8.1.