לדלג לתוכן

8.3 Warp divergence ו scoreboard stalls תרגול

תרגול - Warp divergence ו-scoreboard stalls

בתרגול הזה תראו במו עיניכם את שתי המכשלות מהשיעור - התפצלות ה-warp ועצירות ה-scoreboard - חיות ב-SASS וב-Nsight Compute. תתחילו מ-kernel מתפצל: תהדרו אותו, תשלפו את ה-SASS ותזהו בו את שלישיית ה-FSETP / FADD / @P0 FMUL ואת נקודת ההתכנסות, ואז תמדדו ב-ncu את ה-branch efficiency ואת מספר ה-threads הפעילים לפקודה. אחר כך תוכיחו שבדיקת-הגבולות if (idx < n) היא כמעט חינמית (יעילות ~100%). בהמשך תבודדו את שני סוגי ה-scoreboard stall: long scoreboard ב-kernel memory-bound, ו-short scoreboard ב-kernel עתיר-טרנסצנדנטלים. לבסוף תלמדו לקרוא את סימון המחסום [B..:W..:S..] על זוג load-then-use אמיתי, ובבונוס תיתקלו במלכודת מיחזור ה-scoreboards. עבדו לפי הסדר. הריצו הכל על GPU של NVIDIA; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס (החליפו sm_90a בארכיטקטורה שלכם - למשל sm_75 ל-T4).

הכנה

צרו קובץ divergence.cu עם מקרו הבדיקה, ה-kernel המתפצל, וממלא נתונים אקראיים:

#include <cstdio>
#include <cstdlib>
#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)

__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (data[idx] > 0.5f) {
            data[idx] = data[idx] * 4.0f;   // Block A
        } else {
            data[idx] = data[idx] + 2.0f;   // Block B
        }
        data[idx] = data[idx] * data[idx];
    }
}

הcompilation: nvcc -O2 -arch=sm_90a -o divergence divergence.cu.

תרגיל 1 - לקרוא את ה-SASS של ההתפצלות

  1. הדרו את divergence.cu ושלפו את ה-SASS של ה-kernel: cuobjdump --dump-sass ./divergence (או הדרו ל-cubin עם nvcc -cubin -arch=sm_90a divergence.cu -o divergence.cubin ואז nvdisasm -c divergence.cubin).
  2. אתרו את פקודת ה-FSETP - איזה predicate היא קובעת, ומה התנאי (השוואה לאיזה קבוע)?
  3. אתרו את פקודת ה-FADD ואת פקודת ה-FMUL המותנית (@P0). איזו מהן מבוצעת עבור כל ה-threads, ואיזו רק עבור חלקם? קשרו כל אחת ל-Block A / Block B בקוד המקור.
  4. זהו את פקודת ההתכנסות - הפקודה שאחריה כל ה-threads שוב מריצים יחד (הריבוע data[idx]*data[idx]).
  5. הסבירו במשפט: מדוע ה-FADD מבוצעת גם עבור threads שערכם גדול מ-0.5, למרות שהתוצאה שלה נדרסת מיד?

רמז: חפשו בפלט את המחרוזות FSETP, FADD, FMUL. הפקודה המותנית מתחילה ב-@P0. הפקודה הלא-מותנית (בלי @) היא זו שרצה לכולם. זכרו את הפילוסופיה מההרצאה: הcompiler מעדיף לבזבז FADD זול מאשר להוסיף predication נוסף.

תרגיל 2 - למדוד branch efficiency והתפצלות ב-ncu

  1. הוסיפו main שמקצה מערך של N = 1<<20 floats, ממלא אותו בערכים אקראיים ב-[0,1) (עם rand()/RAND_MAX), מעתיק ל-device, ומשיק את divergent_kernel<<<(N+255)/256, 256>>>(d_data, N). עטפו כל קריאת CUDA ב-CUDA_CHECK.
  2. הריצו את ה-branch efficiency: ncu --metrics smsp__sass_average_branch_targets_threads_uniform.pct ./divergence.
  3. הריצו את מדד ה-threads הפעילים לפקודה: ncu --metrics smsp__thread_inst_executed_per_inst_executed.ratio ./divergence. כמה threads פעילים בממוצע מתוך 32?
  4. הסבירו את שני המספרים. מדוע ה-branch efficiency יוצאת סביב 50% ולא 0% (רמז: יש שני ענפים ב-kernel, לא אחד)? מדוע ה-active threads יוצאים סביב 16?

רמז: הנתונים אקראיים אחידים, ולכן התנאי הפנימי data > 0.5 מפצל כל warp כמעט תמיד לחצי-חצי. אבל בדיקת-הגבולות החיצונית idx < n אחידה כמעט תמיד. ה-branch efficiency היא ממוצע על שני הענפים. ה-.ratio של active threads הוא מספר בין 0 ל-32.

תרגיל 3 - להוכיח שבדיקת-הגבולות כמעט חינמית

  1. כתבו bounds_only.cu: kernel שכל מה שהוא עושה הוא בדיקת-גבולות וכתיבה אחידה, בלי תנאי תלוי-נתונים:
    __global__ void bounds_only(float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) out[idx] = idx * 2.0f;
}
  2. השיקו אותו עם n שאיננו כפולה של 32, כדי להבטיח שיהיה warp אחד שרוכב על הגבול - למשל n = 1000003, blockDim = 256.
  3. מדדו: ncu --metrics smsp__sass_average_branch_targets_threads_uniform.pct ./bounds_only.
  4. חשבו ידנית את ה-branch efficiency הצפויה: כמה warps רצים בסך הכל, כמה מתוכם מתפצלים, ומה היחס? השוו למספר שמדדתם.
  5. נסחו את המסקנה: האם "יש לי if ב-kernel" פירושו בהכרח התפצלות?

