1.6 היררכיית הזיכרון בחומרה הרצאה

בשיעורים הקודמים של פרק 1 פירקנו את ה-GPU (מעבד גרפי / כרטיס מסך) לרכיבי החישוב שלו: ראינו שה-SM (מעבד רב-זרמי - Streaming Multiprocessor) הוא המקבילה האמיתית ל-core של ה-CPU, הכרנו את ה-CUDA Cores (ליבות CUDA) ואת ה-Tensor Cores (ליבות טנזור), ואת מתזמני ה-warp (warp - קבוצת 32 threads) שמזינים אותם. אבל ליבה שאין לה מה לחשב היא רק סיליקון חם. השיעור הזה עוסק בצד השני של המשוואה - מהיכן מגיעים הביטים שהליבות מעבדות, ולאן חוזרות התוצאות. נצייר את פירמידת הזיכרון הפיזית של ה-GPU: מהמהיר והקטן ביותר (קובץ האוגרים) ועד לאיטי והגדול ביותר (זיכרון ה-GPU), עם המספרים המדויקים של ה-H100. זהו שיעור על החומרה - על הטרנזיסטורים ועל הפשרות הפיזיקליות שמכתיבות את המבנה. את המקבילה הלוגית שלו - כיצד מודל התכנות של CUDA חושף את ההיררכיה הזו דרך registers, shared memory ו-global memory - נלמד בפרק 3, ואת ההשלכות על ביצועים דרך מודל ה-Roofline נעמיק בפרק 7. כאן אנחנו בונים את היסודות שעליהם שני הפרקים האלה יישענו.

למה בכלל צריך היררכיה - SRAM מול DRAM¶

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

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

זיכרון סטטי - SRAM (Static RAM). תא בנוי מ-6 טרנזיסטורים (6T) ששומרים ביט בלולאת היזון חוזר יציבה. הוא מהיר מאוד ולא צריך רענון, אבל הוא גדול פיזית ויקר - ששה טרנזיסטורים לכל ביט בודד. מזה בנויים קובץ האוגרים והמטמונים.
זיכרון דינמי - DRAM (Dynamic RAM). תא בנוי מטרנזיסטור אחד וקבל אחד (1T1C) ששומר ביט כמטען חשמלי. הוא צפוף וזול - אפשר לדחוס מיליארדי תאים - אבל איטי, והמטען דולף כך שצריך לרענן אותו כל הזמן. מזה בנוי זיכרון ה-GPU.

הפשרה הזו היא כל הסיפור. הגלוסרי מנסח זאת בחדות: תאי ה-DRAM של זיכרון ה-GPU הם "איטיים יותר אך קטנים יותר מ-SRAM שבו משתמשים באוגרים ובזיכרון הcache". לא ניתן לקבל את שניהם. אילו בנינו 80 GiB מ-SRAM, השבב היה בגודל של שולחן ועולה כמו מטוס; אילו בנינו את קובץ האוגרים מ-DRAM, הליבות היו רעבות רוב הזמן וממתינות לרענון. הפתרון הוא היררכיה: שכבה זעירה וזריזה של SRAM צמוד לליבות, שכבות ביניים, ובתחתית ים עצום ואיטי של DRAM. ככל שמתרחקים מהליבות, כל מדרגה גדֵלה בסדר גודל ומאטה בסדר גודל. תפקיד המתכנת - שנלמד בפרק 3 ובפרק 8 - הוא לשמור נתונים "חמים" קרוב לפסגה כמה שיותר זמן.

הפסגה - קובץ האוגרים - register file¶

בקצה המהיר ביותר של הפירמידה יושב קובץ האוגרים (register file). הגלוסרי מגדיר אותו כ"מאגר הביטים העיקרי בין המניפולציות של הליבות עליהם" - כלומר המחסן שבו יושבים האופרנדים בזמן שהליבה מחשבת עליהם. זה בדיוק כמו האוגרים של ה-CPU, אבל גדול בהרבה ומשותף בין אלפי threads.

מאפייני הביצועים שלו:

