2.2 היררכיית ה threads blocks ו grids הרצאה

בשיעור 2.1 עמדנו על התחתית של ההיררכיה: ה-thread הבודד, האטום של מודל התכנות, שמחזיק אוגרים משלו ולעיתים מונה-פקודות משלו, ואת ה-warp - קבוצת 32 threads שהחומרה מתזמנת יחד, פרט מימוש שאינו חלק רשמי מהמודל. עכשיו נטפס שתי רמות מעלה. מעל ה-thread יושב ה-block (בלוק) - קבוצת threads שמתוזמנת יחד על SM אחד ויכולה לשתף פעולה; ומעל הבלוקים יושב ה-grid (רשת) - כל הthreads של kernel launch אחד, פרוסים על פני כל ה-GPU. בשיעור הזה נבין את שלוש הרמות ואת המיפוי שלהן לחומרה, נוכיח מדוע חוק העצמאות של הבלוקים הוא הלב של ההרחבה השקופה (transparent scaling), נראה למה אין מחסום (barrier) בין בלוקים ומדוע ניסיון לחסום בלוק אחד על אחר מזמין deadlock, ונשלוט באריתמטיקת האינדקסים החד-ממדית והדו-ממדית שממנה כל thread מגלה על איזה נתון לעבוד. בשיעור 2.3 נשתמש בכל זה כדי לכתוב kernels אמיתיים.

שלוש הרמות של ההיררכיה - thread, block ו-grid¶

מודל התכנות של CUDA בנוי על היררכיה מקוננת בת שלוש רמות, וזו אחת משתי אבני היסוד שלו (השנייה היא היררכיית הזיכרון, שנעמיק בה בפרק 3). כל רמה מקבצת את הרמה שמתחתיה:

                 thread block grid  (grid - all threads of the launch)
      +------------------------------------------------------+
      |   block            block            block            |
      | +---------+      +---------+      +---------+         |
      | | warp    |      | warp    |      | warp    |   ...   |
      | | warp    |      | warp    |      | warp    |         |
      | +---------+      +---------+      +---------+         |
      | (<=1024      (<=1024        (<=1024                   |
      |  threads)     threads)       threads)                 |
      +------------------------------------------------------+
        thread  <  (warp)  <  block/CTA  <  grid

thread - יחידת ההרצה האטומית. לכל thread מזהה ייחודי מבוסס-אינדקס בתוך הבלוק שלו, וממנו הוא מחשב על איזה איבר לעבוד.
warp - קבוצת 32 threads שהחומרה מתזמנת יחד. חשוב לזכור מ-2.1: ה-warp הוא פרט מימוש, לא רמה רשמית במודל התכנות. הוא לא מופיע בקוד שלנו כאובייקט, אבל הוא קובע ביצועים.
block (בלוק) - קבוצת threads שמתוזמנת יחד על SM אחד. גודל מרבי של 1024 threads על החומרה הנוכחית, ונבחר כמעט תמיד ככפולה של גודל ה-warp (32), כדי שהthreads ייארזו לתוך warps שלמים. בלוק מכיל warp אחד או יותר.
grid (רשת) - אוסף כל הבלוקים שנוצרים מkernel launch אחד. הlaunch אחד = grid אחד = כל הthreads שמריצים את אותו ה-kernel. ה-grid יכול להיות חד-ממדי, דו-ממדי או תלת-ממדי (1D / 2D / 3D), בהתאם למבנה הבעיה.

המיפוי לחומרה מדויק ומחייב: thread רץ על core (ליבה, compute pipeline), block מתוזמן כולו על SM אחד, ו-grid נפרס על פני כל ה-SMs של הכרטיס. ב-H100, למשל, ה-grid מתחלק על פני 132 ה-SMs. שימו לב לאסימטריה: thread ל-core הוא יחסי-רבים (SM מריץ עשרות warps), אבל בלוק ל-SM הוא כלל ברזל של מיקום - כל threads הבלוק על אותו SM, אף פעם לא מפוצלים בין שניים.

הבלוק וה-CTA - block ו-Cooperative Thread Array¶

