3.3 Shared memory ו tiling הרצאה

בשיעור 2.3 כתבנו את כפל המטריצות (matmul) הנאיבי וגילינו אמת מרה: ה-kernel רץ, נותן תוצאה נכונה, ובכל זאת מנצל שבריר זעום מכוח ה-GPU. הסיבה, כפי שראינו, היא שהוא memory-bound - הוא מבצע בערך פעולת נקודה-צפה אחת (FLOP) לכל קריאה מ-global memory (זיכרון גלובלי), ורוחב הפס האריתמטי של ה-CUDA Cores גבוה בהרבה מרוחב הפס אל ה-HBM3. בפרק 3 עד כה מיפינו את היררכיית הזיכרון: ראינו שבין ה-registers הפרטיים לכל thread לבין ה-global memory הענק והאיטי יושבת שכבת ביניים - shared memory (זיכרון משותף), מהירה מאוד ומשותפת לכל ה-threads שב-block. השיעור הזה הוא לב האופטימיזציה של הקורס כולו: נלמד את טכניקת ה-tiling (חלוקה-לtiles), התבנית שממנפת את ה-shared memory כדי להפוך את ה-matmul הנאיבי מ-kernel זוחל ל-kernel מהיר, ונקפיץ את arithmetic intensity מ-1 לערך של עשרות. בדרך נבין למה נחוצים בדיוק שני מחסומי __syncthreads(), איך ה-shared memory שאנחנו צורכים מגביל את ה-occupancy (תפוסה), ונציץ בהתנגשויות הבנקים (bank conflicts) שאורבות מתחת לפני השטח.

למה ה-matmul הנאיבי מבזבז את ה-GPU - the memory wall¶

נזכיר את ה-kernel הנאיבי מ-2.3, בדיוק כפי שהופיע ב-glossary:

__global__ void mm(const float* A, const float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; k++) {
            sum += A[row * N + k] * B[k * N + col];   // two global reads, two FLOPs
        }
        C[row * N + col] = sum;
    }
}

בואו נספור את התעבורה מול הזיכרון בכל איטרציה של הלולאה הפנימית. אנחנו קוראים איבר אחד מ-A ואיבר אחד מ-B - שתי קריאות מ-global memory, 8 בתים - ומבצעים כפל וחיבור, כלומר שתי FLOPs. היחס הוא FLOP אחת לכל קריאה מ-global memory (שתי FLOPs חלקי שתי קריאות), או במונחי בתים: 0.25 FLOP לכל בית. זו arithmetic intensity נמוכה באופן קיצוני.

הבעיה החמורה יותר היא שאין כאן שום ניצול-חוזר (reuse). שקלו את השורה row של C: כל N האיברים שבה קוראים את אותה שורה של A. בגרסה הנאיבית, כל אחד מ-N ה-threads שמחשבים את השורה הזו קורא את השורה של A מ-global memory מחדש, בנפרד. אותו נתון נשלף מה-HBM3 שוב ושוב. מטריצה N x N מבצעת 2*N^3 FLOPs אבל מייצרת בסדר גודל של 2*N^3 קריאות global - כמעט אפס ניצול-חוזר.

התוצאה מנוסחת ב-glossary בחדות: "לעולם לא תנצלו כך את כל ה-GPU, כי רוחב הפס האריתמטי של ה-CUDA Cores ב-FLOPs/s גבוה בהרבה מרוחב הפס בין ה-GPU RAM ל-SMs." ב-H100, שיא ה-FP32 (ללא Tensor Cores) הוא כ-66.9 TFLOPS, בעוד רוחב הפס של ה-HBM3 הוא בסדר גודל של 3.35 TB/s. היחס בין השניים הוא כ-20 FLOPs לכל בית שנקרא. ה-matmul הנאיבי מספק 0.25 - פי 80 בערך פחות מהנקודה שבה החישוב מתחיל להיות צוואר הבקבוק. אנחנו מרעיבים את יחידות החישוב.

הפתרון היחיד הוא להעלות את arithmetic intensity: לקרוא כל נתון מ-global memory פעם אחת, להניח אותו במחסן המהיר של ה-SM, ולסחוט ממנו כמה שיותר FLOPs לפני שנזרוק אותו. המחסן הזה הוא ה-shared memory, והשיטה היא tiling.

