3.1 היררכיית הזיכרון הלוגית הרצאה

בפרק 1 הכרנו את החומרה מלמטה: בשיעור 1.6 פירקנו את מדרג הזיכרון הפיזי של ה-GPU - קובץ האוגרים (register file) המהיר שבתוך כל SM, cache ה-L1 שחולק את אותו SRAM עם ה-shared memory (ב-H100 - 256 KiB לכל SM), ובתחתית ה-GPU RAM האיטי (ב-H100 - 80 GiB של HBM3). בפרק 2 בנינו את היררכיית ה-threads מלמעלה: thread בודד, שנארז ב-block, שנארז ב-grid (רשת). השיעור הזה פותח את פרק הזיכרון ומחבר את שני הקצוות: הוא מציג את היררכיית הזיכרון הלוגית של מודל התכנות של CUDA, ומראה שהיא לא שרירותית אלא מראָה מדויקת של היררכיית ה-threads. לכל רמה בהיררכיית ה-threads יש בלוק זיכרון משלה, שמשותף לכל הthreads באותה רמה: register לכל thread, shared memory לכל block, global memory לכל grid. הרעיון המכריע שילווה אותנו בכל הפרק הוא שהזיכרון הזה מנוהל בידי המתכנת (scratchpad), לא בידי החומרה - בשונה מהותי מcache ה-CPU השקוף. בשיעור זה נבנה את המפה הלוגית; בשיעורים 3.2 עד 3.4 נצלול לכל רמה בקוד רץ.

עיקרון-העל - זיכרון שמשקף את היררכיית ה-threads¶

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

   thread hierarchy               memory hierarchy
   ─────────────────             ─────────────────
   thread            ───────►    registers      (private to thread)
      │                             │
   block (CTA)       ───────►    shared memory  (shared by block)
      │                             │
   grid              ───────►    global memory  (shared by the whole grid)

ככל שעולים בהיררכיה, הקבוצה גדֵלה, ההיקף (scope) של הזיכרון מתרחב, הקיבולת גדֵלה - והמהירות יורדת. register פרטי ל-thread בודד והוא המהיר ביותר; shared memory משותף ל-block שלם ואיטי יותר; global memory משותף לכל ה-grid, ענק ואיטי בהרבה. שימו לב שההכלה היא חד-כיוונית וברורה: מה ש-thread רואה ב-register שלו איש מלבדו לא רואה; מה ש-block כותב ב-shared memory שלו רואים רק חבריו ל-block; ורק את ה-global memory רואים כולם.

ההבדל העמוק ביותר מ-CPU הוא לא במבנה אלא בניהול. בcache של CPU החומרה מחליטה, באופן שקוף ואוטומטי, מה נשאר קרוב ומה מפונה - המתכנת לא נוגע בזה. ב-GPU, לעומת זאת, ה-shared memory הוא scratchpad מפורש: אתם, המתכנתים, מחליטים אילו נתונים לטעון אליו, מתי, ומתי לפַנות. כפי שראינו בשיעור 1.6, אפילו בשפה עילית כמו CUDA C האזור הזה של L1/shared הוא "mostly programmer-managed". זו לא נוחות - זו אחריות: המהירות של kernel טוב תלויה כמעט כולה במידה שבה ניצלתם נכון את ה-scratchpad הזה.

רגיסטרים - Registers (פרטיים ל-thread)¶

בתחתית ההיררכיה, המהירים והקטנים ביותר, נמצאים ה-registers (רגיסטרים, אוגרים). הם מאחסנים את המידע שבו מטפל thread בודד, ופיזית הם יושבים בקובץ האוגרים - register file של ה-SM. כפי שראינו בשיעור 1.6, קובץ האוגרים מהיר בערך בסדר גודל (פי 10) מcache ה-L1, ובנוי מ-SRAM שמסוגל "לעמוד בקצב" של ליבות החישוב, כך שהוא לא צוואר בקבוק לחשבון עצמו.

בקוד CUDA C רגיל, ה-registers כמעט בלתי-נראים. משתנה סקלרי מקומי בתוך kernel - int i, float sum - יושב בדרך כלל ב-register, בלי שתכתבו על כך מילה:

