7.3 רוחבי פס ו peak rate פתרון
פתרון - רוחבי פס ו-peak rate¶
נעבור תרגיל אחרי תרגיל עם החישובים, הקוד, הפקודות והפלט הצפוי. כל המספרים כאן הם מ-H100 (SXM) עם CUDA 12.4; רוחב הפס וה-FLOP/s המושגים משתנים מעט בין הרצות ובין דורות דרייבר, והם ערכים אופייניים - העיקר הוא היחסים והאחוזים מהשיא. מי שעובד על T4 ב-Colab יראה מספרים אחרים לגמרי, אבל אותה תמונה של "אחוז מהשיא".
פתרון תרגיל 1 - חישוב קצב השיא מהמפרט¶
קצב השיא של FP32 (non-Tensor). מהמפרט: 132 SM, 128 ליבות FP32 ל-SM, תדר 1.98 GHz, ו-FMA = 2 FLOP:
FLOP per cycle = 132 x 128 x 2 = 33,792 FLOP/cycle
peak rate = 33,792 x 1.98e9
= 6.691e13 FLOP/s
= 66.9 TFLOPS (FP32, non-Tensor)
זה בדיוק המספר מה-whitepaper של ה-H100.
קצב השיא של BF16 על ה-Tensor Cores - חישוב הפוך מ-989 TFLOP/s:
FLOP per cycle (total) = 989e12 / 1.98e9 = 499,494 FLOP/cycle
FLOP per SM per cycle = 499,494 / 132 = 3,784 FLOP/SM/cycle
FMA per SM per cycle = 3,784 / 2 = 1,892 BF16 FMA/SM/cycle
השוואה בין היחידות:
כלומר כל SM מבצע כמעט פי 15 יותר FMA של BF16 על ה-Tensor Cores מאשר FP32 על ה-CUDA Cores - וזו בדיוק הקפיצה מ-66.9 ל-989 TFLOP/s.
מלכודת ה-FMA (סעיף 4): אם שוכחים ש-FMA = 2 FLOP ומכפילים ב-1 במקום ב-2, מקבלים ל-FP32:
זו טעות של פי 2 (50% נמוך מדי). זו הטעות הכי נפוצה בכל חישוב קצב שיא.
למה זה עבד: קצב השיא הוא כפל טהור של ארבעה מספרי מפרט קבועים - אין מדידה, רק (SM x ליבות x FLOP-לליבה x תדר). החישוב ההפוך של BF16 מוכיח ש-989 עקבי עם אותה נוסחה, ומחזיר את גורם ה-FMA=2 לגלוי. איך להכליל: אותה נוסחה עובדת לכל כרטיס ולכל (יחידה x דיוק); רק צריך את הספירה הנכונה של יחידות הביצוע ואת גורם ה-FLOP-לפעולה (2 ל-FMA, ולפעולות Tensor Core גדולות עוד יותר). תמיד ציינו איזה דיוק - "השיא" בלי דיוק חסר משמעות.
פתרון תרגיל 2 - מדידת רוחב פס מושג של העתקה device-to-device¶
// bwpeak.cu (exercise 2 part)
#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 copyKernel(const float* __restrict__ in,
float* __restrict__ out, size_t n) {
size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
size_t stride = (size_t)gridDim.x * blockDim.x;
for (; i < n; i += stride)
out[i] = in[i];
}
int main() {
const size_t N = 256ull << 20; // 256M floats = 1 GiB per array
const size_t bytesArr = N * sizeof(float);
float *d_in, *d_out;
CUDA_CHECK(cudaMalloc(&d_in, bytesArr));
CUDA_CHECK(cudaMalloc(&d_out, bytesArr));
CUDA_CHECK(cudaMemset(d_in, 1, bytesArr));
int block = 256, grid = 4096;
copyKernel<<<grid, block>>>(d_in, d_out, N); // warmup
CUDA_CHECK(cudaDeviceSynchronize());
cudaEvent_t a, b;
CUDA_CHECK(cudaEventCreate(&a));
CUDA_CHECK(cudaEventCreate(&b));
const int ITERS = 50;
CUDA_CHECK(cudaEventRecord(a));
for (int k = 0; k < ITERS; ++k)
copyKernel<<<grid, block>>>(d_in, d_out, N);
CUDA_CHECK(cudaEventRecord(b));
CUDA_CHECK(cudaEventSynchronize(b));
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, a, b));
double sec = (ms / 1e3) / ITERS;
double movedBytes = 2.0 * bytesArr; // read + write
double gbps = movedBytes / sec / 1e9;
printf("copy: %.0f GB/s (%.1f%% of 3350 GB/s HBM3 peak)\n",
gbps, gbps / 3350.0 * 100.0);
CUDA_CHECK(cudaEventDestroy(a)); CUDA_CHECK(cudaEventDestroy(b));
CUDA_CHECK(cudaFree(d_in)); CUDA_CHECK(cudaFree(d_out));
return 0;
}
$ nvcc -arch=sm_90 -O3 bwpeak.cu -o bwpeak
$ ./bwpeak
copy: 3040 GB/s (90.7% of 3350 GB/s HBM3 peak)
כ-91% מהשיא - מצוין. סריקת תצורות ההפעלה (סעיף 5) מראה שברגע שיש מספיק blocks כדי להרוות את הכרטיס, רוחב הפס יציב למדי:
block=128 grid=8192 : 3010 GB/s (89.9%)
block=256 grid=4096 : 3040 GB/s (90.7%)
block=512 grid=2048 : 3035 GB/s (90.6%)
block=64 grid=1024 : 2360 GB/s (70.4%) <-- too few threads, not saturating
התצורה האחרונה נופלת כי אין מספיק warps דרים כדי להסתיר את latency ה-HBM (הסתרת latency - שיעור 7.4); בשאר התצורות ה-kernel רווי וגורף ~90%.
למה זה עבד: ההעתקה לא מבצעת שום חישוב, ולכן היא memory-bound מובהקת וגורפת כמעט את כל רוחב הפס. ספירת הבייטים היא 2 x N x 4 כי כל אלמנט נקרא וגם נכתב - זה החלק שהכי קל לטעות בו. החימום וה-50 האיטרציות מבטיחים שלא מודדים overhead הפעלה ראשונה. איך להכליל: זהו ה-benchmark הקנוני לרוחב פס (STREAM copy). ל-triad (a=b+s*c) הבייטים הם 3 x N x 4 (שתי קריאות וכתיבה). אחוז מהשיא של 85%-92% הוא הסימן שה-kernel רווי-זיכרון; פחות מזה מרמז שה-launch לא מרווה או שהגישות לא מאוחדות (coalescing).
פתרון תרגיל 3 - נקודת המפנה לשני דיוקים¶
H100 BF16 TC : 989 TFLOP/s / 3.35 TB/s = 989e12 / 3.35e12 = 295 FLOP/byte
H100 FP8 TC : 1979 TFLOP/s / 3.35 TB/s = 1979e12 / 3.35e12 = 591 FLOP/byte
שני הערכים תואמים לטבלה של שיעור 7.2 (295 ו-~592; ההפרש הזעיר הוא עיגול).
מדוע FP8 כפול מ-BF16 (סעיף 3): רוחב פס הזיכרון זהה לשני הדיוקים (אותו HBM3, 3.35 TB/s) - ה-FP8 רק מכפיל את רוחב הפס האריתמטי (מ-989 ל-1979). נקודת המפנה היא היחס ביניהם, ולכן היא מוכפלת יחד עם המונה. המשמעות: ב-FP8, kernel חייב לבצע 591 פעולות לכל בית כדי להישאר compute-bound, מול "רק" 295 ב-BF16. דיוק נמוך יותר קונה יותר חישוב, אבל מייקר את הכניסה למועדון ה-compute-bound.
סיווג הכפלת-סקלר (סעיף 4):
y[i] = a * x[i] -> read 4B + write 4B = 8 bytes per 1 FLOP
arithmetic intensity = 1 / 8 = 0.125 FLOP/byte
0.125 << 295 (BF16) and 591 (FP8)
עמוק, עמוק משמאל לנקודת המפנה - memory-bound באופן מוחלט. שום דבר לא יעזור לו חוץ מרוחב פס גבוה יותר; ה-Tensor Cores של הכרטיס ישבו סרק.
בונוס (סעיף 5):
A100 BF16 : 312 / 2 = 156 FLOP/byte
H100 BF16 : 989 / 3.35 = 295 FLOP/byte
H100 FP8 : 1979 / 3.35 = 591 FLOP/byte
B200 FP4 : 9000 / 8 = 1125 FLOP/byte
ascending order: 156 < 295 < 591 < 1125
המגמה: נקודת המפנה מטפסת עם הדורות (A100 -> H100 -> B200) ועם ירידת הדיוק (BF16 -> FP8 -> FP4), כי רוחב הפס האריתמטי גדל הרבה יותר מהר מרוחב פס הזיכרון. ככל שהחומרה חדשה יותר, כך קשה יותר להישאר compute-bound.
למה זה עבד: נקודת המפנה היא פשוט מנה של שני מספרים שכבר חישבנו - היחידות (FLOP/s)/(בייט/s) מצטמצמות ל-FLOP/בייט, אותה יחידה של ציר ה-x. איך להכליל: כדי לסווג כל kernel, חשבו את ה-arithmetic intensity שלו (FLOP חלקי בייטים שנעים מ-HBM) והשוו לנקודת המפנה של הכרטיס-והדיוק שאתם מריצים. גדול ממנה = compute-bound, קטן = memory-bound.
פתרון תרגיל 4 - למה הביצוע בפועל תמיד מתחת לשיא¶
שלוש הסיבות (סעיף 1):
- לחץ אוגרים (register pressure). kernel שצורך הרבה אוגרים מוריד תפוסה, ואם הוא מכריח גלישה (spill) - כל spill הוא גישת DRAM איטית (ראינו זאת בשיעור 1.6). פחות warps דרים משמעו פחות הסתרת latency.
- מגבלת רוחב פס זיכרון. אם ה-kernel memory-bound, הליבות מבזבזות את רוב זמנן בהמתנה ל-HBM; קצב החישוב שלהן לעולם לא יתקרב לשיא כי אין להן מספיק נתונים להאכיל אותן.
- מחסומי סנכרון (barriers). כל
__syncthreads()מכריח warps מהירים להמתין לאיטיים; הזמן שהם ממתינים הוא זמן שבו יחידות הביצוע יושבות סרק.
אלה בדיוק שלושת החשודים שהגלוסרי מונה כסיבה לכך שקצב השיא "מושג לעתים נדירות".
סיווג ה-copyKernel (סעיף 2): הוא memory-bound מובהק - הוא מזיז נתונים בלי לחשב עליהם. הוא "נוגע" בגג הזיכרון המשופע (גורף ~90% מרוחב הפס), אבל גג החישוב שלו כמעט אפס, כי הוא מבצע ~0 פעולות ממשיות: אין FMA, אין כפל, רק load ו-store.
אימות ב-ncu (סעיף 3):
$ ncu --metrics \
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\
sm__throughput.avg.pct_of_peak_sustained_elapsed \
./bwpeak
copyKernel(const float *, float *, unsigned long)
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed 90.6 %
sm__throughput.avg.pct_of_peak_sustained_elapsed 6.8 %
ה-DRAM throughput גבוה (90.6%) וה-SM throughput נמוך (6.8%) - החתימה המדויקת של kernel memory-bound, בדיוק כפי שסיווגנו: נוגע בגג הזיכרון, לא בגג החישוב.
מדוע רק 90% ולא 100% (סעיף 4): אפילו ה-kernel האופטימלי מפסיד את ה-10% האחרונים לoverheads פיזיות של ה-DRAM: מחזורי רענון (refresh) של תאי ה-HBM, מעברי שורות (row activation) בבנקי הזיכרון, ו-overhead התחלה/סיום של כל הפעלה. השיא של 3.35 TB/s מניח זרימה רציפה מושלמת שאף פעם לא באמת קורית. 90% הוא ה-speed of light המעשי.
למה זה עבד: שני מדדי ה-pct_of_peak_sustained_elapsed הם בדיוק "אחוז מהשיא" של שני הגגות - האחד לזיכרון, האחר לחישוב - ולכן הם קוראים ישירות את מיקום ה-kernel על ה-Roofline. איך להכליל: הריצו כל kernel עם שני המדדים האלה; הגבוה מביניהם אומר לכם באיזה גג אתם נתקעים, וזה מה שצריך לייעל. אם שניהם נמוכים - ה-kernel חסום-תקורה (overhead-bound) ולא מרווה את הכרטיס כלל.
פתרון תרגיל 5 - מדידת FLOP/s מושג ואחוז מהשיא¶
// bwpeak.cu (exercise 5 part)
#define ITERS_INNER 1000000
__global__ void fmaBench(float x, float* sink, int iters) {
// eight parallel FMA chains - enough independent work to saturate the pipeline
float a0=x, a1=x+1, a2=x+2, a3=x+3, a4=x+4, a5=x+5, a6=x+6, a7=x+7;
const float c = 1.0000001f, d = 0.5f;
for (int i = 0; i < iters; ++i) {
a0 = a0*c + d; a1 = a1*c + d; a2 = a2*c + d; a3 = a3*c + d;
a4 = a4*c + d; a5 = a5*c + d; a6 = a6*c + d; a7 = a7*c + d;
}
sink[threadIdx.x] = a0+a1+a2+a3+a4+a5+a6+a7; // prevents the loop from being eliminated
}
הקוד ב-main (בנוסף לחלק של תרגיל 2):
const int block2 = 256;
const int grid2 = 132 * 8; // 1056 blocks -> 270,336 threads
const size_t threads = (size_t)grid2 * block2;
float* d_sink;
CUDA_CHECK(cudaMalloc(&d_sink, block2 * sizeof(float)));
fmaBench<<<grid2, block2>>>(1.0f, d_sink, ITERS_INNER); // warmup
CUDA_CHECK(cudaDeviceSynchronize());
cudaEvent_t c0, c1;
CUDA_CHECK(cudaEventCreate(&c0)); CUDA_CHECK(cudaEventCreate(&c1));
CUDA_CHECK(cudaEventRecord(c0));
fmaBench<<<grid2, block2>>>(1.0f, d_sink, ITERS_INNER);
CUDA_CHECK(cudaEventRecord(c1));
CUDA_CHECK(cudaEventSynchronize(c1));
float ms2 = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms2, c0, c1));
// 8 FMA per iteration, x2 FLOP per FMA
double flop = (double)threads * ITERS_INNER * 8.0 * 2.0;
double tflops = flop / (ms2 / 1e3) / 1e12;
printf("fma: %.1f TFLOP/s (%.1f%% of 66.9 TFLOPS FP32 peak)\n",
tflops, tflops / 66.9 * 100.0);
CUDA_CHECK(cudaFree(d_sink));
$ nvcc -arch=sm_90 -O3 bwpeak.cu -o bwpeak
$ ./bwpeak
copy: 3040 GB/s (90.7% of 3350 GB/s HBM3 peak)
fma: 61.4 TFLOP/s (91.8% of 66.9 TFLOPS FP32 peak)
ספירת ה-FLOP: 270,336 threads x 1,000,000 iters x 8 FMA x 2 = 4.325e12 FLOP. חלוקה בזמן (~70 ms) נותנת ~61 TFLOP/s, שהם ~92% מ-66.9. אימות ב-ncu (סעיף 4):
$ ncu --metrics \
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\
sm__throughput.avg.pct_of_peak_sustained_elapsed \
./bwpeak
fmaBench(float, float *, int)
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed 0.1 %
sm__throughput.avg.pct_of_peak_sustained_elapsed 92.3 %
בדיוק ההפך מה-copyKernel: SM throughput גבוה (92.3%) ו-DRAM throughput אפסי (0.1%) - החתימה של kernel compute-bound. ה-92% מאשרים את חישוב הידני.
למה זה עבד: שמונה שרשראות ה-FMA המקבילות נותנות לחומרה מספיק עבודה בלתי-תלויה כדי להסתיר את latency של ה-FMA ולהחזיק את ה-CUDA Cores עסוקים כמעט תמיד; כל האופרנדים חיים באוגרים, אז אין גישות זיכרון להאט. הכתיבה ל-sink מונעת מהcompiler למחוק את הלולאה כ"קוד מת". איך להכליל: זו התבנית למדידת גג החישוב - שרשראות מקבילות של פעולה בודדת (כאן FMA), ספירת FLOP מדויקת (x2 ל-FMA), חלוקה בזמן. אם קיבלתם הרבה פחות מ-90%, כנראה אין מספיק ILP (שרשראות מקבילות) או שהcompiler ייעל חלק מהחישוב החוצה.
פתרון תרגיל 6 (בונוס) - מיפוי ה-Roofline של שני ה-kernels¶
זוגות ה-(arithmetic intensity, ביצוע מושג) של שני ה-kernels:
kernel | AI [FLOP/byte] | achieved perf | roof it touches
--------------+----------------+-------------------+-----------------
copyKernel | ~0 (no compute) | 3040 GB/s | memory roof (diagonal)
fmaBench | very high | 61.4 TFLOP/s | compute roof (horizontal)
מיקום על ה-Roofline של H100 (גג חישוב 66.9 TFLOPS ל-FP32, שיפוע 3.35 TB/s):
performance [TFLOP/s]
66.9 |·······································____________ <- compute roof 66.9
| ,-'' [*] fmaBench (61.4, ~92%)
| ,-''
| ,-''
| ,-''
| ,-''
| ,-''
| ,-''
| [*] copyKernel sits here on the left: AI~0, on the diagonal roof
|,-''_________________________________________________ AI [FLOP/byte]
0
ה-copyKernel יושב בקצה השמאלי הקיצוני, על הגג המשופע (memory-bound); ה-fmaBench יושב על הגג האופקי (compute-bound). טבלת ה-2x2 מ-ncu מזהה זאת בלי לצייר בכלל:
kernel | DRAM % of peak | SM % of peak | classification
-----------+----------------+--------------+---------------
copyKernel | 90.6 % | 6.8 % | memory-bound
fmaBench | 0.1 % | 92.3 % | compute-bound
האלכסון של הטבלה מספר את כל הסיפור: כל kernel גבוה בדיוק בעמודה של הגג שחוסם אותו ונמוך בשנייה.
למה זה עבד: שני מדדי ה-pct_of_peak_sustained_elapsed הם המנה achieved/peak של שני הגגות; ה-kernel memory-bound גורף את גג הזיכרון ומשאיר את החישוב סרק, וההפך ל-kernel compute-bound. איך להכליל: זו האבחנה שנפעיל בכל פרויקט - מריצים עם שני המדדים, מסתכלים איזה גבוה, ומייעלים את המשאב שנוגעים בו. אם שניהם נמוכים, ה-kernel לא מרווה את הכרטיס והבעיה היא overhead או launch, לא הגג עצמו.