המילה "block" (או "thread block") היא השם ברמת מודל התכנות. לאותו דבר בדיוק יש שם נוסף ברמה הנמוכה, ברמת ה-PTX וה-SASS (שפות הביניים והמכונה של ה-GPU): Cooperative Thread Array, בקיצור CTA - מערך תהליכונים משתף-פעולה. שני השמות מתייחסים לאותה ישות; פשוט אחד הוא ההפשטה שאנחנו כותבים בה, והשני הוא איך שהחומרה והcompiler מדברים עליה. כשתקראו PTX או פלט של cuobjdump ותיתקלו ב-CTA, דעו שמדובר בבלוק.

התכונה הקריטית של הבלוק היא ערובת ה-SM המשותף: כל threads בלוק מתוזמנים בו-זמנית על אותו SM יחיד. המיקום המשותף הזה (co-residency) איננו פרט טכני שולי - הוא בדיוק מה שמאפשר שיתוף פעולה בתוך בלוק, ומכאן נגזרות שתי היכולות הבלעדיות של הרמה הזו:

shared memory (זיכרון משותף) - כל threads הבלוק רואים אותו בלוק זיכרון מהיר שיושב פיזית ב-L1 של ה-SM (ב-H100, מתוך 256 KiB של L1/shared מאוחד; עד 228 KiB מהם ניתנים להקצאה כ-shared memory לבלוק). זהו ערוץ תקשורת מהיר בין הthreads, בסדרי גודל מהיר מ-global memory.
barriers (מחסומים) - הפקודה __syncthreads() היא מחסום ברמת הבלוק: כל threads הבלוק חייבים להגיע אליו לפני שאף אחד ממשיך. זו נקודת סנכרון שמאפשרת, למשל, לטעון נתונים ל-shared memory ואז להבטיח שכולם מוכנים לפני שקוראים אותם.

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

        why are all the block's threads on the same SM?  ->  because only that way can you have:

   thread 0 ─┐
   thread 1 ─┤   __shared__ float tile[...]   (memory on the SM)
   thread 2 ─┤   ...everyone reads/writes to it...
     ...     ─┤
   thread N ─┘   __syncthreads();             (barrier on the SM)
                 ^^^^^^^^^^^^^^^^ possible only when everyone is on the same SM

חוק העצמאות וההרחבה השקופה - independence ו-transparent scaling¶

וכעת לכלל היחיד, החשוב ביותר, שקובע נכונות ברמת ה-grid: הבלוקים חייבים לרוץ באופן עצמאי, כך שכל סדר הרצה של הבלוקים יהיה תקף. הזמן (runtime) חופשי להריץ את הבלוקים בכל סדר - החל מסדרתי לחלוטין (בלוק אחרי בלוק) ועד מקבילי לחלוטין (כולם בבת אחת) - ותוכנית CUDA נכונה חייבת להישאר נכונה תחת כל שזירה (interleaving) של הבלוקים. אסור להניח שבלוק 5 רץ אחרי בלוק 3, אסור להניח ששניהם רצים בו-זמנית, אסור להניח כלום על הסדר.

האי-דטרמיניזם הזה נראה כמו מגבלה, אבל הוא בדיוק המנוע של התכונה המרכזית של CUDA - הרחבה שקופה (transparent scaling). מטרת התכנון של המודל היא שקוד שנכתב נכון "יתרחב בשקיפות ככל שהתקני ה-GPU מתרחבים במשאבי ההרצה המקבילית שלהם" - כלומר, שאותו קוד בדיוק יֵרוץ מהר יותר על כרטיס גדול יותר בלי שכתבנו אותו מחדש. ניסוח פרובוקטיבי מהמקור: המודל תוכנן כדי למנוע מהמתכנת לכתוב קוד ש"ייכשל בלהיות מהיר יותר כשהמשתמש קונה GPU חדש של NVIDIA".

הדוגמה הממחישה: אותו grid בדיוק, עם אותם בלוקים, רץ נכון בשני קצוות:

          grid with 8 blocks  (B0..B7)

  GPU with one SM (sequential):   [B0][B1][B2][B3][B4][B5][B6][B7]   one after another
                           SM0 ─────────── over time ───────────►

  GPU with 4 SMs (parallel):   SM0: [B0][B4]
                           SM1: [B1][B5]     four at once,
                           SM2: [B2][B6]     then four more
                           SM3: [B3][B7]

  H100 with 132 SMs:         all eight blocks at once, with SMs to spare

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

