לדלג לתוכן

8.2 ניצולת SM, pipe ו issue efficiency תרגול

תרגול - ניצולת - SM, pipe ו-issue efficiency

בתרגול הזה תמדדו את כל ארבע שכבות היררכיית הניצולת במו ידיכם, על אותו קובץ מקור אחד. תתחילו מהתופעה המכוננת - kernel של block יחיד שמראה 100% ניצולת GPU ב-nvidia-smi אבל מתחת ל-1% ניצולת SM ב-ncu - ותתקנו אותה בשינוי תצורת הlaunch בלבד. משם תרדו שכבה: תפרופלו kernel חישובי ותקראו את ניצולת pipes ה-fma, ה-tensor וה-lsu כדי לראות איזה pipe שולט. אחר-כך תיצרו ביד kernel עם יעילות הנפקה נמוכה - שרשרת FFMA תלויה עם מעט warps - ותקשרו את זה ל-occupancy ולחוק Little משיעור 7.4. לבסוף תלכו את כל מסלול הדיבוג top-down על kernel אחד ותכתבו אבחנה. עבדו לפי הסדר - כל תרגיל בונה על קודמו. הריצו הכל על GPU אמיתי עם Nsight Compute מותקן; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס (132 SM), אבל כל תרגיל ירוץ על כל GPU של NVIDIA - רק המספרים המוחלטים ישתנו (על T4 החליפו 132 ב-40 SM).

הכנה

צרו קובץ profile.cu. הוא מכיל שני kernels - אחד חישובי (שרשרת FFMA, שולט בpipe ה-fma) ואחד זיכרוני (העתקה, שולט ב-lsu) - ו-main שבוחר kernel ותצורת launch מהפרמטרים בשורת-הפקודה. כך כל התרגילים משתמשים באותו בינארי ורק משנים ארגומנטים.

#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>

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

// pure compute load: dependent FFMA chain. dominates the fma pipe.
// the chain is intentionally dependent - so latency hiding (occupancy) affects issue efficiency.
__global__ void compute_fma(const float* in, float* out, int n, int iters) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float x = in[idx];
        const float a = 1.0000001f, b = 0.9999999f;
        #pragma unroll 1
        for (int i = 0; i < iters; i++) x = x * a + b;   // dependent FFMA
        out[idx] = x;
    }
}

// memory load: one load, one store, almost no arithmetic. dominates the lsu.
__global__ void mem_stream(const float* in, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) out[idx] = in[idx] * 2.0f;
}

int main(int argc, char** argv) {
    const char* mode = argc > 1 ? argv[1] : "compute";
    int blocks  = argc > 2 ? atoi(argv[2]) : 132 * 8;   // default: flood all the SMs
    int threads = argc > 3 ? atoi(argv[3]) : 256;
    int iters   = argc > 4 ? atoi(argv[4]) : 4096;

    const int N = 132 * 8 * 256;                          // fixed size, enough for every configuration
    float *in, *out;
    CUDA_CHECK(cudaMalloc(&in,  (size_t)N * sizeof(float)));
    CUDA_CHECK(cudaMalloc(&out, (size_t)N * sizeof(float)));
    CUDA_CHECK(cudaMemset(in, 0, (size_t)N * sizeof(float)));

    // warmup (ncu will skip it with -s 1)
    if (!strcmp(mode, "mem")) mem_stream<<<blocks, threads>>>(in, out, N);
    else                      compute_fma<<<blocks, threads>>>(in, out, N, iters);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    cudaEvent_t s, e;
    CUDA_CHECK(cudaEventCreate(&s));
    CUDA_CHECK(cudaEventCreate(&e));
    CUDA_CHECK(cudaEventRecord(s));
    if (!strcmp(mode, "mem")) mem_stream<<<blocks, threads>>>(in, out, N);
    else                      compute_fma<<<blocks, threads>>>(in, out, N, iters);
    CUDA_CHECK(cudaEventRecord(e));
    CUDA_CHECK(cudaEventSynchronize(e));
    float ms = 0.0f;
    CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
    printf("mode=%s blocks=%d threads=%d iters=%d -> %.3f ms\n",
           mode, blocks, threads, iters, ms);

    CUDA_CHECK(cudaFree(in));
    CUDA_CHECK(cudaFree(out));
    return 0;
}

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

nvcc -O2 -arch=sm_90a -o profile profile.cu