__global__ void demo(const float* g, float* out, int n) {
    int   i   = blockIdx.x * blockDim.x + threadIdx.x;  // register
    float acc = 0.0f;                                   // register
    for (int k = 0; k < 4; ++k) {
        acc += g[i * 4 + k];   // g[...] comes from global; acc accumulates in register
    }
    if (i < n) out[i] = acc;   // write back to global
}

הרגיסטרים נהיים מפורשים רק כשיורדים מתחת ל-CUDA C. כפי שראינו בשיעור על שרשרת הכלים, ב-PTX (הייצוג הווירטואלי) הם מופיעים כרגיסטרים וירטואליים בני 32 סיביות (למשל הצהרה כמו .reg .f32 %f<7>; שמכריזה על שבעה אוגרי נקודה-צפה), וב-SASS (אסמבלר זורם, קוד המכונה של ה-GPU) הם הופכים לאוגרים פיזיים - R0, R7 וכן הלאה. את ההקצאה הפיזית עושה ה-ptxas, הcompiler שמתרגם PTX ל-SASS.

כאן נכנס אחד המתחים המרכזיים של כל הקורס - registers מול occupancy (תפוסה). קובץ האוגרים של כל SM הוא משאב סופי (ב-H100, כ-65,536 אוגרים בני 32 סיביות לכל SM, ולכל היותר 255 אוגרים ל-thread). ככל שכל thread צורך יותר registers, כך פחות threads יכולים לשבת (resident) על ה-SM בו-זמנית, כי צריך לחלק ביניהם את אותו קובץ אוגרים. פחות warps תושבים פירושו פחות יכולת להסתיר latency (השהיה) - יש פחות warps להחליף אליהם כשאחד תקוע בהמתנה לזיכרון. לכן היעד של הcompiler הוא דווקא להגביל את מספר האוגרים ל-thread כדי שיותר blocks ייכנסו ל-SM. הכלל הלא-אינטואיטיבי: יותר registers ל-thread לא שווה מהיר יותר; לעיתים דווקא הגבלת האוגרים מעלה את התפוקה הכוללת.

החריג ל"רגיסטר הוא פרטי ל-thread". ברירת המחדל של CUDA היא שאוגר שייך ל-thread אחד בלבד. אבל על חומרה מודרנית זה לא כלל מוחלט: פקודות ה-MMA (matrix-multiply-accumulate) של ה-Tensor Cores (ליבות טנזור) חולקות את אוגרי הקלט והפלט שלהן בין כל 32 הthreads של ה-warp. כפי שראינו בפרק על ה-Tensor Cores, פקודת wmma יחידה מבוצעת בשיתוף בידי warp שלם, והtile של המטריצה פרוס על פני האוגרים של 32 הthreads. אז "register פרטי ל-thread" הוא המצב הרגיל - אבל לא אמת מוחלטת בכל מסלולי החומרה.

זיכרון משותף - Shared Memory (משותף ל-block)¶

הרמה שמעל הרגיסטרים היא ה-shared memory (זיכרון משותף), והיא מתאימה בדיוק לרמת ה-block בהיררכיית ה-threads. הוא קטן בהרבה מ-global memory אבל מהיר בהרבה ממנו - גם ב-throughput (תפוקה) וגם ב-latency. פיזית, כפי שהדגשנו בשיעור 1.6, ה-shared memory יושב באותו SRAM של cache ה-L1 בתוך כל SM (ב-H100 - 256 KiB לכל SM, שמתחלקים בין L1 ל-shared). ההיקף שלו הוא ה-block: כל הthreads של אותו block - מה שנקרא רשמית cooperative thread array (CTA) - רואים אותו וחולקים אותו, ואף thread מ-block אחר לא נוגע בו.

מכריזים על shared memory במילת המפתח __shared__:

__global__ void stage(const float* g, float* out) {
    __shared__ float tile[256];      // 256 floats = 1 KiB, shared by the block

    int t = threadIdx.x;
    tile[t] = g[blockIdx.x * blockDim.x + t];  // 1. load from global to shared
    __syncthreads();                            // 2. barrier: everyone finished loading

    float s = 0.0f;                             // 3. compute on shared
    for (int k = 0; k < blockDim.x; ++k) s += tile[k];

    if (t == 0) out[blockIdx.x] = s;            // 4. write back to global
}

