9.1 פרופיילינג Nsight Systems ו Nsight Compute פתרון
פתרון - Nsight Systems ו-Nsight Compute¶
נעבור תרגיל אחרי תרגיל עם הקוד, הפקודות והפלט הצפוי. כל המספרים כאן הם מ-H100 (SXM) עם CUDA 12.4 ו-Nsight 2024.x; הזמנים המוחלטים, רוחבי הפס וחלק ממספרי הדוגמה בפלט של nsys/ncu משתנים בין הרצות, דרייברים ודורות, והם ערכים אופייניים - העיקר הוא היחסים, הצורה של הדוחות, וספירות ה-sectors וההתנגשויות שהן דטרמיניסטיות. מי שעובד על T4 ב-Colab יראה מספרים מוחלטים אחרים לגמרי, אבל בדיוק אותה תמונה ואותם קטעי דוח.
פתרון תרגיל 1 - nsys על אפליקציה overhead-bound¶
// overhead.cu
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#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 tiny_add(float* x, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) x[i] += 1.0f;
}
int main() {
const int N = 1024;
float* d;
CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));
CUDA_CHECK(cudaMemset(d, 0, N * sizeof(float)));
for (int k = 0; k < 10000; ++k)
tiny_add<<<(N + 255) / 256, 256>>>(d, N); // 10,000 tiny dispatches
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaFree(d));
return 0;
}
$ nvcc -arch=sm_90 -O3 overhead.cu -o overhead
$ nsys profile -o rep_overhead ./overhead
Generating '/tmp/rep_overhead.nsys-rep'
$ nsys stats rep_overhead.nsys-rep
** CUDA API Summary (cuda_api_sum):
Time(%) Total Time(ns) Num Calls Avg(ns) Name
------- -------------- --------- --------- ------------------
97.4 104,882,001 10,000 10,488.2 cudaLaunchKernel
2.3 2,510,443 1 2,510,443 cudaDeviceSynchronize
0.2 215,880 1 215,880 cudaMalloc
** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):
Time(%) Total Time(ns) Instances Avg(ns) Name
------- -------------- --------- ------- ---------------
100.0 18,400,120 10,000 1,840 tiny_add(float*, int)
ה-CPU בילה ~105 מילישנייה בקריאות cudaLaunchKernel בעוד שה-GPU עבד בסך הכל ~18 מילישנייה. יחס של כמעט פי 6 לטובת הoverhead - האפליקציה overhead-bound מובהקת.
למה זה עבד: כל cudaLaunchKernel עולה ~10 מיקרו-שנייה של overhead CPU (שיעור 7.4), ללא קשר לכמה עבודה ה-kernel עושה. 10,000 dispatches כפול 10,488 ננו-שנייה בממוצע = ~105 מילישנייה, שכולן overhead טהורה. ה-GPU בטל את רוב הזמן ומחכה לקבל עבודה - בציר הזמן של ה-GUI זה נראה כרצף פערים בין kernels זעירים. איך להכליל: כשזמן ה-CUDA API עולה על זמן ה-GPU, אתם overhead-bound, וזה בדיוק המקרה שבו CUDA Graphs (שיעור 5.5) עוזר - הוא אוסף את 10,000 הdispatches לdispatch host יחיד ומצניח את הoverhead. תמיד השוו את שתי הטבלאות (API מול Kernel) לפני שמתחילים לאפטמז את ה-kernel עצמו.
פתרון תרגיל 2 - מציאת חוסר חפיפה בין copy ל-compute¶
הגרסה הגרועה (הכל בטור על ה-default stream) מול הגרסה עם streams ו-pinned memory:
// pipeline.cu (essentials - the good version)
const int CHUNKS = 8, CN = 1 << 22; // 8 chunks of 4M floats
float *h; CUDA_CHECK(cudaHostAlloc(&h, CHUNKS*CN*sizeof(float), cudaHostAllocDefault));
float *d; CUDA_CHECK(cudaMalloc(&d, CHUNKS*CN*sizeof(float)));
cudaStream_t st[CHUNKS];
for (int c = 0; c < CHUNKS; ++c) CUDA_CHECK(cudaStreamCreate(&st[c]));
for (int c = 0; c < CHUNKS; ++c) {
size_t off = (size_t)c * CN;
CUDA_CHECK(cudaMemcpyAsync(d+off, h+off, CN*sizeof(float),
cudaMemcpyHostToDevice, st[c]));
heavy_kernel<<<CN/256, 256, 0, st[c]>>>(d+off, CN); // overlaps with the copy of another chunk
CUDA_CHECK(cudaMemcpyAsync(h+off, d+off, CN*sizeof(float),
cudaMemcpyDeviceToHost, st[c]));
}
CUDA_CHECK(cudaDeviceSynchronize());
$ nsys stats rep_serial.nsys-rep # the bad version (default stream)
** CUDA GPU MemOps + Kernels timeline (shortened):
HtoD chunk0 [ 0.0 ms - 1.4 ms ]
krnl chunk0 [ 1.4 ms - 3.2 ms ]
DtoH chunk0 [ 3.2 ms - 4.6 ms ]
HtoD chunk1 [ 4.6 ms - 6.0 ms ] <- starts only after chunk 0 fully finished
...
Total: 36.8 ms (all serial, zero overlap)
$ nsys stats rep_overlap.nsys-rep # the good version (streams + pinned)
HtoD chunk0 [ 0.0 ms - 1.4 ms ]
krnl chunk0 [ 1.4 ms - 3.2 ms ]
HtoD chunk1 [ 1.4 ms - 2.8 ms ] <- overlaps with the kernel of chunk 0!
krnl chunk1 [ 3.2 ms - 5.0 ms ]
HtoD chunk2 [ 2.8 ms - 4.2 ms ] <- overlaps again
...
Total: 17.1 ms (full overlap - about 2.1x faster)
למה זה עבד: בגרסה הגרועה כל פעולה מחכה לקודמת - העתקה, חישוב, העתקה, ורק אז האצווה הבאה - כי ה-default stream מסנכרן הכל. בגרסה הטובה, שלושת התנאים יחד (cudaHostAlloc ל-pinned memory, cudaMemcpyAsync, ו-stream נפרד לכל אצווה) מאפשרים למנוע ה-DMA להעתיק אצווה אחת בזמן שה-SMs מחשבים אצווה אחרת. חותמות הזמן ב-nsys חופפות - זה האישור החזותי. איך להכליל: חפיפת copy-compute דורשת תמיד את שלושת הדברים; cudaMemcpy רגיל על זיכרון pageable לעולם לא יחפוף. ב-nsys מזהים חוסר חפיפה כשטווחי הזמן של ה-MemOps וה-Kernels לא חופפים - וזו בעיה מערכתית שרק nsys יראה, לא ncu.
פתרון תרגיל 3 - ncu --set full וקריאת הקטעים¶
נשתמש ב-SAXPY כדוגמה הקנונית לmemory-bound:
// saxpy.cu
__global__ void saxpy(float a, const float* __restrict__ x,
float* __restrict__ y, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i]; // 2 FLOP, 12 byte -> AI = 1/6
}
$ nvcc -arch=sm_90 -O3 -lineinfo saxpy.cu -o saxpy
$ ncu --set full -k saxpy --launch-count 1 -o rep_saxpy ./saxpy
$ ncu -i rep_saxpy.ncu-rep --page details
saxpy(float, const float *, float *, int)
Section: GPU Speed Of Light Throughput
----------------------------------------------------------------------
Compute (SM) Throughput [%] 3.10
Memory Throughput [%] 91.24 <- memory close to peak
DRAM Throughput [%] 91.24
----------------------------------------------------------------------
-> Kernel is memory-bound. Compute is idle; DRAM near peak.
Section: GPU Speed Of Light Roofline Chart
Arithmetic Intensity [FLOP/byte] 0.167 (= 1/6)
Achieved Performance [GFLOP/s] ~1100
-> point on the diagonal roof, deep to the left of the ridge (295)
Section: Occupancy
Theoretical Occupancy [%] 100.00
Achieved Occupancy [%] 89.7
Registers Per Thread [reg] 16
Section: Warp State (Stall Reasons)
Stall Long Scoreboard [cycles/inst] 8.42 <- dominant
Stall Wait [cycles/inst] 0.31
למה זה עבד: SAXPY מבצע 2 FLOP על כל 12 בייט (טעינת x, טעינת y, כתיבת y) - arithmetic intensity של 1/6, שהוא 0.167 FLOP/byte, הרבה מתחת לנקודת הרכס 295 של ה-H100 (שיעור 7.2). לכן קטע ה-Speed of Light מראה ניצולת זיכרון ~91% מול ניצולת חישוב ~3%, וה-roofline מצייר את הנקודה על הגג האלכסוני - memory-bound חד-משמעי. סיבת העצירה הדומיננטית היא long scoreboard (8.42 מחזורים להוראה) - המתנה ל-LDG שחוזר מ-HBM, בדיוק כפי ש-8.3 חוזה לקוד memory-bound. ה-occupancy גבוה (~90%) כי SAXPY משתמש רק ב-16 registers לthread. איך להכליל: --set full נותן בהרצה אחת את כל התמונה. הסדר לקריאה: קודם Speed of Light/roofline (חישוב או זיכרון?), אז occupancy (מספיק warps?), אז stalls (מה עוצר?). שלושתם עקביים - memory-bound => DRAM בשיא => long scoreboard.
פתרון תרגיל 4 - אישור בעיית coalescing עם ncu¶
נשתמש ב-rowsum_bad מול rowsum_good משיעור 8.4:
$ ncu --metrics \
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio \
./rowsum
rowsum_bad(const float *, float *, int, int)
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio 32.0
rowsum_good(const float *, float *, int, int)
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio 4.0
ובקטע המלא, ה-Memory Workload וה-Stalls של שתי הגרסאות:
$ ncu --set full -k "rowsum_bad|rowsum_good" ./rowsum
rowsum_bad:
Memory Throughput [%] 94.1 Sectors/Req 32.0
Stall Long Scoreboard 11.8 cyc/inst
Duration 4.81 ms
Roofline: on the diagonal roof, far from the ridge
rowsum_good:
Memory Throughput [%] 88.6 Sectors/Req 4.0
Stall Long Scoreboard 2.1 cyc/inst
Duration 0.36 ms (13.3x faster)
Roofline: on the diagonal roof, closer to the usable memory peak
למה זה עבד: rowsum_bad נותן ל-threads שכנים כתובות במרחק N*4 בייט, וכל warp מפעיל 32 sectors נפרדים לבקשה - מתוכם רק 1/8 שימושי. rowsum_good נותן כתובות רצופות ולכן 4 sectors לבקשה, כולם שימושיים. שלושת המדדים מספרים את אותו סיפור: 32 sectors => יותר LDG => יותר long scoreboard stalls (11.8 מול 2.1 מחזורים) => יותר זמן. ה-roofline של שתיהן על הגג האלכסוני (שתיהן memory-bound), אבל הרעה מבזבזת פי 8 בייטים. איך להכליל: sectors_per_request הוא המדד הישיר ביותר ל-coalescing, וקטע ה-Stalls מאשר אותו מזווית שנייה - גישה לא מאוחדת מייצרת יותר long scoreboard stalls כי היא מפעילה יותר עסקאות global. כשרואים 32 sectors לבקשה, מיד יודעים לחפש היכן threadIdx.x לא רץ על הממד הרציף (שיעור 8.4).
פתרון תרגיל 5 - אישור bank conflict עם ncu¶
נשתמש ב-transposeNaive (tile[32][32]) מול transposePadded (tile[32][33]) משיעור 8.4:
$ ncu --metrics \
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum \
./transpose
transposeNaive(const float *, float *, int, int)
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum 16,252,928
transposePadded(const float *, float *, int, int)
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum 0
בקטע המלא:
$ ncu --set full -k transposeNaive ./transpose
Section: Memory Workload Analysis
L1/TEX Hit Rate [%] 98.7
Shared Memory Bank Conflicts 16,252,928 <- the problem, in shared
Global Load Sectors/Req 4.0 <- global is clean!
Global Store Sectors/Req 4.0
Section: Warp State (Stall Reasons)
Stall Short Scoreboard [cyc/inst] 6.9 <- waiting for LDS from shared
Stall Long Scoreboard [cyc/inst] 1.2
למה זה עבד: הבעיה כולה ב-shared (מיליוני התנגשויות בקריאת tile[threadIdx.x][threadIdx.y]), בעוד שה-global נקי לגמרי - 4 sectors לבקשה גם בקריאה וגם בכתיבה, כי ה-threadIdx.x רץ על הממד הרציף בשני הצדדים (שיעור 8.4). לכן סיבת העצירה הדומיננטית היא short scoreboard (LDS נשאר על ה-SM) ולא long - זה בדיוק מבדיל bank conflict מבעיית coalescing. אחרי הריפוד ל-[32][33], ה-stride 33 זר ל-32, ההתנגשויות צונחות ל-0, וה-short scoreboard יורד. איך להכליל: bank conflict הוא תמיד תופעת shared ולכן מופיע כ-short scoreboard על LDS/STS, בעוד coalescing הוא תופעת global שמופיעה כ-long scoreboard על LDG/STG. סיבת העצירה עצמה מרמזת באיזה זיכרון לחפש - זכרו את המלכודת מ-8.3 שמיחזור scoreboards עלול לבלבל, ולכן תמיד הצליבו עם מדד ההתנגשויות הישיר.
פתרון תרגיל 6 - קורלציית מקור-SASS¶
$ nvcc -arch=sm_90 -O3 -lineinfo rowsum.cu -o rowsum
$ ncu --set full --import-source yes -k rowsum_bad -o rep_src ./rowsum
$ ncu -i rep_src.ncu-rep --page source
# Source | Address | SASS | Sampling(%)
------------------+---------+----------------------------+-------------
for(col=0;col<N;) | | |
acc += A[row*N | 0x1e00 | LDG.E.SYS R6, [R2] | ########## 74.3
+col]; | 0x1e10 | FADD R4, R4, R6 | # 3.1
s[row] = acc; | 0x1f40 | STG.E.SYS [R8], R4 | # 1.9
הצלבה עם ה-SASS הגולמי (בדיוק סימון ה-scoreboards משיעור 8.3):
[B------:R-:W3:-:S02] /*1e00*/ LDG.E.SYS R6, [R2] ; // writes scoreboard 3
[B---3--:R-:W-:Y:S04] /*1e10*/ FADD R4, R4, R6 ; // blocked until 3 clears
למה זה עבד: תצוגת ה-Source (שאפשרית רק בזכות -lineinfo, שמטמיע את מיפוי SASS-למקור בלי להאט - להבדיל מ--G) חושפת שהשורה acc += A[row*N+col] אוכלת 74% מדגימות ה-stall. הוראת ה-SASS שלה, LDG.E.SYS, כותבת scoreboard 3 וממתינה לזיכרון שחוזר מ-HBM; ה-FADD שצורך את הערך חוסם על scoreboard 3 (B---3--) - זהו ה-long scoreboard stall. השרשרת המלאה סגורה: השורה החמה במקור -> הוראת LDG ב-SASS -> long scoreboard stall -> נקודה על הגג האלכסוני ב-roofline -> memory-bound. איך להכליל: כשקטע ה-Stalls אומר "long scoreboard", תצוגת ה-Source מראה איזו שורה בדיוק אחראית, וכמעט תמיד זו גישת global לא-מאוחדת (LDG) או, אם ה-stall הוא short, גישת shared עם conflict (LDS). המיפוי הזה הוא אותה יכולת nvdisasm משיעור 4.4, עכשיו עם מוני stall מוצמדים לכל שורה.
פתרון תרגיל 7 (בונוס) - לולאת האופטימיזציה המלאה¶
הרצנו את הלולאה המלאה על pipeline.cu הגרוע. הטבלה מסכמת את שלוש האיטרציות:
iteration | tool | what we found | fix | speedup
--------+-------+----------------------------------+--------------------------+--------
1 | nsys | copy and compute serialized, zero overlap | streams + pinned (5.3) | x2.1
2 | ncu | kernel memory-bound, 32 sec/req | fix coalescing (8.4) | x11.0
3 | nsys | now overhead-bound: many dispatches | CUDA Graphs (5.5) | x1.4
--------+-------+----------------------------------+--------------------------+--------
cumulative end-to-end speedup: ~x32
הפלט בכל צעד:
# step 1 - nsys: the systemic problem (missing overlap)
Total (serial): 36.8 ms -> Total (overlap): 17.1 ms
# step 2 - ncu on the kernel after closing the overlap
Roofline: memory-bound; Sectors/Req: 32.0; Long Scoreboard: 11.8 cyc
after fixing the access pattern: Sectors/Req: 4.0; Duration: 11x faster
# step 3 - nsys again: the new bottleneck
cudaLaunchKernel jumps to 40% of the time -> overhead-bound
after CUDA Graphs: overhead plummets, another 1.4x
למה זה עבד: בכל איטרציה השתמשנו בכלי הנכון למפלס הנכון. nsys ("איפה") גילה קודם את הבעיה המערכתית - חוסר חפיפה - שהיא בעיה ש-ncu לעולם לא היה רואה, כי הוא מסתכל על kernel בודד ולא על ציר הזמן. רק אחרי שסגרנו את החפיפה עברנו ל-ncu ("למה") כדי לתקן את ה-kernel עצמו, וכשהצוואר הזה נפתר, הצוואר החדש (dispatch overhead) שוב הצריך את nsys. איך להכליל: זו מתודולוגיית ה-Theory of Constraints משיעור 7.1 - מצא, הרם, חזור על הצוואר החדש - כשכל איטרציה מתחילה ב-nsys ("איפה") ורק אז יורדת ל-ncu ("למה") אם הצוואר הוא kernel. אם היינו מתחילים מ-ncu, היינו מבזבזים שעות על אופטימיזציה של kernel בזמן שרוב הזמן בכלל הלך לחוסר החפיפה - צוואר שהיה נשאר בלתי-נראה. תמיד מהמערכתי אל הפרטני, לא ההפך. זהו בדיוק זרימת העבודה שנפעיל בפרויקטים 9.3-9.4.