הכרזת זיכרון משותף - shared¶

ה-shared memory מוכרז בתוך ה-kernel במילת המפתח __shared__. יש שתי צורות: סטטית (בגודל קבוע הידוע בזמן קומפילציה) ודינמית (בגודל שנקבע בזמן הlaunch).

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

#define TILE_WIDTH 16
__global__ void k(...) {
    __shared__ float As[TILE_WIDTH][TILE_WIDTH];   // 16*16*4 = 1024 bytes per block
    __shared__ float Bs[TILE_WIDTH][TILE_WIDTH];
    ...
}

שימו לב היטב: As ו-Bs אינם משתנים פרטיים לכל thread כמו משתנה רגיל בתוך ה-kernel. הם משותפים לכל ה-threads של ה-block - זו בדיוק הנקודה. thread מספר (0,0) יכול לכתוב ל-As[0][0] ו-thread מספר (5,3) יקרא את הערך הזה. ה-block כולו חולק לוח כתיבה אחד. פיזית, כפי שראינו ב-3.1, ה-shared memory יושב ב-L1 data cache של ה-SM, ולכן הגישה אליו מהירה בסדר גודל מ-global memory בשתי המידות - גם ב-latency (השהיה) וגם ב-throughput (תפוקה).

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

extern __shared__ float smem[];       // size unknown at compile time

__global__ void kDyn(int tile, ...) {
    float* As = smem;                 // first half
    float* Bs = smem + tile * tile;   // second half
    // As[ty*tile + tx] instead of As[ty][tx] - manual two-dimensional indexing
    ...
}

// at launch: the third parameter in the triple angle brackets is the number of bytes
size_t smemBytes = 2 * tile * tile * sizeof(float);
kDyn<<<grid, block, smemBytes>>>(tile, ...);

הפרמטר השלישי ב-<<<grid, block, smemBytes>>> הוא מספר הבתים של ה-shared memory הדינמי שיוקצה לכל block (ברירת המחדל, כשהוא מושמט, היא 0). כל ההקצאה הדינמית מגיעה כבלוק זיכרון אחד ורציף שמצביע smem מצביע לתחילתו; אם צריך שני מערכים לוגיים, מחלקים אותו ידנית כפי שעשינו כאן. המחיר של הדינמיות הוא ויתור על האינדוקס הדו-ממדי הנוח As[ty][tx] - חייבים לשטח אותו ל-As[ty*tile + tx] בעצמנו.

מאפיין	סטטי (`__shared__ float As[T][T]`)	דינמי (`extern __shared__`)
גודל ידוע ב	קומפילציה	זמן launch
מהיכן הגודל	מהמערך עצמו	הפרמטר השלישי של `<<< >>>` (בתים)
אינדוקס	דו-ממדי נוח `As[y][x]`	ידני ומשוטח `As[y*T+x]`
כמה מערכים	כמה שרוצים	בלוק אחד רציף, מחלקים ידנית
מתי להשתמש	גודל tile קבוע (רוב המקרים)	tile בגודל זמן-ריצה, kernel גנרי

תבנית ה-staging - load, sync, compute, sync¶

ה-glossary מתאר את התבנית הקנונית של kernel מהיר בארבעה שלבים, וזו התבנית שנחזור עליה שוב ושוב בקורס:

טעינה של נתח נתונים מ-global memory אל shared memory, בשיתוף פעולה בין כל ה-threads של ה-block.
מחסום - __syncthreads() - כדי לוודא שכל ה-block סיים לטעון לפני שמישהו מתחיל לחשב.
חישוב מתוך ה-shared memory המהיר, שוב ושוב, כמה שיותר.
מחסום שני לפני שמאתחלים מחדש את אותו ה-shared memory לנתח הבא, ולבסוף כתיבה חזרה ל-global memory.