התבנית הזו היא לב הליבה של כתיבת kernels מהירים, וכדאי לחקוק אותה: global -> shared -> חישוב -> global. טוענים בלוק נתונים פעם אחת מה-global האיטי אל ה-shared המהיר, מבצעים עליו כמה שיותר פעולות אריתמטיות בזמן שהוא "חם" ב-SRAM, ורק בסוף כותבים תוצאה חזרה ל-global. הניסוח של הגלוסרי חד: הניהול הקפדני של ה-scratchpad הזה - טעינת נתונים כך שיאפשרו את המספר המרבי של פעולות אריתמטיות לפני שצריך לטעון נתונים חדשים - הוא "המפתח לאמנות של תכנון kernels בעלי ביצועים גבוהים ב-CUDA".

הסיבה שהתבנית עובדת היא כלכלית. גישה ל-global עולה יקר; גישה ל-shared זולה בסדרי גודל. אם טענו איבר פעם אחת ל-shared ואז השתמשנו בו 20 פעם, "פרסנו" (amortized) את עלות הטעינה האיטית על פני 20 פעולות מהירות. המדד לזה - arithmetic intensity (עצימות אריתמטית), יחס הפעולות האריתמטיות לבתים שנטענו - הוא בדיוק מה שה-shared memory בא להגדיל. את המסגרת הפורמלית (מודל ה-Roofline) נפתח בפרק על ביצועים; כאן מספיק להבין את התפקיד: ה-shared memory הוא שכבת הביניים המבוימת (staging tier) שממירה תעבורת global איטית לחישוב מהיר.

זיכרון גלובלי - Global Memory (משותף ל-grid)¶

בקצה העליון של ההיררכיה נמצא ה-global memory (זיכרון גלובלי). ההגדרה של הגלוסרי מדויקת ושווה שינון: הוא גלובלי גם בהיקפו וגם בזמן-חייו. בהיקף - הוא נגיש לכל thread בכל ה-grid. בזמן-חיים - הוא חי כל עוד התוכנית רצה, ולא נעלם כש-kernel מסתיים. פיזית הוא ה-GPU RAM (ב-H100 - 80 GiB של HBM3), התחתית של המדרג הפיזי: הכי גדול, הכי איטי, והכי רחב-היקף. את ה-global memory מקצים מהhost (ה-CPU) דרך מקצי הזיכרון של ה-CUDA Runtime API או ה-Driver API - זהו בדיוק ה-cudaMalloc שהכרנו בשיעור 0.3.

מכיוון שכל הthreads בכל ה-grid ניגשים אליו, ה-global memory מצריך תיאום (synchronization), וכאן שני מנגנונים נפרדים לחלוטין שאסור לבלבל ביניהם:

פעולות אטומיות - atomics. פעולה אטומית (כמו atomicAdd) מבטיחה שקריאה-שינוי-כתיבה על כתובת ב-global תתבצע ללא הפרעה. אטומיים עובדים בין כל הthreads ב-grid, בדיוק כמו זיכרון משותף בין ליבות CPU. זהו המנגנון היחיד לתיאום בטוח בין blocks שונים.
מחסומים - barriers. מחסום כמו __syncthreads() מסנכרן threads, אבל רק בתוך block יחיד (CTA). הוא ממתין עד שכל threads ה-block הגיעו לנקודה, ואז משחרר את כולם. מחסום כזה אינו מסנכרן בין blocks שונים.

וכאן האזהרה הקריטית ביותר של השיעור: סדר הביצוע בין blocks אינו דטרמיניסטי (indeterminate). החומרה מתזמנת blocks ל-SMs בסדר שאינו מובטח, ולא כל ה-blocks בהכרח תושבים בו-זמנית - block עשוי להמתין בתור עד ש-block אחר יסיים. מכאן נובע כלל ברזל:

   legal:      every block updates a shared global counter with atomicAdd,
               without assuming anything about which one runs before which.

   forbidden:  block 0 spin-waits until block 1 writes a flag to global.
               If block 1 hasn't been scheduled to an SM yet (or is
               waiting for block 0 to free up space), this is a deadlock.

