7.3 רוחבי פס ו peak rate הרצאה

בשיעור 7.1 בנינו את מודל ה-Roofline ובשיעור 7.2 הכרנו את שני הצירים שלו - את ה-arithmetic intensity (עצימות אריתמטית) על ציר ה-x ואת הביצוע המושג על ציר ה-y - ואת שני ה"גגות" שחוסמים כל kernel: גג חישוב אופקי וגג זיכרון משופע. אבל עד עכשיו ציירנו את הגגות בלי לנקוב במספר מדויק של אף אחד מהם. השיעור הזה נותן לגגות את הגובה והשיפוע שלהם. נגדיר שלושה מושגים שכל הפרק נשען עליהם: רוחב הפס של הזיכרון (memory bandwidth) שקובע את שיפוע גג הזיכרון, רוחב הפס האריתמטי (arithmetic bandwidth) שקובע את גובה גג החישוב, וקצב השיא (peak rate) - "מהירות האור" של הכרטיס, המספר התיאורטי שמחושב ישירות מהמפרט ומשמש מכנה בכל מדד ניצול. נחשב את קצב השיא של ה-H100 יד ביד מהמספרים שראינו בפרק 1, נחבר את שני רוחבי הפס לנקודת המפנה (ridge point) שכבר פגשנו ב-7.2, ולבסוף נצא מהתיאוריה למעבדה: נכתוב kernel שמודד את רוחב הפס וה-FLOP/s שהושגו בפועל ונחשב אחוז מהשיא. זה הכלי שנשתמש בו בכל פרויקט מכאן והלאה כדי לענות על השאלה "כמה טוב ה-kernel שלי באמת".

רוחב פס של זיכרון - memory bandwidth

רוחב הפס של הזיכרון הוא הקצב המרבי שבו ניתן להעביר נתונים בין שתי רמות סמוכות בהיררכיית הזיכרון, נמדד בבייטים לשנייה. הגדרה זו נשמעת פשוטה אבל טומנת בחובה דקות אחת חשובה: ל-GPU אין רוחב פס אחד, יש לו רבים - אחד לכל זוג שכבות סמוכות בפירמידה שבנינו בשיעור 1.6. יש רוחב פס בין קובץ האוגרים ל-shared memory, בין ה-L1 ל-L2, בין ה-L2 לזיכרון ה-GPU, ומעבר לכרטיס - דרך ה-PCIe אל הhost ודרך ה-NVLink אל כרטיסים אחרים.

מבין כל רוחבי הפס האלה, אחד שולט על מודל ה-Roofline של רוב ה-kernels: רוחב הפס בין זיכרון ה-GPU (ה-HBM / global memory) לבין קבצי האוגרים שב-SM. הסיבה, כפי שהגלוסרי מנסח אותה, נחרצת: "מערכי העבודה (working sets) של רוב ה-kernels נכנסים רק בזיכרון ה-GPU, ולא גבוה יותר בהיררכיה" - כלומר לא במלואם ב-cache, ב-shared memory או באוגרים. לכן כל בית שה-kernel מעבד חייב לזרום, במוקדם או במאוחר, דרך הpipe הזה, והוא זה שקובע את שיפוע גג הזיכרון ב-Roofline. זהו הpipe שכיווננו אליו כבר בסוף שיעור 1.6.

המספרים המובילים של הדור הנוכחי, לפי דור ה-HBM שכל כרטיס נושא:

כרטיס	סוג זיכרון	רוחב פס זיכרון
A100 80GB SXM (Ampere)	HBM2e	2 TB/s
H100 SXM (Hopper)	HBM3	3.35 TB/s
B200 (Blackwell)	HBM3e	8 TB/s

שימו לב ש-TB/s כאן הוא במונחים עשרוניים (10^12 בייט לשנייה), כמקובל בשיווק רוחבי פס - להבדיל מ-GiB של קיבולת שהוא בינארי. ה-H100 שלנו מספק 3.35e12 בייט בשנייה בין ה-HBM3 לשבב. מהיר במונחים אנושיים, אבל כפי שנראה מיד - איטי בסדר גודל מקצב הצריכה של הליבות. הפער הזה בין הpipe לליבות הוא כל הסיפור של ה-Roofline.

רוחב פס אריתמטי - arithmetic bandwidth

רוחב הפס האריתמטי הוא קצב השיא שבו המערכת יכולה לבצע עבודה אריתמטית - התקרה התיאורטית של פעולות חישוב לשנייה, נמדדת ב-FLOP/s (ומכפיליו TFLOP/s, PFLOP/s). זהו גובה גג החישוב ב-Roofline: הקו האופקי שאף kernel לא יכול לחצות.

