לדלג לתוכן

4.1 PTX שפת הביניים הווירטואלית תרגול

תרגול - PTX - שפת הביניים הווירטואלית

בתרגול הזה תפתחו את "המותן הצר" ותראו אותו במו עיניכם: תפיקו PTX מ-kernel שכתבתם, תקראו את הצהרות ה-registers, את פעולת ה-fma ואת ה-registers המובנים, ותקשרו כל שורת PTX לשורת ה-CUDA C++ שהולידה אותה. אחר כך תשנו את הקוד ותראו את ה-PTX משתנה, תכתבו inline PTX משלכם, ולבסוף תוכיחו לעצמכם את התאימות-קדימה של PTX מול הקיבעון של SASS. רוב התרגילים כאן דורשים רק את הcompiler nvcc ולא GPU בפועל (הפקת PTX וקריאתו היא פעולת compilation), אבל תרגילים 4 ו-6 דורשים כרטיס NVIDIA להרצה - אם אין לכם, הריצו אותם ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס שלנו (sm_90a); על כרטיס אחר החליפו את הארכיטקטורה בהתאם. עבדו לפי הסדר - כל תרגיל בונה על הקודם.

הכנה

צרו קובץ square.cu עם ה-kernel מההרצאה:

#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call)                                                   \
    do {                                                                   \
        cudaError_t err_ = (call);                                         \
        if (err_ != cudaSuccess) {                                         \
            fprintf(stderr, "CUDA error %s at %s:%d\n",                    \
                    cudaGetErrorString(err_), __FILE__, __LINE__);         \
            exit(EXIT_FAILURE);                                            \
        }                                                                  \
    } while (0)

__global__ void square(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        out[i] = in[i] * in[i] + 1.5f;
}

בררו קודם את הארכיטקטורה של הכרטיס שלכם (אם יש), והשתמשו בה במקום sm_90a לכל אורך התרגול:

nvidia-smi --query-gpu=name,compute_cap --format=csv

תרגיל 1 - הפקת PTX וקריאת שלוש המשפחות

  1. הפיקו את ה-PTX של ה-kernel: nvcc -arch=sm_90a -ptx square.cu -o square.ptx.
  2. פתחו את square.ptx ומצאו את שורות ה-.reg. כמה משפחות registers הוצהרו, ומאיזה סוג כל אחת (.f32, .b32, .b64, .pred)? מה המשמעות של הסימון %f<3>?
  3. מצאו את הוראת ה-fma.rn.f32. איזה קבוע מופיע בה, ומה ערכו העשרוני? מה המשמעות של הסיומת rn?
  4. מצאו את שלוש ההוראות mov.u32 ... %tid.x / %ntid.x / %ctaid.x. ודאו שהן שם.

רמז: הסימון %f<3> מכריז על משפחת שמות %f0, %f1, %f2 (שלושה registers של float). הקבוע ב-fma הוא 0f3FC00000, וזהו 1.5 בייצוג ביטים של IEEE 754 float (הקידומת 0f = ליטרל float גולמי). הסיומת rn היא round-to-nearest-even, מצב העיגול של ברירת המחדל. אם אינכם רואים את שורות ה-mov של האינדקסים, ודאו שלא ביקשתם אופטימיזציה שהעלימה את ה-kernel - הן תמיד שם כשמשתמשים ב-threadIdx/blockIdx.

תרגיל 2 - מיפוי PTX לשורות ה-CUDA C++

זהו תרגיל קריאה (על נייר / בעורך). ליד כל שורת PTX בגוף ה-kernel, כתבו לאיזו שורת CUDA C++ היא שייכת.

  1. אתרו את ההוראה שמחשבת את האינדקס הגלובלי i. לאיזו הוראת PTX אחת התרגמו שלושת ה-mov בשילוב עם הכפל-חיבור? (רמז לשם ההוראה: mad.)
  2. אתרו את שתי ההוראות שמממשות את if (i < n). שימו לב: הcompiler הפך את התנאי - לאיזה תנאי, ולמה?
  3. אתרו את ld.global.f32 ואת st.global.f32. איזו שורת C++ כל אחת מממשת?
  4. ספרו: כמה גישות ל-global memory יש בגוף ה-kernel, מול כמה פעולות אריתמטיות (fma)? מה זה אומר על arithmetic intensity החישוב?

רמז: mad.lo.s32 %r1, %r4, %r3, %r5 הוא בדיוק i = blockIdx.x * blockDim.x + threadIdx.x (הכפל-חיבור על שלמים). setp.ge.s32 %p1, %r1, %r2 מחשב את התנאי ההפוך i >= n ואז @%p1 bra מדלג על הגוף - זה הדפוס הרגיל של קומפילציית if. שתי גישות global (ld אחת, st אחת) מול פעולה אריתמטית אחת = יחס נמוך = memory-bound.

תרגיל 3 - שינוי הקוד וצפייה בשינוי ה-PTX

  1. שנו את השורה ל-out[i] = in[i] * in[i] + 3.0f; (במקום 1.5f). הפיקו PTX מחדש. איזה קבוע מופיע עכשיו ב-fma? (תצטרכו את הייצוג ההקסדצימלי של 3.0f.)
  2. שנו את השורה ל-out[i] = in[i] * 2.0f; (כפל בלבד, בלי חיבור). האם עדיין מופיעה fma, או mul.f32? הסבירו.
  3. שנו את השורה ל-out[i] = in[i] * in[i] + in[i];. עכשיו כל שלושת האופרנדים של ה-fma הם registers ולא קבוע - ודאו זאת ב-PTX.
  4. הוסיפו לולאה קטנה בתוך ה-kernel, למשל for (int k = 0; k < 4; ++k) acc += in[i+k];. חפשו ב-PTX את מבנה הענף/הקפיצה (bra) ואת ה-registers של ה-predicate (.pred) שנוספו.