ההשלכה המעשית: אי אפשר לסנכרן את כל ה-grid מתוך kernel רגיל. אם אתם צריכים ש"כל ה-blocks יסיימו שלב א' לפני שמישהו יתחיל שלב ב'", הדרך הבטוחה היא לסיים את ה-kernel ולהשיק kernel שני - גבול ה-kernel הוא נקודת סנכרון גלובלית טבעית. (קיים גם מנגנון grid.sync() של cooperative groups, אבל הוא דורש launch מיוחדת ותנאים שמבטיחים שכל ה-blocks תושבים; נעסוק בו בפרק מתקדם.) את התיאום התוך-block, לעומת זאת, עושים בזול עם __syncthreads().

זיכרון מקומי - Local Memory (המלכוד של ה-scope)¶

עכשיו למלכוד שמפיל מתחילים רבים. לצד register, shared ו-global קיים מרחב רביעי שנקרא local memory (זיכרון מקומי) - והשם שלו מטעה לחלוטין. "local" כאן הוא scope (היקף), לא שכבת מהירות. local memory הוא פרטי ל-thread בדיוק כמו register, אבל פיזית הוא יושב ב-GPU RAM האיטי - אותו DRAM/HBM שבו יושב ה-global memory. במילים אחרות: "local" נשמע קרוב ומהיר, אבל בפועל הוא רחוק ואיטי כמו global.

מתי משתנה של thread נופל ל-local memory במקום ל-register? בשלושה מקרים עיקריים:

גלישת רגיסטרים - register spilling. כש-thread צריך יותר registers ממה שזמין בקובץ האוגרים, הcompiler "שופך" (spills) את העודף ל-local memory שב-GPU RAM. כפי שראינו בשיעור 1.6, הגלישה הזו מגיעה עם "קנס ביצועים משמעותי".
מערכים מקומיים באינדוקס דינמי. מערך אוטומטי בתוך kernel שניגשים אליו באינדקס שאינו ידוע בזמן קומפילציה לא יכול לשבת באוגרים (לאוגרים אין כתובות שאפשר לחשב בזמן ריצה), ולכן הוא נוחת ב-local memory.
מבנים או מערכים גדולים מדי מכדי להיכנס לתקציב האוגרים.

הנה דוגמה שכמעט מבטיחה גלישה - מערך גדול באינדוקס דינמי:

__global__ void spills(const int* idx, float* out, int n) {
    float scratch[64];               // large; likely to spill to local memory
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    for (int k = 0; k < 64; ++k) scratch[k] = 0.0f;
    scratch[idx[i] % 64] += 1.0f;    // dynamic index -> not possible in registers
    out[i] = scratch[idx[i] % 64];
}

הסכנה בגלישה היא שהיא שקטה: ה-kernel נכון לחלוטין, פשוט איטי, כי כל גישה ל-scratch נעשית עכשיו מול DRAM ולא מול אוגר. בתרגול נחשוף גלישה כזו במפורש - נספור אוגרים עם nvcc --ptxas-options=-v ונראה שורת lmem (local memory) גדֵלה. שימו לב שאין מילת מפתח __local__ שכותבים ידנית: local memory תמיד נוצר בעקיפין, בהחלטת הcompiler. הלקח: "local" הוא scope פרטי ל-thread, לא הבטחת מהירות.

זיכרון קבוע - Constant Memory (לקריאה בלבד)¶

מרחב חמישי, שנעסוק בו לעומק בשיעור 3.4, הוא constant memory (זיכרון קבוע). גם הוא בהיקף ה-grid כולו, אבל עם שלוש תכונות ייחודיות: הוא לקריאה בלבד מנקודת מבט ה-kernel (הhost כותב אליו לפני הlaunch), הוא ממוטמן (cached) בcache constant ייעודי בכל SM, והוא ממוטב לשידור (broadcast). המשמעות של broadcast: כשכל 32 הthreads של warp קוראים את אותה כתובת בדיוק, החומרה משרתת את כולם בקריאה אחת, בעלות של גישה יחידה. מכריזים עליו עם __constant__, וגודלו מוגבל (כ-64 KB בסך הכל). זה הופך אותו מושלם לפרמטרים שכל הthreads קוראים אותם יחד - מקדמים, קבועי מסנן, מטריצות טרנספורמציה. את הפרטים, כולל מה קורה כשthreads באותו warp קוראים כתובות שונות (אז השידור מתפצל וההטבה נעלמת), נראה בשיעור 3.4.

