לדלג לתוכן

8.5 Register pressure ואיזון משאבים תרגול

תרגול - Register pressure ואיזון משאבים

בתרגול הזה תעברו את כל הלולאה של איזון המשאבים במו ידיכם, על שני kernels אמיתיים. תתחילו בקריאת הפלט של -Xptxas -v - מספר האוגרים, ה-shared memory ובייטי הדליפה - ותלמדו לזהות kernel בריא מ-kernel לחוץ במבט. אחר כך תכפו דליפה על kernel כבד-אוגרים ותמדדו כמה היא מאטה. משם תיישמו __launch_bounds__ ו--maxrregcount כדי לשנות את מספר האוגרים, ותמדדו את ה-occupancy (עם cudaOccupancyMaxActiveBlocksPerMultiprocessor) ואת ה-runtime. השיא הוא תרגיל 4: תסרקו את מספר האוגרים ותמצאו את הערך שממקסם throughput - שכמעט אף פעם אינו הערך שממקסם occupancy. עבדו לפי הסדר; כל תרגיל בונה על קודמו. הריצו על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן (שיעור 0.2). ה-H100 הוא כרטיס הייחוס - המספרים המוחלטים ישתנו על כרטיס אחר, אבל המגמות (ובעיקר סף ה-32 האוגרים ל-occupancy מלא, שהוא ספציפי ל-H100) יישארו.

הכנה

צרו קובץ regpressure.cu עם מקרו-הבדיקה, kernel "קליל" (מעט אוגרים) ו-kernel "כבד" (הרבה אוגרים חיים). הקומפילציה (החליפו sm_90a בארכיטקטורה שלכם - למשל sm_75 ל-T4):

nvcc -O2 -arch=sm_90a -Xptxas -v -o regpressure regpressure.cu

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

// light kernel - few registers, easily full occupancy
__global__ void light(const float* __restrict__ in, float* __restrict__ out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) out[idx] = in[idx] * 2.0f + 1.0f;
}

// heavy kernel - 64 live accumulators at once, strongly tends toward many registers
__global__ void heavy(const float* __restrict__ in, float* __restrict__ out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n) return;
    float x = in[idx];
    float a[64];
    #pragma unroll
    for (int i = 0; i < 64; i++) a[i] = x + (float)i;
    #pragma unroll
    for (int k = 0; k < 48; k++) {
        #pragma unroll
        for (int i = 0; i < 64; i++)
            a[i] = a[i] * a[i] + a[(i + 1) & 63];   // all accumulators stay live
    }
    float s = 0.0f;
    #pragma unroll
    for (int i = 0; i < 64; i++) s += a[i];
    out[idx] = s;
}

תרגיל 1 - קריאת הפלט של ptxas -v

  1. קמפלו את regpressure.cu עם -Xptxas -v (השורה בהכנה). ה-nvcc ידפיס בלוק ptxas info נפרד לכל אחד משני ה-kernels.
  2. עבור כל kernel רשמו שלושה מספרים: Used ___ registers, כמות smem (אם מופיעה), וכמות spill stores / spill loads.
  3. לפי טבלת ה-occupancy מההרצאה (H100: 32 אוגרים = 100%, 64 = 50%, 128 = 25% ...), אמדו את ה-occupancy המקסימלי שכל kernel מאפשר מבחינת אוגרים בלבד.
  4. איזה מהשניים "בריא" ואיזה "לחוץ"? האם ל-heavy יש כבר דליפה במצב הזה (בלי שום דגל הגבלה)?

רמז: הפלט הולך ל-stderr, אז אם אתם מפנים אותו לקובץ השתמשו ב-2>. שלושת המספרים החשובים הם registers, spill stores ו-spill loads. אם spill stores/loads שווים 0 - אין דליפה. חשבו occupancy כ-floor(65536 / ceil(32*R/256)*256) / 64. ה-light צריך לצאת עם ~16 אוגרים; ה-heavy עם עשרות רבות של אוגרים (המספר המדויק תלוי בגרסת ה-compiler).

תרגיל 2 - לכפות דליפה ולמדוד את ההאטה

  1. הוסיפו ל-main קוד שמריץ את heavy על מערך גדול (למשל n = 4M), עם block של 256 threads, ומודד את הזמן עם cudaEvent (חימום + מינימום על 20 חזרות), כפי שעשינו בשיעורים קודמים.
  2. קמפלו בלי הגבלת אוגרים ומדדו את הזמן. רשמו גם את מספר האוגרים ובייטי הדליפה מ--Xptxas -v.
  3. קמפלו שוב עם -maxrregcount=32, שכופה תקרה של 32 אוגרים ל-thread. הסתכלו שוב ב--Xptxas -v: כמה בייטי spill stores / spill loads יש עכשיו? כמה אוגרים נרשמו?
  4. מדדו את הזמן של הגרסה עם -maxrregcount=32 והשוו לזמן מסעיף 2. מה קרה ל-runtime, למרות שה-occupancy אמור היה לעלות?
  5. הסבירו את הפרדוקס: איך ייתכן ש-occupancy גבוה יותר הביא runtime גרוע יותר?

רמז: תצטרכו שתי קומפילציות נפרדות (או שני קבצי הרצה) - אחת בלי הדגל ואחת עם -maxrregcount=32. הדליפה יורדת ל-local memory שהוא בעצם GPU RAM איטי (ראו ההרצאה). הביטו במספרי הדליפה: הם קפצו מ-0 (או מעט) לכמה מאות בייטים. כל ערך שדלף הפך גישת-אוגר-כמעט-חינם לגישת-זיכרון-של-מאות-מחזורים. זו הדגמה שהורדת אוגרים נמוך מדי היא הרסנית.

