8.5 Register pressure ואיזון משאבים פתרון
פתרון - Register pressure ואיזון משאבים¶
הערה: מספרי האוגרים, בייטי הדליפה, ה-occupancy והזמנים בפלטים למטה הם דוגמאות ממכונת H100 SXM עם CUDA 12, וישתנו אצלכם לפי הכרטיס, גרסת ה-compiler ומצב המערכת. במיוחד, מספר האוגרים ש-ptxas מקצה ל-heavy תלוי חזק בגרסת ה-nvcc - אל תיבהלו אם קיבלתם 96 במקום 128. מה שלא ישתנה הוא המבנה, סדרי הגודל והמגמות: light קליל, heavy כבד, הגבלה נמוכה מדי גורמת לדליפה, ונקודת ה-throughput המיטבית איננה בקצה ה-occupancy המרבי. סף ה-32 האוגרים ל-occupancy מלא הוא ספציפי ל-H100.
פתרון תרגיל 1 - קריאת הפלט של ptxas -v¶
הקומפילציה:
הפלט הצפוי (ptxas.txt):
ptxas info : Compiling entry function '_Z5lightPKfPfi' for 'sm_90a'
ptxas info : Function properties for _Z5lightPKfPfi
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 16 registers, 372 bytes cmem[0]
ptxas info : Compiling entry function '_Z5heavyPKfPfi' for 'sm_90a'
ptxas info : Function properties for _Z5heavyPKfPfi
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 128 registers, 372 bytes cmem[0]
הקריאה, kernel אחר kernel:
light: 16 registers, 0 smem, 0 spill
16 registers -> 32*16=512, rounded to 512, 65536/512=128 -> capped at 64 warps
-> max occupancy (from registers) = 100% ["healthy"]
heavy: 128 registers, 0 smem, 0 spill
128 registers -> 32*128=4096, 65536/4096=16 warps
-> max occupancy (from registers) = 16/64 = 25% ["under pressure"]
ל-heavy אין דליפה במצב הזה (spill stores/loads = 0): ה-ptxas הצליח להכניס את כל 64 הצוברים לתוך 128 אוגרים, מתחת לתקרת ה-255. הוא "לחוץ" לא כי הוא דולף, אלא כי מחיר ה-128 אוגרים הוא occupancy של 25% בלבד.
למה זה עבד: שלושת המספרים ב--Xptxas -v הם כל מה שצריך כדי למקם kernel על טבלת ה-occupancy. light עם 16 אוגרים רחוק מכל תקרה - occupancy מלא. heavy עם 128 אוגרים כבר חתך את ה-occupancy לרבע, אבל בלי דליפה - כי 128 עדיין מתחת ל-255. זיהוי "בריא מול לחוץ" הוא בדיוק המבט הזה: אוגרים גבוהים = occupancy נמוך; דליפה גדולה מ-0 = בעיה נפרדת וחמורה יותר.
איך להכליל: זו הבדיקה הראשונה על כל kernel חם, לפני כל אופטימיזציה. אם האוגרים גבוהים, ה-occupancy כנראה חסום-אוגרים ואפשר לשקול הגבלה; אם יש דליפה, זו כבר סכנה מיידית. תמיד הפרידו בין שתי השאלות: "כמה אוגרים" (משפיע על occupancy) ו"כמה דליפה" (משפיע ישירות על מהירות ה-thread).
פתרון תרגיל 2 - לכפות דליפה ולמדוד את ההאטה¶
תוספת ה-main למדידה:
static float timeKernel(void (*launch)(const float*, float*, int),
const float* d_in, float* d_out, int n, int reps) {
int threads = 256, blocks = (n + threads - 1) / threads;
cudaEvent_t s, e;
CUDA_CHECK(cudaEventCreate(&s)); CUDA_CHECK(cudaEventCreate(&e));
launch<<<blocks, threads>>>(d_in, d_out, n); // warmup
CUDA_CHECK(cudaDeviceSynchronize());
float best = 1e30f;
for (int r = 0; r < reps; r++) {
CUDA_CHECK(cudaEventRecord(s));
launch<<<blocks, threads>>>(d_in, d_out, n);
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;
}
int main(void) {
const int n = 4 * 1024 * 1024;
float *d_in, *d_out;
CUDA_CHECK(cudaMalloc(&d_in, n * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_out, n * sizeof(float)));
CUDA_CHECK(cudaMemset(d_in, 0, n * sizeof(float)));
float ms = timeKernel(heavy, d_in, d_out, n, 20);
printf("heavy: %.3f ms\n", ms);
CUDA_CHECK(cudaFree(d_in)); CUDA_CHECK(cudaFree(d_out));
return 0;
}
(ה-launch כאן מקבל מצביע-ל-kernel דרך wrapper; פשוט יותר לקרוא ל-heavy<<<...>>> ישירות אם מעדיפים.) שתי הקומפילציות והמדידות:
nvcc -O2 -arch=sm_90a -Xptxas -v -o rp_free regpressure.cu # no limit
nvcc -O2 -arch=sm_90a -maxrregcount=32 -Xptxas -v -o rp_32 regpressure.cu
./rp_free ; ./rp_32
הפלט הצפוי של ה--Xptxas -v עבור heavy בשתי הגרסאות, ואחריו הזמנים:
# no limit:
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 128 registers
# with -maxrregcount=32:
488 bytes stack frame, 456 bytes spill stores, 612 bytes spill loads
ptxas info : Used 32 registers
heavy (128 reg, 25% occ) : 1.83 ms
heavy ( 32 reg, 100% occ): 4.71 ms
התוצאה מפתיעה במבט ראשון: הגרסה עם ה-occupancy הגבוה פי ארבעה (100% מול 25%) רצה פי ~2.5 יותר לאט. הפרדוקס נפתר בשורת הדליפה: כשכפינו 32 אוגרים, 64 הצוברים לא נכנסו, וה-ptxas דלף מאות בייטים ל-local memory. כל צובר שדלף הפך גישת-אוגר (כמעט חינם) לגישת GPU RAM (מאות מחזורים), בתוך הלולאה החמה שרצה 48 x 64 פעמים לכל thread. ה-occupancy הגבוה יותר לא התקרב לפצות על אסון הדליפה.
למה זה עבד: -maxrregcount=32 היא תקרה קשיחה. כשה-kernel באמת צריך יותר אוגרים מהתקרה, ה-compiler לא יכול להמציא אוגרים - הוא דולף ל-local memory, שהוא GPU RAM איטי (ההרצאה). כאן ראינו את הצד ההרסני של "להוריד אוגרים כדי להעלות occupancy": ירדנו נמוך מדי, נכנסנו לדליפה, וההאטה מהדליפה בלעה בקלות את הרווח מה-occupancy.
איך להכליל: לעולם אל תורידו את תקרת האוגרים בלי לבדוק ב--Xptxas -v אם נכנסתם לדליפה. spill stores/loads שקפצו מ-0 למאות בייטים הם אזהרה שה-occupancy שהרווחתם עלול לעלות ביוקר. הכלל: occupancy גבוה יותר שווה משהו רק אם הוא לא בא במחיר של דליפה.
פתרון תרגיל 3 - למדוד occupancy ו-runtime עם launch_bounds¶
הגרסה עם ה-__launch_bounds__ (גוף זהה ל-heavy):
__global__ void __launch_bounds__(256, 8)
heavy_lb(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];
}
float s = 0.0f;
#pragma unroll
for (int i = 0; i < 64; i++) s += a[i];
out[idx] = s;
}
הקריאות ל-reportOccupancy והמדידות:
reportOccupancy("heavy", (const void*)heavy, 256);
reportOccupancy("heavy_lb", (const void*)heavy_lb, 256);
הפלט הצפוי (כולל -Xptxas -v על heavy_lb):
# ptxas for heavy_lb:
496 bytes stack frame, 464 bytes spill stores, 620 bytes spill loads
ptxas info : Used 32 registers
heavy blocks/SM=2 occupancy=25.0%
heavy_lb blocks/SM=8 occupancy=100.0%
summary table:
kernel registers/thread occupancy runtime
heavy 128 25% 1.83 ms
heavy_lb 32 100% 4.68 ms
כן, heavy_lb הגיע ל-occupancy של 100% שכיוונו אליו - בדיוק כפי שהחשבון של ההרצאה חוזה (__launch_bounds__(256, 8) -> 65536/(256*8) = 32 אוגרים -> occupancy מלא). אבל ה-runtime שלו גרוע יותר (4.68 מול 1.83 ms), מאותה סיבה בדיוק כמו בתרגיל 2: כדי לרדת ל-32 אוגרים ה-compiler נאלץ לדלוף. __launch_bounds__ הוא בסך הכל דרך מסודרת יותר (לכל kernel) להשיג את מה ש--maxrregcount עשה גלובלית - ואותה מלכודת חלה עליו.
למה זה עבד: cudaOccupancyMaxActiveBlocksPerMultiprocessor מחזיר את ה-occupancy התיאורטי מהאוגרים וה-smem, ולכן הוא מאשש שהגענו ל-100% התיאורטי. אבל occupancy תיאורטי גבוה איננו throughput גבוה: אותה הורדה ל-32 אוגרים שהעלתה את ה-occupancy היא בדיוק זו שהכניסה לדליפה. שני הכלים (__launch_bounds__ ו--maxrregcount) מחליפים דליפה ב-occupancy - וכאן ההחלפה הפסידה.
איך להכליל: העדיפו __launch_bounds__ על -maxrregcount בקוד production, כי הוא מכוון לכל kernel בנפרד ונשמר יחד עם הקוד. אבל תמיד מדדו runtime, לא רק occupancy: cudaOccupancyMaxActiveBlocksPerMultiprocessor אומר לכם כמה warps ייכנסו, לא כמה מהר ה-kernel ירוץ. המספר השני הוא היחיד שקובע.
פתרון תרגיל 4 - למצוא את מספר האוגרים שממקסם throughput¶
הסקריפט מההכנה, והפלט הצפוי לצד מספרי האוגרים וה-occupancy שנאספו מ--Xptxas -v ומתרגיל 3:
maxrreg actual registers occupancy runtime
24 24 100% 6.02 ms <- heavy spilling
28 28 100% 5.31 ms <- spilling
32 32 100% 4.71 ms <- spilling
40 40 80% 3.24 ms <- light spilling
48 48 66% 2.15 ms
64 64 50% 1.72 ms <- minimum (no spilling)
96 96 33% 1.78 ms
128 128 25% 1.83 ms
168 168 18% 2.05 ms
255 168* 18% 2.05 ms *the kernel doesn't need more than 168
עקומת ה-runtime היא "בטן" קלאסית: יורדת מהקצה השמאלי (occupancy מרבי אך דליפה), מגיעה למינימום סביב 64 אוגרים, ואז עולה שוב לאט לכיוון הקצה הימני (בלי דליפה אך occupancy נמוך). שני הקצוות:
left edge (24-40 registers): spilling to local memory -> GPU RAM accesses in the hot loop -> slow
right edge (128-255 registers): occupancy 25% and below -> not enough warps to hide latency -> slow
middle (~64 registers): no spilling, occupancy 50% - enough to hide latency. the fastest.
מספר האוגרים שממקסם throughput כאן הוא 64 (occupancy 50%), ולא הקצה של ה-occupancy המרבי (24-32 אוגרים, 100%). ההפרש בין המהיר (1.72 ms) לאיטי-מ-occupancy-מלא (4.71 ms) הוא פי ~2.7 - כולו לרעת ה-occupancy הגבוה.
למה זה עבד: זה הביטוי המדיד המלא של המסר. הקצה השמאלי לוקה בדליפה; הקצה הימני לוקה בהסתרת-latency חלשה (חוק Little, 7.4 - מעט warps מדי). המיטב באמצע: מספיק אוגרים כדי להימנע מדליפה, מספיק warps כדי להסתיר latency. 50% occupancy מספיק בהחלט להסתרה (זכרו: כבר ~20% מסתירים שרשרת load אופיינית), ומעבר לזה אין מה להרוויח, רק אוגרים לאבד.
איך להכליל: זו הלולאה שסוגרת את הפרק - שנו אוגרים, מדדו runtime, חזרו. ה-throughput-optimum כמעט תמיד ב-occupancy בינוני, לא מרבי. אל תרדפו אחרי 100% occupancy; רדפו אחרי runtime מינימלי, וברוב ה-kernels הכבדים הוא נמצא בנקודה שבה יש מספיק warps להסתרה וגם מספיק אוגרים למניעת דליפה.
פתרון תרגיל 5 (בונוס) - אימות עם ncu וקריאת סטטיסטיקות local memory¶
הרצת הפרופיילר על שתי הגרסאות:
הפלט הרלוונטי הצפוי (מקוצר לשורות המכריעות):
# rp_free (no limit, 128 registers)
Registers Per Thread register/thread 128
Achieved Occupancy % 23.9
Local Memory (per thread) byte/thread 0
Compute (SM) Throughput % 71.4
Warp State (Stall) - top: Stall Wait, Stall Not Selected
# rp_32 (-maxrregcount=32, 32 registers)
Registers Per Thread register/thread 32
Achieved Occupancy % 88.2
Local Memory (per thread) byte/thread 488
Compute (SM) Throughput % 29.6
Warp State (Stall) - top: Stall Long Scoreboard
הקריאה:
metric rp_free (128) rp_32 (32) conclusion
Registers Per Thread 128 32 matches ptxas -v
Achieved Occupancy 24% 88% we raised occupancy...
Local Memory 0 488 B ...but started spilling
SM Throughput 71% 30% ...and throughput crashed
dominant stall reason Wait/NotSel Long Scoreboard spilling = waiting on GPU RAM
כל השורות מספרות את אותו סיפור. "Registers Per Thread" של ה-ncu תואם למספר מ--Xptxas -v (128 מול 32) - אימות צולב נחמד. הגרסה המוגבלת אכן העלתה את ה-Achieved Occupancy (מ-24% ל-88%), אבל ה-"Local Memory" קפץ מ-0 ל-488 בייט ל-thread - זו הדליפה - וה-Compute Throughput צנח מ-71% ל-30%. וסיבת ה-stall הדומיננטית בגרסה שדולפת היא Stall Long Scoreboard: המתנה ל-load מהזיכרון הגלובלי, בדיוק לאן שהדליפה הלכה (local memory יושב ב-GPU RAM).
למה זה עבד: ה-ncu מודד את מה שהחשבון של ההרצאה חזה. ה-occupancy התיאורטי (100%) וה-Achieved (88%) קרובים; פערים קטנים נובעים מ-tail effects ומאיזון blocks. אבל השורה המכרעת היא ה-throughput: למרות ה-occupancy הגבוה, ה-SM Throughput נמוך יותר, כי כל warp מבזבז את זמנו בהמתנה לדליפה במקום לחשב. Stall Long Scoreboard הוא טביעת-האצבע של הדליפה.
איך להכליל: ncu --set full הוא בורר-הסכסוכים הסופי. כשאתם מתלבטים אם הגבלת-אוגרים עזרה, הביטו בשלושת המדדים יחד: Registers Per Thread (כמה אוגרים), Achieved Occupancy (כמה warps בפועל), ו-Local Memory / SM Throughput (האם דלפתם וכמה מהר בסוף). occupancy גבוה עם Stall Long Scoreboard דומיננטי ו-Local Memory גדול מ-0 הוא החתימה של register pressure שנפתר לא-נכון - בדיוק מה שראינו לאורך כל התרגול.