לאורך התרגול נשתמש ב-ncu כדי לפרופל רק את הlaunch השנייה (המדודה) ולדלג על החימום, עם -s 1 -c 1 (דלג 1, ספור 1). ריצת ncu דורשת הרשאות; על מכונת ענן/Colab ייתכן שתצטרכו sudo או דגל --target-processes all.

תרגיל 1 - ניצולת GPU מול ניצולת SM

  1. הריצו ידנית ./profile compute 1 256 40000 (block יחיד, iters גדול כדי שירוץ מספיק זמן). במקביל, בטרמינל שני, הריצו nvidia-smi dmon -s u והביטו בעמודת ה-sm %.
  2. עכשיו מדדו את ניצולת ה-SM האמיתית עם ncu:
    ncu -s 1 -c 1 \
  --metrics sm__cycles_active.avg.pct_of_peak_sustained_elapsed \
  ./profile compute 1 256 40000
  3. רשמו את שני המספרים זה-לצד-זה. חשבו ידנית מה צפוי: block יחיד רץ על SM אחד מתוך 132 - כמה זה באחוזים?
  4. תקנו: הריצו שוב עם grid שמציף את כל ה-SM - ./profile compute 1056 256 40000 (1056 = 132 x 8) - ומדדו שוב את ניצולת ה-SM. השוו.

רמז: block לעולם לא מתפצל בין SM, לכן block יחיד = SM אחד עובד, 131 בטלים. הצפי הוא 1/132 = 0.76%. nvidia-smi יראה sm % גבוה כי הוא רק שואל "היה SM פעיל בחלון הדגימה", לא "כמה SM היו פעילים". שימו לב לסיומת elapsed במדד - היא כוללת את המחזורים הבטלים במכנה, ולכן חושפת את ה-SM המובטלים (בניגוד ל-active שהיינו משתמשים בו ל-pipe).

תרגיל 2 - איזה pipe שולט - pipe utilization

  1. פרופלו את ה-kernel החישובי וקראו את שלושת pipes הביצוע במדידה אחת:
    ncu -s 1 -c 1 --metrics \
  sm__inst_executed_pipe_fma.avg.pct_of_peak_sustained_active,\
  sm__inst_executed_pipe_tensor.avg.pct_of_peak_sustained_active,\
  sm__inst_executed_pipe_lsu.avg.pct_of_peak_sustained_active \
  ./profile compute 1056 256 4096
  2. חזרו על אותה מדידה בדיוק על ה-kernel הזיכרוני: ./profile mem 1056 256.
  3. סדרו את שש התוצאות בטבלה 2x3 (compute מול mem, על פני fma/tensor/lsu). איזה pipe שולט בכל kernel?
  4. הסבירו במשפט: מדוע pipe_tensor הוא 0% בשני המקרים, למרות שה-H100 מלא ב-Tensor Cores?

רמז: compute_fma הוא לולאת x = x*a + b - כל איטרציה היא הוראת FFMA אחת על ה-CUDA Cores, ולכן pipe ה-fma יגיע גבוה וה-lsu יהיה זניח (load אחד בהתחלה, store אחד בסוף). mem_stream הפוך - load ו-store לכל thread כמעט בלי חשבון, ולכן ה-lsu שולט. ה-tensor נשאר 0 כי אף אחד מהשניים לא מנפיק הוראת HMMA/WMMA - pipe ה-Tensor מונע רק מ-wmma/mma או מספריות כמו cuBLAS (שיעור 6.1). זכרו: ניצולת pipe ממוצעת רק על מחזורים פעילים, אז קראו אותה יחד עם ניצולת ה-SM מתרגיל 1.

תרגיל 3 - יעילות הנפקה נמוכה ו-occupancy

  1. הריצו את ה-kernel החישובי עם מעט warps מאוד - warp אחד לכל block, block אחד בערך לכל SM:
    ncu -s 1 -c 1 --metrics \
  smsp__issue_active.avg.pct_of_peak_sustained_active,\
  sm__warps_active.avg.pct_of_peak_sustained_active \
  ./profile compute 132 32 20000

    המדד הראשון הוא יעילות ההנפקה (Issue Slot Utilization); השני הוא ה-occupancy שהושג בפועל (Achieved Occupancy).
  2. עכשיו הגדילו בהדרגה את מספר ה-warps התושבים והריצו שוב אחרי כל שינוי - ./profile compute 132 128 20000, ואז ./profile compute 528 256 20000. רשמו את שני המדדים בכל צעד.
  3. שרטטו טבלה: occupancy מול issue efficiency לאורך שלושת הצעדים. מה הקשר ביניהם?
  4. חוק Little בעיפרון: השרשרת x = x*a + b היא FFMA תלוי, ו-latency של FFMA הוא בערך 4 מחזורים. כמה warps כשירים לכל scheduler צריך בערך כדי להסתיר את זה ולמלא כל מחזור? (רמז: זו בדיוק המדידה שראיתם).