כמו אצל רוחב הפס של הזיכרון, גם כאן הדקות המרכזית היא שאין מספר אחד לכרטיס. לכל צירוף של (יחידת ביצוע x דיוק מספרי) יש רוחב פס אריתמטי משלו. אותו H100 מספק מספרים שונים לחלוטין תלוי במה בדיוק אתם מריצים:

• ליבות CUDA מול Tensor Cores. ה-CUDA Cores הם היחידות הכלליות שראינו בשיעור 1.3; ה-Tensor Cores (שיעור 1.5) הם יחידות ייעודיות לכפל מטריצות. הגלוסרי נותן כלל אצבע חד: היחס בין רוחב הפס של ה-Tensor Core לזה של ה-CUDA Core הוא בערך 100:1 - שני סדרי גודל. ביצועי השיא של GPU מודרני נשלטים כמעט לחלוטין על ידי ה-Tensor Cores.
• דיוק מספרי. בתוך ה-Tensor Cores, כל חצייה של הדיוק בערך מכפילה את רוחב הפס. על B200: BF16 נותן 2250 TFLOP/s, ה-FP8 מכפיל ל-4500, וה-FP4 מכפיל שוב ל-9000 (תשעה PFLOP/s - המספר המוביל של Blackwell). על H100: BF16 נותן 989, ו-FP8 מכפיל ל-1979.
• ממשי מול שלם. ברוב ה-GPU, פעולות על מספרים ממשיים (FP) מקבלות יותר רוחב פס מפעולות שלמות (INT).

הטבלה המלאה, שנחזור אליה מיד כשנחשב נקודות מפנה:

מערכת	רוחב פס אריתמטי (TFLOP/s)	דיוק / יחידה
H100 SXM FP32 (non-Tensor)	66.9	CUDA Core
A100 80GB BF16 TC	312	Tensor Core
H100 SXM BF16 TC	989	Tensor Core
B200 BF16 TC	2250	Tensor Core
H100 SXM FP8 TC	1979	Tensor Core
B200 FP8 TC	4500	Tensor Core
B200 FP4 TC	9000	Tensor Core

הביטו בשורה הראשונה מול השאר: אותו H100 עצמו נותן 66.9 TFLOP/s ב-FP32 על ה-CUDA Cores, אבל 989 ב-BF16 על ה-Tensor Cores - פי ~15 - ו-1979 ב-FP8. המסקנה המעשית: כשאתם מדברים על "השיא של H100" אתם חייבים לומר באיזו יחידה ובאיזה דיוק, אחרת המספר חסר משמעות.

קצב השיא - peak rate

קצב השיא הוא הקצב המרבי התיאורטי שבו מערכת חומרה יכולה להשלים עבודה - "הגבול העליון המוחלט של ביצועי ה-GPU כשכל יחידת ביצוע פועלת בקצב מלא ביעילות מושלמת". קצב השיא הוא הכללה של רוחב הפס האריתמטי: רוחב הפס האריתמטי הוא בדיוק קצב השיא כשמיישמים אותו על פעולות אריתמטיות. מהנדסי NVIDIA מכנים אותו בחיבה "מהירות האור" (speed of light) - גבול פיזיקלי שכל ביצוע ממשי נמדד מולו.

שלוש תכונות מגדירות אותו:

1. הוא מחושב ישירות מהמפרט הקבוע של הכרטיס - מספר ה-SM, מספר הליבות ל-SM, פעולות לליבה למחזור, ותדר השעון. אין בו מדידה, רק כפל.
2. הוא מניח יעילות מושלמת - אפס לחץ אוגרים, אפס מגבלת רוחב פס, אפס עצירות על מחסומי סנכרון (barriers). אף אחד מאלה לא קורה במציאות.
3. הוא המכנה בכל מדד ניצול (utilization). ניצול = מושג / שיא. לכן קצב השיא הוא "השופט האולטימטיבי של ניצול ה-GPU" - היעד שכל אופטימיזציה מכוונת אליו.

בואו נחשב אותו ל-H100 ב-FP32, בדיוק מהמספרים שספרנו בפרק 1. הנוסחה הכללית:

peak rate = (number of SMs) x (cores per SM) x (operations per core per cycle) x (clock frequency)

והמספרים של ה-H100:

number of SMs:                132
FP32 cores per SM:            128
operations per core per cycle: 2      <-- FMA = 2 real operations!
clock frequency (boost, subsystem): 1.98 GHz  (1980 MHz)

FLOP per cycle  = 132 x 128 x 2        = 33,792 FLOP/cycle
peak rate       = 33,792 x 1.98e9      = 6.69e13 FLOP/s
                = 66.9 TFLOPS   (FP32, non-Tensor)

