1.3 יחידות ביצוע CUDA cores, SFU ו LSU תרגול

תרגול - יחידות ביצוע - CUDA cores, SFU ו-LSU

בתרגול הזה תלמדו לזהות באיזה pipe משתמש כל קוד, לא לפי תחושת בטן אלא לפי ראיות. תקמפלו kernel שקורא ל-expf ו-sqrtf ותצודו את הוראות ה-MUFU.* ב-SASS עם cuobjdump, תחשבו ביד את יחס תפוקת ה-FP64:FP32 של H100 ותאמתו מול ה-whitepaper, תסווגו כמה kernels לפי הpipe שהם מלחיצים, ולבסוף תוכיחו במדידה עם ncu ש-kernel memory-bound משאיר את יחידות החשבון בטלות. עבדו לפי הסדר - כל תרגיל בונה אינטואיציה לבא אחריו. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס שלנו, אבל התרגילים ירוצו על כל GPU של NVIDIA (התאימו את sm_90a).

הכנה

צרו קובץ pipes.cu עם המאקרו לבדיקת שגיאות וכמה kernels שנשתמש בהם לאורך התרגול. שימו לב שאין כאן main מלא לכל תרגיל - חלק מהתרגילים רק מהדרים ובודקים SASS ולא מריצים.

#include <cstdio>
#include <cmath>

#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)

// transcendental kernel: routed to SFU (MUFU.*)
__global__ void transcendental(const float* x, float* y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float v = x[i];
        y[i] = expf(v) + sqrtf(v) + sinf(v);   // EX2 + RSQ/SQRT + SIN
    }
}

// memory-bound kernel: pure copy, almost no arithmetic
__global__ void copyKernel(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i];             // one load, one store, zero ALU
}

// compute-bound kernel: lots of FMA on the same value, one memory access
__global__ void computeKernel(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float v = in[i], acc = 0.0f;
        #pragma unroll
        for (int k = 0; k < 256; ++k) acc = acc * 1.001f + v;  // heavy FFMA
        out[i] = acc;
    }
}

הקומפילציה הבסיסית (החליפו sm_90a בארכיטקטורה שלכם - למשל sm_75 ל-T4):

nvcc -O2 -arch=sm_90a -c pipes.cu -o pipes.o

כדי לברר את ה-compute capability:

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

תרגיל 1 - ציד הוראות MUFU ב-SASS

1. הדרו את pipes.cu לאובייקט: nvcc -O2 -arch=sm_90a -c pipes.cu -o pipes.o.
2. הוציאו את ה-SASS של ה-kernel הטרנסצנדנטי: cuobjdump --dump-sass pipes.o.
3. סננו רק את הוראות ה-SFU: העבירו את הפלט דרך grep MUFU.
4. זהו איזו הוראת MUFU.* מתאימה לכל קריאה: expf, sqrtf, sinf. אילו הוראות ראיתם, וכמה מכל אחת?
5. הדרו שוב עם -use_fast_math והשוו את פלט ה-MUFU - מה השתנה במספר ההוראות סביב כל טרנסצנדנטית?

רמז: expf לא מיושם בהוראת "exp" - חפשו MUFU.EX2 (חזקת בסיס 2) עם כפל ב-log2(e) לפניו. sqrtf המדויק בנוי סביב MUFU.RSQ (הופכי שורש) ועידון Newton-Raphson; sinf נותן MUFU.SIN. עם -use_fast_math הטרנסצנדנטיות מתקצרות משמעותית כי הקוד המתקן (הענפים והעידון) נעלם.

תרגיל 2 - יחס התפוקה FP64:FP32 ביד

נתון על H100 SXM: 132 SM, 128 נתיבי FP32 ל-SM, 64 נתיבי FP64 ל-SM, שעון boost של בערך 1.98 GHz. כל נתיב מבצע FMA (כפל+חיבור, שתי פעולות נקודה צפה) במחזור.

1. חשבו את תפוקת השיא ב-FP32 ב-TFLOPS: 132 x 128 x 2 x 1.98e9.
2. חשבו את תפוקת השיא ב-FP64 ב-TFLOPS: 132 x 64 x 2 x 1.98e9.
3. חשבו את היחס FP64:FP32. הסבירו במילה אחת מאיפה היחס נובע.
4. אמתו את שני המספרים מול טבלת המפרט ב-whitepaper של H100 (חפשו "NVIDIA H100 Tensor Core GPU Architecture"). מצאו את השורות "FP32" ו-"FP64". שימו לב לשורה נוספת "FP64 Tensor Core" - הסבירו למה היא לא רלוונטית לחישוב הזה.

