3.1 היררכיית הזיכרון הלוגית פתרון
פתרון - היררכיית הזיכרון הלוגית¶
הערה: מספרי האוגרים, גדלי ה-stack frame ופלטי ה-ptxas שלהלן הם דוגמאות מ-H100 עם sm_90a וגרסת CUDA מסוימת. הערכים המדויקים משתנים בין ארכיטקטורות וגרסאות compiler - מה שקבוע הוא הדפוס (איזה kernel גולש ל-local memory, מי מגביל את התפוסה), לא הספרה האחרונה. הריצו אצלכם וראו את המספרים שלכם.
פתרון תרגיל 1 - מיפוי מרחבי הזיכרון¶
המיפוי המלא, סימון-סימון:
A coeff[4] -> constant (__constant__, read-only, written from the host)
B tile[128] -> shared (__shared__, shared by the block)
C i -> register (local scalar)
D acc -> register (local accumulating scalar)
E buf[8] -> local (local array, dynamic index)
F in -> global (pointer from cudaMalloc)
G coeff[k]/tile[..] -> constant + shared (read from two spaces)
H buf[k % n] -> local (write to array in local memory)
I out -> global (pointer from cudaMalloc)
(2) שני משתנים עם scope זהה ומיקום מנוגד: acc (D) ו-buf (E). שניהם פרטיים ל-thread בודד (אותו scope בדיוק), אבל acc הוא סקלר ולכן יושב ב-register file (ה-SRAM המהיר ביותר ב-SM), בעוד buf הוא מערך שניגשים אליו באינדקס דינמי (buf[k % n], buf[i % 8]) ולכן אינו יכול לשבת באוגרים - הוא נוחת ב-local memory שב-GPU RAM האיטי. אותו scope, קצוות מנוגדים של מדרג המהירות.
(3) coeff מול in: שניהם בהיקף grid (כל הthreads רואים אותם), אבל coeff הוא constant memory - לקריאה בלבד מה-kernel, נכתב מהhost לפני הlaunch, ממוטמן וממוטב לשידור כשכל ה-warp קורא את אותה כתובת. in הוא global memory רגיל - קריא וכתיב, הוקצה ב-cudaMalloc, וללא הטבת השידור.
למה זה עבד: ההבחנה נשענת על שני חוקים פשוטים - מסמן המשתנה (__shared__, __constant__) קובע מרחב במפורש, ומבנה הגישה קובע register מול local: סקלר או מערך באינדקס קבוע בזמן קומפילציה נשאר ב-register, מערך באינדקס דינמי יורד ל-local.
איך להכליל: לכל משתנה ב-kernel שאלו שתי שאלות בסדר: "יש עליו מסמן מרחב?" (אם כן - זה המרחב), ואם לא - "האם ניגשים אליו באינדקס שידוע בזמן קומפילציה?" (אם כן register, אם לא local). מצביעים שמגיעים כארגומנטים מצביעים כמעט תמיד ל-global.
פתרון תרגיל 2 - לגרום לגלישה ל-local memory ולראות אותה¶
הקומפילציה עם -v מדפיסה את תקציב הזיכרון של כל kernel:
פלט מייצג (המספרים שלכם עשויים להיות שונים):
ptxas info : Compiling entry function '_Z5tightPKfPfi' for 'sm_90a'
ptxas info : Function properties for _Z5tightPKfPfi
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 '_Z5heavyPKiPKfPfi' for 'sm_90a'
ptxas info : Function properties for _Z5heavyPKiPKfPfi
256 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 32 registers, 380 bytes cmem[0]
הקריאה: ל-tight יש 0 bytes stack frame - כל המשתנים שלו נשארו באוגרים, אין local memory. ל-heavy יש 256 bytes stack frame - זהו בדיוק המערך scratch[64] (64 floats x 4 בתים = 256 בתים) שנחת ב-local memory כי הגישה scratch[idx[i] % 64] היא באינדקס דינמי. הגדילו את scratch ל-scratch[128] ותראו את ה-stack frame קופץ ל-512 בתים.
למה זה עבד: stack frame הוא ה-local memory שהcompiler הקצה ל-thread ב-GPU RAM. מערך שניגשים אליו באינדקס שאינו ידוע בזמן קומפילציה לא יכול לחיות באוגרים (לאוגר אין כתובת שאפשר לחשב בזמן ריצה), ולכן ה-ptxas שם אותו ב-local. השורות spill stores/spill loads הן תופעה קרובה אך נפרדת - גלישה של סקלרים כשנגמרים האוגרים; שתיהן נוחתות באותו local memory האיטי.
איך להכליל: --ptxas-options=-v (או -Xptxas -v) הוא הכלי הראשון לכל חקירת ביצועים - הוא חושף registers, stack frame ו-spills לכל kernel. כלל אצבע: stack frame או spill גדולים מ-0 הם דגל אדום שדורש בדיקה, כי כל גישה אליהם היא מול DRAM איטי.
פתרון תרגיל 3 - shared memory הוא פרטי ל-block¶
הקוד המלא של ה-main:
int blockSize = 4, numBlocks = 3, n = blockSize * numBlocks;
int *d_out, h_out[12];
CUDA_CHECK(cudaMalloc(&d_out, n * sizeof(int)));
privacy<<<numBlocks, blockSize>>>(d_out);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaMemcpy(h_out, d_out, n * sizeof(int), cudaMemcpyDeviceToHost));
for (int i = 0; i < n; ++i) printf("out[%2d] = %d\n", i, h_out[i]);
CUDA_CHECK(cudaFree(d_out));
הפלט:
out[ 0] = 0
out[ 1] = 0
out[ 2] = 0
out[ 3] = 0
out[ 4] = 1
out[ 5] = 1
out[ 6] = 1
out[ 7] = 1
out[ 8] = 2
out[ 9] = 2
out[10] = 2
out[11] = 2
הthreads 0-3 (block 0) ראו 0; threads 4-7 (block 1) ראו 1; threads 8-11 (block 2) ראו 2. אף thread ב-block 2 לא ראה 0 או 1.
למה זה עבד: המשתנה __shared__ int marker אינו משתנה יחיד גלובלי - הוא מוקצה מחדש בעותק נפרד לכל block, ב-L1/SRAM של ה-SM שעליו ה-block רץ. כל block כותב לעותק שלו, וקורא רק אותו. לו ה-shared היה חוצה blocks, כל 12 הthreads היו רואים ערך אחיד (של הכותב האחרון).
איך להכליל: זו הוכחה תפעולית להיקף ה-block של __shared__. אותה טכניקה - "כל block מטביע את זהותו ואז כולם מדווחים מה ראו" - מאמתת גבולות scope של כל מרחב זיכרון. אם צריכים לשתף בין blocks, חייבים לרדת ל-global memory ולתאם עם atomics.
פתרון תרגיל 4 - תיאום חוקי מול deadlock בין blocks¶
(1) גרסה A - חוקית ונכונה. atomicAdd(counter, 1) מבטיח שכל הגדלה של המונה תתבצע ללא הפרעה (קריאה-שינוי-כתיבה אטומית), גם כשthreads מ-blocks שונים מגדילים בו-זמנית. הקוד אינו מניח שום סדר בין ה-blocks - התוצאה הסופית (מספר האיברים החיוביים) זהה בכל סדר ביצוע. זה בדיוק דפוס התיאום החוקי: atomics עובדים על פני כל ה-grid בלי הנחות סדר.
(2) גרסה B - deadlock. תרחיש קונקרטי: נניח ל-GPU יש קיבולת ל-N blocks תושבים בו-זמנית, ואנחנו משיקים grid עם יותר מ-N blocks. החומרה מתזמנת קודם את block 0 (שנכנס למקום פנוי), אבל ה-block האחרון (gridDim.x - 1) עדיין לא תוזמן כי אין מקום - הוא ממתין בתור שיתפנה SM. אלא ש-block 0 תקוע בלולאת while (*flag == 0) וממתין ל-block האחרון שיכתוב את הדגל. block 0 לא יסתיים ולא יפַנה מקום, ה-block האחרון לא יתחיל, והדגל לעולם לא ייכתב. deadlock. הכשל נובע ישירות מכך שסדר הביצוע בין blocks אינו דטרמיניסטי ולא כל ה-blocks בהכרח תושבים יחד.
(3) הדרך הנכונה - שני kernels. גבול ה-kernel הוא נקודת סנכרון גלובלית טבעית. מפצלים את הלוגיקה: kernel ראשון עושה את "שלב הכותב", הhost ממתין לסיומו (cudaDeviceSynchronize או פשוט הlaunch הבאה באותו stream), ואז kernel שני עושה את "שלב הקורא":
producer<<<grid, block>>>(d_data); // all blocks finish here
consumer<<<grid, block>>>(d_data); // starts only after the previous one has fully finished
למה זה עבד: ה-GPU מבטיח שכל ה-blocks של kernel אחד סיימו לפני שה-kernel הבא באותו stream מתחיל - זה הסנכרון הגלובלי היחיד שאפשר לסמוך עליו. spin-wait בין blocks באותו kernel מסתמך על הנחת סדר שהחומרה לא נותנת.
איך להכליל: הכלל: תיאום בתוך block - __syncthreads(); תיאום נתונים בין blocks - atomics; סנכרון מלא בין כל ה-grid - סיום kernel וlaunch של kernel חדש (או grid.sync() של cooperative groups עם launch מיוחדת שמבטיחה תושבוּת). לעולם אל תבנו ממתין-מסתובב בין blocks של אותו kernel.
פתרון תרגיל 5 - registers מול occupancy¶
ניקח את המספרים המייצגים מתרגיל 2: heavy צורך 32 אוגרים ל-thread, tight צורך 16.
עבור heavy (R = 32):
max resident threads by register limit = 65536 / 32 = 2048
H100 physical max per SM = 2048
=> the register file is not the limiter: 2048 = 2048, full occupancy reached.
עבור tight (R = 16):
65536 / 16 = 4096 > 2048
=> definitely not the limiter; the limit is the max thread count (2048), not registers.
כדי לראות מתי האוגרים כן מגבילים, נניח kernel כבד שצורך 64 אוגרים ל-thread:
65536 / 64 = 1024 < 2048
=> registers cut occupancy in half: only 1024 resident threads instead of 2048.
fewer resident warps -> less ability to hide latency.
למה זה עבד: קובץ האוגרים של ה-SM (65,536 אוגרים בני 32 סיביות ב-H100) הוא משאב סופי שמתחלק בין כל הthreads התושבים. כשצריכת האוגרים ל-thread גבוהה, פחות threads נכנסים, והתפוסה יורדת. זה בדיוק המתח registers-מול-occupancy מההרצאה: יותר אוגרים ל-thread אינו מהיר יותר - הוא עלול דווקא לחתוך את מספר ה-warps הזמינים להסתרת latency.
איך להכליל: חסם התפוסה מהאוגרים הוא min(2048, 65536 / R) ל-SM (בקירוב; ההקצאה בפועל בגרנולריות של warp, 32 threads). אם kernel אטי ותפוסתו נמוכה, בדקו את R בפלט ה-ptxas; אפשר לכפות תקרה עם -maxrregcount=N או __launch_bounds__ כדי להעלות תפוסה - במחיר גלישה אפשרית ל-local, ולכן תמיד מודדים.
פתרון תרגיל 6 (בונוס) - קריאת PTX והרגיסטרים הווירטואליים¶
הפקת ה-PTX:
גוף ה-kernel tight נראה בערך כך (קוצר ונוקה לשם הבהירות):
.visible .entry _Z5tightPKfPfi(
.param .u64 in, .param .u64 out, .param .u32 n)
{
.reg .pred %p<2>;
.reg .f32 %f<4>;
.reg .b32 %r<6>;
.reg .b64 %rd<7>;
ld.param.u64 %rd1, [in];
ld.param.u64 %rd2, [out];
ld.param.u32 %r1, [n];
mov.u32 %r2, %ntid.x; // blockDim.x
mov.u32 %r3, %ctaid.x; // blockIdx.x
mov.u32 %r4, %tid.x; // threadIdx.x
mad.lo.s32 %r5, %r3, %r2, %r4;// i = blockIdx*blockDim + threadIdx
setp.ge.s32 %p1, %r5, %r1; // i >= n ?
@%p1 bra DONE;
cvta.to.global.u64 %rd3, %rd1;
mul.wide.s32 %rd4, %r5, 4;
add.s64 %rd5, %rd3, %rd4;
ld.global.f32 %f1, [%rd5]; // f1 = in[i]
fma.rn.f32 %f2, %f1, 0f40000000, 0f3F800000; // in[i]*2.0 + 1.0
cvta.to.global.u64 %rd6, %rd2;
add.s64 %rd6, %rd6, %rd4;
st.global.f32 [%rd6], %f2; // out[i] = f2
DONE:
ret;
}
הזיהוי:
- רגיסטרים וירטואליים מסוג
.f32(הנקודה הצפה):%f1מחזיק אתin[i],%f2מחזיק את התוצאהin[i]*2.0f + 1.0f. שימו לב שכל החישוב הצטמצם לפקודתfma.rn.f32אחת (fused multiply-add), עם הקבועים0f40000000(שהוא 2.0f) ו-0f3F800000(שהוא 1.0f). - רגיסטרים מסוג
.b32/.s32לחישוב האינדקס:%r2..%r4מקבלים אתblockDim,blockIdx,threadIdx, ו-%r5מחזיק אתiדרךmad.lo.s32. - גישות זיכרון:
ld.global.f32אחת (הקריאהin[i]) ו-st.global.f32אחת (הכתיבהout[i]) - סך הכל שתי גישות ל-global מול פעולה אריתמטית אחת (ה-fma).
למה זה עבד: ה-PTX חושף בדיוק את מה שהיה בלתי-נראה ב-CUDA C - הרגיסטרים הווירטואליים בני 32 סיביות. הצהרות ה-.reg הן ההכרזות עליהם, ופקודות ה-ld.global/st.global מסמנות במפורש כל מגע עם global memory. יחס של שתי גישות זיכרון לפעולת חשבון אחת הוא בדיוק החתימה של חישוב memory-bound.
איך להכליל: קריאת PTX היא הדרך לראות את היררכיית הזיכרון "מבפנים": .reg = רגיסטרים, ld/st.global = global, ld/st.shared = shared, ld.const = constant, ו-ld/st.local = local (הסימן לגלישה). כשרוצים להבין למה kernel איטי, ספירת פקודות ה-ld.global מול הפעולות האריתמטיות נותנת אומדן מיידי לarithmetic intensity - הכלי שנפַתח פורמלית בפרק על ביצועים.