זהו בדיוק המספר שרשום ב-whitepaper של ה-H100. שימו לב לגורם שמפילים אותו הכי הרבה תלמידים: ה-FMA (Fused Multiply-Add) - הפעולה a*b+c - נספרת כשתי פעולות ממשיות (כפל וחיבור), כי היא אכן מבצעת שתיים בהוראה אחת. שכחה של הגורם הזה חוצה את התוצאה לחצי (33.5 TFLOPS) ונותנת מספר שגוי לחלוטין. זכרו: FMA = 2 FLOP.

אותה נוסחה עובדת ל-Tensor Cores, רק שיחידת "פעולות לליבה למחזור" גדולה בהרבה. נראה זאת הפוך: NVIDIA מפרסמת ל-H100 BF16 Tensor Core את המספר 989 TFLOP/s. נאמת שהוא עקבי ושה-FMA=2 שם:

FLOP per cycle (total)  = 989e12 / 1.98e9      ~= 499,500 FLOP/cycle
FLOP per SM per cycle   = 499,500 / 132        ~= 3,784 FLOP/SM/cycle
FMA per SM per cycle    = 3,784 / 2            ~= 1,892 BF16 FMA/SM/cycle

כלומר כל SM מבצע כ-1,892 פעולות FMA של BF16 בכל מחזור על ארבעת ה-Tensor Cores שלו (שיעור 1.5) - עשרות מונים יותר מ-128 ה-FP32 FMA שה-CUDA Cores מבצעים. זה בדיוק היחס בן ~15 שראינו בטבלה, וזו הסיבה שכל kernel שמכבד את עצמו רוצה לרוץ על ה-Tensor Cores.

וקריטי: קצב השיא הוא תיאורטי ומושג לעתים נדירות. הוא מניח אפס עצירות. ה-kernel האמיתי נופל ממנו בגלל לחץ אוגרים, מגבלת רוחב פס הזיכרון, ומחסומי סנכרון - בדיוק שלושת החשודים שהגלוסרי מונה. כמה הוא נופל? זה מה שנמדוד בסעיף הביצוע-בפועל.

נקודת המפנה וטבלת ה-Roofline - ridge point

עכשיו ששני רוחבי הפס מוגדרים במספרים, אפשר לחבר אותם. נקודת המפנה (ridge point), שהגדרנו ב-7.2 כמאזן-המכונה (machine balance), היא היחס ביניהם:

ridge point [FLOP/byte] = arithmetic bandwidth [FLOP/s] / memory bandwidth [byte/s]

היחידות מצטמצמות יפה: (FLOP/s) חלקי (בייט/s) = FLOP/בייט - בדיוק היחידה של ה-arithmetic intensity על ציר ה-x. נקודת המפנה היא ה-arithmetic intensity הקריטי: kernel משמאלה (arithmetic intensity נמוכה יותר) חסום-זיכרון (memory-bound), kernel מימינה חסום-חישוב (compute-bound). זו הנקודה שבה שני הגגות נפגשים.

נחשב ל-H100 BF16:

989 TFLOP/s / 3.35 TB/s = 989e12 / 3.35e12 = 295 FLOP/byte

כלומר על H100 עם Tensor Cores ב-BF16, kernel חייב לבצע לפחות 295 פעולות ממשיות לכל בית שהוא מושך מ-HBM כדי בכלל להישאר compute-bound. זה מספר עצום. הטבלה המלאה - זו שראינו ב-7.2, עכשיו עם המקורות שלה גלויים:

מערכת	רוחב פס אריתמטי (TFLOP/s)	רוחב פס זיכרון (TB/s)	נקודת מפנה (FLOP/בייט)
A100 80GB BF16 TC	312	2	156
H100 SXM BF16 TC	989	3.35	295
B200 BF16 TC	2250	8	281
H100 SXM FP8 TC	1979	3.35	~591
B200 FP8 TC	4500	8	562
B200 FP4 TC	9000	8	1125

הסתכלו על העמודה הימנית מלמעלה למטה. נקודת המפנה מטפסת בפראות מ-156 (A100 BF16) עד 1125 (B200 FP4). למה? כי רוחב הפס האריתמטי גדל הרבה יותר מהר מרוחב פס הזיכרון בין הדורות, ובתוך דור - ככל שיורדים בדיוק (BF16 -> FP8 -> FP4). המשמעות מטרידה: ככל שהחומרה חדשה ומהירה יותר, כך קשה יותר ל-kernel להישאר compute-bound. כדי לנצל את ה-9000 TFLOP/s של B200 FP4 צריך 1125 פעולות לכל בית - ורוב ה-kernels בעולם האמיתי (שכבות רשת, אלמנט-אחר-אלמנט) לא מתקרבים לזה. הם נופלים משמאל לנקודת המפנה ונשארים memory-bound. זו מגבלת הביצועים המרכזית של החומרה המודרנית, ולכן חצי מהקורס עוסק בהעלאת ה-arithmetic intensity (tiling, fusion, שימוש חוזר ב-shared memory).