אין מחסום חוצה-בלוקים - global memory, atomics ו-deadlock¶

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

אז איך בלוקים כן מתאמים ביניהם, כשצריך? הדרך היחידה היא דרך global memory (זיכרון גלובלי) בעזרת פעולות אטומיות (atomics). thread בבלוק אחד כותב לכתובת ב-global memory, וthread בבלוק אחר קורא אותה; כדי שהעדכונים לא ידרסו זה את זה, משתמשים בפקודות כמו atomicAdd, שמבצעות קריאה-שינוי-כתיבה כפעולה אחת בלתי-ניתנת-לחלוקה. זו תקשורת איטית (global memory) ומבוססת-atomics, לא מבוססת-מחסום.

ומכאן מלכוד ה-deadlock, שחובה להפנים: מכיוון שסדר הבלוקים נשלט על ידי הנהג (driver) ואינו מובטח, חסימת בלוק אחד בהמתנה לבלוק אחר יכולה בקלות להוביל ל-deadlock. דמיינו kernel שבו בלוק 1 מסתובב בלולאת המתנה (busy-wait) עד שבלוק 0 יכתוב דגל ל-global memory:

// incorrect code - may cause a deadlock, do not do this
__global__ void bad(volatile int* flag) {
    if (blockIdx.x == 1) {
        while (flag[0] == 0) { /* waiting for block 0 */ }  // busy-wait
    } else if (blockIdx.x == 0) {
        flag[0] = 1;                                  // releases block 1
    }
}

על GPU קטן עם SM אחד, ייתכן שהנהג יתזמן קודם את בלוק 1. בלוק 1 יתפוס את ה-SM ויסתובב לנצח בהמתנה, ובלוק 0 לעולם לא יקבל הזדמנות לרוץ ולשחרר אותו - קיפאון מוחלט. הכלל: אסור לכתוב קוד שבו בלוק אחד ממתין באופן פעיל להתקדמות של בלוק אחר. אם צריך תלות כזו בין שלבים, מפצלים לשני kernel launches נפרדים - הסיום של ה-kernel הראשון הוא נקודת הסנכרון החוצה-בלוקים היחידה שאפשר לסמוך עליה.

אריתמטיקת האינדקסים - blockIdx, blockDim, threadIdx ו-gridDim¶

השקנו kernel אחד באלפי עותקים; איך כל עותק יודע על איזה נתון לעבוד? מהמזהים המובנים. כל thread רואה ארבעה משתני dim3 מובנים שהחומרה ממלאת עבורו (ל-dim3 שלושה שדות: .x, .y, .z):

משתנה	משמעות
`threadIdx`	האינדקס של ה-thread בתוך הבלוק שלו
`blockDim`	מספר הthreads בכל בלוק (מימדי הבלוק)
`blockIdx`	האינדקס של הבלוק בתוך ה-grid
`gridDim`	מספר הבלוקים ב-grid (מימדי ה-grid)

בעולם החד-ממדי, האינדקס הגלובלי הקנוני - הביטוי שראינו כבר בשיעור 0.3 ושחוזר בכל תוכנית CUDA:

int i = blockIdx.x * blockDim.x + threadIdx.x;

הפירוק במספרים, עם blockDim.x = 256: הthread הראשון של בלוק 0 מקבל 0*256 + 0 = 0; הthread האחרון של בלוק 0 מקבל 0*256 + 255 = 255; הthread הראשון של בלוק 1 מקבל 1*256 + 0 = 256; וכן הלאה. כל האיברים מכוסים פעם אחת בדיוק, בלי חפיפה ובלי חורים. גודל ה-grid מחושב בעיגול כלפי מעלה (n + blockSize - 1) / blockSize, ותמיד נדרש שומר גבולות if (i < n) כדי להגן על הthreads העודפים שנוצרים כשה-grid לא מחלק בדיוק את n.

אינדוקס דו-ממדי עם dim3 - 2D indexing¶

כשהבעיה עצמה דו-ממדית - תמונה, מטריצה, לוח - נוח למפות אליה grid דו-ממדי, כך שהאינדקסים של הthread משקפים ישירות שורה ועמודה. משתמשים בטיפוס dim3 כדי לבנות תצורת הרצה דו-ממדית. שימו לב לאמנה שקל להתבלבל בה: .x הוא הממד המהיר (עמודות), .y הוא האיטי (שורות).