מלכוד ה-global - הבלבול המפורסם¶

לפני הטבלה המסכמת, צריך לפרק פעם נוספת את הבלבול שהזכרנו כבר בשיעור 0.3, כי הוא חוזר ומכשיל: למילת המפתח __global__ אין שום קשר ל-global memory.

__global__ היא מסמן (qualifier) של פונקציה. היא מסמנת kernel - פונקציה שמושקת מהhost ומורצת על הdevice. זה שם של סוג פונקציה.
global memory הוא מרחב זיכרון - ה-DRAM שיושב על ה-GPU, ה-80 GiB של ה-H100. זה שם של מקום.

השניים חולקים את המילה "global" במקרה גמור. אחד ממעצבי CUDA המוקדמים אמר בבדיחות הדעת שהשם נבחר "כדי למקסם בלבול אצל המפתחים". שימו לב שאף לא אחד ממסמני הפונקציות של CUDA מתאר מרחב זיכרון: __global__ = kernel, __device__ = פונקציה שרצה על הdevice ונקראת מהdevice, __host__ = פונקציה של הhost. לעומתם, מסמני המשתנים הם אלה שמתארים זיכרון: __shared__, __constant__, __device__ (על משתנה גלובלי - global scope ב-global memory). זכרו את ההפרדה: מסמן על פונקציה מדבר על היכן הפונקציה רצה; מסמן על משתנה מדבר על היכן הזיכרון יושב.

טבלת מרחבי הזיכרון - the memory spaces¶

הנה כל ההיררכיה הלוגית בטבלה אחת. זו מפת הדרכים של כל פרק הזיכרון - חזרו אליה בכל פעם שאתם לא בטוחים היכן משתנה חי:

מרחב	היקף (scope)	זמן-חיים	מיקום פיזי	מהירות	מילת מפתח
register	thread בודד	ה-thread	register file ב-SM	הכי מהיר	(אין; משתנה מקומי)
local	thread בודד	ה-thread	GPU RAM	איטי (כמו global)	(אין; גלישה/מערך)
shared	block (CTA)	ה-block	L1/SRAM ב-SM	מהיר מאוד	`__shared__`
constant	grid כולו	התוכנית	GPU RAM + cache constant	מהיר בקריאת-שידור	`__constant__`
global	grid כולו	התוכנית	GPU RAM (HBM)	איטי, גדול	`__device__` / `cudaMalloc`

שתי שורות שמלמדות את כל הסיפור: register ו-local חולקים היקף זהה (thread בודד) אבל מיקום פיזי הפוך - האחד ה-SRAM המהיר ביותר, השני ה-DRAM האיטי ביותר; המרחק ביניהם הוא כל ההבדל בין kernel מהיר ל-kernel שגולש. וshared ו-global חולקים את אותה מטרה - שיתוף נתונים - בשתי רמות היררכיה שונות: shared משתף בזול בתוך block, global משתף ביוקר על פני grid, וה-scratchpad שביניהם הוא מה שממיר את היוקר לזול.

מפה זו לוגית בלבד. את המספרים הפיזיים המדויקים - קיבולת קובץ האוגרים, פיצול L1/shared, רוחב הפס של ה-HBM3 - ראינו בשיעור 1.6, והם שנותנים לכל שורה כאן את ה"מחיר" האמיתי שלה. בשיעורים 3.2 עד 3.4 נרד לכל אחד מהמרחבים האלה בקוד: 3.2 ל-global וניהול הזיכרון, 3.3 ל-shared ולתבנית ה-tiling, 3.4 ל-constant ולזיכרונות המיוחדים.

התמונה המלאה - מהלוגי לפיזי על H100¶