תרגיל 3 - למדוד occupancy ו-runtime עם launch_bounds

  1. הוסיפו לקוד פונקציה שמדפיסה את ה-occupancy התיאורטי של kernel נתון באמצעות ה-Runtime API:
void reportOccupancy(const char* name, const void* kernel, int threadsPerBlock) {
    int maxBlocks = 0;
    CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &maxBlocks, kernel, threadsPerBlock, 0));
    cudaDeviceProp p; CUDA_CHECK(cudaGetDeviceProperties(&p, 0));
    int activeWarps = maxBlocks * (threadsPerBlock / 32);
    int maxWarps    = p.maxThreadsPerMultiProcessor / 32;
    printf("%-16s  blocks/SM=%d  occupancy=%.1f%%\n",
           name, maxBlocks, 100.0 * activeWarps / maxWarps);
}
  1. צרו גרסה שנייה של ה-kernel הכבד עם __launch_bounds__, למשל:
__global__ void __launch_bounds__(256, 8)
heavy_lb(const float* __restrict__ in, float* __restrict__ out, int n) { /* body identical to heavy */ }
  1. קראו ל-reportOccupancy עבור heavy (המקורי) ועבור heavy_lb, שניהם עם threadsPerBlock = 256. מדדו גם את מספר האוגרים של שניהם ב--Xptxas -v.
  2. מדדו את ה-runtime של שני ה-kernels על אותו קלט. רשמו טבלה: kernel, אוגרים ל-thread, occupancy, runtime.
  3. האם heavy_lb הגיע ל-occupancy הגבוה יותר שכיוונתם אליו? האם ה-runtime שלו טוב או רע יותר מ-heavy? מה זה מלמד?

רמז: __launch_bounds__(256, 8) על H100 מכוון ל-65536 / (256*8) = 32 אוגרים ל-thread, כלומר occupancy מלא. שימו לב אם זה גרם לדליפה (הסתכלו ב--Xptxas -v על heavy_lb). ה-cudaOccupancyMaxActiveBlocksPerMultiprocessor מחזיר occupancy תיאורטי מהאוגרים וה-smem; זה בדיוק המספר מטבלת ההרצאה. אם ה-runtime של הגרסה עם ה-occupancy הגבוה גרוע יותר - זו התובנה של השיעור.

תרגיל 4 - למצוא את מספר האוגרים שממקסם throughput

  1. כתבו סקריפט bash שסורק ערכי -maxrregcount שונים, מקמפל את regpressure.cu מחדש לכל ערך, ומריץ את מדידת ה-heavy:
for R in 24 28 32 40 48 64 96 128 168 255; do
  nvcc -O2 -arch=sm_90a -maxrregcount=$R -o rp_$R regpressure.cu 2>/dev/null
  printf "maxrreg=%3d  " "$R"
  ./rp_$R
done
  1. לכל ערך, אספו שלושה מספרים: אוגרים ל-thread בפועל (מ--Xptxas -v), occupancy (מתרגיל 3), ו-runtime.
  2. שרטטו (או תארו במילים) את ה-runtime כפונקציה של מספר האוגרים. איפה המינימום?
  3. סמנו את שני הקצוות: בקצה הנמוך (מעט אוגרים) - מה מאט? בקצה הגבוה (הרבה אוגרים) - מה מאט?
  4. מצאו את מספר האוגרים שממקסם throughput. האם הוא זהה למספר שממקסם occupancy (הקצה הנמוך)? הסבירו למה לא.

רמז: הצפי הוא עקומת "U" (או "בטן"): מעט מדי אוגרים -> דליפה -> איטי; יותר מדי אוגרים -> occupancy נמוך -> latency לא מוסתר -> איטי; ובאמצע נקודת-מתיקות. הנקודה המיטבית כמעט אף פעם אינה בקצה ה-occupancy המרבי (24-32 אוגרים), כי שם יש דליפה. זו בדיוק המסקנה: מודדים throughput, לא occupancy. ודאו שהמדידה כוללת חימום ומינימום על כמה חזרות.

תרגיל 5 (בונוס) - אימות עם ncu וקריאת סטטיסטיקות local memory

  1. הריצו את הפרופיילר על שתי הגרסאות מתרגיל 2 (בלי הגבלה, ועם -maxrregcount=32):
ncu --set full --kernel-name heavy ./regpressure
  1. אתרו בפלט את המדדים: "Registers Per Thread", "Achieved Occupancy", ו-"Local Memory" (או סטטיסטיקות ה-spill). איך הם משתנים בין שתי הגרסאות?
  2. הסתכלו בסעיף סיבות ה-stall (Warp State). בגרסה שדולפת, האם עולה Stall Long Scoreboard (המתנה לזיכרון גלובלי, שאליו הלכה הדליפה)?
  3. הצליבו: הגרסה עם ה-occupancy הגבוה (המוגבל ל-32) - האם ה-throughput שלה (Compute / Memory Throughput) באמת גבוה יותר? קשרו למסקנה מתרגיל 4.
  4. (רשות) הריצו גם על heavy_lb מתרגיל 3, וודאו שה"Registers Per Thread" ו-"Achieved Occupancy" מתאימים לחישוב התיאורטי מתרגיל 3.

רמז: ncu מדווח "Registers Per Thread" ישירות - השוו למספר מ--Xptxas -v (הם צריכים להתאים). "Achieved Occupancy" הוא ה-occupancy בפועל בזמן ריצה, שעשוי להיות נמוך מהתיאורטי. הדליפה מופיעה כתעבורה ל-local memory ולרוב מרימה את Stall Long Scoreboard, כי local memory יושב ב-GPU RAM. זה מאשש מהמדידה את כל הסיפור של ההרצאה: אותו kernel, occupancy גבוה יותר, אבל throughput נמוך יותר בגלל הדליפה.