8.1 Occupancy ומצבי warp פתרון
פתרון - Occupancy ומצבי warp¶
הערה: כל הזמנים המוחלטים, ה-occupancy הנמדד וספירות ה-spill בפלטים למטה הם דוגמאות ממכונת H100 SXM עם CUDA 12, וישתנו אצלכם לפי הכרטיס, גרסת הדרייבר ומצב המערכת. מה שלא ישתנה הוא המבנה, סדרי הגודל והמגמות. תרגילי החשבון (1, 5) מדויקים ואינם תלויים בחומרה (בהנחת מגבלות ה-H100). הקוד מתקמפל עם nvcc -O2 -arch=sm_90a (החליפו ל-sm_75 על T4 וכדומה). לאורך הפתרון נשתמש בקובץ occ.cu הבא, שמשרת את תרגילים 2-4:
#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)
extern __shared__ float smem[];
// memory-dependent work; smemFloats controls how much shared memory is consumed
__global__ void work(const float* in, float* out, int n, int iters, int smemFloats) {
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;
if (smemFloats > 0) {
int idx = threadIdx.x % smemFloats;
smem[idx] = acc;
__syncthreads();
acc = smem[idx];
}
if (t < n) out[t] = acc;
}
פתרון תרגיל 1 - חישוב occupancy ביד¶
נתונים: 32 threads/block, 8 אוגרים/thread, 12 KB smem/block. מגבלות H100 SM: 64 warps, 32 blocks, 65,536 אוגרים, 228 KB smem.
(1) warps per block = 32 / 32 = 1 warp
(2) register limit = 65,536 / (32 x 8) = 65,536 / 256 = 256 blocks
(3) shared memory limit = 228 KB / 12 KB = 19 blocks
(4) warp limit = 64 / 1 = 64 blocks
hardware limit = 32 blocks
max blocks = min(256, 19, 32, 64) = 19 blocks <- bottleneck: shared memory
occupancy = 19 blocks x 1 warp / 64 = 0.297 ≈ 30%
הצוואר הוא shared memory - היא נגמרת ב-19 blocks. סעיף 6, להעלאה ל-50% (= 32 warps = 32 blocks, כי יש warp אחד ל-block): צריך שמגבלת ה-smem תרשה לפחות 32 blocks, כלומר 228 / smem >= 32 -> smem <= 7.125 KB. נוריד ל-6 KB:
shared memory limit = 228 / 6 = 38 blocks
max blocks = min(256, 38, 32, 64) = 32 blocks <- now the bottleneck: hardware
occupancy = 32 / 64 = 50%
למה זה עבד: הרעיון הוא שכל משאב מגביל בנפרד את מספר ה-blocks, וה-SM יכול להכיל רק את המינימום. חישבנו ארבע מגבלות ולקחנו את הקטנה. הורדת ה-smem מ-12 ל-6 KB הכפילה את מגבלת ה-smem מ-19 ל-38, אבל אז מגבלת החומרה (32) הפכה לצוואר החדש - ולכן קיבלנו בדיוק 50%, לא יותר. שינוי מספר האוגרים לא היה עוזר כלל, כי הם מעולם לא היו הצוואר (256 >> 19).
איך להכליל: תמיד חשבו את כל המגבלות וזהו את המינימום; שחררו רק את הצוואר, ואז בדקו שוב - כי צוואר חדש תמיד עלול לצוץ. הנוסחה: occupancy = min(מגבלות) x warps_per_block / max_warps_per_SM.
פתרון תרגיל 2 - אימות עם ה-Occupancy API¶
הוסיפו את ה-main הבא ל-occ.cu:
int main(void) {
cudaDeviceProp p;
CUDA_CHECK(cudaGetDeviceProperties(&p, 0));
printf("%s (CC %d.%d)\n", p.name, p.major, p.minor);
printf("maxThreads/SM=%d regs/SM=%d smem/SM=%zu KB warpSize=%d\n",
p.maxThreadsPerMultiProcessor, p.regsPerMultiprocessor,
p.sharedMemPerMultiprocessor / 1024, p.warpSize);
int blockSize = 32;
size_t dynSmem = 12 * 1024; // 12 KB
int blocks = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocks, work, blockSize, dynSmem));
int maxWarps = p.maxThreadsPerMultiProcessor / p.warpSize;
int activeWarps = blocks * (blockSize / p.warpSize);
printf("blockSize=%d dynSmem=12KB -> %d blocks, %d warps, occ=%.1f%%\n",
blockSize, blocks, activeWarps, 100.0 * activeWarps / maxWarps);
int minGrid = 0, bestBlock = 0;
CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(
&minGrid, &bestBlock, work, dynSmem, 0));
printf("suggested blockSize=%d (minGrid=%d)\n", bestBlock, minGrid);
return 0;
}
בדיקת מספר האוגרים בפועל, וההרצה:
nvcc -O2 -arch=sm_90a --ptxas-options=-v -o occ occ.cu
# ptxas info : Used 18 registers, ... (the real number; different from the example's 8)
./occ
הפלט הצפוי (H100 SXM):
NVIDIA H100 80GB HBM3 (CC 9.0)
maxThreads/SM=2048 regs/SM=65536 smem/SM=228 KB warpSize=32
blockSize=32 dynSmem=12KB -> 19 blocks, 19 warps, occ=29.7%
suggested blockSize=128 (minGrid=2112)
ה-API מאשש: 19 blocks, 19 warps, 29.7% - בדיוק החישוב הידני. ה-suggested blockSize=128 מלמד ש-32 threads/block הוא בחירה גרועה: עם block של 128, אותם 12 KB smem עדיין מגבילים ל-19 blocks לפי smem, אבל מגבלת ה-warps (64/4 = 16 blocks) נותנת 16 blocks x 4 warps = 64 warps = 100% occupancy.
למה זה עבד: ה-API מגלם את אותו חישוב min-על-המגבלות שעשינו ביד, אבל מתחשב באוגרים האמיתיים של ה-kernel (18 כאן, לא 8) ובגרנולריות ההקצאה של Hopper. במקרה הזה האוגרים לא הצוואר בכל מקרה, אז התוצאה זהה (19). אילו ה-smem השמור לדרייבר או עיגול ל-128 בייט היו קופצים, ה-API היה יכול להחזיר 18 - וזה היה הפוסק המדויק. cudaOccupancyMaxPotentialBlockSize הוא heuristic מוכן לבחירת block size שממקסם occupancy - שימושי כברירת-מחדל שפויה.
איך להכליל: אל תחשבו occupancy ביד בקוד ייצור - קראו ל-cudaOccupancyMaxActiveBlocksPerMultiprocessor. חלקו את התוצאה כדי לקבל occupancy, והשתמשו ב-cudaOccupancyMaxPotentialBlockSize כדי לקבל נקודת-פתיחה טובה ל-block size. תמיד ודאו שה-dynSmem שאתם מעבירים ל-API זהה לזה שתעבירו ב-<<<...>>>.
פתרון תרגיל 3 - אימות ב-ncu¶
הlaunch על grid אמיתי והרצת הפרופיילר:
הפלט הצפוי (block=32, dynSmem=12 KB):
Section: Occupancy
------------------------------- ----------- ------------
Block Limit SM block 32
Block Limit Registers block 113
Block Limit Shared Mem block 19 <- bottleneck
Block Limit Warps block 64
Theoretical Active Warps per SM warp 19
Theoretical Occupancy % 29.69
Achieved Occupancy % 27.31
Achieved Active Warps Per SM warp 17.48
------------------------------- ----------- ------------
ה-Block Limit Shared Mem = 19 הוא הנמוך - הצוואר שחישבנו, ו-Theoretical Occupancy = 29.69% (= 19/64). ה-Achieved = 27.31% נמוך מעט יותר. עכשיו עם 6 KB:
Block Limit Shared Mem block 38
Theoretical Active Warps per SM warp 32
Theoretical Occupancy % 50.00
הצוואר עבר ל-Block Limit SM = 32, וה-occupancy קפץ ל-50% - בדיוק כמו החישוב הידני מתרגיל 1 סעיף 6.
למה זה עבד: ncu מודד את מגבלות ה-block ישירות מהחומרה ומציג אותן שורה-שורה, כך שהמשאב-המגביל גלוי מיד. ה-Achieved נמוך מ-Theoretical בגלל אפקט זנב: ה-blocks האחרונים בכל SM מסיימים בזמנים שונים, וברגעים אלה פחות warps תושבים, אז הממוצע על פני המחזורים הפעילים יורד. אם ה-grid קטן מ-132 blocks, חלק מה-SMs בכלל לא מקבלים עבודה וה-Achieved צונח עוד יותר.
איך להכליל: ncu --section Occupancy הוא הדרך המהירה ביותר לענות "מי הצוואר של ה-kernel שלי". קראו את ארבע שורות ה-Block Limit, זהו את הנמוכה, ושחררו רק אותה. פער גדול בין Theoretical ל-Achieved = חשד לזנב או ל-grid קטן מדי - הגדילו את ה-grid או אזנו את העבודה.
פתרון תרגיל 4 - סריקה: occupancy מול ביצועים¶
לולאת הסריקה (block קבוע 128, סך העבודה קבוע):
static float timeKernel(int grid, int block, size_t dynSmem, int smemFloats,
const float* in, float* out, int n, int iters) {
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
work<<<grid, block, dynSmem>>>(in, out, n, iters, smemFloats); // warmup
CUDA_CHECK(cudaDeviceSynchronize());
float best = 1e30f;
for (int r = 0; r < 20; r++) {
CUDA_CHECK(cudaEventRecord(s));
work<<<grid, block, dynSmem>>>(in, out, n, iters, smemFloats);
CUDA_CHECK(cudaEventRecord(e));
CUDA_CHECK(cudaEventSynchronize(e));
float ms; CUDA_CHECK(cudaEventElapsedTime(&ms, s, e));
if (ms < best) best = ms;
}
CUDA_CHECK(cudaEventDestroy(s)); CUDA_CHECK(cudaEventDestroy(e));
return best;
}
// inside main, after allocating d_in/d_out of size n:
int block = 128, maxWarps = 64;
int smemKB[] = {0, 2, 4, 8, 16, 32, 48};
printf("%8s %10s %10s\n", "smem_KB", "occ_%", "time_ms");
for (int i = 0; i < 7; i++) {
size_t dynSmem = (size_t)smemKB[i] * 1024;
int smemFloats = smemKB[i] * 256; // 1KB = 256 floats
int blocks = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocks, work, block, dynSmem));
double occ = 100.0 * blocks * (block / 32) / maxWarps;
int grid = 132 * 16; // fixed amount of work, fills the GPU
float ms = timeKernel(grid, block, dynSmem, smemFloats, d_in, d_out, n, 512);
printf("%8d %10.1f %10.3f\n", smemKB[i], occ, ms);
}
הפלט הצפוי (H100 SXM):
smem_KB occ_% time_ms
0 100.0 3.084
2 100.0 3.086
4 100.0 3.089
8 100.0 3.091
16 87.5 3.098
32 43.8 3.142
48 25.0 3.588
וסריקת גודל ה-block (smem=0):
block occ_% time_ms
32 50.0 4.401
64 100.0 3.140
128 100.0 3.084
256 100.0 3.083
512 100.0 3.091
1024 100.0 3.126
התוצאה המכוננת: ה-occupancy יורד מ-100% ל-25% ככל שה-smem ל-block גדל, אבל הזמן כמעט לא משתנה בטווח רחב - מ-100% עד ~44% occupancy הזמן קבוע בערך על 3.09-3.14 ms. רק ב-25% occupancy (16 warps ל-SM) הזמן קופץ, כי ירדנו מתחת לסף חוק Little וה-latency כבר לא מוסתר לגמרי. בסריקת ה-block, המהיר ביותר הוא 128-256 threads (100% occupancy), אבל 512 ו-1024 - גם הם 100% - איטיים במקצת, ו-32 threads (רק 50%) הכי איטי. שיא ה-occupancy אינו נקודה בודדת של שיא ביצועים; יש רמה שלמה של occupancy שנותנים את אותו זמן.
למה זה עבד: ברגע שיש מספיק warps כדי להסתיר את ה-latency (חוק Little, ~24-30 warps), warps נוספים לא קונים כלום - ה-latency כבר מוסתר, וה-throughput רווי. לכן הזמן קבוע לאורך טווח occupancy רחב. הקפיצה ב-25% מראה מה קורה כשיורדים מתחת לסף: פתאום אין מספיק מלאי, ה-scheduler רעב, וה-latency מתגלה. block=32 איטי כי הוא נחסם ל-50% (מגבלת 32 blocks בחומרה) ומבזבז חצי מהמלאי.
איך להכליל: אל תרדפו אחרי 100% occupancy. מצאו את הסף שבו הזמן מתייצב, ותנו לעצמכם מרווח מעליו - זהו ה-occupancy ה"מספיק". מעבר לו, השקיעו את המשאבים (אוגרים, smem) בטובת כל thread, לא במספר ה-threads. בחרו block size מתוך הרמה השטוחה (כאן 128-256), לא את הקיצון.
פתרון תרגיל 5 - זיהוי מצבי ה-warp בדוח ncu¶
סעיף 1 - warps פעילים.
active/scheduler = 4.46; per SM (4 schedulers) = 4.46 x 4 = 17.84 warps
occupancy = 17.84 / 64 ≈ 27.9% ≈ 28%
סעיף 2 - כשירים ונבחרים. בממוצע 0.61 warps כשירים לכל scheduler, ורק 0.27 נבחרים (Issued). "נבחר" תמיד קטן-או-שווה ל"כשיר" כי ה-scheduler בוחר את הנבחר מתוך מאגר הכשירים, ולכל היותר warp אחד למחזור - אי אפשר לבחור warp שאינו כשיר.
סעיף 3 - No Eligible = 73%. ב-73% מהמחזורים הפעילים ל-scheduler אין אף warp כשיר להנפיק. באותם מחזורים כל ה-warps התושבים (הפעילים) נמצאים במצב תקוע - הם פעילים אבל אף אחד לא עומד בארבעת תנאי הכשירות. ה-scheduler יושב בטל.
סעיף 4 - Stall Long Scoreboard. ה-warps תקועים בעיקר בהמתנה ל-load מהזיכרון הגלובלי. ה-Long Scoreboard הוא ה-scoreboard שעוקב אחרי תוצאות מ-GPU RAM (זיכרון גלובלי), בניגוד ל-Short Scoreboard שעוקב אחרי shared memory ואריתמטיקה. 11.9 מחזורים מתוך 16.55 מבוזבזים בהמתנה הזו - רוב זמן ה-stall.
סעיף 5 - אבחון. יש רק ~17.8 warps פעילים על ה-SM (28% occupancy). חוק Little דורש ~24-30 warps כדי להסתיר latency של זיכרון (תוצאת Volkov, שיעור 7.4). אנחנו מתחת לסף: אין מספיק warps כדי לכסות את ה-Long Scoreboard, ולכן ה-scheduler רעב ב-73% מהמחזורים. כאן, ובאופן חריג, התשובה כן להעלות occupancy - אבל לא כמטרה, אלא כאמצעי: אנחנו מתחת לרצפת חוק Little, וצריך עוד warps רק כדי להגיע אליה. אילו היינו כבר ב-40 warps עם No Eligible גבוה, המסקנה הייתה הפוכה - הבעיה במקום אחר (למשל latency שאי אפשר להסתיר, שדורש תיקון אלגוריתמי).
למה זה עבד: הדוח מתרגם ישירות את ארבעת המצבים - Active/Eligible/Selected - ואת סיבות ה-stall. No Eligible גבוה יחד עם occupancy נמוך יחד עם Long Scoreboard דומיננטי הם שלישיית האבחון הקלאסית של kernel memory-latency-bound עם מלאי warps לא מספיק.
איך להכליל: קראו כל דוח ncu כסיפור: כמה warps פעילים (occupancy), כמה מהם כשירים (No Eligible), ועל מה התקועים תקועים (סיבת ה-stall). השוו את מספר ה-warps הפעילים לסף חוק Little - זה מה שקובע אם התרופה היא "עוד occupancy" או "תקן את הצוואר האמיתי".
פתרון תרגיל 6 (בונוס) - register pressure: כשעלייה ב-occupancy מזיקה¶
kernel כבד-אוגרים (unroll שמחזיק הרבה ערכים חמים):
__global__ void heavy(const float* in, float* out, int n) {
int t = blockIdx.x * blockDim.x + threadIdx.x;
float r[40];
#pragma unroll
for (int i = 0; i < 40; i++) r[i] = in[(t + i) % n];
#pragma unroll
for (int i = 0; i < 40; i++)
r[i] = r[i] * r[(i + 1) % 40] + r[(i + 7) % 40]; // dependent, keeps r[] hot
float acc = 0.0f;
for (int i = 0; i < 40; i++) acc += r[i];
if (t < n) out[t] = acc;
}
שלוש הקומפילציות ובדיקת ה-spills:
nvcc -O2 -arch=sm_90a --ptxas-options=-v -maxrregcount=255 -o h255 heavy.cu
nvcc -O2 -arch=sm_90a --ptxas-options=-v -maxrregcount=64 -o h64 heavy.cu
nvcc -O2 -arch=sm_90a --ptxas-options=-v -maxrregcount=32 -o h32 heavy.cu
פלט -v צפוי (block=128), וזמני הריצה:
version regs spill stores/loads theoretical occ. time_ms
-maxrreg=255 82 0 / 0 37.5% 1.90
-maxrreg=64 64 24 / 64 50.0% 2.35
-maxrreg=32 32 160 / 410 100.0% 4.65
הגרסה המהירה ביותר היא בעלת ה-occupancy הנמוך ביותר (37.5%, 1.90 ms), והאיטית ביותר היא זו של 100% occupancy (4.65 ms). ncu על h32 יראה Stall Long Scoreboard גבוה בהרבה מ-h255 - כי כל אחד מ-48 האוגרים העודפים שנשפכו ל-local memory הופך לגישת זיכרון גלובלי.
למה זה עבד: המשאבים על ה-SM קבועים - 65,536 אוגרים לחלוקה בין כל ה-threads התושבים. כדי לדחוס 64 warps (2048 threads, 100% occupancy), כל thread מוגבל ל-65,536/2048 = 32 אוגרים. אבל ה-kernel צריך 82; 50 האוגרים העודפים נשפכים ל-local memory (זיכרון גלובלי איטי), ויוצרים מבול של loads/stores שמתקעים את ה-warps על ה-Long Scoreboard. ה-occupancy הגבוה שהרווחנו לא שווה כלום מול הזיכרון הנוסף שיצרנו. הגרסה של 37.5% occupancy נותנת לכל thread 82 אוגרים מלאים, בלי spill - וזה מנצח בגדול.
איך להכליל: זו בדיוק הסיבה ש-GEMM מהיר על Hopper/Blackwell רץ ב-occupancy חד-ספרתי - הוא מעדיף לתת לכל thread הרבה אוגרים ו-tiles ענקיים על פני הרבה warps. occupancy הוא אמצעי להסתרת latency, לא מטרה: ברגע שיש מספיק warps לחוק Little, הפסיקו להוסיף - כל warp נוסף גוזל אוגרים ו-smem מהאחרים ועלול לגרום ל-spilling שהורס הכל.