מהירות: כ-סדר גודל (פי 10 בערך) מהיר יותר מ-L1 data cache. הוא בנוי מזיכרון מהיר במיוחד ש"יכול לעמוד בקצב של ליבות החישוב" - כלומר הוא אינו צוואר הבקבוק של ה-ALU. אופרנד באוגר זמין כמעט מיידית לפייפליין.
רוחב האוגר: קובץ האוגרים מאורגן כאוגרים ברוחב 32-bit.
הקצאה מחדש דינמית של טיפוסים: אותם אוגרים פיזיים "מוקצים מחדש דינמית בין טיפוסי נתונים שונים, כמו שלמים של 32-bit, מספרים ממשיים של 64-bit, וקבוצות של מספרים ממשיים ברוחב 16-bit או פחות". המשמעות המעשית: ערך של 64-bit (למשל double) תופס זוג אוגרים סמוכים; ערכים קטנים של 16-bit או פחות (למשל half) יכולים להיות ארוזים יחד באוגר אחד. אין "אוגרי double" נפרדים - יש מאגר אחיד של תאי 32-bit שמתפרש לפי הצורך.

הקשר ל-PTX ול-ptxas. האוגרים הפיזיים האלה הם ש"מגבים את האוגרים הווירטואליים בייצוג הביניים PTX (Parallel Thread eXecution)". כזכור מ-0.4, PTX חושף מספר בלתי-מוגבל של אוגרים וירטואליים; הcompiler ptxas, שממיר PTX ל-SASS, הוא זה שמבצע את ההקצאה בפועל - הוא ממפה את האוגרים הווירטואליים אל מספר סופי של אוגרים פיזיים, ומייעל את השימוש עבור ה-thread blocks. שורת ה-PTX .reg .f32 %f<7>; שראינו בשיעור 0.4 היא בדיוק הנחיה כזו: "ה-kernel צורך שבעה אוגרי 32-bit ממשיים".

המספרים של H100. הגלוסרי עצמו אינו נותן מספר לגודל קובץ האוגרים, אבל זו ידיעה סטנדרטית ומתועדת של NVIDIA: לכל SM ב-H100 יש 65,536 אוגרים של 32-bit, כלומר 256 KiB של קובץ אוגרים לכל SM (65,536 x 4 בייט). על פני 132 ה-SM של ה-H100 זה מצטבר ל-33 MiB של אוגרים. תקרת האוגרים לכל thread היא 255.

לחץ אוגרים ותפוסה - register pressure ו-occupancy. כאן מגיעה שרשרת הסיבתיות החשובה ביותר של השיעור. קובץ האוגרים הוא סופי, והוא מחולק בין כל ה-threads שדרים בו-זמנית על ה-SM. ה-H100 יכול להחזיק עד 2048 threads דרים (64 warps) לכל SM. נעשה את החשבון:

תקציב אוגרים ל-SM:  65,536 אוגרים
תפוסה מלאה:          2048 threads דרים

65,536 / 2048  =  32 אוגרים ל-thread  ->  תפוסה 100%

מכאן נובע ישירות:

אוגרים ל-thread	threads מקסימלי ל-SM	תפוסה - occupancy
32	2048	100%
64	1024	50%
128	512	25%
255	256	~12.5%

זהו לחץ האוגרים (register pressure): ככל שכל thread צורך יותר אוגרים, כך פחות warps יכולים לדור בו-זמנית על ה-SM, כי המאגר הסופי חייב להתחלק ביניהם. פחות warps דרים משמעו תפוסה נמוכה יותר, ומכאן יכולת מופחתת להסתיר latency - יש פחות warps שאפשר לעבור אליהם כשאחד תקוע בהמתנה לזיכרון. זו הסיבה ש-ptxas שואף להגביל את מספר האוגרים ל-thread: לפעמים פחות אוגרים = יותר blocks על ה-SM = תפוקה כוללת גבוהה יותר. נחזור לזה לעומק כשנדבר על occupancy בפרק 6.

המדרגה השנייה - L1 ו-shared memory¶

מתחת לאוגרים יושב אזור ה-L1 data cache, שהוא גם ה-shared memory (זיכרון משותף). הגלוסרי מתאר אותו כ"הזיכרון הפרטי של ה-SM": לכל SM יש משלו, והוא מחולק בין ה-thread blocks שמתוזמנים עליו.

מאפיינים טכניים:

