4.4 כלי בינארי לניתוח artifacts הרצאה
לאורך פרק 4 פירקנו את שרשרת הcompilation לגורמים: בשיעור 4.1 הכרנו את PTX כשפת ביניים וירטואלית, בשיעור 4.2 ירדנו ל-SASS - האסמבלי המקורי של החומרה, ובשיעור 4.3 ראינו איך nvcc (compiler-מנהל של CUDA) עוטף את הכל לתוך fat binary (קובץ שמן) יחיד שנושא PTX ו/או SASS. עד עכשיו ייצרנו artifacts (תוצרי compilation). השיעור הזה סוגר את הפרק מהכיוון ההפוך: כיצד לפתוח artifact קיים ולקרוא מה בדיוק יש בתוכו. נכיר את CUDA Binary Utilities (כלי-העזר לקבצים בינאריים של CUDA) - ארגז הכלים לבדיקה סטטית של קבצי הרצה, קבצי אובייקט וספריות. נלמד מתי לשלוף PTX, מתי לרדת ל-SASS, איך לבנות control-flow graph (גרף זרימת בקרה), איך לפענח שמות סימבולים מעוותים, ואיך לגזום קוד device מיותר כדי לכווץ בינארי. אלה כלי הבדיקה הסטטיים שילוו אתכם בכל שאר הקורס, ובמיוחד בפרקים 8 ו-9 כשנתחיל לפרופל ולייעל kernels (קרנלים - פונקציות GPU) אמיתיים. ההבחנה המרכזית: כלים אלה קוראים את הבינארי כפי שהוא נח על הדיסק, בלי להריץ אותו - ניתוח סטטי, בניגוד ל-profiling (פרופיילינג) הדינמי של פרק 9.
למה בכלל צריך ארגז כלים ייעודי - static inspection¶
הבינארי שיצא מ-nvcc הוא קובץ ELF תואם host ABI, כמו שראינו בשיעור 0.4. אפשר להצביע עליו עם כל כלי ELF גנרי - file, readelf, objdump, nm. הבעיה: קוד ה-device אינו יושב ב-ELF כמו קוד host רגיל. הוא ארוז בתוך מקטע (section) בשם .nv_fatbin כמבנה נתונים אטום שנקרא fatbin - מיכל שבתוכו יושבים אחד או יותר קבצי cubin (קוד SASS ארוז כ-ELF פנימי משלו) וגם בלוקי PTX כטקסט דחוס. כלי ELF גנרי רואה את המקטע הזה כערימת בייטים חסרת פשר. הוא לא יודע לפרש את מבנה ה-fatbin, לא יודע לחלץ ממנו את ה-cubin, ובוודאי לא יודע לפרק את ה-SASS שבפנים לאסמבלי קריא.
כאן נכנס ארגז הכלים הייעודי. CUDA Binary Utilities הם אוסף כלים שמגיע עם ה-CUDA Toolkit, וכולם חולקים ידע אחד: הם מבינים את פורמט ה-fatbin ואת פורמט ה-cubin. חלוקת העבודה בין הכלים היא הדבר החשוב להפנים כבר עכשיו:
Question Tool
------------------------------------------ -----------------------
"Is this even a CUDA file? Which section?" readelf (generic, host)
"Which device architectures are embedded?" cuobjdump --list-elf / --list-ptx
"Show me the PTX / SASS inside the binary" cuobjdump -ptx / -sass
"Extract the cubin to a separate file" cuobjdump -xelf all
"Disassemble the cubin in depth, with CFG and sources" nvdisasm
"What's the real name of symbol _Z6vecAdd...?" cu++filt
"Drop device architectures I don't need" nvprune
שני צירים מארגנים את הכלים. הציר הראשון הוא הרזולוציה: readelf נותן מבט host גס על הקובץ; cuobjdump נכנס פנימה ומראה מה יש ב-fatbin; nvdisasm צולל לתוך cubin בודד ומנתח את ה-SASS שלו לעומק. הציר השני הוא הכיוון: רוב הכלים קוראים (readelf, cuobjdump, nvdisasm, cu++filt), ורק אחד כותב - nvprune, שמייצר בינארי חדש רזה יותר.
הבינארי לדוגמה - the running example¶
כדי שלא נדבר בהפשטה, נבנה artifact אחד ונחזור אליו לאורך כל השיעור. זהו אותו vecAdd שהכרנו, עם מאקרו בדיקת השגיאות CUDA_CHECK שנשתמש בו לאורך הקורס:
// vecadd.cu
#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)
__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
c[i] = a[i] + b[i];
}
int main() {
const int n = 1 << 20;
size_t bytes = n * sizeof(float);
float *da, *db, *dc;
CUDA_CHECK(cudaMalloc(&da, bytes));
CUDA_CHECK(cudaMalloc(&db, bytes));
CUDA_CHECK(cudaMalloc(&dc, bytes));
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
vecAdd<<<numBlocks, blockSize>>>(da, db, dc, n);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaFree(da));
CUDA_CHECK(cudaFree(db));
CUDA_CHECK(cudaFree(dc));
return 0;
}
נהדר אותו ל-H100 (Hopper, compute capability 9.0) כך שיטמיע גם SASS של sm_90 וגם PTX של compute_90 - זה ייתן לנו את שני ייצוגי ה-device באותו קובץ, מה שיהפוך את הבדיקה למעניינת. נוסיף -lineinfo כדי שנוכל למפות SASS בחזרה לשורות המקור בהמשך:
nvcc -lineinfo \
-gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_90,code=compute_90 \
vecadd.cu -o vecadd
השורה הראשונה של -gencode מטמיעה SASS ל-Hopper; השנייה מטמיעה PTX של compute_90 לתאימות קדימה. הדגל -lineinfo מוסיף מידע דיבאג קליל (מיפוי הוראה-לשורה) בלי לכבות אופטימיזציות - בניגוד ל--G, שמכבה אופטימיזציות ומייצר SASS שונה לגמרי שלא מייצג את קוד הייצור. זכרו את ההבחנה הזו: לניתוח ביצועים תמיד -lineinfo, לעולם לא -G.
המבט הגנרי - readelf¶
לפני שנשלוף כלי מיוחד, שווה לזכור שהבינארי הוא ELF רגיל. readelf הוא כלי binutils סטנדרטי של לינוקס, לא חלק מ-CUDA, והוא נותן את המבט הראשון:
[23] .nvFatBinSegment PROGBITS 0000000000015c80 00015c80
[24] .nv_fatbin PROGBITS 0000000000015ca0 00015ca0
שני המקטעים מספרים את הסיפור: .nv_fatbin הוא ה-fatbin עצמו - ערימת הבייטים שנושאת את ה-cubin ואת ה-PTX. .nvFatBinSegment הוא מבנה קטן שקוד האתחול של ה-Runtime קורא כדי לרשום את ה-fatbin מול הדרייבר בזמן טעינת התוכנית. readelf מאשר לנו שיש כאן תוכן device ומראה את גודלו וגבולותיו, אבל כאן נגמרת יכולתו. הוא לא יודע לפרש את מה שבפנים - עבור זה נעבור ל-cuobjdump. שווה גם לזכור ש-readelf -h יאמת שזה בכלל ELF, ו-nm vecadd | grep vecAdd יראה את סימבול ה-host שמפעיל את ה-kernel (הפ'ונקציה שקוראת ל-cudaLaunchKernel מאחורי התחביר <<<>>>).
הכלי המרכזי - cuobjdump¶
cuobjdump הוא נקודת הכניסה לכל בדיקה. הוא פותח בינארי host שלם (או קובץ cubin ישיר) ומראה מה יש בו. נתחיל תמיד בשאלה "מה בכלל בפנים", עם דגלי הרשימה:
cuobjdump --list-elf vecadd # which cubins (SASS) are embedded
cuobjdump --list-ptx vecadd # which PTX blocks are embedded
הפלט מאשר בדיוק את מה שביקשנו מ-nvcc: cubin אחד ל-sm_90 ובלוק PTX אחד ל-compute_90. אילו היינו מהדרים עם כמה -gencode, היינו רואים כאן שורה לכל ארכיטקטורה - וזו הדרך המהירה לענות על השאלה "האם הבינארי הזה בכלל תומך ב-GPU שלי". עכשיו נשלוף תוכן. ארבעת הדגלים המרכזיים:
| דגל | מה הוא מציג |
|---|---|
-ptx |
את קוד ה-PTX (שפת הביניים הווירטואלית) |
-sass |
את קוד ה-SASS (אסמבלי מקורי, מפורק להוראות) |
-elf |
את מבנה ה-ELF של ה-cubin (מקטעים, סימבולים, headers) |
-res-usage |
שימוש משאבים לכל kernel: registers, shared, stack |
נציץ ב-PTX. שימו לב שהפעולה c[i] = a[i] + b[i] הפכה ל-add.f32, ושהקוד מסתיים בקוד יציאה:
//
// Generated by NVIDIA NVVM Compiler
// Compiler Build ID: CL-...
// Cuda compilation tools, release 12.x
// Based on NVVM 7
//
.visible .entry _Z6vecAddPKfS0_Pfi(
.param .u64 _Z6vecAddPKfS0_Pfi_param_0,
.param .u64 _Z6vecAddPKfS0_Pfi_param_1,
.param .u64 _Z6vecAddPKfS0_Pfi_param_2,
.param .u32 _Z6vecAddPKfS0_Pfi_param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<4>;
.reg .b32 %r<6>;
.reg .b64 %rd<11>;
mov.u32 %r1, %ctaid.x;
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r1, %r2, %r3;
setp.ge.s32 %p1, %r4, %r5;
@%p1 bra $L__BB0_2;
...
add.f32 %f3, %f2, %f1;
st.global.f32 [%rd10], %f3;
$L__BB0_2:
ret;
}
עכשיו ה-SASS - האסמבלי המקורי שירוץ בפועל על ה-SM (מעבד רב-זרמי) של ה-H100:
code for sm_90
Function : _Z6vecAddPKfS0_Pfi
.headerflags @"EF_CUDA_SM90 EF_CUDA_PTX_SM(EF_CUDA_SM90)"
/*0000*/ IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
/*0010*/ S2R R4, SR_CTAID.X ;
/*0020*/ IMAD.MOV.U32 R5, RZ, RZ, c[0x0][0x0] ;
/*0030*/ S2R R3, SR_TID.X ;
/*0040*/ IMAD R4, R4, R5, R3 ;
/*0050*/ ISETP.GE.AND P0, PT, R4, c[0x0][0x228], PT ;
/*0060*/ @P0 EXIT ;
/*0070*/ MOV R7, 0x4 ;
/*0080*/ IMAD.WIDE R2, R4, R7, c[0x0][0x210] ;
/*0090*/ IMAD.WIDE R8, R4, R7, c[0x0][0x218] ;
/*00a0*/ LDG.E R2, [R2.64] ;
/*00b0*/ LDG.E R0, [R8.64] ;
/*00c0*/ IMAD.WIDE R6, R4, R7, c[0x0][0x220] ;
/*00d0*/ FADD R9, R2, R0 ;
/*00e0*/ STG.E [R6.64], R9 ;
/*00f0*/ EXIT ;
/*0100*/ BRA `(.L_x_0) ;
שימו לב לשלושה דברים. ראשית, ה-FADD היחיד - זו פעולת החיבור שלנו, בדיוק כמו שה-add.f32 ב-PTX ניבא. שנית, S2R R4, SR_CTAID.X ו-SR_TID.X הם חישוב ה-blockIdx.x וה-threadIdx.x, וה-IMAD שאחריהם הוא בדיוק blockIdx.x * blockDim.x + threadIdx.x. שלישית, ISETP.GE ואחריו @P0 EXIT - זה ה-if (i < n) שלנו: אם האינדקס גדול-או-שווה ל-n, ה-thread יוצא מיד. ה-SASS הזה הוא התרגום הישיר של שש שורות ה-CUDA C++ שלנו, וזו בדיוק היכולת שנשתמש בה בפרק 8 כדי לראות אם הcompiler עשה מה שציפינו.
הדגל -res-usage הוא כנראה הדגל המעשי ביותר של cuobjdump לצורכי ביצועים, כי הוא עונה על שאלת ה-occupancy (תפוסה) בלי להריץ כלום:
Function _Z6vecAddPKfS0_Pfi:
REG:18 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:544 TEXTURE:0 SURFACE:0 SAMPLER:0
REG:18 אומר שכל thread צורך 18 registers מקובץ האוגרים של ה-SM. זה מספר קריטי: ל-H100 יש 65,536 registers ל-SM, כך שתקרת ה-threads מבחינת registers היא 65536 / 18 ≈ 3640 threads - הרבה מעל תקרת ה-2048 threads הפיזית של ה-SM, אז ה-kernel הקטן הזה אינו register-bound ויכול להגיע ל-occupancy מלא. את החישוב הזה נעשה שוב ושוב בפרק 8; כאן חשוב רק לראות ש--res-usage נותן את הקלט לו בלי הרצה.
עוד יכולת של cuobjdump היא חילוץ: הדגל -xelf all כותב כל cubin מוטמע לקובץ נפרד על הדיסק, ו--xptx all עושה אותו דבר ל-PTX. זה הגשר לכלי הבא:
הצולל לעומק - nvdisasm¶
cuobjdump הראה לנו SASS, אז למה צריך עוד כלי? כי nvdisasm עושה דברים ש-cuobjdump לא. חלוקת העבודה: cuobjdump הוא הכלי הרחב שפותח בינארי host שלם ומראה מה בפנים; nvdisasm הוא המפרק הממוקד שעובד על cubin בודד ומנתח אותו לעומק. הוא לא יודע לפתוח בינארי host - צריך קודם לחלץ את ה-cubin (עם cuobjdump -xelf all, כמו שעשינו), ורק אז להאכיל אותו:
הדגל -c מפרק את ה-SASS. עד כאן דומה ל-cuobjdump -sass, אבל שלוש היכולות הבאות ייחודיות ל-nvdisasm והן הסיבה שהוא קיים.
היכולת הראשונה - גרף זרימת בקרה - control-flow graph. הדגל -cfg פולט את גרף הזרימה של ה-kernel בפורמט DOT (שפת התיאור של Graphviz), כך שאפשר להפוך אותו לתמונה. הדגל -bbcfg עושה זאת ברזולוציה של basic blocks (בלוקים בסיסיים - רצפי הוראות בלי הסתעפות):
nvdisasm -cfg vecadd.1.sm_90.cubin > vecadd.dot
dot -Tpng vecadd.dot -o vecadd.png # requires graphviz
digraph "_Z6vecAddPKfS0_Pfi" {
rankdir = TB;
node [shape = record];
blk_0 [label = "{BB0 | S2R ... | ISETP.GE ... | @P0 EXIT}"];
blk_1 [label = "{BB1 | LDG ... | FADD ... | STG ... | EXIT}"];
blk_0 -> blk_1;
blk_0 -> blk_exit;
}
הגרף מראה בדיוק את מבנה ה-if: בלוק אחד מחשב את האינדקס ובודק את התנאי, ומשם יש שתי קשתות - אחת ל-EXIT מוקדם (כש-i >= n) ואחת לבלוק שעושה את הטעינה, החיבור והכתיבה. עבור kernel עם לולאות ו-branches מסובכים, ה-CFG הזה הוא הדרך היחידה השפויה להבין את זרימת הבקרה, ובפרק 8 נשתמש בו כדי לאתר warp divergence (התפצלות ה-warp).
היכולת השנייה - מיפוי למקור - source correlation. אם הידרנו עם -lineinfo (כמו שעשינו), הדגל -g מכניס לצד כל הוראת SASS את שורת המקור שממנה היא נגזרה:
//## File "vecadd.cu", line 17
/*0040*/ IMAD R4, R4, R5, R3 ;
//## File "vecadd.cu", line 18
/*0050*/ ISETP.GE.AND P0, PT, R4, c[0x0][0x228], PT ;
/*0060*/ @P0 EXIT ;
//## File "vecadd.cu", line 19
/*....*/ FADD R9, R2, R0 ;
עכשיו כל הוראת מכונה מקושרת לשורת ה-CUDA C++ שכתבנו: שורה 17 היא חישוב האינדקס, שורה 18 היא ה-if, שורה 19 היא החיבור. זו בדיוק החוליה שמאפשרת ל-Nsight Compute בפרק 9 להראות לכם "השורה הזו בקוד שלך אחראית ל-40% מה-stalls".
היכולת השלישית - ניתוח סטטי מדויק. nvdisasm יודע גם -poff (להצמיד offset לכל הוראה), -hex (להראות את הקידוד הבינארי לצד האסמבלי), ו--plr (register liveness - אילו registers חיים בכל נקודה). אלה כלים לניתוח מיקרו-ארכיטקטורלי עמוק שנשתמש בהם רק כשנרד לכוונון היד-על-הדופק בפרק 8.
פענוח שמות - cu++filt¶
שמתם לב שכל הכלים קוראים ל-kernel שלנו _Z6vecAddPKfS0_Pfi ולא vecAdd? זה name mangling (עיוות שמות) של C++: הcompiler מקודד את שם הפונקציה ואת טיפוסי הפרמטרים לתוך סימבול יחיד, כדי לתמוך ב-overloading. cu++filt הוא הגרסה של CUDA ל-c++filt הקלאסי - הוא מפענח את השם בחזרה:
הנה פענוח הקידוד, כי כדאי להבין אותו ולא רק להריץ כלי: _Z הוא הקידומת של סימבול C++ מעוות, 6vecAdd הוא השם באורך 6 תווים, ואז הפרמטרים - PKf הוא pointer-to-const-float (const float), S0_ הוא הפניה חוזרת (substitution) לאותו טיפוס בדיוק כדי לחסוך מקום, Pf הוא pointer-to-float (float), ו-i הוא int. שימו לב שהטיפוס S0_ הוא בדיוק המנגנון שמצמצם את const float* השני. השימוש המעשי: כשאתם קוראים SASS או פלט profiler של kernel עם template (תבנית) מסובך, השמות המעוותים בלתי-קריאים לחלוטין, ו-cu++filt הוא מה שהופך אותם חזרה למשהו אנושי. אפשר גם להזרים אליו פלט שלם דרך pipe והוא יפענח כל סימבול שהוא מזהה:
גזום קוד device - nvprune¶
כל הכלים עד כה קראו. nvprune הוא היחיד שכותב: הוא לוקח קובץ אובייקט host (.o) או ספרייה סטטית (.a) שמכילה קוד device לכמה ארכיטקטורות, וזורק את הארכיטקטורות שאתם לא צריכים - ומקטין את הקובץ. הוא לא עובד על קובץ הרצה מקושר סופי, אלא על אובייקטים וספריות שעדיין לא קושרו. השימוש הטיפוסי הוא כיווץ ספרייה ענקית כמו cuBLAS הסטטית, שמגיעה מ-NVIDIA עם SASS לכל דור GPU קיים:
# the original library carries SASS for sm_70, sm_80, sm_90, sm_100, and more
nvprune --arch sm_90 libcublas_static.a -o libcublas_sm90.a
התוצאה היא ספרייה שנושאת רק קוד ל-Hopper. אם אתם בונים קונטיינר שירוץ אך ורק על H100, אין טעם לשאת SASS ל-Turing ול-Ampere; nvprune יכול לחתוך את גודל הבינארי הסופי בעשרות אחוזים. אפשר גם לגזום לפי -gencode מלא כדי לשמור על PTX לתאימות קדימה:
זהירות: nvprune הוא אופטימיזציית גודל, לא אופטימיזציית ביצועים - הקוד שנשאר רץ בדיוק אותו דבר. וכמובן, אם תגזמו ארכיטקטורה שבסוף כן תריצו עליה, התוכנית תיכשל בטעינה או תיפול ל-JIT איטי מ-PTX (אם נשאר PTX). גזמו רק מה שאתם בטוחים שלא צריך.
תהליך העבודה המלא - the end-to-end workflow¶
עכשיו כשמכירים את כל הכלים, נחבר אותם לתהליך אחד. השאלה המנחה תמיד קובעת את הכלי. הנה עץ ההחלטה שכדאי לזכור:
I have a suspicious binary. What do I do?
|
+-- "Is this even CUDA?" -> readelf -S | grep nv_fatbin
|
+-- "Which architectures are inside?" -> cuobjdump --list-elf / --list-ptx
|
+-- "How many registers / shared?" -> cuobjdump -res-usage
|
+-- "Show me PTX" -> cuobjdump -ptx
|
+-- "Show me SASS" -> cuobjdump -sass (fast, top-level view)
|
+-- "Deep SASS analysis:"
| first cuobjdump -xelf all (extracts cubin)
| then nvdisasm -cfg (flow graph)
| nvdisasm -g (source mapping, requires -lineinfo)
|
+-- "The names are unreadable" -> ... | cu++filt
|
+-- "The binary is too large" -> nvprune --arch sm_90
וזה החיבור לשאר הקורס. כל הכלים בשיעור הזה הם סטטיים - הם קוראים את הבינארי כפי שהוא נח, בלי להריץ אותו. הם עונים על "מה יש כאן ומה נוצר", אבל לא על "כמה מהר זה רץ ולמה". על השאלה השנייה עונים כלי ה-profiling הדינמיים של פרק 9, שרצים תוך כדי הרצת ה-kernel: Nsight Systems לתמונה מערכתית של כל ה-timeline, ו-Nsight Compute לניתוח מיקרו-ארכיטקטורלי של kernel בודד. שני העולמות משלימים זה את זה בדיוק: כשה-profiler הדינמי מצביע על שורת SASS איטית, אתם חוזרים לכלים הסטטיים של השיעור הזה - nvdisasm -g כדי לראות איזו שורת מקור זו, ו-nvdisasm -cfg כדי להבין את מבנה הבקרה סביבה. ה-source correlation שראינו הוא בדיוק החוליה ש-Nsight Compute מנצל כדי לתלות מדידות דינמיות על הוראות סטטיות. ארגז הכלים הזה הוא לכן לא נספח לפרק 4 - הוא הבסיס שעליו נשענת כל מתודולוגיית האופטימיזציה של פרקים 8 ו-9.
מלכודות נפוצות - common gotchas¶
כמה טעויות חוזרות ששווה לסמן מראש:
-Gמזהם את ה-SASS. אם תהדרו עם-G(debug מלא) ותסתכלו על ה-SASS, תראו קוד לא-מאופטם עם spills מיותרים שאינו מייצג את הייצור. לניתוח ביצועים תמיד-lineinfo, שמוסיף מיפוי-למקור בלי לכבות אופטימיזציות.nvdisasmלא פותח בינארי host. הוא עובד רק על cubin. אם תריצוnvdisasm vecaddעל קובץ ההרצה הוא ייכשל; צריך קודםcuobjdump -xelf allכדי לחלץ את ה-cubin.- בינארי עם PTX בלבד אין לו SASS להראות. אם הידרתם
code=compute_90בלבד (PTX ולא SASS), אזcuobjdump -sassיחזיר ריק - ה-SASS ייוצר רק בזמן ריצה ב-JIT. תראו SASS סטטי רק אם הטמעתםcode=sm_90. - קידוד ה-SASS לא מתועד רשמית. כפי שראינו בשיעור 4.2, NVIDIA לא מפרסמת את הסמנטיקה המדויקת של כל הוראת SASS ולא את הקידוד הבינארי שלה. הכלים האלה הם הדרך הרשמית היחידה לקרוא SASS; אין assembler רשמי שכותב SASS מהיד.
--list-elfמול--list-ptx. שני דגלים שונים לשני ייצוגים שונים. אם--list-elfריק אבל--list-ptxמלא, הבינארי נושא רק PTX ויעבור JIT בזמן ריצה - סימן לתאימות-קדימה, לא לתקלה.
סיכום¶
- ה-CUDA Binary Utilities הם ארגז כלי הבדיקה הסטטיים של artifacts - הם קוראים את הבינארי כפי שהוא נח על הדיסק, בלי להריץ אותו, בניגוד ל-profiling הדינמי של פרק 9.
readelfנותן מבט host גנרי ומאשר קיום.nv_fatbin, אבל לא יודע לפרש את התוכן הפנימי - שם צריך את הכלים המודעים-CUDA.cuobjdumpהוא נקודת הכניסה:--list-elf/--list-ptxלרשימת ארכיטקטורות,-ptx/-sass/-elfלהצגת התוכן,-res-usageלשימוש registers ו-shared, ו--xelf allלחילוץ ה-cubin.nvdisasmהוא המפרק הממוקד שעובד על cubin בודד ומוסיף מה ש-cuobjdump לא: גרף זרימת בקרה (-cfg/-bbcfg), מיפוי הוראות לשורות מקור (-g, דורש-lineinfo), וניתוח סטטי עמוק.cu++filtמפענח שמות C++ מעוותים -_Z6vecAddPKfS0_Pfiהופך ל-vecAdd(float const*, float const*, float*, int); אפשר להזרים אליו פלט שלם דרך pipe.nvpruneהוא הכלי היחיד שכותב: הוא גוזם קוד device לארכיטקטורות מיותרות מ-.oומ-.aומכווץ בינארי - אופטימיזציית גודל, לא ביצועים.- חלוקת העבודה:
cuobjdumpרחב (פותח בינארי host שלם, "מה בפנים"),nvdisasmעמוק (מפרק cubin בודד, "מה זה עושה"); הם מרכיבים - חלץ עם cuobjdump, נתח עם nvdisasm. - הכלל: תמיד
-lineinfoולעולם לא-Gלניתוח ביצועים;nvdisasmלא פותח בינארי host; בינארי עם PTX בלבד לא יראה SASS סטטי. - הכלים הסטטיים האלה הם הבסיס למתודולוגיית האופטימיזציה של פרקים 8 ו-9: כש-Nsight Compute מצביע על שורת SASS איטית, ה-source correlation של
nvdisasm -gהוא מה שמחזיר אתכם לקוד המקור.