8.1 Occupancy ומצבי warp הרצאה
בפרק 7 בנינו את התיאוריה של הביצועים: מודל ה-Roofline, ההבחנה בין compute-bound לmemory-bound, וה-arithmetic intensity. בשיעור 7.4 סגרנו את המשולש עם הoverhead וחוק Little, וראינו את השורה התחתונה: כדי להסתיר latency צריך מספיק warps תושבים על ה-SM, ושרשרת load אופיינית של כ-416 מחזורים דורשת בערך 13 warps כדי להיעלם לגמרי. אבל השארנו שאלה פתוחה: מאיפה מגיעים ה-warps האלה, כמה מהם באמת יושבים על ה-SM שלי בזמן ריצה, ומה מגביל את המספר? הפרק הזה - הפרק הראשון של האופטימיזציה בפועל - פותח בדיוק שם. נלמד את ארבעת המצבים שבהם warp יכול להימצא בכל מחזור-שעון (פעיל, כשיר, תקוע, נבחר), נגדיר מהו מחזור פעיל (active cycle), ונגיע ל-occupancy (תפוסה) - היחס בין ה-warps הפעילים למקסימום האפשרי. נחשב occupancy תיאורטי ביד על דוגמת H100 מלאה, נראה כיצד לחזות אותו בקוד עם ה-Occupancy API, ונאמת אותו ב-ncu. והכי חשוב - נפרק את המיתוס הגדול של הפרק: occupancy גבוה יותר הוא לא תמיד טוב יותר. חוק Little אמר כמה warps צריך; כאן נראה שמעבר לזה, occupancy נוסף לא רק שלא קונה כלום - הוא לפעמים מזיק, ושה-GEMM המהיר בעולם על Hopper רץ ב-occupancy חד-ספרתי.
ארבעת מצבי ה-warp - Warp Execution State¶
ה-glossary מתאר את מצב ה-warps שמריצים kernel בעזרת ארבעה שמות-תואר, וההערה הקריטית הראשונה היא שהם לא בלעדיים (non-exclusive): warp יכול להיות פעיל וגם תקוע בו-זמנית, פעיל וגם כשיר, וכן הלאה. אלה לא ארבעה תאים נפרדים אלא ארבע תכונות שמתארות, מחזור אחר מחזור, מאילו warps ה-warp scheduler יכול לבחור להנפיק פקודה. נעבור עליהם אחד-אחד.
פעיל - active. warp נחשב פעיל מהרגע שה-threads שלו מתחילים לרוץ ועד הרגע שכל ה-threads בו יצאו מה-kernel. במילים אחרות, active פירושו תושב (resident) על ה-SM עם משאבים מוקצים - אוגרים בקובץ-האוגרים, מקום ב-shared memory - מרגע הlaunch ועד הסיום. אוסף ה-warps הפעילים הוא המאגר שממנו ה-scheduler בוחר מועמדים בכל מחזור. הדקות המכרעת: warp פעיל לא בהכרח מבצע פקודות. הוא יכול להיות פעיל ותקוע ובטל בו-זמנית. ב-H100 (compute capability 9.0) המקסימום הוא 64 warps פעילים ל-SM = 2048 threads ל-SM (64 x 32).
כשיר - eligible. warp כשיר הוא warp פעיל שמוכן להנפיק את הפקודה הבאה שלו. זו קבוצת המועמדים המיידיים של ה-scheduler. כדי להיות כשיר, כל ארבעת התנאים הבאים חייבים להתקיים יחד:
warp eligible <=> all four hold:
1. the next instruction has already been fetched
2. the required execution unit is free (Tensor Core / CUDA Core / LSU)
3. all instruction dependencies are resolved (source registers ready)
4. no sync barrier is blocking progress
אם ולו תנאי אחד לא מתקיים - ה-warp לא כשיר, כלומר הוא תקוע.
תקוע - stalled. warp תקוע הוא warp פעיל שלא יכול להנפיק את הפקודה הבאה שלו בגלל תלות שלא נפתרה או התנגשות משאבים. תקוע הוא בעצם המשלים של כשיר בתוך קבוצת הפעילים. שלוש הסיבות:
- תלות ביצוע (execution dependency) - ממתין לתוצאה של פקודה אריתמטית קודמת.
- תלות זיכרון (memory dependency) - ממתין לתוצאה של פעולת זיכרון קודמת.
- התנגשות צנרת (pipeline conflict) - יחידת הביצוע הדרושה כבר תפוסה על-ידי warp אחר.
החומרה עוקבת אחרי תלויות בעזרת שני סוגי scoreboard: ה-short scoreboard לתוצאות מ-shared memory ומאריתמטיקה, וה-long scoreboard לתוצאות מהזיכרון הגלובלי (GPU RAM). הסטטיסטיקות האלה יופיעו לנו ב-ncu כ-Stall Short Scoreboard ו-Stall Long Scoreboard, ונחזור אליהן בסוף השיעור. וכאן המקום למחיקת התפיסה השגויה החשובה ביותר בכל השיעור: warps תקועים אינם דבר רע. להפך - אוסף גדול של warps תקועים בו-זמנית הוא בדיוק מה שדרוש כדי להסתיר latency של פעולות ארוכות כמו load מהזיכרון או פקודת HMMA של ה-Tensor Core, שרצה עשרות מחזורים. ה-warp התקוע של warp אחד הוא בדיוק החור שעבודה של warp אחר ממלאת. זו הנקודה שקישרנו אליה כל שיעור 7.4.
נבחר - selected. warp נבחר הוא warp כשיר שה-scheduler בחר להנפיק לו פקודה במחזור הנוכחי. התהליך, בכל מחזור: ה-warp scheduler מסתכל על מאגר ה-warps הכשירים, בוחר אחד אם יש כאלה, ומנפיק לו פקודה. מכאן מדד חשוב - issue efficiency (יעילות ההנפקה): החלק מתוך המחזורים הפעילים שבהם warp נבחר ופקודה הונפקה.
הנה כל התמונה כרשת של חריצים (slots) על פני מחזורים. כל scheduler מנפיק לכל היותר warp אחד למחזור:
cycle n n+1 n+2 n+3 n+4
warp 0 stalled eligible selected stalled stalled (active the whole time)
warp 1 selected stalled stalled eligible selected
warp 2 eligible selected stalled stalled eligible
warp 3 stalled stalled eligible selected stalled
│ │ │
│ │ └── warp 0 selected: eligible -> got an instruction
│ └── only warp 1,3 eligible; warp 2 selected
└── all four active, but only 1,2 eligible; warp 1 selected
שימו לב לזרימה: כל הארבעה פעילים בכל מחזור (הם תושבים), אבל רק חלק מהם כשירים בכל מחזור נתון, ורק אחד נבחר לכל scheduler. מעבר המצבים הטבעי: ה-threads מתחילים -> ה-warp נעשה פעיל -> אם ארבעת התנאים מתקיימים הוא כשיר, אחרת תקוע -> ה-scheduler בוחר warp כשיר אחד, שהופך לנבחר -> הפקודה מונפקת -> בפקודה הבאה ה-warp עלול שוב להיתקע על תלות חדשה. כל זה קורה בזמן שה-warp נשאר פעיל עד ש-threads שלו יוצאים מה-kernel.
המחזור הפעיל - Active Cycle¶
לפני ה-occupancy צריך מונח ביניים קטן אבל קריטי: מחזור פעיל. ה-glossary מגדיר בפשטות: מחזור פעיל הוא מחזור-שעון שבו ל-SM יש לפחות warp פעיל אחד תושב. ה-warp התושב יכול להיות באותו מחזור כשיר או תקוע - כך או כך המחזור נספר כפעיל.
active cycle ? at least one warp resident on the SM
├── yes, and at least one eligible -> active cycle
├── yes, but all are stalled -> active cycle (!)
└── no, no warp is resident -> inactive cycle (excluded)
הדקות שה-glossary מדגיש: מחזור פעיל סופר כל מחזור שבו לפחות warp אחד תושב, בלי קשר לשאלה אם הוא תקוע בהמתנה לתלות. כלומר מחזורים תקועים עדיין נספרים כמחזורים פעילים - "פעיל" עוסק בתושבוּת, לא בהנפקה בפועל. אל תבלבלו "מחזור פעיל" עם "מחזור שבו הונפקה פקודה". ההבחנה הזו היא הלב של ההבדל בין תפוסה תיאורטית לנמדדת שנראה מיד: ה-occupancy הנמדד ממוצע על פני מחזורים פעילים. אם ה-SM עמד ריק לגמרי בחצי מזמן ה-kernel (בין גלים של blocks), אותם מחזורים ריקים בכלל לא נכנסים למכנה של ה-occupancy הנמדד - ולכן occupancy נמדד גבוה לא מבטיח שה-SM עבד לאורך כל הזמן.
תפוסה - Occupancy: תיאורטי מול נמדד¶
עכשיו ההגדרה המרכזית. occupancy הוא היחס בין מספר ה-warps הפעילים למספר המרבי של warps פעילים על הdevice.
ב-H100, המכנה הוא 64. אם על ה-SM שלכם יושבים 32 warps פעילים, ה-occupancy הוא 32/64 = 50%. פשוט. אבל יש שתי מדידות שונות של אותו יחס, וההבחנה ביניהן קריטית:
- תפוסה תיאורטית (theoretical occupancy) - החסם העליון ל-occupancy, הנגזר מתצורת הlaunch ומיכולות הdevice. זהו תקרה סטטית: בהינתן כמה אוגרים ו-shared memory כל block צורך, וכמה threads יש ב-block, החומרה יכולה להכניס לכל היותר כך-וכך warps ל-SM. מחושב לפני הריצה, ללא הרצה בכלל.
- תפוסה נמדדת (achieved occupancy) - ה-occupancy האמיתי במהלך הריצה, ממוצע על פני המחזורים הפעילים. זה מה שבאמת קרה - נמדד על-ידי מונים בחומרה.
התפוסה הנמדדת כמעט תמיד נמוכה מהתיאורטית, ומכמה סיבות: זנב לא-מאוזן של blocks בסוף ה-kernel (כשה-blocks האחרונים מסיימים בקצב שונה), חוסר-איזון בין ה-SMs, ומחזורים שבהם פחות warps תושבים בפועל. הפער בין השתיים הוא מדד אבחוני בפני עצמו: תפוסה תיאורטית גבוהה אבל נמדדת נמוכה מרמזת על בעיית איזון או על גל-זנב (tail effect) - למשל grid קטן מדי שלא ממלא את כל 132 ה-SMs של ה-H100.
הדוגמה המלאה - חישוב תפוסה על H100¶
נעשה עכשיו את החישוב במלואו, ביד, על מגבלות ה-H100 האמיתיות. זו הדוגמה שכדאי לזכור. מגבלות המשאבים של SM בודד ב-H100:
| משאב | מגבלה ל-SM |
|---|---|
| warps פעילים מקסימלי | 64 |
| thread blocks מקסימלי | 32 |
| אוגרים (קובץ-האוגרים) | 65,536 אוגרי 32-ביט |
| shared memory | 228 KB |
וה-kernel שנשיק, עם פרמטרי הlaunch הבאים:
- מספר ה-threads ל-block: 32 (כלומר
32/32 = 1 warp ל-block) - אוגרים ל-thread: 8
- כמות shared memory ל-block: 12 KB
הרעיון: כל אחד מהמשאבים מגביל בעצמו כמה blocks נכנסים ל-SM, ואנחנו לוקחים את המינימום - המשאב שנגמר ראשון הוא צוואר-הבקבוק. נחשב ארבע מגבלות:
(1) register limit:
registers per block = 32 threads x 8 registers = 256 registers
blocks = 65,536 / 256 = 256 blocks
(2) shared memory limit:
blocks = 228 KB / 12 KB = 19 blocks
(3) hardware limit (blocks per SM):
blocks = 32
(4) warp limit:
warps per block = 1; 64 / 1 = 64 blocks
bottleneck = min(256, 19, 32, 64) = 19 blocks
הצוואר הוא shared memory: היא נגמרת ב-19 blocks, הרבה לפני שאזל קובץ-האוגרים (שהיה מרשה 256) או מגבלת החומרה (32). עם block אחד = warp אחד, זה נותן 19 warps פעילים מתוך מקסימום 64:
ה-kernel הזה shared-memory-bound, ומקורק ל-occupancy תיאורטי של כ-30%. שימו לב מה זה אומר: אם היינו רוצים occupancy גבוה יותר, מגדיל האוגרים בכלל לא יעזור (הם לא הצוואר) - צריך להוריד את ה-12 KB ל-block. לו היינו מורידים ל-6 KB, המגבלה הייתה 228/6 = 38, וזה כבר נחסם על-ידי מגבלת החומרה min(38, 32) = 32 blocks = 32 warps = 50% occupancy. זהו כל האופי של אופטימיזציית occupancy: מזהים את המשאב-המגביל ומשחררים אותו - לא נוגעים בשאר.
הנוסחה הכללית שכדאי לחקוק:
max blocks per SM = min( blocks by registers,
blocks by shared memory,
max blocks in hardware,
blocks by warp limit )
theoretical occupancy = (max blocks x warps-per-block) / max-warps-per-SM
הערה על דיוק: החומרה מקצה אוגרים ו-shared memory בגרנולריות מסוימת (למשל shared memory ב-Hopper מוקצה ביחידות של 128 בייט, ויש כמות shared memory שמורה לדרייבר). לכן ה-API הרשמי עשוי לפעמים להחזיר 18 במקום 19. החשבון ידני מעגל את הפרטים האלה; ה-API הוא הפוסק המדויק. את שניהם נראה עכשיו.
חישוב התפוסה ב-API - Occupancy Calculator¶
לא צריך לעשות את החשבון הזה ביד בכל פעם. ה-CUDA Runtime מגלם אותו בשתי פונקציות מרכזיות שמתחשבות בכל הגרנולריות והמגבלות של ה-compute capability שלכם:
cudaOccupancyMaxActiveBlocksPerMultiprocessor- "בהינתן ה-kernel הזה, גודל block זה, וכמות ה-shared memory הדינמית הזו, כמה blocks ייכנסו ל-SM?"cudaOccupancyMaxPotentialBlockSize- "איזה גודל block ימקסם את ה-occupancy עבור ה-kernel הזה?" - נותן heuristic מוכן לבחירת block size.
הנה תוכנית מלאה שמחשבת occupancy תיאורטי עבור kernel, ומדפיסה גם את גודל ה-block המומלץ:
#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)
// example kernel: 12 KB static shared memory (= 3072 floats)
__global__ void myKernel(const float* in, float* out, int n) {
__shared__ float tile[3072]; // 12 KB
int t = blockIdx.x * blockDim.x + threadIdx.x;
if (t < n) {
tile[threadIdx.x % 3072] = in[t];
__syncthreads();
out[t] = tile[threadIdx.x % 3072] * 2.0f;
}
}
int main(void) {
int dev = 0;
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, dev));
int blockSize = 32;
size_t dynSmem = 0; // no dynamic shared memory here
int blocksPerSM = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksPerSM, myKernel, blockSize, dynSmem));
int activeWarps = blocksPerSM * (blockSize / prop.warpSize);
int maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;
double occ = (double)activeWarps / maxWarps;
printf("device : %s (CC %d.%d)\n", prop.name, prop.major, prop.minor);
printf("blocks per SM : %d\n", blocksPerSM);
printf("active warps / SM : %d (max %d)\n", activeWarps, maxWarps);
printf("theoretical occ. : %.1f%%\n", occ * 100.0);
// which block size maximizes occupancy?
int minGrid = 0, bestBlock = 0;
CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(
&minGrid, &bestBlock, myKernel, 0, 0));
printf("suggested block : %d threads (minGrid %d)\n", bestBlock, minGrid);
return 0;
}
קומפילציה והרצה על H100:
פלט צפוי (H100 SXM, CUDA 12):
device : NVIDIA H100 80GB HBM3 (CC 9.0)
blocks per SM : 19
active warps / SM : 19 (max 64)
theoretical occ. : 29.7%
וה-API מאשש את החשבון הידני: 19 blocks, 19 warps, כ-30% occupancy, shared-memory-bound. שימו לב שה-myKernel כאן צורך 12 KB shared memory כ-tile סטטי, בדיוק כמו בדוגמה. מספר האוגרים בפועל תלוי בגוף ה-kernel ובאופטימיזציות ה-ptxas; כדי לראות אותו הריצו nvcc --ptxas-options=-v (יידפיס Used N registers), וכדי לכפות תקרת אוגרים השתמשו ב-__launch_bounds__ או ב--maxrregcount. אם תרצו לשנות את ה-occupancy, שחקו עם blockSize וב-dynSmem וראו את blocksPerSM זז.
יותר תפוסה הוא לא תמיד טוב יותר¶
וכאן מגיעה הנקודה החשובה ביותר בכל השיעור, זו שמפרידה בין מי שקרא spec למי שמבין ביצועים. האינטואיציה הנאיבית אומרת: occupancy גבוה = הרבה warps = הרבה מקביליות = מהיר יותר, אז נמקסם occupancy. זה שגוי. ה-glossary אומר זאת מפורשות: ברגע שה-occupancy מספיק להסתרת latency, הגדלה נוספת שלו עלולה דווקא לפגוע בביצועים.
הסיבה נעוצה בחשבון המשאבים שהרגע עשינו, רק במהופך. המשאבים על ה-SM קבועים: 65,536 אוגרים ו-228 KB shared memory, לחלוקה בין כל ה-warps התושבים. אם תדחסו יותר warps ל-SM (occupancy גבוה), כל thread מקבל פחות אוגרים ופחות shared memory:
64 warps x 32 threads = 2048 threads on the SM
65,536 registers / 2048 threads = 32 registers per thread (at 100% occupancy)
versus:
16 warps x 32 = 512 threads
65,536 / 512 = 128 registers per thread (at 25% occupancy)
ב-100% occupancy כל thread מוגבל ל-32 אוגרים. אם ה-kernel שלכם צריך 80 אוגרים כדי להחזיק את הנתונים החמים שלו באוגרים, דחיסה ל-32 תכריח גלישת אוגרים (register spilling) אל ה-local memory (שהוא בעצם זיכרון גלובלי איטי) - וזה יהרוס את הביצועים הרבה יותר משכל ה-warps הנוספים הרוויחו. אותו סיפור עם shared memory: occupancy גבוה מכריח פחות shared memory ל-block, מה שמגביל את גודל ה-tile ב-GEMM או בקונבולוציה.
חוק Little נותן את הסף התחתון: ראינו ש-13 warps (כ-20% occupancy) מסתירים שרשרת load אופיינית, וש-~24-30 warps מסתירים עומס מעורב. מעל הסף הזה, occupancy נוסף לא קונה הסתרת latency נוספת - כל ה-latency כבר מוסתר - אבל הוא כן גובה את המחיר של פחות אוגרים ל-thread. זה בדיוק "אזור התשואה השלילית".
הדוגמה הנגדית המדהימה שה-glossary נותן: kernels של GEMM בעל-ביצועים על ארכיטקטורות Hopper ו-Blackwell רצים לעיתים קרובות ב-occupancy חד-ספרתי (אחוזים בודדים). איך ייתכן? כי הם לא צריכים הרבה warps כדי לרווות את ה-Tensor Cores. זכרו את יחס ה-100:1 בין רוחב-הפס של ה-Tensor Core ל-CUDA Core (שיעור 7.3): מספר קטן של warps שמזין את ה-Tensor Cores ב-wgmma/HMMA ברצף מספיק כדי להחזיק אותם רוויים לגמרי. ה-kernels האלה משתמשים במעט warps אבל נותנים לכל אחד הרבה אוגרים ו-shared memory ענק (tiles גדולים, double-buffering, warp-specialization עם warps של producer ו-consumer). אילו הם היו רודפים אחרי occupancy גבוה, הם היו נאלצים לצמצם את ה-tiles ואת האוגרים - וההאצה הייתה צונחת.
המסקנה הפילוסופית, בלשון ה-glossary: occupancy מודד איזה חלק ממשימותיו המקביליות המרביות ה-GPU מטפל בו-זמנית, וזה אינו כשלעצמו יעד אופטימיזציה ברוב ה-kernels. במקום זה, מייעלים את המשאב שהוא באמת הצוואר:
- ל-kernel compute-bound -> ממקסמים את ניצול צנרות החישוב (pipe / Tensor Core utilization).
- ל-kernel memory-bound -> ממקסמים את ניצול רוחב-הפס של הזיכרון.
ה-occupancy הוא אמצעי להסתרת latency מספקת, לא מטרה. הוא רלוונטי רק כדי לוודא שיש מספיק warps כדי לספק את חוק Little - וברגע שהגעתם לסף, מפסיקים להסתכל עליו ועוברים לצוואר האמיתי.
איפה כל זה מופיע ב-ncu¶
כל המונחים של השיעור הם בדיוק מה שתראו כשתריצו את NSight Compute על kernel. נריץ:
שלוש הסקציות הרלוונטיות. ראשית, סקציית ה-Occupancy - זו מאשרת את החשבון הידני שלנו, כולל שורת ה-block limit לכל משאב, כך שרואים מיד מי הצוואר:
Section: Occupancy
------------------------------- ----------- ------------
Block Limit SM block 32
Block Limit Registers block 256
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.84
Achieved Active Warps Per SM warp 17.82
------------------------------- ----------- ------------
שימו לב: Block Limit Shared Mem = 19 הוא הנמוך מכולם, בדיוק המשאב-המגביל שחישבנו. ה-Theoretical Occupancy הוא 29.69% (= 19/64), וה-Achieved מעט נמוך יותר (27.84%) - בגלל אפקטי זנב ואיזון, בדיוק כפי שאמרנו. שנית, סקציית ה-Scheduler Statistics מראה את מצבי ה-warp הממוצעים לכל scheduler למחזור:
Section: Scheduler Statistics
---------------------------- ----------- ------------
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 <- warp scheduler starving
One or More Eligible % 26.90
---------------------------- ----------- ------------
זו קריאה ישירה של המצבים: בממוצע יש 4.46 warps פעילים לכל scheduler (19 warps חלקי 4 schedulers ≈ 4.75, קרוב), אבל רק 0.61 כשירים, ורק 0.27 נבחרים למחזור. ה-No Eligible = 73% הוא הדגל האדום: ב-73% מהמחזורים ל-scheduler אין אף warp כשיר להנפיק - הוא רעב. עם 19 warps בלבד, אין מספיק מלאי כדי להסתיר את ה-latency. שלישית, סקציית ה-Warp State Statistics מספרת למה ה-warps תקועים - כאן נכנסות סיבות ה-stall שהזכרנו:
Section: Warp State Statistics
------------------------------------- ----------- ------------
Warp Cycles Per Issued Instruction cycle 16.55
Stall Long Scoreboard cycle 11.90 <- waiting for a global load
Stall Short Scoreboard cycle 1.40
Stall Wait cycle 1.80
------------------------------------- ----------- ------------
ה-Stall Long Scoreboard הדומיננטי אומר שרוב ה-warps תקועים בהמתנה ל-load מהזיכרון הגלובלי (ה-long scoreboard מהסעיף הראשון). האבחון המלא: תפוסה נמוכה (30%), scheduler רעב (73% No Eligible), והתקיעות היא latency של זיכרון גלובלי. הפתרון כאן דווקא כן יהיה להעלות occupancy - אבל רק כי אנחנו מתחת לסף של חוק Little, לא כמטרה בפני עצמה. כפי שנראה בשיעורים 8.2 ו-8.5, המדדים האלה הם המצפן שמכוון אותנו למשאב הנכון לתקן.
סיכום¶
- מצב ה-warp מתואר בארבעה שמות-תואר לא-בלעדיים: פעיל (תושב מרגע הlaunch עד שכל ה-threads יצאו), כשיר (פעיל וגם מוכן להנפיק - ארבעה תנאים: פקודה נשלפה, יחידה פנויה, תלויות נפתרו, אין barrier), תקוע (פעיל אך לא יכול להנפיק - תלות ביצוע/זיכרון/צנרת), ונבחר (הכשיר שה-scheduler הנפיק לו במחזור זה).
- ה-warps התקועים הם דבר טוב, לא רע: אוסף warps תקועים בו-זמנית הוא בדיוק מה שמסתיר את ה-latency של loads ושל פקודות Tensor Core ארוכות; החומרה עוקבת אחרי התקיעות עם short scoreboard (shared/אריתמטיקה) ו-long scoreboard (זיכרון גלובלי).
- מחזור פעיל הוא מחזור שבו יש לפחות warp אחד תושב על ה-SM - גם אם הוא תקוע; "פעיל" עוסק בתושבוּת ולא בהנפקה, וה-occupancy הנמדד ממוצע על פני המחזורים הפעילים בלבד.
- ה-occupancy = warps פעילים / מקסימום warps פעילים; יש תיאורטי (חסם עליון מתצורת הlaunch וממגבלות הdevice, מחושב לפני הריצה) ונמדד (מה שקרה בפועל, ממוצע על מחזורים פעילים, כמעט תמיד נמוך מהתיאורטי).
- בדוגמת ה-H100 (64 warps, 32 blocks, 65,536 אוגרים, 228 KB smem ל-SM) עם kernel של 32 threads/block, 8 אוגרים/thread ו-12 KB smem/block: מגבלות ה-blocks הן אוגרים 256, smem 19, חומרה 32 ->
min = 19-> smem-bound -> תפוסה ~30%. - לוקחים את המינימום על פני מגבלות האוגרים, ה-shared memory, מגבלת החומרה ומגבלת ה-warps; המשאב שנגמר ראשון הוא הצוואר, ורק אותו משחררים - לא נוגעים בשאר.
- ה-Occupancy API מחשב את זה מדויק:
cudaOccupancyMaxActiveBlocksPerMultiprocessorנותן blocks ל-SM (ומשם occupancy), ו-cudaOccupancyMaxPotentialBlockSizeמציע block size ממקסם-occupancy; החישוב הידני עשוי לסטות ב-block אחד בגלל גרנולריות הקצאה ו-smem שמור. - ה-occupancy גבוה יותר הוא לא תמיד טוב יותר: המשאבים ל-SM קבועים, אז יותר warps = פחות אוגרים/smem ל-thread, מה שעלול לגרום ל-register spilling; ברגע שיש מספיק warps להסתרת latency (חוק Little), occupancy נוסף לא קונה כלום ולעיתים מזיק.
- חישובי GEMM מהירים על Hopper/Blackwell רצים ב-occupancy חד-ספרתי כי מעט warps מספיקים לרווֹת את ה-Tensor Cores (יחס 100:1 מול CUDA Cores); הם מעדיפים tiles גדולים ואוגרים רבים ל-thread על פני occupancy גבוה.
- ה-occupancy הוא אמצעי להסתרת latency מספקת, לא מטרה: ל-kernel compute-bound ממקסמים ניצול צנרות החישוב, ל-kernel memory-bound ממקסמים ניצול רוחב-הפס - ואת ה-occupancy מסתכלים רק כדי לוודא שעברנו את סף חוק Little.
- ב-ncu כל זה נמדד: סקציית Occupancy מראה את ה-block limit לכל משאב ואת התיאורטי מול הנמדד, Scheduler Statistics מראה warps פעילים/כשירים/נבחרים ו-
No Eligible, ו-Warp State Statistics מראה את סיבות ה-stall (למשלStall Long Scoreboard).