int width = 1920, height = 1080;             // example image
dim3 block(16, 16);                          // 256 threads per block, 16x16
dim3 grid((width  + block.x - 1) / block.x,  // blocks across width (columns)
          (height + block.y - 1) / block.y); // blocks across height (rows)
myKernel<<<grid, block>>>(img, width, height);

בתוך ה-kernel, כל thread מחשב את השורה והעמודה שהוא אחראי עליהן, ואז ניגש לזיכרון ליניארי בעזרת row * width + col (כי מטריצות ותמונות נשמרות ברצף שורה-אחר-שורה, row-major):

__global__ void myKernel(unsigned char* img, int width, int height) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;   // .x -> column
    int row = blockIdx.y * blockDim.y + threadIdx.y;   // .y -> row
    if (row < height && col < width) {                 // 2D bounds guard
        int idx = row * width + col;                   // row-major flattening
        img[idx] = 255 - img[idx];                     // example: pixel inversion
    }
}

הנה המיפוי הוויזואלי של grid דו-ממדי על מטריצה. כל בלוק הוא tile של 16x16 threads, והבלוקים מרצפים את המישור:

              col = blockIdx.x*blockDim.x + threadIdx.x  ─────►
        +----------------+----------------+----------------+
        |  block (0,0)   |  block (1,0)   |  block (2,0)   |
   r    |  16x16 threads |  16x16 threads |  16x16 threads |
   o    +----------------+----------------+----------------+
   w    |  block (0,1)   |  block (1,1)   |  block (2,1)   |
   |    |  16x16 threads |  16x16 threads |  16x16 threads |
   v    +----------------+----------------+----------------+
   (blockIdx.y, threadIdx.y)     idx = row * width + col

שומר הגבולות כאן הוא if (row < height && col < width) - שני התנאים, כי ה-grid מעוגל כלפי מעלה בשני הממדים, וייווצרו threads עודפים גם מימין למטריצה וגם מתחתיה. הרחבה ל-3D זהה ברוחה: מוסיפים .z ומחשבים depth באותה תבנית. את ה-kernel הזה לכפל מטריצות נראה במלואו בשיעור 2.3.

כמה בלוקים נכנסים ל-SM - occupancy ומגבלות משאבים¶

ראינו שבלוק שלם יושב על SM אחד. שאלה מעשית: כמה בלוקים יכולים לשבת על אותו SM בו-זמנית? התשובה קובעת את ה-occupancy (תפוסה) - כמה warps תושבים יש על ה-SM ביחס למקסימום - וזה בתורו קובע כמה טוב ה-SM יכול להסתיר latency (השהיה), כפי שראינו בפרק 1. ככל שיותר בלוקים תושבים, יותר warps זמינים להחלפה כשאחד נתקע.

מספר הבלוקים לכל SM חסום על ידי המשאבים הסופיים של ה-SM, וכל אחד מהם יכול להיות הגורם המגביל:

קובץ האוגרים (register file): ב-H100, 65,536 אוגרים בני 32 סיביות ל-SM. כל thread צורך מספר אוגרים קבוע (שה-nvcc קובע בקומפילציה), וסך האוגרים של כל הthreads התושבים לא יכול לחרוג מהתקציב.
חריצי warp (warp slots): ה-SM יכול לעקוב אחר מספר מוגבל של warps תושבים - ב-H100, 64 warps (כלומר 2,048 threads).
בייטים של shared memory: כל בלוק שמכריז על __shared__ צורך נתח מה-L1/shared, וסך הצריכה חסום (עד 228 KiB לבלוק, ופחות מכך אם רוצים כמה בלוקים תושבים).
מגבלה קשיחה של בלוקים ל-SM: ב-Hopper, עד 32 בלוקים תושבים ל-SM.

החשבון נעשה בזמן קומפילציה על ידי nvcc: הcompiler יודע כמה אוגרים ו-shared memory כל בלוק צורך, וכך אפשר לחשב מראש כמה בלוקים ייכנסו. בואו נחשב דוגמה מלאה ל-H100, עם kernel שבו כל בלוק הוא 256 threads (8 warps) ומשתמש ב-8 KiB shared memory לבלוק:

assume: blockSize = 256 threads = 8 warps ;  shared = 8 KiB/block