הפקודה __syncthreads() היא מחסום (barrier) ברמת ה-block: כל ה-threads של ה-block חייבים להגיע אליה לפני שאף אחד מהם ממשיך הלאה. זה בדיוק המנגנון שהוזכר ב-2.2 - שיתוף פעולה בתוך block אפשרי כי כל ה-threads שלו מתוזמנים יחד על אותו SM. שימו לב לאזהרה קריטית: __syncthreads() חייב להיות במסלול שכל ה-threads של ה-block מגיעים אליו. אם תקראו לו בתוך if שרק חלק מה-threads נכנסים אליו, ה-block ייתקע לנצח (deadlock), כי המחסום ממתין ל-threads שלעולם לא יגיעו.

   the staging pattern on a timeline (one block, several warps)
   ─────────────────────────────────────────────────────►  time

   warp 0:  load ─┐            compute ────┐            load ...
   warp 1:  load ─┤            compute ────┤            load ...
   warp 2:  load ─┤            compute ────┤            load ...
   warp 3:  load ─┘            compute ────┘            load ...
                  │                        │
              __syncthreads()          __syncthreads()
              (load barrier)           (reuse barrier)

הרעיון העמוק: המחיר של קריאה מ-global memory משולם פעם אחת בשלב הטעינה, ואז מתחלק (amortized) על פני עשרות פעולות אריתמטיות מהירות מתוך ה-shared memory. זה בדיוק מה שמעלה את arithmetic intensity.

ה-matmul עם tiling - the tiled kernel¶

עכשיו נחבר הכל ל-matmul המרובע עם tiling. הרעיון הגיאומטרי: מחלקים את מטריצת התוצאה C לtiles בגודל TILE_WIDTH x TILE_WIDTH. כל block אחראי על tile אחד של C. כדי לחשב את הtile, ה-block מתקדם לאורך A ו-B בשלבים (phases), tile בכל פעם: בכל שלב הוא טוען tile אחד של A וtile אחד של B ל-shared memory, ומצטבר את המכפלות החלקיות.

        B: tiles along the columns
        ┌────┬────┬────┐
        │ B0 │    │    │
        ├────┼────┼────┤
        │ B1 │    │    │
        ├────┼────┼────┤
   A    │ B2 │    │    │
 ┌───┬──┴────┴────┴────┘
 │A0 │A1│A2│        C
 ├───┼──┼──┤   ┌────┬────┬────┐
 │   │  │  │   │ C  │    │    │   C = sum_t (A_t * B_t)
 └───┴──┴──┘   └────┴────┴────┘   each block accumulates one tile of C

הנה ה-kernel המלא. שימו לב לשני מחסומי ה-__syncthreads():

#define TILE_WIDTH 16

__global__ void mmTiled(const float* A, const float* B, float* C, int N) {
    __shared__ float As[TILE_WIDTH][TILE_WIDTH];
    __shared__ float Bs[TILE_WIDTH][TILE_WIDTH];

    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TILE_WIDTH + ty;   // row in C this thread is responsible for
    int col = blockIdx.x * TILE_WIDTH + tx;   // column in C

    float sum = 0.0f;
    int numTiles = (N + TILE_WIDTH - 1) / TILE_WIDTH;

    for (int t = 0; t < numTiles; t++) {
        // cooperative load: each thread loads one element of A and one of B
        int aCol = t * TILE_WIDTH + tx;
        int bRow = t * TILE_WIDTH + ty;
        As[ty][tx] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
        Bs[ty][tx] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;

        __syncthreads();                       // barrier 1: the tile is fully loaded

        // compute from shared memory only
        for (int k = 0; k < TILE_WIDTH; k++)
            sum += As[ty][k] * Bs[k][tx];

        __syncthreads();                       // barrier 2: before overwriting the tile in the next stage
    }

    if (row < N && col < N)
        C[row * N + col] = sum;
}

והlaunch, עם block דו-ממדי בגודל הtile ו-grid דו-ממדי שמכסה את C:

dim3 block(TILE_WIDTH, TILE_WIDTH);            // 16x16 = 256 threads
dim3 grid((N + TILE_WIDTH - 1) / TILE_WIDTH,
          (N + TILE_WIDTH - 1) / TILE_WIDTH);
mmTiled<<<grid, block>>>(d_A, d_B, d_C, N);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());

הריפוד ב-0.0f בטעינה (ה-? :) מטפל במטריצות ש-N שלהן אינו מתחלק ב-TILE_WIDTH: הtiles החלקיים בשוליים מתמלאים באפסים, שאינם תורמים לסכום. כך ה-kernel נכון לכל N, לא רק לכפולות של גודל הtile.

החשבון של הניצול-חוזר. ניקח TILE_WIDTH = 16, כלומר block של 256 threads. בכל שלב:

ה-block מבצע 256 קריאות מ-global memory כדי לטעון את As (thread אחד לכל איבר), ועוד 256 כדי לטעון את Bs. סך הכל 512 קריאות global.
לאחר מכן, כל אחד מ-256 ה-threads מבצע את הלולאה הפנימית TILE_WIDTH = 16 פעמים, וכל איטרציה היא כפל-חיבור = 2 FLOPs. כלומר 256 * 16 * 2 = 8192 FLOPs.
היחס: 8192 FLOPs חלקי 512 קריאות = 16 FLOPs לכל קריאה מ-global memory.

זו קפיצה מ-1 (הנאיבי) ל-TILE_WIDTH (המרוצף). הסיבה פשוטה: כל איבר שנטען ל-shared memory מנוצל TILE_WIDTH פעמים לפני שנזרק. איבר של As[ty][k] נקרא בידי כל TILE_WIDTH ה-threads שבשורה ty של ה-block; איבר של Bs[k][tx] נקרא בידי כל TILE_WIDTH ה-threads שבעמודה tx. ל-TILE_WIDTH = 32 הarithmetic intensity מוכפלת ל-32 FLOPs לקריאה.

מהי המשמעות למהירות? להלן מדידות אופייניות ל-N = 4096 על H100 (הערכים משתנים בין הרצות ובין דורות חומרה - קחו אותם כסדרי גודל, לא כמספרים מקודשים):

גרסה	GFLOP/s (בקירוב)	אחוז משיא FP32 (~66.9 TFLOPS)
נאיבי (2.3)	~300	~0.5%
מרוצף, `TILE_WIDTH = 16`	~3,000	~4.5%
מרוצף, `TILE_WIDTH = 32`	~4,500	~7%
cuBLAS (Tensor Cores)	~50,000+	-

הקפיצה מהנאיבי למרוצף היא בסדר גודל של פי 10 עד פי 15 - וזו רק ההתחלה. גם ה-kernel המרוצף שלנו רץ על CUDA Cores, ולכן, כפי שה-glossary מזכיר, הוא אינו תקרת הביצועים: ה-matmul המהיר ביותר רץ על Tensor Cores, שרוחב הפס האריתמטי שלהם גבוה בהרבה. ל-cuBLAS (שהוא, אגב, column-major) ולריצוף ברמת ה-register נגיע בפרויקט 3.5 ובפרק על Tensor Cores.

למה שני מחסומים - the two-barrier rule¶

הנקודה העדינה ביותר בכל ה-kernel היא שני מחסומי ה-__syncthreads(). שניהם הכרחיים, וכל אחד מגן מפני סכנה (hazard) אחרת. הבה נבין בדיוק מפני מה.

מחסום 1 (מחסום הטעינה), לפני החישוב. ה-threads טוענים את הtile ל-As ו-Bs, ואז ניגשים לחשב. אבל ה-threads של ה-block אינם רצים בצעד אחיד - הם מפוזרים על פני כמה warps, ו-warps מתוזמנים בזמנים שונים. בלי המחסום, thread ב-warp מהיר עלול להתחיל את הלולאה הפנימית ולקרוא את As[ty][k] בעוד thread ב-warp אחר עדיין לא כתב את חלקו לtile. זו סכנת קריאה-לפני-כתיבה (read-before-write): קוראים נתון שעוד לא נטען. המחסום הראשון מבטיח שהtile מלא לגמרי לפני שמישהו קורא ממנו.

מחסום 2 (מחסום הניצול-חוזר), אחרי החישוב. זהו המחסום שמתחילים שוכחים, וזה הבאג הקלאסי. אחרי הלולאה הפנימית, השלב הבא בלולאת ה-t יכתוב tile חדש אל אותם As ו-Bs (אותו זיכרון פיזי). בלי המחסום השני, thread ב-warp מהיר שכבר סיים את החישוב שלו ימשיך לסיבוב הבא של הלולאה ויתחיל לדרוס את As[ty][tx] בנתון של הtile החדש - בעוד thread ב-warp איטי עדיין קורא את As[ty][k] של הtile הישן בלולאה הפנימית שלו. זו סכנת כתיבה-לפני-קריאה (write-after-read): דורסים נתון שעוד לא סיימו לקרוא. התוצאה היא ערכי C שגויים, ובאופן מרושע - לא דטרמיניסטיים: לפעמים נכון, לפעמים לא, תלוי בתזמון ה-warps.

   The bug from omitting the second barrier:

   slow warp:   ...still in the inner loop, reading As[ty][k]  (tile t)
   fast warp:   already at stage t+1, writing As[ty][tx] = A[...]   (tile t+1)
                              └──► overwrote the data the slow warp is still reading!