רמז: היחס אמור לצאת עגול ופשוט. הוא נובע ישירות ממספר הנתיבים - חצי נתיבי FP64 מול FP32. שורת ה-"FP64 Tensor Core" מתייחסת לpipe אחר (ה-Tensor Core), לא לנתיבי ה-FP64 הרגילים שעליהם מדובר כאן.

תרגיל 3 - איזה pipe כל kernel מלחיץ

לכל אחד מקטעי ה-kernel הבאים, קבעו מהו הpipe הדומיננטי (FP32 / FP64 / INT32 / SFU / LSU) והסבירו בשורה למה. אל תריצו - נתחו את הקוד.

// A
out[i] = a[i] * b[i] + c[i];                       // float

// B
double s = 0.0;
for (int k = 0; k < N; ++k) s += A[k] * B[k];      // double

// C
out[i] = expf(a[i]) / (1.0f + expf(-a[i]));        // sigmoid-ish

// D
for (int k = 0; k < 64; ++k) out[base + k] = in[base + k];   // block copy

// E
int h = (idx * 2654435761u) >> 16;                 // hashing, integers
out[h & mask] = idx;

1. סווגו כל אחד מ-A עד E.
2. לאילו שניים מהם המספר "128 CUDA cores" מטעה במיוחד, ולמה?
3. עבור C, הסבירו למה הוא עלול להיחסם למרות שהוא נראה "קליל".

רמז: חפשו את הפעולה שחוזרת הכי הרבה. double מפנה לpipe ה-FP64 (חצי נתיבים). expf מפנה ל-SFU הנדיר. העתקה טהורה מפנה ל-LSU. חשבון שלמים מפנה לנתיבי ה-INT32 (גם הם חצי מ-FP32 ב-H100).

תרגיל 4 - להוכיח שה-ALU יושב בטל

כאן נמדוד בפועל. נשתמש ב-ncu כדי להשוות ניצול pipes בין ה-copyKernel (memory-bound) ל-computeKernel (compute-bound). הוסיפו main שמקצה מערכים, ממלא אותם, ומריץ כל kernel פעם אחת על מיליון איברים (השתמשו בשלד vecadd מהתרגול 0.3 והחליפו את שם ה-kernel).

1. הדרו עם מידע קווים ל-ncu: nvcc -O2 -lineinfo -arch=sm_90a pipes.cu -o pipes.
2. פרופלו את שני ה-kernels וקבלו ניצול pipes:

ncu --metrics \
  sm__inst_executed_pipe_lsu.avg.pct_of_peak_sustained_active,\
  sm__inst_executed_pipe_fma.avg.pct_of_peak_sustained_active \
  ./pipes

1. השוו: ב-copyKernel, מי גבוה - pipe ה-LSU או pipe ה-FMA? ב-computeKernel, מי גבוה?
2. הוסיפו את המדד gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed וראו אצל מי רוחב הפס של ה-HBM כמעט רווי.
3. נסחו במשפט: מה בדיוק "יושב בטל" ב-copyKernel, ולמה זה מבזבז את רוב ה-GPU.

רמז: copyKernel הוא load ו-store בלבד - אפס FMA. תצפו לראות ניצול LSU/DRAM גבוה וניצול FMA שואף לאפס. זו בדיוק ההוכחה שהוראות זיכרון לא עוברות דרך ה-CUDA Cores: הpipe החישובי בטל בזמן שה-LSU וה-HBM עמוסים.

תרגיל 5 (בונוס) - צוואר בקבוק ה-SFU

1. הריצו את ה-transcendental על מיליון איברים ופרופלו את pipe ה-SFU (הטרנסצנדנטי נקרא xu ב-ncu):

ncu --metrics \
  sm__inst_executed_pipe_xu.avg.pct_of_peak_sustained_active,\
  sm__inst_executed_pipe_fma.avg.pct_of_peak_sustained_active \
  ./pipes

1. החליפו במקור את expf/sinf בגרסאות המהירות __expf/__sinf (intrinsics), הדרו והריצו שוב את הפרופיילינג. מה קרה לזמן ה-kernel ולניצול ה-xu?
2. הסבירו: אם ה-SFU רווי ב-100% אבל ה-FMA נמוך, איזה שינוי אלגוריתמי יעזור - להוסיף עוד נתיבי חישוב, או להפחית טרנסצנדנטיות?

רמז: הצורה המהירה __expf היא הוראת MUFU.EX2 בודדת בלי קוד תיקון, אז היא מפחיתה הוראות אבל עדיין מעמיסה על אותו pipe xu נדיר. אם ה-xu רווי, הוספת FP32 לא תעזור בכלל - צריך פשוט פחות קריאות טרנסצנדנטיות, או לפרוס אותן על יותר warps כדי להסתיר את ה-latency.