בואו נניח את המפה הלוגית של השיעור הזה על גבי המדרג הפיזי של שיעור 1.6, כדי לראות למה כל רמה "עולה" מה שהיא עולה. הנה ה-H100 מלמעלה למטה, עם המרחב הלוגי לצד המבנה הפיזי והמספרים הקונקרטיים:

   logical space     physical structure (H100)        relative speed
   ─────────────     ──────────────────────────       ──────────────
   register    ◄──   register file in every SM         fastest
                     (~65,536 32-bit registers/SM)     ("keeps pace with the cores")
                          │  ~10x faster than L1
   shared      ◄──   L1 / SRAM in every SM             very fast
                     (256 KiB per SM, 33 MiB aggregate) (~10x slower than the cores)
                          │
                     L2 cache (shared by all SMs)
                          │
   local  ─┐
   global ─┴──►      GPU RAM - HBM3                    slow, huge
   constant          (80 GiB aggregate)                (DRAM, ~3.35 TB/s bandwidth)

שלוש תובנות שהתמונה הזו מבהירה. ראשית, local ו-global יושבים באותו DRAM פיזי - זו הסיבה שגלישה ל-local כואבת כמו גישה ל-global, למרות שה-scope שלה פרטי ל-thread. שנית, shared ו-L1 חולקים את אותם 256 KiB של SRAM לכל SM - כשאתם מקצים יותר __shared__, נשאר פחות L1 אוטומטי, וזה tradeoff שנכוונן בשיעור 3.3. שלישית, ההפרשים אינם ליניאריים: קובץ האוגרים מהיר בערך פי 10 מ-L1, וה-L1 עצמו איטי בערך פי 10 מהליבות - כך שגישה ל-register מהירה בסדרי גודל מגישה ל-DRAM. כל תכנון kernel טוב הוא במהותו מאבק לדחוף כמה שיותר מהעבודה למעלה בפירמידה הזו: לשמור נתונים חמים באוגרים, לביים ב-shared, ולמזער את המגע עם ה-DRAM למטה.

סיכום¶

היררכיית הזיכרון הלוגית של CUDA היא מראָה של היררכיית ה-threads: register לכל thread, shared memory לכל block, global memory לכל grid; ככל שעולים, הקבוצה וההיקף גדֵלים והמהירות יורדת.
הזיכרון הזה מנוהל בידי המתכנת (scratchpad) ולא בידי החומרה - זה ההבדל המהותי מcache CPU שקוף, וזו האחריות שקובעת את ביצועי ה-kernel.
register פרטי ל-thread, יושב בקובץ האוגרים של ה-SM, הכי מהיר; יותר registers ל-thread שווה occupancy נמוכה יותר; החריג הוא MMA של Tensor Cores, שחולק אוגרים על פני ה-warp.
shared memory משותף ל-block, יושב פיזית ב-L1, והוא שכבת הביום; אמנות ה-kernel המהיר היא למקסם את מספר הפעולות האריתמטיות לכל טעינה אליו (התבנית global -> shared -> חישוב -> global).
global memory גלובלי בהיקף (כל ה-grid) ובזמן-חיים (כל התוכנית), יושב ב-GPU RAM; מתאמים אותו עם atomics (בין כל הthreads) ועם barriers (רק בתוך block); סדר הביצוע בין blocks אינו דטרמיניסטי.
מכיוון שאין סדר מובטח בין blocks, אסור ש-block ימתין ל-block אחר בתוך אותו kernel; סנכרון גלובלי אמיתי מושג על ידי סיום ה-kernel וlaunch של kernel חדש.
local memory הוא scope פרטי ל-thread, לא שכבה מהירה: פיזית הוא ב-GPU RAM האיטי ונוצר מגלישת registers או ממערכים מקומיים באינדוקס דינמי - והגלישה שקטה ויקרה.
constant memory לקריאה בלבד, ממוטמן וממוטב לשידור כשכל ה-warp קורא אותה כתובת (טיפול מלא ב-3.4).
__global__ מסמן kernel (פונקציה שמושקת מהhost ורצה על הdevice) ואין לו שום קשר ל-global memory; מסמני פונקציה מתארים היכן הקוד רץ, מסמני משתנה (__shared__, __constant__) מתארים היכן הזיכרון יושב.