רמז: כשיש warp יחיד לכל scheduler, אחרי כל הנפקת FFMA ה-warp נתקע ~4 מחזורים עד שהתוצאה מוכנה (תלות), ואין warp אחר להנפיק ממנו - אז המתזמן יושב בטל ויעילות ההנפקה צונחת. כשמוסיפים warps, בכל מחזור יש warp כשיר אחר, החורים מתמלאים, ויעילות ההנפקה עולה. זה בדיוק חוק Little משיעור 7.4: כדי להסתיר latency של 4 מחזורים ב-throughput של הנפקה למחזור צריך בערך 4 warps כשירים לכל scheduler. שימו לב: occupancy נמוכה היא הסיבה, יעילות הנפקה נמוכה היא הסימפטום הנמדד.

תרגיל 4 - מסלול הדיבוג המלא - top-down

  1. בחרו kernel אחד "חולה" - הריצו ./profile compute 1 256 40000.
  2. הלכו את כל ארבע השכבות עליו, בסדר, במדידה אחת:
    ncu -s 1 -c 1 --metrics \
  sm__cycles_active.avg.pct_of_peak_sustained_elapsed,\
  sm__inst_executed_pipe_fma.avg.pct_of_peak_sustained_active,\
  smsp__issue_active.avg.pct_of_peak_sustained_active \
  ./profile compute 1 256 40000
  3. כתבו פסקת אבחנה קצרה (3-4 משפטים) שעונה: איפה בעץ ההחלטה מהשיעור אתם עוצרים, מהו הצוואר, ומה השורה הבאה שתשנו בקוד. הסבירו במפורש מדוע אסור להיתפות למספר ה-fma הגבוה.
  4. יישמו את התיקון שגזרתם, מדדו שוב, והראו איזו שכבה השתפרה.

רמז: המדד הראשון (ניצולת SM) יהיה זעיר (~0.76%), ולכן עוצרים כבר בשכבה 2 של העץ - אין טעם אפילו להסתכל על ה-pipe וה-issue. ה-fma עלול להראות מספר יפה כי הוא ממוצע רק על ה-SM היחיד שעובד ומחריג את 131 הבטלים - זו בדיוק מלכודת "המחזורים הפעילים" מההרצאה. השורה שתשנו היא תצורת הlaunch: <<<1, 256>>> -> <<<1056, 256>>>.

תרגיל 5 (בונוס) - למה nvidia-smi ו-ncu לא מסכימים

  1. הריצו את ה-kernel של block יחיד בלולאה ארוכה ברקע כך שירוץ כמה שניות רצוף (למשל עטפו את הlaunch בלולאת host, או הגדילו את iters ל-2000000 והריצו ./profile compute 1 256 2000000 &).
  2. בזמן שהוא רץ, קראו nvidia-smi --query-gpu=utilization.gpu --format=csv -l 1 (דגימה לשנייה).
  3. השוו את המספר הזה למדד sm__cycles_active...elapsed של ncu על אותו kernel.
  4. כתבו הסבר: מדוע nvidia-smi מראה ~100% בעוד ניצולת ה-SM האמיתית מתחת ל-1%? מהו בדיוק ההבדל במה שכל אחד מהם מודד, ולמה זה הופך את nvidia-smi לצעד ראשון בדיבוג אבל לא לצעד אחרון?

רמז: nvidia-smi מדווח ניצולת GPU/kernel - מדד גס, כמעט בינארי, שאומר רק "היה kernel פעיל בחלון הדגימה". הוא לא מבחין בין SM אחד עובד ל-132 עובדים. ncu מודד את ניצולת ה-SM האמיתית ברמת מחזור-שעון ממוצעת על כל 132 ה-SM, ולכן רואה את 131 הבטלים. זה בדיוק הפער מההרצאה: הקומה העליונה משקרת, הקומה השנייה חושפת.