רמז: 3.0f בייצוג ביטים הוא 0f40400000. כשמכפילים ב-2.0f בלי חיבור, הcompiler יפלוט mul.f32 ולא fma (אין מה לצבור). כשכל האופרנדים registers, ה-fma תיראה fma.rn.f32 %fD, %fA, %fB, %fC בלי ליטרל 0f.... לולאה מוסיפה תווית (label) והוראת bra מותנית בפרדיקט - זה מבנה הבקרה של PTX.

תרגיל 4 - כתיבת inline PTX משלכם

זהו תרגיל מעשי שדורש הרצה על GPU. תכתבו inline PTX שקורא register מיוחד ותוודאו שהוא עובד.

  1. הוסיפו לקובץ פונקציית device שקוראת את %laneid (מיקום הthread בתוך ה-warp):
__device__ __forceinline__ unsigned lane_id() {
    unsigned id;
    asm volatile("mov.u32 %0, %%laneid;" : "=r"(id));
    return id;
}
  1. כתבו kernel קטן שכל thread כותב את ה-lane_id() שלו למערך פלט, launch block יחיד של 64 threads, והעתיקו את התוצאה לhost והדפיסו.
  2. הריצו. הthreads 0-31 (warp 0) צריכים להדפיס 0..31, וthreads 32-63 (warp 1) צריכים להדפיס שוב 0..31. הסבירו למה ה-laneid "מתאפס" כל 32.
  3. בונוס: כתבו inline PTX שני שסופר ביטים דולקים עם popc.b32 (ראו ההרצאה), והדפיסו popcount(threadIdx.x) לכמה threads.

רמז: שימו לב ל-%%laneid (שני סימני אחוז) - סימן % בודד שמור ל-placeholders כמו %0. האילוץ "=r" הוא פלט 32 סיביות. ה-laneid הוא threadIdx.x % 32 בבלוק חד-ממדי, ולכן הוא מתאפס בכל גבול warp. עבור popc: asm("popc.b32 %0, %1;" : "=r"(r) : "r"(x)); - כאן אין צורך ב-volatile כי הפעולה טהורה.

תרגיל 5 - PTX ISA version מול target

זהו תרגיל קריאה קצר על ההנחיות בראש קובץ ה-PTX.

  1. הפיקו PTX פעמיים, פעם עם -arch=sm_90a ופעם עם -arch=sm_80 (Ampere). השוו את שורת .target בשני הקבצים.
  2. הסתכלו על שורת .version. מה היא מייצגת - את הכרטיס, או את גרסת הcompiler שהפיק? האם היא השתנתה בין שתי ההפקות?
  3. הסבירו במילים שלכם: איזו משתי ההנחיות (.version או .target) קובעת על אילו GPU פיזיים אפשר יהיה להריץ את הקוד?

רמז: .target הוא ה-compute capability של הקוד (sm_90a מול sm_80) - הוא זה שקובע תאימות חומרה. .version הוא גרסת ה-PTX ISA של הcompiler שהפיק את הקובץ, והוא תלוי בגרסת ה-CUDA Toolkit המותקנת, לא בכרטיס. אותה גרסת compiler תפלוט אותו .version לשני הטרגטים.

תרגיל 6 (בונוס) - תאימות-קדימה: PTX מנצח SASS

זהו תרגיל מעשי שדורש הרצה על GPU, ומדגים בדיוק למה PTX שווה את הטרחה. הריצו על H100 (או על כל כרטיס - בחרו טרגט ישן מהכרטיס שלכם).

  1. קמפלו את square.cu כך שיטמיע רק SASS של Ampere (בלי PTX):
nvcc -gencode arch=compute_80,code=sm_80 -o sq_sass square.cu
  1. הריצו ./sq_sass על ה-H100 (הוסיפו main שמשיק את ה-kernel ובודק עם CUDA_CHECK). רשמו את השגיאה המדויקת שהמאקרו תופס.
  2. עכשיו קמפלו כך שיטמיע PTX של Ampere (ייצוג ביניים נייד):
nvcc -gencode arch=compute_80,code=compute_80 -o sq_ptx square.cu
  1. הריצו ./sq_ptx על ה-H100. הפעם היא צריכה לעבוד. הסבירו את ההבדל בין שתי ההרצות, והצביעו על השלב המדויק שבו ה-PTX הפך ל-SASS של Hopper.

רמז: SASS של sm_80 לא רץ על SM של Hopper, ולכן sq_sass נכשל עם no kernel image is available for execution on the device. sq_ptx מטמיע PTX של compute_80, וכשאין SASS מתאים ה-hardware driver מבצע לו JIT ל-SASS של sm_90 בזמן הטעינה - ולכן הוא רץ. זהו בדיוק מנגנון התאימות-קדימה: ה-JIT קורה ב-driver בזמן ריצה, לא ב-nvcc בזמן הבנייה.