ציור ה-Roofline המלא של H100 BF16, עם המספרים מהשיעור הזה:

  performance [TFLOP/s]
  989 |············································_________________  <- compute roof = arithmetic bandwidth (989)
      |                                     ,-''
      |                                 ,-''
      |                             ,-''  <- memory roof slope = memory bandwidth (3.35 TB/s)
      |                         ,-''
      |                     ,-''
      |                 ,-''
      |             ,-''
      |         ,-''      |
      |     ,-''          |
      |_,-''______________|__________________________________________ arithmetic intensity
      0                  295                                        [FLOP/byte]
                    ridge point
              memory-bound  |  compute-bound

מדידת הביצוע בפועל - achieved bandwidth ו-achieved FLOP/s

כל המספרים עד כה היו תיאורטיים - חושבו מהמפרט. עכשיו נצא למעבדה ונמדוד את מה שהושג בפועל (achieved), כי היחס בין השניים - achieved / peak - הוא כל מה שחשוב. נתחיל ברוחב הפס.

הדרך הקנונית למדוד רוחב פס מושג היא kernel בסגנון STREAM - העתקה פשוטה מ-device ל-device, שלא עושה שום חישוב ולכן memory-bound מובהק, וגורף כמעט את כל רוחב הפס:

// bwcopy.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)

// device-to-device copy with a grid-stride loop
__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;
    // warmup, so we don't measure first-launch overhead
    copyKernel<<<grid, block>>>(d_in, d_out, N);
    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;        // one read + one write
    double gbps      = movedBytes / sec / 1e9;
    printf("achieved: %.0f GB/s  (%.1f%% of 3350 GB/s HBM3 peak)\n",
           gbps, gbps / 3350.0 * 100.0);

    CUDA_CHECK(cudaFree(d_in));
    CUDA_CHECK(cudaFree(d_out));
    return 0;
}

שתי נקודות שקובעות אם המדידה נכונה. ראשית, ספירת הבייטים: ההעתקה out[i] = in[i] נוגעת בכל אלמנט פעמיים - קריאה אחת מ-in וכתיבה אחת ל-out - ולכן הבייטים שזזו הם 2 x N x 4, לא N x 4. שכחה של ה-2 חוצה את התוצאה. שנית, החימום (warmup): ההפעלה הראשונה של kernel כוללת overhead JIT/הקצאה; מודדים תמיד אחרי חימום ועל פני הרבה איטרציות. הפלט האופייני על H100:

$ nvcc -arch=sm_90 -O3 bwcopy.cu -o bwcopy
$ ./bwcopy
achieved: 3040 GB/s  (90.7% of 3350 GB/s HBM3 peak)

כ-91% מהשיא - וזה מצוין. kernel memory-bound שמגיע ל-85%-92% מרוחב הפס התיאורטי הוא כמעט אופטימלי; אף פעם לא מגיעים ל-100% בגלל overheads רענון ה-DRAM, מעברי שורות, ותחילת/סיום. 91% הוא ה"speed of light" המעשי של רוחב הפס.

באותה רוח מודדים FLOP/s מושג: מריצים kernel עתיר-FMA, סופרים כמה FLOP הוא ביצע (מספר ה-FMA x 2), מחלקים בזמן, ומחלקים ב-66.9 TFLOPS. kernel חישובי טהור (בלי גישות זיכרון בלולאה הפנימית) יכול להגיע ל-90%+ מ-66.9; kernel אמיתי מעורב יגיע להרבה פחות.

מדידה מקצועית עם ncu. במקום לספור ידנית, ה-Nsight Compute מודד את שני האחוזים ישירות מונים חומרתיים. השאילתה הקנונית ל"אחוז מהשיא" של שני הגגות:

ncu --metrics \
  gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\
  sm__throughput.avg.pct_of_peak_sustained_elapsed \
  ./bwcopy

המדד הראשון הוא אחוז מרוחב פס הזיכרון (הגג המשופע) והשני הוא אחוז מגג החישוב. עבור copyKernel נצפה לראות את הראשון סביב 90% והשני נמוך מאוד - החתימה המובהקת של kernel memory-bound. עבור kernel עתיר-FMA היחס יתהפך. שני המדדים האלה, יחד, הם הקריאה הישירה ביותר של מיקום ה-kernel על ה-Roofline של השיעור הזה.

