לדלג לתוכן

8.1 Occupancy ומצבי warp תרגול

תרגול - Occupancy ומצבי warp

בתרגול הזה תעברו את כל שרשרת ה-occupancy מקצה לקצה: קודם תחשבו occupancy תיאורטי ביד ותשחזרו את דוגמת ה-30% חסומת-shared-memory מההרצאה, אחר כך תאמתו אותה עם ה-Occupancy API בקוד, ואז תסתכלו לחומרה בעיניים ותאשרו את אותו מספר ב-ncu. משם תעשו את הניסוי החשוב באמת - סריקת גודל block ו-shared memory, ציור של occupancy מול זמן-ריצה, והגילוי במו ידיכם ששיא ה-occupancy אינו שיא הביצועים. לבסוף תקראו דוח ncu אמיתי ותזהו בו כל אחד מארבעת מצבי ה-warp. עבדו לפי הסדר - כל תרגיל בונה על קודמו. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אבל התרגילים ירוצו על כל GPU של NVIDIA (רק המספרים המוחלטים ומגבלות ה-SM ישתנו - את שלכם תשלפו מ-deviceQuery או מ-cudaGetDeviceProperties).

הכנה

צרו קובץ occ.cu עם מקרו הבדיקה ו-kernel נסיוני שצורך shared memory מדיד. נשתמש ב-shared memory דינמי כדי שנוכל לשנות את גודלו בזמן הlaunch בלי לקמפל מחדש:

nvcc -O2 -arch=sm_90a -o occ occ.cu      # replace with sm_75 on T4

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

// simple kernel that consumes dynamic shared memory (its size is set at launch)
extern __shared__ float smem[];
__global__ void work(const float* in, float* out, int n, int iters) {
    int t = blockIdx.x * blockDim.x + threadIdx.x;
    float acc = 0.0f;
    for (int k = 0; k < iters; k++)
        acc += in[(t + k) % n] * 1.0009f;   // dependent work, memory latency
    smem[threadIdx.x % (blockDim.x)] = acc;
    __syncthreads();
    if (t < n) out[t] = smem[threadIdx.x % (blockDim.x)];
}

תרגיל 1 - חישוב occupancy ביד

בלי קוד. שחזרו את דוגמת ההרצאה על H100. נתונות מגבלות ה-SM: 64 warps, 32 blocks, 65,536 אוגרים, 228 KB shared memory. ה-kernel: 32 threads/block, 8 אוגרים/thread, 12 KB shared memory/block.

  1. חשבו כמה warps יש ב-block.
  2. חשבו את מגבלת ה-blocks לפי אוגרים: 65,536 / (threads x regs).
  3. חשבו את מגבלת ה-blocks לפי shared memory: 228 KB / smem-per-block.
  4. חשבו את מגבלת ה-blocks לפי מקסימום ה-warps: 64 / warps-per-block.
  5. קחו את המינימום על פני כל המגבלות (כולל מגבלת החומרה 32). מי הצוואר? מהו ה-occupancy התיאורטי באחוזים?
  6. מה תשנו כדי להעלות את ה-occupancy ל-50%? הראו את החשבון החדש. (רמז: תגעו רק בצוואר.)

רמז: הנוסחה max blocks = min(אוגרים, shared memory, חומרה, warps), ואז occupancy = max_blocks x warps_per_block / 64. שימו לב שאם הצוואר הוא shared memory, שינוי מספר האוגרים בכלל לא ישפיע - חייבים להוריד את ה-smem ל-block. בסעיף 6 בדקו שוב את כל המגבלות אחרי השינוי - צוואר חדש עלול לצוץ (מגבלת החומרה של 32 blocks).

תרגיל 2 - אימות עם ה-Occupancy API

  1. הוסיפו ל-occ.cu פונקציית main שקוראת ל-cudaGetDeviceProperties ומדפיסה את maxThreadsPerMultiProcessor, regsPerMultiprocessor, sharedMemPerMultiprocessor, ו-warpSize.
  2. קראו ל-cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks, work, blockSize, dynSmem) עם blockSize = 32 ו-dynSmem = 12*1024 (12 KB). הדפיסו את blocks.
  3. חשבו והדפיסו את ה-occupancy: blocks * (blockSize/warpSize) / maxWarpsPerSM.
  4. השוו את התוצאה לחישוב הידני מתרגיל 1. אם יש הבדל של block אחד - הסבירו למה (גרנולריות הקצאה, shared memory שמור).
  5. קראו ל-cudaOccupancyMaxPotentialBlockSize(&minGrid, &bestBlock, work, dynSmem, 0) והדפיסו את bestBlock. האם ה-block המומלץ שווה ל-32?