רמז: gridDim = ceil(1000003/256) = 3907, כלומר 3907 x 256 = 1,000,192 threads = 31,256 warps. רק ה-warp שרוכב על idx = 1,000,003 מתפצל. היעילות היא (31256-1)/31256. אתם אמורים לקבל ~99.99%.

תרגיל 4 - long scoreboard מול short scoreboard

  1. כתבו שני kernels בקובץ stalls.cu:
  2. mem_bound: memory-bound טהור - out[i] = in[i] על מערך גדול (הרבה LDG/STG, מעט חישוב).
  3. transcendental: עתיר-SFU - כל thread מריץ לולאה של הרבה sinf/expf/sqrtf על ערך יחיד באוגר, עם גישת זיכרון מינימלית. למשל:
    __global__ void transcendental(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float x = in[i];
        for (int k = 0; k < 64; k++) x = sinf(x) + sqrtf(fabsf(x)) + expf(x*0.01f);
        out[i] = x;
    }
}
  4. פרפלו כל אחד עם פירוק סיבות העצירה: ncu --section WarpStateStats ./stalls. (או הריצו ncu --set full וקִראו את הטבלה "Warp State Statistics".)
  5. עבור mem_bound - מהי סיבת העצירה הדומיננטית? עבור transcendental - מהי?
  6. שלפו את ה-SASS של transcendental (cuobjdump --dump-sass) ואתרו את פקודות ה-MUFU (למשל MUFU.SIN, MUFU.EX2, MUFU.RSQ). קשרו אותן ל-short scoreboard.

רמז: LDG/STG יוצאים מה-SM אל ה-HBM ולכן גורמים ל-Stall Long Scoreboard; MUFU.* הן פעולות SFU תוך-SM עם latency משתנה ולכן גורמות ל-Stall Short Scoreboard. sinf/expf/sqrtf מתהדרות ל-MUFU.SIN/MUFU.EX2/MUFU.RSQ בתוספת חישובי range-reduction.

תרגיל 5 - לקרוא את סימון המחסום על זוג load-then-use

  1. כתבו kernel קטן שטוען שני ערכים משני מצביעים בלתי-תלויים ואז מחבר/מכפיל אותם:
    __global__ void load_use(const int* a, const int* b, int* c, int i) {
    c[i] = a[i] * b[i];   // two independent LDGs, then use both
}
  2. שלפו את ה-SASS עם סימון המחסום: nvcc -cubin -arch=sm_90a stalls.cu -o k.cubin && nvdisasm -c k.cubin (או cuobjdump --dump-sass ./stalls).
  3. אתרו את שתי פקודות ה-LDG ואת פקודת השימוש (IMAD/IMUL/STG). לכל אחת, קִראו את השדות בסוגריים [B..:R..:W..:..:S..].
  4. ענו: איזה scoreboard כל LDG מסמן (שדה ה-W)? על איזה scoreboard פקודת השימוש חוסמת (שדה ה-B)? הסבירו כיצד המחסום מבטיח שהכפל לא יונפק לפני ששני ה-loads חזרו.
  5. אם שני ה-LDG השתמשו באותו scoreboard בשדה ה-W - הסבירו מה המשמעות לגבי המתנת פקודת השימוש.

רמז: הפורמט הוא [B<barrier>:R<read>:W<write>:<yield>:S<stall>]. W2 פירושו "מסמן scoreboard 2"; B--2--- פירושו "המתן ל-scoreboard 2 שיתנקה". שדה ה-S הוא מספר מחזורי-העצירה הסטטי. אם שני ה-loads מסמנים את אותו scoreboard, חסימה עליו ממתינה לשניהם.

תרגיל 6 (בונוס) - מלכודת מיחזור ה-scoreboards

  1. הריצו ncu --set full על ה-kernel transcendental מתרגיל 4 והסתכלו בסיבות העצירה. אם ה-kernel מערבב גם MUFU (short) וגם LDG (long), האם ncu מפריד ביניהן נקי?
  2. שלפו את ה-SASS ובדקו: כמה scoreboards שונים באמת בשימוש (הסתכלו על שדות ה-W השונים)? זכרו שיש רק 6.
  3. בנו מקרה שבו load (long) ופעולת MUFU (short) חולקים בטעות את אותו scoreboard, והסבירו מדוע ncu עלול לתייג את העצירה בסיווג הלא-נכון.
  4. נסחו כלל אצבע: מתי סומכים על תווית ה-short/long של ncu, ומתי חוזרים ל-SASS כמקור-האמת?

רמז: ה-glossary מזהיר במפורש שמיחזור scoreboards עלול לגרום ל-ncu לבלבל short ו-long כשהם משתמשים באותו scoreboard. עם 6 scoreboards בלבד ל-warp, הcompiler ממחזר אותם בקוד עמוס. כשהמספרים מפתיעים - cuobjdump --dump-sass, אתרו את הפקודה החמה, וקִראו ישירות את שדה ה-B ואת הפקודה שסימנה את אותו scoreboard בשדה ה-W.