אין להסתמך על סנכרון מרומז בתוך warp. מאז ארכיטקטורת Volta, ה-GPU תומך ב-independent thread scheduling, כך שאפילו threads באותו warp אינם בהכרח בצעד אחיד. שני המחסומים הם היחידים שמבטיחים נכונות. בתרגול תסירו את המחסום השני בכוונה, תראו במו עיניכם תוצאות שגויות ולא יציבות, ואז תתקנו.

אינטראקציה עם התפוסה - occupancy interaction¶

ה-shared memory אינו חינם: הוא משאב מוגבל לכל SM, ולכן כמות ה-shared memory שכל block צורך מגבילה כמה blocks יכולים לשבת (resident) בו-זמנית על אותו SM. זהו לב חישוב ה-occupancy, שנקדיש לו את פרק 8.1; כאן רק נציץ.

ב-H100, כל SM מכיל 256 KiB של L1/shared משותפים (עם חלוקה בת-הגדרה ביניהם), ומאפשר עד 2048 threads תושבים ועד 32 blocks תושבים. ה-occupancy בפועל נקבע על ידי המגביל החמור מכולם מבין: מכסת ה-threads, מכסת ה-blocks, ה-registers לכל thread, וה-shared memory לכל block. נחשב ל-mmTiled שלנו:

עבור TILE_WIDTH = 16 (block של 256 threads, 2 * 16 * 16 * 4 = 2048 בתים = 2 KiB shared לכל block):

מגבלת threads: 2048 / 256 = 8 blocks.
מגבלת blocks: 32 blocks.
מגבלת shared: 256 KiB / 2 KiB = 128 blocks.
המגביל החמור: 8 blocks (מכסת ה-threads). ה-shared memory כלל אינו הצוואר כאן.

עבור TILE_WIDTH = 32 (block של 1024 threads, 2 * 32 * 32 * 4 = 8192 בתים = 8 KiB shared לכל block):

מגבלת threads: 2048 / 1024 = 2 blocks.
מגבלת shared: 256 KiB / 8 KiB = 32 blocks.
המגביל החמור: 2 blocks (שוב מכסת ה-threads).

בשני המקרים ה-shared memory אינו המגביל - הtiles שלנו קטנים. אבל שנו את החשבון: kernel שצורך, נניח, 100 KiB shared לכל block יוגבל ל-2 blocks בלבד לפי ה-shared, גם אם ה-threads וה-registers היו מאפשרים יותר. זה הטרייד-אוף: הtile הגדול יותר משפר את הניצול-החוזר (arithmetic intensity TILE_WIDTH), אבל צורך יותר shared memory ועלול לחנוק את ה-occupancy. הערה מעשית: החלוקה הסטטית __shared__ מוגבלת כברירת מחדל ל-48 KiB לכל block; לחריגה מעבר לכך חייבים shared memory דינמי ובקשה מפורשת דרך cudaFuncSetAttribute עם cudaFuncAttributeMaxDynamicSharedMemorySize.

הצצה להתנגשויות בנקים - bank conflicts preview¶

ה-shared memory מהיר, אבל יש לו מבנה פנימי שיכול להאט אותנו אם ניגשים אליו לא נכון. הזיכרון מחולק ל-32 בנקים (banks), ברוחב 4 בתים כל אחד, כך שמילים בנות 32 סיביות עוקבות נופלות בבנקים עוקבים. ה-32 בנקים מתאימים בדיוק ל-32 ה-threads של warp. הכלל:

אם 32 ה-threads של warp ניגשים ל-32 כתובות שנופלות ב-32 בנקים שונים - הגישה מתבצעת במחזור אחד, מלוא המהירות.
אם כל ה-threads ניגשים לאותה כתובת בדיוק - זהו broadcast (שידור), וגם הוא מהיר, מחזור אחד.
אבל אם שני threads (או יותר) ניגשים לכתובות שונות שנופלות באותו בנק - זו bank conflict (התנגשות בנקים), וה-hardware מסדר (serialize) את הגישות. התנגשות של דרך-k מאיטה את הגישה פי k.