register limit (with 32 registers/thread):
   32 registers/thread * 256 threads = 8,192 registers/block
   65,536 / 8,192            = 8 blocks
warp slot limit:
   64 warps / 8 warps-per-block  = 8 blocks
shared memory limit:
   228 KiB / 8 KiB           = 28 blocks
hard block limit:            = 32 blocks
                              ------------------
   limiting factor = minimum    = 8 resident blocks

שמונה בלוקים תושבים פירושם 8 * 256 = 2,048 threads תושבים - בדיוק המקסימום של ה-SM, כלומר תפוסה מלאה (100% occupancy). עכשיו נראה את כוחו של לחץ האוגרים: אם אותו kernel היה צורך 64 אוגרים לthread במקום 32, מגבלת האוגרים הייתה 65,536 / (64*256) = 4 בלוקים בלבד, כלומר 4 * 256 = 1,024 threads תושבים - 50% תפוסה בלבד. אותו קוד, פי שניים אוגרים, חצי מהתפוסה. זו בדיוק הסיבה שכתיבת kernels "רזים" באוגרים היא מפתח לביצועים. את הטיפול המלא ב-occupancy - איך מודדים אותה, איך מכווננים אותה, ומתי תפוסה גבוהה בכלל עוזרת - נפתח בפרק 8.

סיכום¶

היררכיית הthreads בת שלוש רמות: thread (רץ על core) < block/CTA (מתוזמן כולו על SM אחד) < grid (נפרס על כל ה-SMs); ה-warp הוא פרט מימוש שיושב בין ה-thread לבלוק ואינו רמה רשמית במודל.
בלוק מכיל עד 1024 threads, נבחר ככפולה של 32, ומכיל warp אחד או יותר; grid הוא אוסף כל הבלוקים של kernel launch אחד, ויכול להיות 1D/2D/3D.
"block" הוא שם מודל התכנות ו-"Cooperative Thread Array" (CTA) הוא אותו דבר ברמת PTX/SASS; ערובת ה-SM המשותף (כל threads הבלוק על SM אחד) היא מה שמאפשר shared memory ו-barriers, ולכן הבלוק הוא היחידה הקטנה ביותר של שיתוף פעולה שנחשפת למתכנת.
חוק העצמאות: הבלוקים חייבים לרוץ באופן עצמאי וכל סדר שזירה שלהם חייב להיות תקף; האי-דטרמיניזם הזה הוא בדיוק מה שמאפשר הרחבה שקופה - אותו grid רץ נכון על GPU עם SM אחד (סדרתי) או עם 132 SMs (מקבילי מלא).
אין מחסום חוצה-בלוקים: __syncthreads() הוא ברמת הבלוק בלבד; בלוקים מתאמים ביניהם רק דרך global memory עם atomics, וחסימת בלוק בהמתנה לבלוק אחר מזמינה deadlock כי הנהג שולט בסדר - תלות בין שלבים פותרים בשני kernel launches נפרדים.
האינדקס הגלובלי החד-ממדי הוא i = blockIdx.x * blockDim.x + threadIdx.x, עם gridDim ו-blockDim כמימדים; תמיד נדרש שומר גבולות if (i < n).
אינדוקס דו-ממדי עם dim3: col = blockIdx.x*blockDim.x + threadIdx.x, row = blockIdx.y*blockDim.y + threadIdx.y, גישה לזיכרון ב-row*width + col, ושומר גבולות כפול if (row < height && col < width); זכרו ש-.x הוא הממד המהיר (עמודות).
מספר הבלוקים שנכנסים ל-SM נקבע על ידי המשאב הסופי הכי מגביל - קובץ האוגרים (65,536 ל-SM ב-H100), חריצי warp (64), בייטים של shared memory (עד 228 KiB), ומגבלת 32 בלוקים ל-SM; nvcc מחשב זאת בזמן קומפילציה.
מספר הבלוקים התושבים קובע את ה-occupancy: בדוגמת 256 threads ו-32 אוגרים/thread מתקבלים 8 בלוקים תושבים = 2,048 threads = תפוסה מלאה, אך הכפלת האוגרים ל-64 מורידה ל-4 בלוקים = 50% - לחץ אוגרים הוא בלם ביצועים מרכזי, ונרחיב עליו בפרק 8.