מהירות: הוא "ממוקם יחד עם, ורק כ-סדר גודל איטי יותר מ, הרכיבים שמבצעים את החישובים (למשל ה-CUDA Cores)". כלומר: על ה-SM, פי 10 בערך איטי יותר מיחידות החישוב - ופי 10 בערך איטי יותר מקובץ האוגרים.
טכנולוגיה: הוא בנוי מ-SRAM - "אותו תא מוליך-למחצה בסיסי שמשמש במטמוני CPU ובאוגרים".
מנגנון הגישה: אליו ניגשים יחידות ה-LSU (Load/Store Units) של ה-SM - יחידות הטעינה והאחסון שהכרנו כשפירקנו את ה-SM.

המספרים של H100 (מדויק): לכל SM יש 256 KiB של L1/shared. על פני 132 ה-SM זה מצטבר ל-33 MiB (256 KiB x 132 = 33,792 KiB = 33 MiB). זהו אותו גודל בדיוק כמו קובץ האוגרים המצטבר - צירוף מקרים נחמד שכדאי לזכור: ב-H100 הן קובץ האוגרים והן ה-L1/shared הם 256 KiB ל-SM. מתוך 256 ה-KiB האלה, ב-Hopper אפשר להקצות עד 228 KiB ל-shared memory במפורש, והשאר ל-L1; החלוקה ניתנת להגדרה, ונרחיב עליה בפרק 3.

המסגור הקריטי - cache מנוהל-מתכנת מול מנוהל-חומרה. זו הנקודה החשובה ביותר בסעיף, וההבדל התפיסתי הגדול בין GPU ל-CPU. הגלוסרי מנסח זאת מפורשות: "ב-CPU הcache הזה מנוהל לחלוטין על ידי החומרה. ב-GPU הcache הזה מנוהל בעיקר על ידי המתכנת, אפילו בשפות עיליות כמו CUDA C". ב-CPU, ה-L1 הוא שקוף (transparent): אתם ניגשים לזיכרון, והחומרה מחליטה לבד מה לשמור בcache ומה לפנות - אין לכם שליטה ישירה. ב-GPU, אזור ה-shared memory הוא לוח שרטוט (scratchpad) שאתם מנהלים ביד: אתם מצהירים __shared__ float tile[256];, טוענים אליו נתונים במפורש מ-global memory, ומחליטים בעצמכם מה יושב שם ומתי. זו אינה החלטה של החומרה - זו החלטה שלכם. הדפוס הקנוני, שנלמד בפרק 3, הוא: טען נתח מ-global memory ל-shared memory, בצע עליו מקסימום חישובים, וכתוב חזרה. השליטה הידנית הזו היא "אמנות תכנון ה-kernels המהירים".

המדרגה השלישית - L2 cache¶

בין ה-L1 הפרטי-לכל-SM לבין זיכרון ה-GPU יושבת שכבה אחת נוספת שהגלוסרי מזכיר רק בקצרה: L2 cache. שווה להציג אותו כהקשר סטנדרטי, כי הוא חלק אינטגרלי מהתמונה.

שני מאפיינים מבחינים את ה-L2 מה-L1:

הוא משותף לכל ה-SM. בניגוד ל-L1 שהוא פרטי לכל SM, ה-L2 הוא device-wide: יחידה אחת גדולה שכל 132 ה-SM חולקים. הוא נקודת המפגש שדרכה SM אחד יכול לראות נתונים שכתב SM אחר.
הוא מנוהל-חומרה. בניגוד ל-L1/shared שהוא בעיקר מנוהל-מתכנת, ה-L2 חוזר להיות cache שקוף וקלאסי בסגנון CPU - החומרה מחליטה מה לשמור בו ומה לפנות. אתם לא מצהירים עליו ולא מנהלים אותו ידנית.

ב-H100 ה-L2 cache הוא בגודל 50 MB (ידיעה סטנדרטית של NVIDIA; הגלוסרי אינו נוקב במספר). כל גישה ל-global memory עוברת דרכו: קריאה שנמצאת ב-L2 (hit) חוזרת מהר בהרבה מקריאה שצריכה לרדת עד ל-DRAM. ה-L2 הוא לכן החיץ שממתן את האיטיות של זיכרון ה-GPU.

התחתית - זיכרון ה-GPU - GPU RAM¶

בתחתית הפירמידה נמצא זיכרון ה-GPU (GPU RAM), הידוע גם כ-VRAM (Video RAM) וכ-global memory (זיכרון גלובלי) בהקשר של מודל התכנות. הגלוסרי מגדיר אותו כ"זיכרון הרמה-התחתונה של ה-GPU... מאגר גדול (מגה-בייטים עד ג'יגה-בייטים) שניתן לכתובת על ידי כל ה-SM של ה-GPU". זו השכבה הגדולה ביותר, האיטית ביותר, והמשותפת ביותר.