רמז: כדי לקבל את מספר האוגרים בפועל של ה-kernel, קמפלו עם nvcc --ptxas-options=-v והסתכלו בשורת Used N registers. ה-API מתחשב באוגרים האמיתיים האלה, לא ב-8 מהדוגמה - אז אם ה-kernel שלכם צורך יותר, מגבלת האוגרים עשויה להשתנות. כדי לשחזר בדיוק את 8 האוגרים אפשר לכפות -maxrregcount=8 (על חשבון spilling). ה-dynSmem שאתם מעבירים ל-API חייב להיות אותו גודל שתעבירו בפועל ב-<<<grid, block, dynSmem>>>.

תרגיל 3 - אימות ב-ncu

  1. הריצו את ה-kernel work על grid אמיתי (למשל work<<<1024, 32, 12*1024>>>(...)), עם n גדול (למשל מיליון) ו-iters בינוני (למשל 256).
  2. הריצו ncu --set full -k work -c 1 ./occ (או --section Occupancy בלבד למהירות).
  3. אתרו בפלט את סקציית ה-Occupancy. רשמו את ארבע שורות ה-Block Limit (SM, Registers, Shared Mem, Warps). מי הנמוך? האם הוא תואם לצוואר שחישבתם?
  4. רשמו את Theoretical Occupancy ואת Achieved Occupancy. למה הנמדד נמוך מהתיאורטי?
  5. הריצו שוב עם dynSmem = 6*1024 (6 KB) והשוו: לאיזה occupancy קפצתם, ומי הצוואר החדש?

רמז: סקציית ה-Occupancy ב-ncu מראה בדיוק את ארבע מגבלות ה-block שחישבתם ביד, וזו הדרך המהירה ביותר לזהות את המשאב-המגביל. ה-Achieved נמוך מהתיאורטי בגלל אפקטי זנב (ה-blocks האחרונים) ואיזון בין ה-SMs; אם ה-grid קטן מ-132 blocks, לא כל ה-SMs בכלל מקבלים עבודה. עברו מ-12 KB ל-6 KB וראו את Block Limit Shared Mem מכפיל את עצמו, עד שמגבלה אחרת (החומרה, 32) הופכת לצוואר.

תרגיל 4 - סריקה: occupancy מול ביצועים

זהו התרגיל המרכזי. תגלו במו ידיכם ששיא ה-occupancy אינו שיא הביצועים.

  1. כתבו לולאה ב-main שסורקת את גודל ה-shared memory הדינמי: 0, 2, 4, 8, 16, 32, 48 KB ל-block (עם block קבוע של 128 threads).
  2. לכל גודל, השתמשו ב-cudaOccupancyMaxActiveBlocksPerMultiprocessor כדי לחשב ולדווח את ה-occupancy התיאורטי.
  3. לכל גודל, מדדו גם את זמן הריצה של ה-kernel work (עם cudaEvent, חימום + מינימום על כמה חזרות). שמרו את ה-iters וכמות העבודה קבועים - רק גודל ה-smem משתנה.
  4. הדפיסו טבלה: smem_KB | occupancy% | time_ms. שרטטו (או תארו) את שתי העקומות. היכן ה-occupancy מקסימלי? היכן הזמן מינימלי? האם זו אותה נקודה?
  5. חזרו על הסריקה עם ציר אחר: קבעו smem=0 וסרקו block size 32, 64, 128, 256, 512, 1024. שוב, occupancy מול זמן. מצאו את גודל ה-block המהיר ביותר והשוו אותו ל-bestBlock שה-API הציע בתרגיל 2.

רמז: כדי שהסריקה תהיה הוגנת, סך העבודה חייב להיות זהה בכל נקודה - קבעו את מספר הגישות הכולל (grid x block x iters) ולא רק את ה-iters. אתם אמורים לראות שה-occupancy יורד ככל שה-smem ל-block גדל (פחות blocks נכנסים), אבל הזמן לא בהכרח גדל - עד סף מסוים הוא כמעט קבוע, כי כבר יש מספיק warps להסתיר את ה-latency (חוק Little). מעבר לסף, occupancy נוסף לא קונה כלום. זו ההוכחה הניסויית ל"מספיק occupancy, לא מקסימלי".