מלכודות ותפיסות שגויות - gotchas

• "ל-GPU יש מספר שיא אחד" - שגוי. יש רוחב פס אריתמטי לכל (יחידה x דיוק): FP32 CUDA (66.9), BF16 Tensor (989), FP8 Tensor (1979)... תמיד ציינו איזה. אותו כלל תקף לרוחב פס זיכרון - יש אחד לכל זוג שכבות; הרלוונטי ל-Roofline הוא HBM<->אוגרים.
• שכחת FMA=2. הטעות הכי נפוצה בחישוב שיא. FMA היא שתי פעולות ממשיות; ספירתה כאחת חוצה את קצב השיא. תמיד x2.
• בלבול TB עשרוני מול GiB בינארי. רוחב פס נמדד ב-TB/s עשרוני (10^12), קיבולת ב-GiB בינארי (2^30). אל תערבבו - זה מזייף אחוזים ב-~7%.
• ספירת בייטים שגויה במדידת רוחב פס. העתקה נוגעת בכל אלמנט פעמיים (קריאה + כתיבה). ל-STREAM triad (a=b+s*c) זה שלושה מערכים. ספרו נכון או קבלו אחוז שקרי.
• "קצב השיא ניתן להשגה" - כמעט אף פעם לא. הוא תיאורטי ומניח אפס עצירות. ה-91% שמדדנו הוא כבר מצוין; מי שמדווח "100% מהשיא" כנראה ספר בייטים לא נכון.
• "תפוסה גבוהה = ניצול גבוה" - לא בהכרח. נראה בשיעור 7.4 ש-kernel יכול לרוץ בתפוסה חד-ספרתית ועדיין להרוות את ה-Tensor Cores. הניצול (achieved/peak) הוא המדד האמיתי, לא התפוסה.

סיכום

• רוחב פס זיכרון הוא הקצב המרבי (בייט/s) בין שתי רמות בהיררכיה; ל-GPU יש רבים, אך הרלוונטי ל-Roofline הוא זה שבין ה-HBM לקבצי האוגרים (כי מערכי העבודה יושבים ב-HBM), והוא קובע את שיפוע גג הזיכרון: H100 3.35 TB/s (HBM3), A100 2 (HBM2e), B200 8 (HBM3e).
• רוחב פס אריתמטי הוא קצב השיא של פעולות חישוב (FLOP/s) וקובע את גובה גג החישוב; יש אחד לכל (יחידת ביצוע x דיוק): Tensor Core מול CUDA Core ~100:1, FP > INT, וכל חצייה של דיוק בערך מכפילה את הקצב (H100 BF16 989, FP8 1979; B200 FP4 9000).
• קצב השיא ("מהירות האור") מחושב מהמפרט הקבוע בהנחת יעילות מושלמת, והוא המכנה בכל מדד ניצול; שיא ה-FP32 של H100: 132 SM x 128 FP32 x 2 FLOP/FMA x 1.98 GHz = 66.9 TFLOPS.
• זכרו FMA=2: הפעולה a*b+c נספרת כשתי פעולות ממשיות; שכחת הגורם הזה חוצה כל חישוב שיא.
• נקודת המפנה = רוחב פס אריתמטי / רוחב פס זיכרון [FLOP/בייט]; היא ה-arithmetic intensity הקריטי שמפריד memory-bound (משמאל) מ-compute-bound (מימין), והיא מטפסת מ-156 (A100 BF16) עד 1125 (B200 FP4).
• הפער בין הגגות הולך וגדל בין דורות ובדיוקים נמוכים - ולכן קשה יותר ויותר להישאר compute-bound על חומרה חדשה, וזו המוטיבציה להעלאת arithmetic intensity לאורך הקורס.
• הביצוע בפועל נמדד עם kernel בסגנון STREAM (רוחב פס: 2 x N x sizeof בייט חלקי זמן, אחוז מ-3.35 TB/s) או kernel עתיר-FMA (FLOP/s, אחוז מ-66.9); ncu מדווח את שני האחוזים ישירות דרך dram_throughput ו-sm__throughput ב-pct_of_peak_sustained_elapsed.
• קצב השיא כמעט אף פעם לא מושג: 85%-92% מרוחב הפס הוא כבר מצוין; הפער נובע מלחץ אוגרים, מגבלת רוחב פס, ומחסומי סנכרון - בדיוק שלושת החשודים שנעמיק בהם בשיעור 7.4 ובפרויקטים.