ב-mmTiled שלנו, לולאת החישוב As[ty][k] * Bs[k][tx] דווקא נוחה: As[ty][k] היא אותה כתובת לכל 32 ה-threads של warp נתון (broadcast), ו-Bs[k][tx] רצה על tx העוקב (בנקים שונים). אבל דפוסי גישה אחרים - למשל קריאת עמודה של tile, שבה threads עוקבים קופצים בקפיצות של TILE_WIDTH - יכולים למפות את כולם לאותו בנק ולגרום להתנגשות של עד 32 דרכים. הטריק הקלאסי לשבירת ההתנגשות הוא ריפוד (padding) של הtile בעמודה מדומה אחת:

__shared__ float As[TILE_WIDTH][TILE_WIDTH + 1];   // +1 shifts the bank mapping

הריפוד ב-+1 מזיז כל שורה בבנק אחד, כך שגישה לעמודה שלמה פוגעת ב-32 בנקים שונים במקום באחד. זה מבזבז מעט shared memory אך יכול לתת שיפור ניכר. את הטיפול המלא בהתנגשויות בנקים - הכללים המדויקים, איך למדוד אותן עם ncu, ומתי הריפוד עוזר - נעשה בפרק 8.4. כאן די לזכור: ה-shared memory אינו קופסה שטוחה; יש לו 32 בנקים, ודפוס גישה גרוע מסדר אותם בטור ומאבד את היתרון.

סיכום¶

ה-matmul הנאיבי memory-bound כי הוא מבצע כ-FLOP אחת לכל קריאה מ-global memory, בעוד ה-H100 דורש כ-20 FLOPs לבית כדי להיות compute-bound; הפתרון הוא להעלות את arithmetic intensity.
ה-shared memory מוכרז ב-__shared__, משותף לכל ה-threads של ה-block, ויושב פיזית ב-L1 של ה-SM; יש צורה סטטית (גודל בקומפילציה) וצורה דינמית (extern __shared__ + הפרמטר השלישי של <<< >>> במספר בתים).
תבנית ה-staging היא: טעינה שיתופית global -> shared, מחסום, חישוב מתוך shared, מחסום שני, ואז המשך או כתיבה חזרה ל-global.
ב-matmul מרוצף כל block צובר tile של C בשלבים, כשבכל שלב נטענים tile של A וtile של B ל-shared memory; כל איבר טעון מנוצל TILE_WIDTH פעמים, וarithmetic intensity קופצת מ-1 ל-TILE_WIDTH FLOPs לכל קריאה.
ל-TILE_WIDTH = 16 מקבלים 16 FLOPs לקריאה (512 קריאות global לעומת 8192 FLOPs בכל שלב של block), והמהירות מזנקת בסדר גודל של פי 10 עד פי 15 לעומת הנאיבי.
נדרשים בדיוק שני __syncthreads(): מחסום הטעינה (מונע קריאה-לפני-כתיבה) לפני החישוב, ומחסום הניצול-חוזר (מונע כתיבה-לפני-קריאה, דריסת הtile) אחריו; השמטת השני היא הבאג הקלאסי שנותן תוצאות שגויות ולא דטרמיניסטיות.
__syncthreads() חייב במסלול שכל ה-threads של ה-block מגיעים אליו, אחרת ה-block נתקע ב-deadlock.
ה-shared memory שכל block צורך מגביל את מספר ה-blocks התושבים על SM ולכן את ה-occupancy; בtiles קטנים המגביל הוא בדרך כלל מכסת ה-threads, אך tile גדול הופך את ה-shared memory לצוואר (הרחבה בפרק 8.1, החלוקה הסטטית מוגבלת ל-48 KiB לכל block).
ה-shared memory מחולק ל-32 בנקים ברוחב 4 בתים; גישה שבה threads שונים פוגעים באותו בנק מסתדרת בטור (bank conflict), וריפוד [TILE_WIDTH + 1] הוא הטריק הקלאסי לשבירתה (טיפול מלא בפרק 8.4).