תרגיל 5 - זיהוי מצבי ה-warp בדוח ncu

בלי קוד. נתונה סקציית Scheduler Statistics מדוח ncu אמיתי של kernel latency-bound:

  Active Warps Per Scheduler         warp         4.46
  Eligible Warps Per Scheduler       warp         0.61
  Issued Warp Per Scheduler          warp         0.27
  No Eligible                           %        73.10
  One or More Eligible                  %        26.90

וסקציית Warp State Statistics:

  Warp Cycles Per Issued Instruction         cycle        16.55
  Stall Long Scoreboard                      cycle        11.90
  Stall Short Scoreboard                     cycle         1.40
  Stall Wait                                 cycle         1.80
  1. כמה warps פעילים יש בממוצע לכל scheduler? אם ל-SM יש 4 schedulers, כמה warps פעילים על ה-SM כולו, ומה ה-occupancy (מתוך 64)?
  2. כמה warps כשירים יש בממוצע? כמה נבחרים? מדוע "נבחר" תמיד קטן-או-שווה ל"כשיר"?
  3. מה אומר No Eligible = 73%? באילו מצבים נמצאים ה-warps ב-73% מהמחזורים האלה?
  4. לפי Stall Long Scoreboard הדומיננטי - על מה בדיוק ה-warps תקועים? איזה scoreboard זה, ולאיזו רמת זיכרון הוא מתייחס?
  5. אבחון: האם ה-kernel הזה latency-bound בגלל occupancy נמוך מדי, או שיש מספיק warps והבעיה במקום אחר? מה הייתם ממליצים - להעלות occupancy, או לא? הצדיקו לפי חוק Little.

רמז: Active הוא כל ה-warps התושבים (פעילים), Eligible הם אלה שעברו את ארבעת התנאים במחזור זה, Issued/Selected הם אלה שקיבלו פקודה. "נבחר" חייב להיות תת-קבוצה של "כשיר" כי ה-scheduler בוחר רק מבין הכשירים, אחד למחזור. No Eligible גבוה = scheduler רעב = לרוב אין מספיק warps או שכולם תקועים על אותה latency ארוכה. Long Scoreboard = המתנה ל-load מהזיכרון הגלובלי. כדי לאבחן אם צריך occupancy, השוו את מספר ה-warps הפעילים לסף של חוק Little (~13-30 warps).

תרגיל 6 (בונוס) - register pressure: כשעלייה ב-occupancy מזיקה

  1. כתבו kernel "כבד-אוגרים" - למשל כזה שמחזיק מערך מקומי גדול של float באוגרים ומבצע עליו הרבה פעולות תלויות (unroll ידני של ~40 משתנים חמים). ודאו עם --ptxas-options=-v שהוא צורך הרבה אוגרים (למשל 80+).
  2. קמפלו שלוש גרסאות עם תקרות אוגרים שונות: -maxrregcount=255 (occupancy נמוך, בלי spilling), -maxrregcount=64 (occupancy בינוני), -maxrregcount=32 (occupancy גבוה, spilling כבד). לכל גרסה רשמו את ה-occupancy התיאורטי (מה-API) ואת Local memory / spill stores/loads מ--v.
  3. מדדו את זמן הריצה של שלוש הגרסאות. איזו הכי מהירה? האם זו בעלת ה-occupancy הגבוה ביותר?
  4. הריצו ncu על הגרסאות והשוו: הגרסה עם occupancy גבוה - האם היא מראה Stall Long Scoreboard גבוה (מה-spills ל-local memory)?
  5. הסבירו במונחי המשאבים הקבועים של ה-SM: מדוע דחיסת יותר warps (occupancy גבוה) הכריחה פחות אוגרים ל-thread, ולמה זה הרס את הביצועים למרות ה-occupancy הגבוה.

רמז: -maxrregcount=32 כופה על כל thread להסתפק ב-32 אוגרים כדי לאפשר 64 warps ל-SM (100% occupancy), אבל אם ה-kernel צריך 80 אוגרים, ה-48 העודפים נשפכים ל-local memory (שהוא זיכרון גלובלי איטי). כל גישה לערך שנשפך היא load/store לזיכרון גלובלי - וזה יוצר בדיוק את ה-Stall Long Scoreboard שהיינו רוצים להימנע ממנו. זו ההדגמה החיה של "occupancy גבוה שמזיק": ה-GEMM המהיר מ-Hopper בוחר במכוון occupancy נמוך כדי לתת לכל thread הרבה אוגרים.