7.2 Compute bound, memory bound ו arithmetic intensity פתרון
פתרון - Compute-bound, memory-bound ו-arithmetic intensity¶
כאן נעבור על הפתרון המלא של כל תרגיל - החישובים, הקוד, הפקודות והפלט. שימו לב: ערכי ה-ncu (מדדי ניצול, בתים, מונים) הם דוגמאות מהרצה על H100 (sm_90a) עם גודל בעיה מסוים, והם ישתנו מעט בין ארכיטקטורות, גדלים וגרסאות compiler - אבל ההיגיון, הסדר-גודל והסיווג (compute-bound מול memory-bound) זהים. מספרי ה-ridge point שבהם נשתמש הם 295 (H100 BF16) ו-281/562/1125 (B200 BF16/FP8/FP4), בדיוק מטבלת ההרצאה.
פתרון תרגיל 1 - חישוב וסיווג של arithmetic intensity¶
1. SAXPY (y = a*x + y, float של 4 בתים, a סקלר באוגר):
FLOPs = 2N (multiply a*x + add +y, per element)
bytes = 12N (load x=4, load y=4, write y=4)
AI = 2N / 12N = 1/6 ≈ 0.167 FLOPs/byte (N cancels out - constant)
0.167 << 295 => memory-bound, and always will be
2. חיבור וקטורים (z = x + y):
FLOPs = N (one add per element)
bytes = 12N (load x=4, load y=4, write z=4)
AI = N / 12N = 1/12 ≈ 0.083 FLOPs/byte
0.083 << 295 => memory-bound, even "worse" than SAXPY
שימו לב: המכנה זהה ל-SAXPY (12N), אבל המונה חצי (חיבור אחד במקום כפל-וחיבור), ולכן הarithmetic intensity חצי. זו הדגמה יפה לכך שפעולות איבר-איבר יושבות עמוק משמאל ל-ridge point.
3. SGEMM (C = A@B + C, מטריצות N x N, float):
FLOPs = 2N³ (N³ MAC operations, each 2 FLOP)
bytes = 16N² (read A,B,C + write C = 4 matrices × N² × 4 bytes)
AI = 2N³ / 16N² = N/8 FLOPs/byte (N does not cancel - linear!)
condition for compute-bound on H100: N/8 >= 295 => N >= 2360
AI classification vs ridge=295
SAXPY 1/6 ≈ 0.167 memory-bound (constant)
vector add 1/12 ≈ 0.083 memory-bound (constant, lower)
SGEMM N=1024 128 memory-bound (still below 295)
SGEMM N=2360 295 exactly on the ridge
SGEMM N=4096 512 compute-bound
למה זה עבד: ה-arithmetic intensity הוא תמיד FLOPs/bytes. בפעולות וקטוריות ה-N מתבטל ומתקבל מספר קבוע וזעום - ולכן הן memory-bound בכל גודל. ב-SGEMM ה-N שורד בגלל היחס O(N^3)/O(N^2), ולכן הarithmetic intensity גדלה לינארית וחוצה את ה-ridge point כשהמטריצה מספיק גדולה. N/8 = 295 נותן את הסף N = 2360.
איך להכליל: התבנית קבועה - כתבו FLOPs, כתבו bytes (רק מה שזז דרך צוואר הבקבוק), חלקו, השוו ל-ridge point של החומרה והדיוק שלכם. הכלל הגס: פעולות איבר-איבר ו-reductions הן memory-bound, כפל מטריצות וקונבולוציות גדולות הן compute-bound. ככל שהגרעין "מרובע יותר" ב-N (יותר reuse לכל בית), כך הarithmetic intensity גבוהה יותר.
פתרון תרגיל 2 - intertoken latency ו-batch לדוגמת ה-LLM¶
מודל 500 מיליארד פרמטרים × 16 ביט, על H100 (memory BW = 3.35 TB/s, arithmetic BW = 989 TFLOP/s BF16, ridge = 295).
1. גודל המודל:
2. intertoken latency ל-batch יחיד (צריך לטעון את כל המשקלים פעם לכל token):
3. arithmetic intensity ל-batch יחיד:
FLOPs = 2 × 500·10^9 = 10^12 FLOP (~2 FLOP per parameter)
bytes = 10^12 (all weights, single use)
AI = 10^12 / 10^12 = 1 FLOP/byte
1 << 295 => clearly memory-bound
4. גודל batch לcompute-bound. הarithmetic intensity ל-batch של B היא B פעמים זו של batch יחיד (אותם משקלים, פי B חישוב):
5. השפעה על throughput:
למה זה עבד: ה-decode נשלט לגמרי בטעינת המשקלים - 1 TB בכל צעד - ולכן ה-latency הוא פשוט גודל חלקי רוחב פס. ל-batch יחיד כל משקל טעון משמש פעם אחת, ולכן AI = 1 והגרעין memory-bound עמוקות. batching מזריק פי B יותר עבודת חישוב על אותה טעינת משקלים, ומכאן שהarithmetic intensity גדלה לינארית וה-throughput איתה - בלי תוספת latency, כי הזמן נקבע ממילא בטעינת המשקלים המשותפת.
איך להכליל: זו הנוסחה המעשית לתכנון הסקת LLM. intertoken latency ≈ (גודל מודל בבתים) / (memory bandwidth), כל עוד אתם memory-bound. גודל ה-batch שממנו נעשים compute-bound שווה בקירוב ל-ridge point של החומרה והדיוק. לכן ב-FP8 (ridge 592) צריך batch גדול יותר מ-BF16 (295) כדי למצות את אותו כרטיס - בדיוק מה שנראה בתרגיל הבא. KV-cache ורצפים ארוכים מסבכים את התמונה, אבל שלד החישוב הזה נשאר.
פתרון תרגיל 3 - מדוע FP8/FP4 מקשים להישאר compute-bound¶
1. חישוב מחדש של ה-ridge points של B200 (memory BW = 8 TB/s לכולם):
BF16: 2250 TFLOP/s / 8 TB/s = 281.25 ≈ 281
FP8: 4500 TFLOP/s / 8 TB/s = 562.5 ≈ 562
FP4: 9000 TFLOP/s / 8 TB/s = 1125
2. מה קורה לרוחבי הפס. בירידת דיוק, ה-arithmetic bandwidth מוכפל בכל צעד (FP8 = ×2 מ-BF16, FP4 = ×2 מ-FP8) כי יחידות החישוב מעבדות פי שניים ערכים "רזים" באותו שטח וזמן. אבל ה-memory bandwidth אינו משתנה - אותו HBM3e של 8 TB/s מזין את הכרטיס ללא קשר לדיוק החישוב. ה-ridge point הוא היחס ביניהם, ולכן כשהמונה מוכפל והמכנה קבוע, הרכס מוכפל.
3. ה-N המינימלי ל-SGEMM compute-bound על B200 (תנאי N/8 >= ridge, כלומר N >= 8 × ridge):
BF16 (ridge 281): N >= 8 × 281.25 = 2250
FP8 (ridge 562): N >= 8 × 562.5 = 4500
FP4 (ridge 1125): N >= 8 × 1125 = 9000
כדי למצות את החומרה ב-FP4 צריך מטריצה בגודל 9000 x 9000 לפחות - פי ארבע בצלע (פי 16 באיברים) מאשר ב-BF16. הדיוק הרזה "רעב" הרבה יותר לarithmetic intensity.
4. הפרדוקס בשורה אחת: דיוק נמוך מכפיל את כוח החישוב אבל משאיר את רוחב פס הזיכרון במקומו, ולכן דורש arithmetic intensity גבוהה יותר כדי להזין את הכוח הזה. זו לא סתירה - היא בדיוק ה"חומת הזיכרון": ככל שהחישוב מזנק קדימה והזיכרון נשאר מאחור, ה-ridge point נודד ימינה, וקשה יותר להיות compute-bound.
למה זה עבד: ridge point = arithmetic BW / memory BW, נקודה. דיוק נמוך פועל אך ורק על המונה (מכפיל arithmetic BW) ולא על המכנה (memory BW קבוע), ולכן הוא מזיז את הרכס ימינה בדיוק ביחס ההכפלה. לכן צריך מטריצות גדולות יותר - arithmetic intensity גבוהה יותר - כדי למצות את החומרה בדיוק נמוך.
איך להכליל: אותה לוגיקה חלה גם בין דורות: Ampere -> Hopper -> Blackwell העלו את החישוב מהר יותר מהזיכרון, ולכן הרכס עלה (156 -> 295 ומעלה). המשמעות המעשית: כשעוברים ל-FP8/FP4 או לכרטיס חדש כדי "להאיץ", חובה להגדיל במקביל את ה-batch ואת גודל הtiles כדי להישאר compute-bound - אחרת פשוט מקבלים יותר יחידות חישוב שיושבות בטלות ומחכות לזיכרון.
פתרון תרגיל 4 - להעלות arithmetic intensity לגרעין memory-bound¶
1. המימוש הלא-מאוחד של d = relu(a*x + b) כשלושה גרעינים נפרדים (float, a ו-b סקלרים):
k1: t1 = a*x read x(4) + write t1(4) = 8 bytes, 1 FLOP (multiply)
k2: t2 = t1 + b read t1(4) + write t2(4) = 8 bytes, 1 FLOP (add)
k3: d = relu(t2) read t2(4) + write d(4) = 8 bytes, 0 FLOP (max, not counted as FLOP)
-------------------------------------------------------------
total: bytes = 24N, FLOPs = 2N
AI = 2N / 24N = 1/12 ≈ 0.083 FLOPs/byte => deeply memory-bound
הביניים t1 ו-t2 נכתבים ל-global memory ונקראים ממנו שוב - בזבוז תעבורה טהור.
2. הגרסה המאוחדת - kernel fusion. גרעין אחד קורא x פעם אחת, מבצע את שלושת השלבים באוגרים, וכותב d פעם אחת:
__global__ void fused(const float* x, float* d, float a, float b, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
float t = a * x[i] + b; // multiply and add, stays in a register
d[i] = t > 0.0f ? t : 0.0f; // relu in a register
}
}
bytes = 8N (read x=4 + write d=4 only; no intermediate in memory)
FLOPs = 2N (same computation)
AI = 2N / 8N = 1/4 = 0.25 FLOPs/byte
improvement: 0.25 / 0.083 = 3x arithmetic intensity
3. טרנספורמציה שנייה - gradient checkpointing. באימון רשת נוירונים, ה-backward pass צריך את ה-activations של ה-forward. במקום לאחסן את כולם ב-global memory ולקרוא שוב (הרבה בתים במכנה), מחשבים אותם מחדש מנקודות ביקורת ספורות בזמן ה-backward. חוסכים תעבורת זיכרון (פחות בתים), משלמים בחישוב חוזר (יותר FLOPs) - ומעלים בכך את ה-arithmetic intensity. לחלופין, דחיסה (compression) מעבירה נתונים דחוסים (פחות בתים) ומשלמת ב-decompress (יותר FLOPs).
4. למה זה "כמעט חינם" ומתי נעצר. בגרעין memory-bound, יחידות החישוב יושבות בטלות בזמן שהזיכרון עובד קשה. ה-FLOPs הנוספים של ה-fusion/recompute/decompress נכנסים לתוך אותו זמן בטל - כמעט בלי לעלות זמן - בעוד החיסכון בבתים מקצר את צוואר הבקבוק האמיתי. ההעלאה משתלמת כל עוד אנחנו משמאל ל-ridge point. ברגע שהעלינו את הarithmetic intensity עד הרכס, החישוב עצמו נעשה צוואר הבקבוק, ומכאן והלאה כל FLOP נוסף כבר עולה זמן אמיתי - ואין יותר טעם.
למה זה עבד: ה-fusion מוחק את הכתיבה-והקריאה של הביניים ל-global memory. אותו חישוב בדיוק (2N FLOPs) מתחלק כעת בפחות בתים (8N במקום 24N), ולכן היחס - ה-arithmetic intensity - עולה פי 3. הזזנו את הנקודה ימינה על ה-Roofline בלי לשנות ולו FLOP אחד של חישוב שימושי.
איך להכליל: כל הטכניקות הגדולות להעלאת ביצועים בגרעינים memory-bound הן וריאציות על "החלף בתים ב-FLOPs": kernel fusion (בלי ביניים בזיכרון), recompute במקום store, דחיסה, ו-Flash Attention (שמחשב את ה-attention בבלוקים בלי לכתוב את מטריצת ה-N x N המלאה ל-DRAM). את בנייתן המעשית של אלה נלמד בפרק 8. הכלל המנחה קבוע: משמאל לרכס, חישוב זול וזיכרון יקר.
פתרון תרגיל 5 - אימות אמפירי עם ncu¶
הקובץ intensity.cu עם גרעין SAXPY (memory-bound) ו-GEMM דרך cuBLAS (compute-bound), כולל מאקרו בדיקת שגיאות:
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#include <cublas_v2.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 saxpy(int n, float a, const float* x, float* y) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i]; // one FFMA per thread => AI = 1/6
}
int main() {
// --- SAXPY: memory-bound ---
const int n = 1 << 26; // 67M elements, clearly exceeds L2
float *dx, *dy;
CUDA_CHECK(cudaMalloc(&dx, sizeof(float) * n));
CUDA_CHECK(cudaMalloc(&dy, sizeof(float) * n));
CUDA_CHECK(cudaMemset(dx, 0, sizeof(float) * n));
CUDA_CHECK(cudaMemset(dy, 0, sizeof(float) * n));
saxpy<<<(n + 255) / 256, 256>>>(n, 2.0f, dx, dy);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
// --- GEMM: compute-bound ---
const int N = 4096; // N/8 = 512 > 295 => compute-bound
float *dA, *dB, *dC;
CUDA_CHECK(cudaMalloc(&dA, sizeof(float) * N * N));
CUDA_CHECK(cudaMalloc(&dB, sizeof(float) * N * N));
CUDA_CHECK(cudaMalloc(&dC, sizeof(float) * N * N));
cublasHandle_t h; cublasCreate(&h);
float alpha = 1.0f, beta = 0.0f;
cublasSgemm(h, CUBLAS_OP_N, CUBLAS_OP_N, N, N, N,
&alpha, dA, N, dB, N, &beta, dC, N);
CUDA_CHECK(cudaDeviceSynchronize());
cublasDestroy(h);
CUDA_CHECK(cudaFree(dx)); CUDA_CHECK(cudaFree(dy));
CUDA_CHECK(cudaFree(dA)); CUDA_CHECK(cudaFree(dB)); CUDA_CHECK(cudaFree(dC));
printf("done\n");
return 0;
}
קומפילציה והרצה תחת הפרופיילר:
הפלט (מקוצר לשורות המכריעות):
saxpy(int, float, const float *, float *)
Compute (SM) Throughput [%] 3.10
Memory Throughput [%] 92.40 <== saturated!
-> Kernel is bound by DRAM bandwidth (memory-bound)
ampere_sgemm_128x128_nn (cuBLAS)
Compute (SM) Throughput [%] 88.70 <== saturated!
Memory Throughput [%] 14.20
-> Kernel is bound by compute (compute-bound)
למה זה עבד: ל-SAXPY arithmetic intensity 1/6, הרחק משמאל ל-ridge point, ולכן ה-DRAM רווי (~92%) בעוד pipelines החישוב כמעט בטלים (~3%) - memory-bound בדיוק כפי שחישבנו. ל-GEMM על 4096 x 4096 arithmetic intensity 512, מעל 295, ולכן התמונה מתהפכת: החישוב רווי והזיכרון רפוי - compute-bound. ncu --set full אף משרטט את דיאגרמת ה-Roofline עם הנקודה של כל גרעין ממוקמת מתחת לגג הנכון, בדיוק הגרף מ-7.1.
איך להכליל: שתי המטריקות Compute (SM) Throughput מול Memory Throughput הן מבחן הסיווג המהיר: מי מהן קרוב ל-100% הוא צוואר הבקבוק. זו הדרך לאמת אמפירית כל השערת Roofline לפני שמשקיעים באופטימיזציה - קודם מודדים באיזה משטר אתם, ואז מטפלים בגורם החוסם הנכון (חוק ה-bottleneck מ-7.1: אתרו, הרימו, חזרו).
פתרון תרגיל 6 (בונוס) - למדוד arithmetic intensity ישירות ב-ncu¶
מדידת המונים הגולמיים שמרכיבים את ה-arithmetic intensity של גרעין ה-SAXPY (n = 1<<26 = 67,108,864):
ncu --metrics \
smsp__sass_thread_inst_executed_op_fadd_pred_on.sum,\
smsp__sass_thread_inst_executed_op_fmul_pred_on.sum,\
smsp__sass_thread_inst_executed_op_ffma_pred_on.sum,\
dram__bytes.sum ./intensity
הפלט (עבור גרעין ה-saxpy):
saxpy(int, float, const float *, float *)
smsp__sass_thread_inst_executed_op_fadd_pred_on.sum inst 0
smsp__sass_thread_inst_executed_op_fmul_pred_on.sum inst 0
smsp__sass_thread_inst_executed_op_ffma_pred_on.sum inst 67,108,864
dram__bytes.sum byte 805,306,368
חישוב ה-arithmetic intensity הנמדד:
FLOPs = fadd + fmul + 2*ffma
= 0 + 0 + 2 × 67,108,864 = 134,217,728 FLOP
bytes = dram__bytes.sum = 805,306,368 byte (= 12 × 67,108,864, as expected)
AI_measured = 134,217,728 / 805,306,368 = 0.1667 ≈ 1/6
המדידה תואמת בדיוק את הערך התיאורטי 1/6 מתרגיל 1.
למה זה עבד: ה-saxpy מתקמפל ל-FFMA אחד לכל thread (כפל-וצבירה מאוחדים), ולכן מונה ה-ffma שווה בדיוק ל-n, ו-fadd/fmul אפס. כל FFMA נספר כ-2 FLOP (כפל + חיבור). ה-dram__bytes.sum יצא 12N בדיוק כי הווקטור (256 MiB לכל מערך) גדול בהרבה מ-L2 (50 MiB ב-H100), כך שאין reuse בcache וכל בית מגיע מ-DRAM. חלוקת המונים נתנה 1/6 מדויק.
איך להכליל: זו הדרך למדוד arithmetic intensity של כל גרעין, גם כזה שקשה לחשב ידנית: סכמו את מוני ה-FLOP (fadd + fmul + fp16/fp64 מתאימים + 2×ffma, ולגרעיני Tensor Cores השתמשו במוני ה-pipe_tensor), חלקו ב-dram__bytes.sum, וקבלו נקודה מדודה על ציר ה-X של ה-Roofline. אם הווקטור קטן וחלקו נשאר ב-L2, ה-dram__bytes.sum יירד וה-arithmetic intensity הנמדדת תעלה מעל התיאוריה - הגדילו את הבעיה עד שהיא חורגת בבירור מ-L2 כדי לאמת את החישוב הידני.