4.1 PTX שפת הביניים הווירטואלית פתרון
פתרון - PTX - שפת הביניים הווירטואלית¶
הערה: ה-PTX שלהלן הופק מ-H100 עם sm_90a וגרסת CUDA מסוימת. שמות ה-registers, סדר ההוראות והתוויות (labels כמו $L__BB0_2) עשויים להשתנות מעט בין גרסאות compiler - מה שקבוע הוא הדפוס: הצהרות .reg, ה-mad לאינדקס, ה-setp/bra ל-if, ה-fma לחישוב, וה-ld/st.global לגישות הזיכרון. הפעילו אצלכם וראו את המספרים שלכם.
פתרון תרגיל 1 - הפקת PTX וקריאת שלוש המשפחות¶
הפקה:
בראש הקובץ יופיעו ההנחיות, ובגוף ה-kernel יופיעו ארבע הצהרות .reg:
.reg .pred %p<2>; // one predicate register (%p1) for the if condition
.reg .f32 %f<3>; // two float registers: %f1 = in[i], %f2 = the result
.reg .b32 %r<6>; // five 32-bit registers: indices and n
.reg .b64 %rd<7>; // six 64-bit registers: pointers and addresses
(2) הוצהרו ארבע משפחות: .pred (בוליאני לתנאי), .f32 (נקודה צפה), .b32 (מספרים שלמים / bit-buckets בני 32) ו-.b64 (מצביעים וכתובות בני 64). הסימון %f<3> פירושו משפחת שמות בגודל 3: %f0, %f1, %f2.
(3) ה-fma:
הקבוע הוא 0f3FC00000, שהוא 1.5 בעשרוני (ייצוג ביטים של IEEE 754 float). הסיומת rn היא round-to-nearest-even - מצב העיגול של ברירת המחדל ב-IEEE 754.
(4) שלוש ה-mov נמצאות מיד לפני חישוב האינדקס:
mov.u32 %r3, %ntid.x; // blockDim.x
mov.u32 %r4, %ctaid.x; // blockIdx.x
mov.u32 %r5, %tid.x; // threadIdx.x
למה זה עבד: הצהרת .reg היא ההכרזה על ה-registers הווירטואליים שה-kernel צורך, מקובצת לפי סוג. הcompiler בחר float ל-in[i] ולתוצאה, שלמים ל-n ולאינדקסים, ו-64 סיביות לכתובות (על .address_size 64), ו-predicate בודד לתנאי.
איך להכליל: תמיד קראו קודם את בלוק ה-.reg - הוא "תקציב" ה-registers הווירטואלי של ה-kernel, ומכתיב מיד אילו סוגי נתונים בשימוש. משפחת .pred מסגירה תנאים; .b64 מרובה מסגירה עבודת מצביעים כבדה.
פתרון תרגיל 2 - מיפוי PTX לשורות ה-CUDA C++¶
המיפוי המלא:
mov.u32 %r3, %ntid.x; \
mov.u32 %r4, %ctaid.x; > int i = blockIdx.x*blockDim.x + threadIdx.x;
mov.u32 %r5, %tid.x; |
mad.lo.s32 %r1, %r4, %r3, %r5; /
setp.ge.s32 %p1, %r1, %r2; \ if (i < n) [inverted to i >= n -> skip]
@%p1 bra $L__BB0_2; /
ld.global.f32 %f1, [%rd5]; // read in[i]
fma.rn.f32 %f2, %f1, %f1, 0f3FC00000; // in[i]*in[i] + 1.5f
st.global.f32 [%rd6], %f2; // out[i] = ...
(1) שלוש ה-mov בשילוב mad.lo.s32 %r1, %r4, %r3, %r5 הם i = blockIdx.x * blockDim.x + threadIdx.x. ההוראה mad (multiply-add) מבצעת %r4 * %r3 + %r5 בהוראה אחת.
(2) setp.ge.s32 %p1, %r1, %r2 מחשב את הפרדיקט i >= n, ו-@%p1 bra $L__BB0_2 קופץ לסוף אם הוא אמת. הcompiler הפך את i < n לתנאי ההפוך i >= n כדי שכישלון התנאי יגרום לקפיצה מעל גוף ה-if - זה הדפוס הסטנדרטי של קומפילציית if המגוננת.
(3) ld.global.f32 %f1, [%rd5] מממשת את הקריאה in[i]; st.global.f32 [%rd6], %f2 מממשת את הכתיבה out[i] = ....
(4) שתי גישות ל-global memory (ld אחת + st אחת) מול פעולה אריתמטית אחת (fma). היחס הנמוך - פחות מפעולה אחת לכל גישת זיכרון - הוא חתימה מובהקת של kernel memory-bound: הזמן נשלט על ידי תעבורת ה-DRAM, לא על ידי החישוב.
למה זה עבד: PTX הוא מספיק קרוב לקוד המקור כדי שכל מבנה שפה גבוהה (חישוב אינדקס, if, קריאה/כתיבה) ממופה לקבוצת הוראות מזוהה. הצירוף setp + @predicate bra הוא תמיד תרגום של הסתעפות.
איך להכליל: כדי לקרוא PTX של kernel זר, אתרו קודם את ה-ld.param (הארגומנטים), אז את ה-mov מה-registers המובנים (חישוב אינדקס), אז את ה-setp/bra (בקרת זרימה), ולבסוף את ה-ld/st.global (ה-I/O). מה שנשאר באמצע הוא החישוב.
פתרון תרגיל 3 - שינוי הקוד וצפייה בשינוי ה-PTX¶
(1) עם + 3.0f הקבוע ב-fma משתנה:
(2) עם in[i] * 2.0f בלבד, אין מה לצבור, ולכן הcompiler פולט כפל ולא fma:
(3) עם in[i]*in[i] + in[i], כל שלושת האופרנדים הם registers - אין ליטרל 0f...:
(4) לולאת for קטנה מוסיפה register של predicate ומבנה קפיצה. בערך:
.reg .pred %p<3>;
...
$L__BB0_3:
ld.global.f32 %f4, [%rd7];
add.f32 %f5, %f5, %f4; // acc += in[i+k]
add.s64 %rd7, %rd7, 4;
add.s32 %r7, %r7, 1;
setp.lt.s32 %p2, %r7, 4; // k < 4 ?
@%p2 bra $L__BB0_3; // loop back
(בפועל הcompiler עשוי לפרוש - unroll - לולאה קצרה וקבועה כזו, ואז לא תראו bra בכלל אלא ארבע add.f32 רצופות. גם זה תשובה נכונה, וזו הזדמנות טובה לראות loop unrolling ב-PTX.)
למה זה עבד: הcompiler בוחר את ההוראה המדויקת לפי הצורה האלגברית: a*b+c הוא fma, a*b לבדו הוא mul, ואופרנד קבוע נכנס כליטרל 0f... בעוד אופרנד משתנה נכנס כ-register. לולאה מולידה setp+bra, אלא אם היא קצרה וקבועה ואז היא נפרשת.
איך להכליל: זו הדרך הזולה ללמוד מה הcompiler עושה: שנו שורה אחת, הפיקו PTX, השוו. כך רואים מתי fma נוצרת, מתי הcompiler מקבע קבועים (constant folding), ומתי הוא פורש לולאות - בלי לנחש.
פתרון תרגיל 4 - כתיבת inline PTX משלכם¶
הקוד המלא:
#include <cstdio>
#include <cstdlib>
#define CUDA_CHECK(call) \
do { \
cudaError_t err_ = (call); \
if (err_ != cudaSuccess) { \
fprintf(stderr, "CUDA error %s at %s:%d\n", \
cudaGetErrorString(err_), __FILE__, __LINE__); \
exit(EXIT_FAILURE); \
} \
} while (0)
__device__ __forceinline__ unsigned lane_id() {
unsigned id;
asm volatile("mov.u32 %0, %%laneid;" : "=r"(id));
return id;
}
__global__ void lanes(unsigned* out) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
out[i] = lane_id();
}
int main() {
const int n = 64;
unsigned *d_out, h_out[64];
CUDA_CHECK(cudaMalloc(&d_out, n * sizeof(unsigned)));
lanes<<<1, n>>>(d_out); // a single block, 64 threads = 2 warps
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaMemcpy(h_out, d_out, n * sizeof(unsigned),
cudaMemcpyDeviceToHost));
for (int i = 0; i < n; ++i)
printf("thread %2d -> laneid %2u\n", i, h_out[i]);
CUDA_CHECK(cudaFree(d_out));
return 0;
}
הcompilation והרצה:
הפלט (מקוצר):
thread 0 -> laneid 0
thread 1 -> laneid 1
...
thread 31 -> laneid 31
thread 32 -> laneid 0
thread 33 -> laneid 1
...
thread 63 -> laneid 31
(3) ה-laneid מתאפס כל 32 כי הוא המיקום של הthread בתוך ה-warp שלו, וב-warp יש בדיוק 32 threads. בבלוק חד-ממדי הוא שווה ל-threadIdx.x % 32: הthreads 0-31 הם warp 0 (laneid 0..31), threads 32-63 הם warp 1 (laneid 0..31 שוב).
(4) ה-popcount:
__device__ __forceinline__ int popcount(unsigned x) {
int r;
asm("popc.b32 %0, %1;" : "=r"(r) : "r"(x));
return r;
}
עבור popcount(threadIdx.x): הthread 7 (בינארי 111) מחזיר 3, הthread 8 (1000) מחזיר 1, הthread 31 (11111) מחזיר 5.
למה זה עבד: בלוק ה-asm volatile מזריק את הוראת ה-PTX mov.u32 %0, %%laneid היישר לזרם ההוראות של ה-kernel. ה-%% מגן על סימן האחוז כדי שהcompiler לא יפרש אותו כ-placeholder, ו-"=r" מקשר את register הפלט למשתנה id. volatile הכרחי כאן: בלעדיו הcompiler עלול היה לחשב את הקריאה פעם אחת ולשכפל, כי הוא לא "יודע" ש-%laneid שונה בין threads.
איך להכליל: התבנית asm volatile("<PTX>" : "=r"(out) : "r"(in)); היא השלד לכל inline PTX. השתמשו ב-volatile לכל דבר עם תופעות לוואי או ערך תלוי-הקשר (registers מיוחדים, גישות זיכרון), וּותרו עליו רק לפעולות טהורות כמו popc. את התכונות המתקדמות באמת (wgmma, tma) מזריקים באותה תבנית בדיוק, רק עם הוראות ארוכות ומורכבות הרבה יותר.
פתרון תרגיל 5 - PTX ISA version מול target¶
(1) שורת ה-.target משתנה בין שתי ההפקות:
(2) שורת ה-.version זהה בשני הקבצים (למשל .version 8.3), כי היא מייצגת את גרסת ה-PTX ISA של הcompiler שהפיק את הקובץ - כלומר את גרסת ה-CUDA Toolkit המותקנת - ולא את הכרטיס. אותו nvcc פולט אותו .version לכל טרגט.
(3) ההנחיה .target היא זו שקובעת על אילו GPU פיזיים אפשר להריץ: היא ה-compute capability המינימלית. GPU בעל compute capability זהה או גבוה יותר יוכל להריץ (דרך JIT); נמוך יותר - לא.
למה זה עבד: אלה שני צירים בלתי-תלויים. .version הוא "באיזו שפה" (איזו גרסת PTX ISA), ותלוי בcompiler. .target הוא "לאיזו חומרה", ותלוי בדגל -arch. אפשר לכתוב PTX חדש (.version גבוה) שמכוון לחומרה ישנה (.target sm_70), ולהפך.
איך להכליל: כשמנתחים PTX זר, .target אומר לכם את רצפת התאימות של הקוד, ו-.version אומר לכם עם איזו ערכת כלים הוא נבנה (שימושי לאבחון "PTX חדש מדי לדרייבר ישן", שגיאה אמיתית שנתקלים בה כשהדרייבר מיושן מכדי ל-JIT את גרסת ה-PTX).
פתרון תרגיל 6 (בונוס) - תאימות-קדימה: PTX מנצח SASS¶
(1)-(2) הטמעת SASS של Ampere בלבד, והרצה על Hopper:
SASS של sm_80 הוא קוד מכונה של Ampere, ואינו רץ על ה-SM של Hopper. המאקרו CUDA_CHECK תופס את השגיאה no kernel image is available for execution on the device (קוד cudaErrorNoKernelImageForDevice).
(3)-(4) הטמעת PTX של Ampere:
הקובץ מכיל PTX של compute_80, שהוא נייד. כשה-runtime מנסה לטעון את ה-kernel ולא מוצא SASS מתאים ל-Hopper, ה-hardware driver (nvidia.ko דרך libcuda.so) מבצע ל-PTX compilation בזמן ריצה (JIT) ל-SASS של sm_90, וה-kernel רץ. ההרצה הראשונה עלולה להיות מעט אטית יותר בגלל זמן ה-JIT (שממוטמן אחר כך).
למה זה עבד: זהו מנגנון התאימות-קדימה במלוא הדרו. SASS מקובע לדור שעבורו הודר; PTX הוא IR נייד שה-driver מתרגם לדור בפועל בזמן ריצה. השלב המדויק שבו PTX הופך ל-SASS של Hopper הוא ה-JIT ב-hardware driver בזמן טעינת ה-module - לא ב-nvcc בזמן הבנייה, ולא ב-ptxas (שרץ רק כשמבקשים code=sm_XX).
איך להכליל: לעמידות, הטמיעו תמיד גם PTX וגם SASS: -gencode arch=compute_90,code=sm_90 ל-SASS מהיר על החומרה הידועה, ועוד -gencode arch=compute_90,code=compute_90 ל-PTX כרשת ביטחון לחומרה עתידית. הקיצור -arch=sm_90 לבדו כבר מטמיע גם SASS של sm_90 וגם PTX של compute_90, ולכן הוא תואם-קדימה דרך JIT. כדי להטמיע SASS בלבד (בלי תאימות-קדימה) יש לבקש זאת במפורש עם -gencode arch=compute_90,code=sm_90 - נוח לפיתוח, מסוכן לשחרור.