טכנולוגיית התאים - DDR מול HBM. הגלוסרי מבחין בין שתי גישות:

כרטיסים צרכניים ו-CPU משתמשים ב-DDR (Double Data Rate) או ב-GDDR, שאינו יושב על אותה פיסת סיליקון (die) עם ה-SM - הוא שבבים נפרדים על גבי הלוח.
כרטיסי מרכזי-נתונים מודרניים כמו ה-H100 משתמשים ב-HBM (High-Bandwidth Memory). ה-HBM ממוקם על מטמון-ביניים משותף (interposer) יחד עם פיסת ה-GPU - ערימות זיכרון סמוכות פיזית לשבב, מחוברות באלפי threads קצרים. התוצאה: "latency מופחת ורוחב פס מוגבר" לעומת DDR שיושב הרחק על הלוח.

המספרים של H100. ה-H100 מאחסן 80 GiB בזיכרון ה-GPU שלו, בטכנולוגיית HBM3. רוחב הפס בין ה-HBM3 לשבב הוא כ-3.35 TB/s (ידיעה סטנדרטית של NVIDIA ל-H100 SXM; הגלוסרי אינו נוקב ברוחב פס). שימו לב לפער העצום מול קובץ האוגרים: ה-DRAM צפוף פי אלפי מונים (80 GiB מול 256 KiB ל-SM), אבל איטי בסדרי גודל.

שני תפקידים מרכזיים. זיכרון ה-GPU עושה שני דברים:

הוא מממש את ה-global memory של CUDA - מרחב הכתובות שכל ה-threads בכל ה-SM רואים. זה מה ש-cudaMalloc מקצה.
הוא סופג register spills (גלישת אוגרים). כשל-kernel דרושים יותר אוגרים ממה שקיים פיזית, האוגרים העודפים "גולשים מקובץ האוגרים" אל local memory (זיכרון מקומי) - שהוא, למרות השם, פשוט אזור מיוחד בזיכרון ה-GPU. זה קושר ישירות חזרה ללחץ האוגרים: כשמפעילים לחץ גבוה מדי, המחיר אינו רק תפוסה נמוכה - הוא גלישה, ואז אוגר "מקומי" עולה כמו גישה ל-DRAM (מאות מחזורים) במקום כמו אוגר (מיידי). זו הסיבה שהגלוסרי מדגיש שגלישה מתרחשת "בקנס ביצועים משמעותי".

הפירמידה המלאה של H100 - the full hierarchy¶

עכשיו נרכיב את הכל לתמונה אחת. מלמעלה (מהיר, קטן, קרוב לליבות) למטה (איטי, גדול, רחוק):

                     H100 physical memory hierarchy

           ^  fast / small / SRAM / close to the cores
          /|\
         / | \    +--------------------------------------------------+
        /  |  \   |  register file                                   |
       /   |   \  |  256 KiB per SM (65,536 x 32-bit), 33 MiB total  |
      /    |    \ |  ~1 cycle, "keeps up with the compute units"  [private per thread] |
     /     |     \+--------------------------------------------------+
    /      |      |  L1 data cache / shared memory  (SRAM)           |
   /       |      |  256 KiB per SM, 33 MiB total, LSU               |
  /        |      |  ~10x slower than the cores   [private per SM, programmer-managed] |
 /         |      +--------------------------------------------------+
/          |      |  L2 cache                       (SRAM)           |
           |      |  50 MB, device-wide, hardware-managed             |
           |      |  ~10-20x slower than L1   [shared by all SMs]    |
           |      +--------------------------------------------------+
          \|/     |  GPU RAM                        (DRAM / HBM3)    |
           v      |  80 GiB, ~3.35 TB/s, includes register spills    |
                  |  ~hundreds of cycles   [global, visible to all threads] |
  slow / large /  +--------------------------------------------------+
  DRAM / far

טבלת ה-latency היחסי (ערכים אופייניים במחזורי SM; הגלוסרי אינו נותן מספרי latency, אלה ערכים נמדדים סטנדרטיים, למשל מ-Luo et al. על Hopper):

