4.1 PTX שפת הביניים הווירטואלית הרצאה
בשיעור 0.4 ציירנו את מפת שרשרת התוכנה המלאה, ובתוכה זיהינו שני ייצוגים של קוד ה-device שקל לבלבל ביניהם: PTX, שפת הביניים הווירטואלית, ו-SASS, האסמבלי הילידי של החומרה. שם הסתפקנו בהיכרות מרחוק. פרק 4 כולו מוקדש לחלק האמצעי הזה של השרשרת - הcompiler ושתי שפות הביניים - והשיעור הנוכחי פותח אותו בהעמקה על PTX (Parallel Thread eXecution - הרצת שרשורים מקבילית). נבין למה PTX הוא בו-זמנית "מכונה וירטואלית" ו"ארכיטקטורת סט הוראות", למה הוא דומה יותר ל-LLVM-IR מאשר ל-x86, ומהו "המותן הצר" שהוא מהווה בין עולם הcompilers לעולם הdevices. נפרק שורות PTX אמיתיות הוראה-הוראה, נראה כיצד להפיק PTX מכל kernel שכתבנו, נלמד לשלב PTX ידני בתוך CUDA C++, ונבין מתי זו הדרך היחידה לנצל חומרה חדשה. בשיעור 4.2 נרד ל-SASS, ובשיעור 4.3 נפרק את הcompiler nvcc עצמו לרכיביו.
שני הפנים של PTX - virtual machine ו-instruction set¶
הדבר הראשון שצריך להבין ב-PTX הוא שהתיעוד הרשמי של NVIDIA מתאר אותו בשתי דרכים שנשמעות סותרות: הוא גם "מכונה וירטואלית" (virtual machine) וגם "ארכיטקטורת סט הוראות" (instruction set architecture / ISA). שני התיאורים נכונים, והם שני צדדים של אותו מטבע. PTX הוא סט ההוראות שבו מתכנתים כנגד מודל של מכונה מופשטת - מכונה שאינה קיימת פיזית באף שבב. בזכות ההפשטה הזו, מתכנת או compiler שמייצרים PTX יכולים לסמוך שהתוכנית תרוץ באותה סמנטיקה על הרבה מכונות פיזיות שונות, כולל מכונות שעדיין לא נבנו.
זו בדיוק הנקודה שמבדילה את PTX מ-ISA אמיתי של מעבד. x86_64, aarch64 או SPARC הם סטי הוראות של חומרה ממשית - כל הוראה בהם מתורגמת ישירות לפעולת סיליקון. PTX, לעומתם, הוא ייצוג ביניים (IR - Intermediate Representation), בדיוק במובן ש-LLVM-IR הוא ייצוג ביניים של compiler. הוא לא רץ על שום חומרה כמות שהוא. רכיבי ה-PTX שבתוך קובץ CUDA עוברים compilation בזמן ריצה (JIT - Just-In-Time) על ידי דרייבר ה-GPU, שמתרגם אותם ל-SASS של המכשיר הספציפי שעליו התוכנית רצה בפועל.
המותן הצר - narrow waist. הדימוי המרכזי שכדאי לזכור הוא ש-PTX הוא ה"מותן הצר" שמפריד בין עולם התוכנה לעולם החומרה. מעל המותן יושבים compilers רבים ושפות רבות - CUDA C++, אבל גם compilers של Triton, של Julia, של Mojo ועוד - וכולם צריכים רק לדעת לפלוט PTX. מתחת למותן יושבים devices רבים - כל דור GPU של NVIDIA - וכל אחד מהם צריך רק לדעת לתרגם PTX ל-SASS שלו. במקום ש-N שפות תצטרכנה להכיר M devices (מכפלה של N כפול M מסלולי compilation), כל צד מכיר רק את PTX. זהו אותו רעיון ארכיטקטוני שהופך את IP למותן הצר של האינטרנט, או את LLVM-IR למותן הצר של שרשראות הבנייה המודרניות.
Triton CUDA C++ Julia Mojo ... <- software world (many compilers and languages)
\ | | /
\ | | /
v v v v
+------------------------------------+
| PTX | <- the narrow waist: stable virtual ISA
| (virtual ISA / IR) |
+------------------------------------+
| | | |
v v v v
Volta Ampere Hopper Blackwell ... <- hardware world (many devices)
מודל המכונה - SIMT מנוסח במדויק¶
מהי אותה "מכונה וירטואלית" שכנגדה כותבים PTX? שם המודל טמון בשם השפה עצמה. במודל המכונה של PTX קיימת יחידת הוראות אחת עבור מספר מעבדים (a single instruction unit for multiple processors). כל מעבד מריץ thread (thread) אחד, אבל כל ה-threads האלה חייבים לבצע את אותן ההוראות - ומכאן השם parallel thread execution, הרצת שרשורים מקבילית. ה-threads מתאמים זה עם זה דרך shared memory (זיכרון משותף), ומשיגים תוצאות שונות זה מזה באמצעות ה-registers (רגיסטרים / אוגרים) הפרטיים שלהם.
שלושת המשפטים האלה הם בדיוק ניסוח מדויק של מודל ה-SIMT (Single Instruction, Multiple Threads) שראינו בפרק הארכיטקטורה. בואו נפרק אותם:
- יחידת הוראות אחת, הרבה מעבדים. זו הסיבה הכלכלית לכל הארכיטקטורה: במקום שכל thread יישא מונה-פקודות (program counter) ומפענח (decoder) משלו כמו ב-CPU, קבוצה שלמה של threads חולקת יחידת הנפקה אחת. ב-H100 היחידה הזו מנפיקה הוראה אחת ל-warp (קבוצה של 32 threads) בכל מחזור.
- אותן הוראות. כל 32 הthreads ב-warp מבצעים את אותה הוראת PTX. כאן נולד warp divergence (התפצלות ה-warp) שראינו בפרק 1: אם התנאי של
ifמפצל את הthreads, החומרה מריצה את שני הענפים בזה אחר זה, כשבכל ענף חלק מהthreads מושבתים. - תיאום דרך shared memory, בידול דרך registers. מה שמבדיל בין הthreads אינו הקוד - הקוד זהה - אלא הנתונים הפרטיים באוגרים. הthread מספר 5 וthread מספר 6 מריצים את אותה הוראת
mov.u32 %r3, %tid.x, אבל כל אחד מקבל לתוך ה-register הפרטי שלו ערךtidאחר (5 ו-6 בהתאמה), וכך אותה הוראה מובילה לתוצאות שונות.
זו הנקודה העדינה והחשובה ביותר במודל: הקוד אחיד, הנתונים פרטיים. אותה זרימת הוראות בדיוק, כפול 32 עותקי registers, שווה 32 חישובים שונים שרצים בו-זמנית.
גרסאות ותאימות קדימה - compute capability ו-JIT¶
ראינו בשיעור 0.4 ש-PTX מגורסת לפי מספר compute capability, שהוא שם נרדף לגרסת ה-SM המינימלית הנתמכת (הטרגט sm_XX). כדאי לחדד את המנגנון, כי הוא הסיבה שבגללה PTX שווה את הטרחה.
קובץ PTX פותח תמיד בשלוש הנחיות (directives) שמצהירות על ההקשר:
.version 8.3 // PTX ISA version (of the compiler that produced it)
.target sm_90a // minimum architecture the code targets (Hopper, H100)
.address_size 64 // address width: 64 bits
ההנחיה .target sm_90a היא זו שמבטאת את ה-compute capability. תאימות קדימה (forward compatibility) משמעה ש-GPU בעל compute capability זהה או גבוה יותר יוכל להריץ את התוכנית, גם אם הוא לא היה קיים כשהקוד הודר. המנגנון שמאפשר זאת הוא ה-JIT: ה-PTX המוטמע בקובץ עובר compilation בזמן ריצה ל-SASS של המכשיר החדש. קובץ שהודר עם PTX של compute_90 יוכל בעתיד לרוץ על Blackwell גם בלי compilation מחדש, כי הדרייבר יתרגם את ה-PTX ל-SASS של Blackwell בזמן הטעינה.
השוו זאת ל-SASS: SASS של sm_90 מקובע ל-Hopper ולעולם לא ירוץ על דור אחר. לכן, מבחינת עמידות, PTX הוא הטרגט העדיף - הוא ה"ביטוח" שהקוד ימשיך לרוץ על חומרה עתידית. זו הסיבה שספריות רבות מטמיעות PTX לצד SASS: SASS לביצועים מיטביים על החומרה הידועה, PTX כרשת ביטחון לחומרה שעוד לא קיימת.
compile time run time
CUDA C++ --nvcc--> PTX (compute_90) --driver JIT--> SASS of the actual device
| ^
| forward-compatible: same PTX ----------+
| translated to Hopper, to Blackwell, ...
חשוב לדייק היכן ה-JIT מתרחש: לא ב-nvcc ולא ב-ptxas בזמן הבנייה, אלא ב-hardware driver (הדרייבר של NVIDIA, libcuda.so ו-nvidia.ko) בזמן טעינת ה-module. זו ההבחנה בין compiler driver ל-hardware driver שהזהרנו עליה בשיעור 0.4: את ה-JIT מ-PTX ל-SASS מבצע ה-hardware driver, לא ה-compiler driver.
אנטומיה של הוראת PTX - קריאת שלוש הדוגמאות¶
PTX הוא קוד קריא, ואחרי מעט תרגול קוראים אותו כמעט כמו פסאודו-קוד. נפרק שלוש שורות אמיתיות, אחת מכל משפחה מרכזית: הצהרת registers, פעולה אריתמטית, וקריאת register מובנה.
1. הצהרת registers וירטואליים.
זו הנחיה לcompiler ה-PTX-ל-SASS (ptxas): ה-kernel צורך שבעה registers וירטואליים מסוג .f32 (floating point בן 32 סיביות). התחביר %f<7> מכריז על משפחת שמות %f0 עד %f6. שימו לב לשתי עובדות: ראשית, אלה registers וירטואליים - PTX מניח מספר בלתי-מוגבל של registers, ורק בשלב ה-SASS הם ממופים למספר הפיזי המוגבל. שנית, כאן טמונה העובדה שהייתה חסרה בשיעור על ה-registers: registers של PTX הם בני 32 סיביות (.f32, .b32, .u32), והם מוקצים דינמית לקבוצות של threads (warps) מתוך קובץ האוגרים של ה-SM. ערך double בן 64 סיביות תופס שני registers כאלה (.f64 ממופה לזוג).
טבלת סוגי ה-registers הנפוצים ב-PTX:
| מסמן | משמעות | דוגמה |
|---|---|---|
.f32 / .f64 |
נקודה צפה 32 / 64 סיביות | %f3, %fd1 |
.b32 / .b64 |
חבילת סיביות (bit-bucket) 32 / 64 | %r1, %rd1 |
.u32 / .s32 |
שלם 32 סיביות ללא/עם סימן | %r2 |
.pred |
דגל בוליאני לתנאים (predicate) | %p1 |
2. פעולת fused multiply-add.
הכפל את התוכן של %f3 ב-%f4, הוסף את הקבוע, ואחסן ב-%f5. כל הערכים הם floating point בני 32 סיביות. שני פרטים שכדאי לפענח: הקבוע 0f3FC00000 הוא ייצוג הקסדצימלי של המספר 1.5 בתקן IEEE 754 (הקידומת 0f מציינת ליטרל של float בייצוג ביטים גולמי). והסיומת rn היא מצב העיגול (rounding mode): rn = round to nearest, ties to even, כלומר עיגול לזוגי הקרוב, שהוא מצב ברירת המחדל של IEEE 754. שם ההוראה, fma, מציין fused - הכפל והחיבור מתבצעים בדיוק מלא יחיד בלי עיגול ביניים, ולכן fma מדויקת יותר (ולא רק מהירה יותר) מ-mul ואחריה add נפרדת.
מצבי העיגול שאפשר לראות כסיומת: rn (nearest even), rz (toward zero), rm (toward minus infinity), rp (toward plus infinity).
3. קריאת registers מובנים - האינדקסים.
שלוש ההוראות האלה מעתיקות לשלושה registers את שלושת הרכיבים שמהם כל kernel מחשב את האינדקס הגלובלי של הthread. אלה registers מיוחדים מובנים (special registers) שהחומרה ממלאת אוטומטית לכל thread:
| register מובנה ב-PTX | המשתנה המקביל ב-CUDA C++ | משמעות |
|---|---|---|
%tid.x |
threadIdx.x |
אינדקס ה-thread בתוך ה-block |
%ntid.x |
blockDim.x |
מספר ה-threads ב-block (מימד ה-CTA) |
%ctaid.x |
blockIdx.x |
אינדקס ה-block בתוך ה-grid |
%nctaid.x |
gridDim.x |
מספר ה-blocks ב-grid |
השם ctaid נגזר מ-CTA (Cooperative Thread Array - מערך threads משתף פעולה), שהוא השם ברמת ה-ISA למה שאנחנו קוראים block, ו-ntid הוא number of threads in the CTA. עם שלושת הערכים האלה, השורה i = blockIdx.x * blockDim.x + threadIdx.x שאנחנו כותבים ב-CUDA C++ הופכת בדיוק להוראת PTX אחת: mad.lo.s32 %r4, %r1, %r2, %r3 (multiply-add על שלמים: %r1 * %r2 + %r3).
מ-CUDA C++ ל-PTX - דוגמה מלאה מקצה לקצה¶
נחבר את הכל. ניקח kernel קטן וממשי ונראה בדיוק לאיזה PTX הוא מתקמפל. ה-kernel מחשב out[i] = in[i] * in[i] + 1.5f - בחרנו את הקבוע 1.5 בכוונה, כדי שנזהה בפלט את אותו 0f3FC00000 שראינו למעלה:
__global__ void square(const float* in, float* out, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x; // global index
if (i < n) // bounds check
out[i] = in[i] * in[i] + 1.5f; // compiles to a single fma
}
מפיקים את ה-PTX ישירות עם דגל של nvcc:
וזהו גוף ה-PTX המתקבל (נוקה וקוצר מעט לשם הבהירות; המספרים המדויקים תלויים בגרסת הcompiler):
.version 8.3
.target sm_90a
.address_size 64
.visible .entry _Z6squarePKfPfi(
.param .u64 _Z6squarePKfPfi_param_0, // in
.param .u64 _Z6squarePKfPfi_param_1, // out
.param .u32 _Z6squarePKfPfi_param_2 // n
)
{
.reg .pred %p<2>;
.reg .f32 %f<3>;
.reg .b32 %r<6>;
.reg .b64 %rd<7>;
ld.param.u64 %rd1, [_Z6squarePKfPfi_param_0]; // in
ld.param.u64 %rd2, [_Z6squarePKfPfi_param_1]; // out
ld.param.u32 %r2, [_Z6squarePKfPfi_param_2]; // n
mov.u32 %r3, %ntid.x; // blockDim.x
mov.u32 %r4, %ctaid.x; // blockIdx.x
mov.u32 %r5, %tid.x; // threadIdx.x
mad.lo.s32 %r1, %r4, %r3, %r5;// i = blockIdx*blockDim + threadIdx
setp.ge.s32 %p1, %r1, %r2; // %p1 = (i >= n)
@%p1 bra $L__BB0_2; // if i>=n, jump to the end
cvta.to.global.u64 %rd3, %rd1; // convert pointer to global space
mul.wide.s32 %rd4, %r1, 4; // byte offset: i * sizeof(float)
add.s64 %rd5, %rd3, %rd4; // address of in[i]
ld.global.f32 %f1, [%rd5]; // f1 = in[i]
fma.rn.f32 %f2, %f1, %f1, 0f3FC00000; // f2 = f1*f1 + 1.5
cvta.to.global.u64 %rd6, %rd2;
add.s64 %rd6, %rd6, %rd4; // address of out[i]
st.global.f32 [%rd6], %f2; // out[i] = f2
$L__BB0_2:
ret;
}
בואו נמפה כל שורת CUDA C++ ל-PTX שלה:
__global__ void square(...)הופך ל-.visible .entry _Z6squarePKfPfi(...). הקידומת.entryמציינת נקודת כניסה של kernel; השם_Z6squarePKfPfiהוא ה-mangled name של C++ (הפיצוח:square(float const*, float*, int)). שימו לב שהארגומנטים הופכים ל-.param, ונטענים לתוך registers בהוראותld.param.int i = blockIdx.x * blockDim.x + threadIdx.xהופך לשלושmovמהאינדקסים המובנים ואזmad.lo.s32אחת - בדיוק כפי שחזינו בסעיף הקודם.if (i < n)הופך ל-setp.ge.s32 %p1, %r1, %r2(חשב את הפרדיקט "i >= n") ואז@%p1 bra $L__BB0_2(אם הפרדיקט אמת, דלג לסוף). שימו לב שהcompiler הפך אתi < nלתנאי ההפוךi >= nכדי לדלג על הגוף - זהו דפוס קומפילציה סטנדרטי.out[i] = in[i] * in[i] + 1.5fהוא לב העניין:ld.global.f32טוענת אתin[i], ה-fma.rn.f32מבצעת אתf1*f1 + 1.5בהוראה אחת (הכפל בעצמו הפך ל-fma עם הקבוע 1.5 כאיבר החיבור), ו-st.global.f32כותבת בחזרה. שתי גישות global מול פעולה אריתמטית אחת - חתימה של חישוב memory-bound, בדיוק כפי שנלמד לכמת בפרק הביצועים.
זו כל היופי של PTX: הוא נמוך מספיק כדי לחשוף כל register וכל גישת זיכרון, אבל גבוה מספיק כדי להישאר קריא ונייד בין דורות.
PTX משורשר - inline PTX ב-CUDA C++¶
לרוב אנחנו נותנים ל-nvcc לייצר את ה-PTX. אבל CUDA C++ מאפשרת גם לשרשר ידנית הוראות PTX לתוך הקוד, בדיוק כפי שאפשר לשרשר assembly של x86 לתוך C. התחביר הוא בלוק asm:
הדוגמה הפשוטה ביותר: קריאת register מיוחד שאין לו משתנה מובנה נוח ב-CUDA C++, למשל %laneid - האינדקס של הthread בתוך ה-warp שלו (0 עד 31):
__device__ __forceinline__ unsigned lane_id() {
unsigned id;
asm volatile("mov.u32 %0, %%laneid;" : "=r"(id));
return id;
}
נפרק את התחביר בקפידה, כי הוא מלא במלכודות:
%0הוא placeholder לאופרנד מספר 0 ברשימת הפלטים/קלטים - הcompiler יחליף אותו ב-register ווירטואלי שהוא בוחר.%%laneid- שימו לב ל-%%הכפול. בתוך בלוקasm, סימן%בודד שמור ל-placeholders כמו%0, ולכן כדי להתייחס ל-register מיוחד של PTX (שנכתב%laneid) צריך לכתוב%%laneid. זו טעות נפוצה מאוד."=r"(id)הוא רשימת הפלטים: האילוץ"=r"מציין register מסוג 32 סיביות לכתיבה (== write-only,r= 32-bit register), והוא נקשר למשתנה C בשםid.volatileאומר לcompiler לא לבצע אופטימיזציות שעלולות להזיז, לשכפל או למחוק את הבלוק. ללאvolatile, אם הcompiler חושב שהפלט לא בשימוש הוא עלול למחוק את ההוראה כליל, או להזיז אותה ולשבש קריאה של register מיוחד שערכו תלוי-הקשר. לכל קריאת register מיוחד השתמשו ב-volatile.
טבלת האילוצים (constraints) הנפוצים לאופרנדים:
| אילוץ | סוג האופרנד |
|---|---|
"r" |
register בן 32 סיביות (int / unsigned) |
"l" |
register בן 64 סיביות (long / מצביע) |
"f" |
register של float |
"d" |
register של double |
"=r" |
פלט 32 סיביות (write-only) |
"+r" |
קלט וגם פלט (read-write) |
דוגמה שנייה, פעולת סיביות: ספירת הביטים הדולקים (population count) בעזרת הוראת popc.b32, עם קלט ופלט:
__device__ __forceinline__ int popcount(unsigned x) {
int r;
asm("popc.b32 %0, %1;" : "=r"(r) : "r"(x)); // r = the number of set bits in x
return r;
}
כאן %0 הוא הפלט r ו-%1 הוא הקלט x; מפריד ה-: השני מפריד בין רשימת הפלטים לרשימת הקלטים. הפעם ויתרנו על volatile כי הפעולה טהורה (הפלט תלוי רק בקלט המפורש), ואפשר לתת לcompiler לבצע אופטימיזציות.
מתי inline PTX הוא הדרך היחידה. ברוב המקרים כתיבת PTX ביד היא נדירה, בדיוק כפי שכתיבת inline assembly של x86 נדירה - שמורה לפינות הרגישות ביותר לביצועים במנועי מסדי נתונים או בליבת מערכת ההפעלה. אבל יש מצב שבו אין ברירה: תכונות חומרה חדשות שהcompiler עדיין לא יודע לפלוט. נכון ל-2025, inline PTX הוא הדרך היחידה לנצל תכונות ספציפיות ל-Hopper כמו ההוראות wgmma (warpgroup matrix-multiply-accumulate - כפל-צבירת מטריצות ברמת warpgroup, ארבעה warps) ו-tma (Tensor Memory Accelerator - מאיץ העברות הזיכרון של ה-Tensor Cores). כך בדיוק בנויים kernels מתקדמים כמו Flash Attention 3 וכמו ה-kernels של Machete (w4a16): הליבה החמה שלהם היא בלוקי inline PTX של wgmma ו-tma, כי אין להם עדיין ביטוי ב-CUDA C++ טהור. זו הסיבה המדויקת שבגללה מתכנת CUDA מתקדם חייב לדעת לקרוא ולכתוב PTX.
כלים לצפייה ב-PTX - nvcc, cuobjdump ו-Godbolt¶
יש שלוש דרכים עיקריות לראות את ה-PTX של קוד, וכל אחת מתאימה למצב אחר:
1. nvcc --ptx - מהמקור. מפיק PTX ישירות מקובץ .cu, בלי לבנות קובץ הרצה מלא. זו הדרך המהירה ביותר לראות מה הcompiler עושה לקוד שכתבתם עכשיו:
2. cuobjdump -ptx - מקובץ מהודר. שולף את ה-PTX שכבר מוטמע בתוך fat binary או ספרייה קיימת. שימושי כשאין לכם את קוד המקור, רק את הקובץ הבינארי (למשל כדי לבדוק אילו ארכיטקטורות ספרייה מסחרית מטמיעה):
cuobjdump -ptx ./square # extracts PTX from the executable file
cuobjdump -ptx libcublas.so # also works on a library
זכרו מהמפה בשיעור 0.4: cuobjdump הוא אחד מ-CUDA Binary Utilities, והוא פותח את ה-fat binary ומראה מה מוטמע בו. שימו לב להבדל: nvcc --ptx מהדר מחדש מהמקור, ואילו cuobjdump -ptx קורא PTX שכבר קיים בבינארי - אם הבינארי הודר בלי PTX (רק SASS), לא יהיה מה לשלוף.
3. Godbolt (Compiler Explorer) - זה לצד זה. האתר godbolt.org יודע להציג CUDA C++, PTX ו-SASS זה לצד זה, ולצבוע איזו שורת מקור הפכה לאילו הוראות. זהו הכלי הטוב ביותר ללמידה ולניסוי מהיר: משנים שורה ב-C++ ורואים מיד איך ה-PTX משתנה, בלי להתקין דבר. נשתמש בו רבות בהמשך הפרק.
טיפ מעשי: כדי לראות גם כמה registers פיזיים ה-kernel ישתמש בסופו של דבר, הוסיפו את הדגל -Xptxas -v בזמן הcompilation. הוא מדפיס שורת Used N registers - אבל שימו לב שזהו כבר מספר ה-registers של ה-SASS (אחרי הקצאה פיזית על ידי ptxas), לא של ה-PTX הווירטואלי. את ההבחנה הזו נעמיק בשיעור 4.2 על SASS.
סיכום¶
- PTX (Parallel Thread eXecution) הוא בו-זמנית "מכונה וירטואלית" ו-"ISA": סט הוראות כנגד מודל מכונה מופשט, שמאפשר לאותה תוכנית לרוץ באותה סמנטיקה על devices רבים, כולל כאלה שעוד לא קיימים.
- שלא כמו x86 או aarch64, PTX אינו קוד מכונה אמיתי אלא ייצוג ביניים (IR) כמו LLVM-IR, שעובר JIT ל-SASS על ידי ה-hardware driver בזמן ריצה.
- PTX הוא ה"מותן הצר" בין עולם הcompilers (CUDA C++, Triton, Julia...) לעולם הdevices (Volta, Ampere, Hopper, Blackwell): כל צד מכיר רק את PTX במקום כל השפות את כל הdevices.
- מודל המכונה הוא SIMT מנוסח במדויק: יחידת הוראות אחת להרבה מעבדים, כל ה-threads מריצים אותן הוראות, מתאמים דרך shared memory, ומשיגים תוצאות שונות דרך registers פרטיים - הקוד אחיד, הנתונים פרטיים.
- registers של PTX הם וירטואליים ובני 32 סיביות (
.f32,.b32,.u32), ומוקצים דינמית ל-warps מתוך קובץ האוגרים של ה-SM;.reg .f32 %f<7>מכריז על שבעה כאלה. fma.rn.f32היא fused multiply-add בדיוק מלא, והסיומתrnהיא מצב עיגול IEEE 754 "round to nearest even"; הקבוע0f3FC00000הוא 1.5 בייצוג ביטים של float.- ה-registers המובנים
%tid.x,%ntid.x,%ctaid.x,%nctaid.xמקבילים ל-threadIdx.x,blockDim.x,blockIdx.x,gridDim.x, ומחישוב האינדקסi = blockIdx*blockDim + threadIdxנולדת הוראתmad.lo.s32אחת. - inline PTX נכתב בבלוק
asm volatile("...":פלטים:קלטים);%%מציין register מיוחד,"=r"/"r"הם אילוצי אופרנד, ו-volatileמונע אופטימיזציות שימחקו או יזיזו את הבלוק. - נכון ל-2025, inline PTX הוא הדרך היחידה לנצל תכונות Hopper כמו
wgmmaו-tma, כפי שנעשה ב-Flash Attention 3 וב-kernels של Machete - זו הסיבה המעשית לדעת PTX. - רואים PTX בשלוש דרכים:
nvcc --ptxמהמקור,cuobjdump -ptxמבינארי מהודר, ו-Godbolt להצגה זה-לצד-זה של C++/PTX/SASS.