שכבה	גודל (H100)	טכנולוגיה	latency אופייני	מנוהל על ידי
קובץ אוגרים	256 KiB / SM	SRAM (6T)	~1 מחזור	`ptxas` (בזמן compilation)
L1 / shared	256 KiB / SM	SRAM	~30 מחזורים	המתכנת (shared) / חומרה (L1)
L2 cache	50 MB	SRAM	~200 מחזורים	חומרה
GPU RAM (HBM3)	80 GiB	DRAM	~450-800 מחזורים	המתכנת (`cudaMalloc`)

שימו לב לשני הצירים שנעים יחד: כל מדרגה כלפי מטה גדֵלה בסדר גודל ומאטה בסדר גודל. זה הביטוי הישיר של פשרת ה-SRAM-מול-DRAM שפתחנו בה.

הנה קטע קוד קצר ששולף את המספרים האלה מהכרטיס עצמו, עם מאקרו בדיקת השגיאות שנלווה אלינו לאורך הקורס:

#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)

int main() {
    cudaDeviceProp p;
    CUDA_CHECK(cudaGetDeviceProperties(&p, 0));
    printf("SM count:              %d\n", p.multiProcessorCount);
    printf("Registers / SM:        %d (= %.0f KiB)\n",
           p.regsPerMultiprocessor, p.regsPerMultiprocessor * 4.0 / 1024.0);
    printf("Shared mem / SM (max): %zu KiB\n",
           p.sharedMemPerMultiprocessor / 1024);
    printf("L2 cache size:         %d MiB\n", p.l2CacheSize / (1024 * 1024));
    printf("Global (HBM) memory:   %.1f GiB\n",
           p.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
    return 0;
}

על H100 הפלט יראה 132 SM, 65,536 אוגרים ל-SM (256 KiB), shared memory עד 228 KiB ל-SM, L2 cache בגודל 50 MiB, וכ-79 GiB global memory. אלה בדיוק המספרים שבפירמידה.

המספר שקובע הכל - רוחב הפס וקשר ל-Roofline¶

מדוע כל זה חשוב כל כך? כי מבין כל המספרים בפירמידה, מספר אחד שולט על ביצועי רוב ה-kernels בעולם האמיתי: רוחב הפס בין זיכרון ה-GPU (התחתית) לבין קבצי האוגרים שב-SM (הפסגה). זהו הpipe שדרכו נתונים חייבים לזרום כדי שהליבות בכלל יוכלו לגעת בהם, וב-H100 הוא רחב "רק" 3.35 TB/s - מהיר במונחים אנושיים, אך איטי לאין ערוך מקצב הצריכה של הליבות.

נעשה חשבון גס להמחשה. ה-H100 מסוגל לכ-66.9 TFLOPS ב-FP32. כדי להאכיל את זה, הליבות "רעבות" ל-66.9e12 אופרנדים בשנייה בערך. אבל ה-HBM3 מספק רק 3.35e12 בייט לשנייה, כלומר כ-8.4e11 ערכי FP32 (4 בייט כל אחד) בשנייה. היחס:

קצב חישוב :  66.9e12 FLOP/s
קצב אספקה :  3.35e12 B/s  /  4 B  =  8.4e11 ערכי FP32/s

לכל ערך FP32 שנטען מ-HBM, החומרה יכולה לבצע בערך
   66.9e12 / 8.4e11  ~  80 פעולות ממשיות

המשמעות: אם ה-kernel שלכם מבצע פחות מ-~80 פעולות אריתמטיות לכל ערך שהוא טוען מ-global memory, הוא memory-bound - הליבות מבזבזות את רוב זמנן בהמתנה ל-HBM, ולא בחישוב. זו בדיוק הכמות שמכונה arithmetic intensity, וזו הצירים של מודל ה-Roofline שנבנה בפרק 7. כל ההיררכיה שציירנו קיימת כדי להילחם במגבלה הזו: shared memory וקבצי האוגרים מאפשרים לכם לטעון ערך פעם אחת מ-HBM האיטי ואז לחשב עליו שוב ושוב מ-SRAM המהיר, ובכך להעלות את ה-arithmetic intensity ולהתרחק מהצוואר. זו הסיבה שההיררכיה קיימת, וזו הסיבה שהיא הנושא של חצי מהקורס.

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

נסגור בכמה נקודות שמבלבלות תלמידים שוב ושוב:

"L1 ב-GPU עובד כמו L1 ב-CPU" - שגוי. ב-CPU כל שכבות הcache שקופות ומנוהלות-חומרה. ב-GPU, אזור ה-L1/shared הוא בעיקר מנוהל-מתכנת - אתם מציבים בו נתונים ביד. רק ה-L2 וה-DRAM חוזרים להתנהג כמו cache CPU קלאסי.
"local memory הוא מהיר כי הוא מקומי" - שגוי. ה-local memory שאליו גולשים אוגרים הוא בזיכרון ה-GPU (DRAM), לא על השבב. "מקומי" מתאר את ה-scope (פרטי ל-thread), לא את המיקום הפיזי. גישה אליו איטית כמו גישה ל-global memory.
גלישת אוגרים שקטה ויקרה. ה-kernel ימשיך לרוץ נכון גם אם אוגרים גולשים - הוא פשוט יהיה איטי בלי להתלונן. תמיד בדקו עם nvcc -Xptxas -v אם יש spill stores/spill loads.
קובץ האוגרים ו-L1 באותו גודל ב-H100 (256 KiB) אך אינם אותו דבר. הם שני מבנים פיזיים נפרדים על ה-SM, עם מנגנוני גישה שונים (הליבות ניגשות ישירות לאוגרים; ה-LSU ניגשות ל-L1).
latency אינו throughput. ה-latency הגבוה של HBM (מאות מחזורים) אינו אומר שה-GPU איטי - ה-GPU מסתיר את ה-latency הזה על ידי מעבר מהיר בין warps (בערך מחזור אחד להחלפה). לכן חשובה התפוסה: יותר warps דרים = יותר עבודה חלופית בזמן שאחד ממתין ל-HBM. הסתרת ה-latency הזו היא הדבק שמחבר את פירמידת הזיכרון לפירמידת החישוב.

סיכום¶

ההיררכיה קיימת בגלל פשרת SRAM מול DRAM: SRAM (אוגרים, cache) מהיר אך גדול-פיזית ויקר; DRAM (זיכרון ה-GPU) צפוף וזול אך איטי - אי אפשר לקבל את שניהם, ולכן בונים מדרגות.
קובץ האוגרים הוא הפסגה: פי 10 מהיר מ-L1, אוגרי 32-bit שמוקצים דינמית (ערך 64-bit = זוג, ערכי 16-bit נארזים), מגבים את האוגרים הווירטואליים של PTX, ומוקצים על ידי ptxas. ב-H100: 65,536 אוגרים (256 KiB) ל-SM.
לחץ אוגרים קובע תפוסה: 65,536/2048 = 32 אוגרים ל-thread לתפוסה מלאה; יותר אוגרים ל-thread = פחות warps דרים = תפוסה נמוכה יותר = פחות הסתרת latency.
L1 / shared memory הוא SRAM פרטי ל-SM, פי 10 איטי מהליבות, 256 KiB ל-SM (33 MiB מצטבר), נגיש דרך ה-LSU, ומחולק בין ה-blocks; המסגור המרכזי: ב-GPU הוא מנוהל-מתכנת (scratchpad ידני), שלא כמו cache CPU השקוף.
L2 cache הוא שכבת ביניים device-wide ומנוהלת-חומרה (50 MB ב-H100), היושבת בין ה-L1 לזיכרון ה-GPU.
זיכרון ה-GPU הוא התחתית: תאי DRAM (HBM3, 80 GiB, ~3.35 TB/s ב-H100), ממומש כ-global memory וסופג גם register spills אל local memory - שלמרות שמו יושב ב-DRAM ואיטי בהתאם.
כרטיסי מרכזי-נתונים משתמשים ב-HBM על interposer משותף (latency נמוך, רוחב פס גבוה) לעומת DDR/GDDR מחוץ ל-die בכרטיסים צרכניים.
רוחב הפס בין זיכרון ה-GPU לקבצי האוגרים הוא הצוואר ששולט ב-Roofline (פרק 7): kernel שמבצע פחות מ-~80 פעולות לכל ערך FP32 שנטען מ-HBM הוא memory-bound; ההיררכיה כולה קיימת כדי להעלות את ה-arithmetic intensity ולהתרחק מהצוואר.
את המקבילה הלוגית (registers / shared / global במודל התכנות) נלמד בפרק 3, את הסתרת ה-latency והתפוסה בפרק 6, ואת ה-coalescing ורוחב הפס בפרק 8 - כולם נשענים על הפירמידה של השיעור הזה.