8.4 Coalescing ו bank conflicts הרצאה

בשיעור 1.6 בנינו את פירמידת הזיכרון של ה-GPU וראינו ששני הצינורות האיטיים שקובעים את גורל ה-kernel הם הדרך אל ה-global memory (זיכרון גלובלי) שיושב ב-HBM, והדרך אל ה-shared memory (זיכרון משותף) שיושב על השבב לצד ה-SM. בשיעור 3.3 השתמשנו ב-shared memory כדי לבצע tiling ולהעלות את ה-arithmetic intensity, ובפרק 7 למדנו לזהות מתי kernel memory-bound וכמה הוא רחוק מרוחב הפס התיאורטי. השיעור הזה נותן את שני הכלים המעשיים הגדולים ביותר לסגור את הפער הזה: coalescing (איחוד גישות) בזיכרון הגלובלי, והימנעות מ-bank conflicts (התנגשות בנקים) בזיכרון המשותף. שתי האופטימיזציות עוסקות באותו רעיון - דפוס הכתובות שכל 32 ה-threads של warp נוגעים בו באותה הוראה - אבל הן פועלות על שני זיכרונות שונים, לפי שני מנגנוני חומרה שונים לחלוטין, וזהו הבלבול הנפוץ ביותר בקורס. נלמד את שתיהן זו לצד זו כדי שההבחנה תהיה חדה: coalescing שייך ל-global, bank conflicts שייכים ל-shared, ואסור לערבב.

שתי אופטימיזציות, שני זיכרונות - coalescing מול bank conflicts

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

	coalescing	bank conflict
הזיכרון	global memory (HBM)	shared memory (על השבב)
יחידת החומרה	בקר ה-DRAM וה-cache lines	32 בנקים של ה-SRAM המשותף
הדפוס הרצוי	כתובות רציפות ל-warp	כתובת בכל בנק אחר
הדפוס הרע	כתובות מפוזרות (stride גדול)	כמה threads באותו בנק
הקנס	יותר transactions של 128B	סריאליזציה, עד פי 32
התיקון	לסדר מחדש את סדר הגישה	padding / transpose של המערך

ההבדל המנטלי החשוב: ב-global memory הבעיה היא שגורפים בייטים מיותרים - כל גישה פיזית מביאה 128 בייט, ואם רק 4 מהם שימושיים, בזבזנו 97%. ב-shared memory הבעיה הפוכה: אין בזבוז בייטים אלא סריאליזציה בזמן - החומרה מחלקת גישה מתנגשת לכמה מחזורי שעון עוקבים. שתיהן פוגעות בביצועים, אבל דרך שני מנגנונים נפרדים. נתחיל מהגלובלי.

איחוד גישות בזיכרון גלובלי - memory coalescing

coalescing הוא, בלשון הגלוסרי, "טכניקת חומרה לשיפור ניצול רוחב הפס של הזיכרון על ידי שירות של כמה קריאות זיכרון לוגיות בגישה פיזית אחת". המנגנון נשען ישירות על מודל ה-SIMT שראינו בשיעור 1.4: כל 32 ה-threads של warp מנפיקים את אותה הוראת load באותו מחזור. אם 32 הכתובות שהם מבקשים רציפות בזיכרון, החומרה יכולה למזג אותן ל-transaction פיזי אחד במקום 32.

המספר הקסום הוא 128 בייט. זהו הגודל האופייני של burst בודד מה-DRAM, ולא במקרה: 32 threads כפול 4 בייט ל-float שווה בדיוק 128 בייט. כלומר warp שלם שקורא מערך של float ברצף מתמלא ב-transaction פיזי אחד:

warp of 32 threads, each reading one float (4B):

thread:   0    1    2    3   ...  31
address: 0x00 0x04 0x08 0x0C ... 0x7C      <- 128 contiguous bytes
          |____________________________|
                 one 128B transaction  (coalesced, perfect)

לעומת זאת, אם ה-threads קופצים ב-stride (צעד) גדול, אותן 32 קריאות מתפזרות על פני הרבה bursts של 128B, שמכל אחד מהם שולפים רק אחד או שניים מהערכים:

stride = 32 floats (128 bytes between thread and thread):

thread 0 -> burst A (4B useful out of 128B)
thread 1 -> burst B (4B useful out of 128B)
thread 2 -> burst C ...
...
32 separate transactions instead of one -> 32x more traffic, 1/32 useful bandwidth

שתי נקודות עקרוניות שהגלוסרי מדגיש. ראשית, זה מנוהל על ידי המתכנת. ב-CPU ה-cache עושה זאת אוטומטית ובאופן שקוף; ב-GPU אתם, המתכנתים, חייבים לסדר את הנתונים ואת דפוס הגישה במודע כך שה-thread ה-i-י ייגע באלמנט ה-i-י. שנית, coalescing חשוב בגלל ה-latency הגבוה של ה-DRAM. ל-DRAM (בין אם GDDR ובין אם HBM) רוחב פס עצום אבל גם latency ארוך - ארוך אף יותר מ-DDR5 של ה-CPU. הסתרת ה-latency הזה (שיעור 7.4) עובדת רק אם כל transaction מנצל את מלוא ה-128B; אחרת מבזבזים גם רוחב פס וגם מחזורים.

ה-kernel הקנוני שממחיש את זה הוא קריאה עם stride משתנה. שימו לב ל-__restrict__ שמבטיח לcompiler שאין aliasing בין המצביעים:

__global__ void strided_read_kernel(const float* __restrict__ in,
                                     float* __restrict__ out,
                                     size_t N, int stride) {
    const size_t t = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
    const size_t T = gridDim.x * (size_t)blockDim.x;

    float acc = 0.f;
    for (size_t j = t * (size_t)stride; j < N;
         j += T * (size_t)stride) {
        float v = in[j];               // fully coalesced when stride == 1
        acc = acc * 1.000000119f + v;  // prevents the compiler from eliminating the read
    }
    if (t < N) out[t] = acc;
}

כאשר stride == 1, threads עוקבים קוראים כתובות עוקבות (t, t+1, t+2, ...) - מאוחד מושלם. כאשר stride == k, thread ו-thread שכן מרוחקים k*4 בייט, וככל ש-k גדל כך פחות מהם נכנסים ל-burst אחד. ההצטברות ל-acc ומ הכתיבה ל-out הן רק כדי שהcompiler לא ימחק את הלולאה כ"קוד מת".

מדד ה-stride - איך ה-coalescing נשבר

הגלוסרי מודד את ה-kernel הזה על Tesla T4 (compute capability 7.5, דור Turing) על מערך של 67,108,864 floats (256 MB), לכל ערך של stride. התוצאות מספרות את כל הסיפור:

צעד - stride	רוחב פס מושג	הערה
1	206.0 GB/s	מאוחד מושלם - קו הבסיס
2	130.5 GB/s	כמעט נחצה
4	68.8 GB/s	נחצה שוב
8	33.8 GB/s	ושוב
16	16.8 GB/s	כאן משתנה האופי
32	15.2 GB/s	הדעיכה מתמתנת
64	13.6 GB/s
128	11.2 GB/s	פי ~18 איטי מ-stride 1

שני דברים בטבלה הזו חשובים מאוד. הראשון: הדעיכה אינה לינארית. בין stride 1 ל-8, כל הכפלה של ה-stride בערך חוצה את רוחב הפס - כי כל הכפלה מכניסה חצי מהבייטים השימושיים לכל burst של 128B, ולכן צריך בערך כפול bursts. אבל מ-stride 16 והלאה האופי משתנה: הדעיכה מתמתנת ומשתטחת. זה קורה כי ברגע שכל thread צורך burst נפרד ממילא, החוק שקובע כבר אינו ה-coalescing אלא החטאות ב-TLB (Translation Lookaside Buffer) וירידה בלוקליות - צוואר בקבוק של תת-מערכת אחרת לגמרי. במילים אחרות, מ-stride 16 ומעלה כבר "איבדנו" את ה-coalescing כמעט לחלוטין, ומה שנשאר הוא רק הידרדרות של locality.

הדבר השני: 11.2 GB/s מול 206 GB/s הוא פי 18 הבדל, מאותו קוד בדיוק - רק סדר הגישה השתנה. זו הסיבה ש-coalescing הוא אחד השיפורים הגדולים ביותר שאפשר לעשות ל-kernel memory-bound. על H100 המספרים המוחלטים גבוהים בהרבה (קריאה מאוחדת קרובה ל-2800-3000 GB/s מתוך 3350 GB/s של ה-HBM3), אבל צורת העקומה זהה: stride 1 גורף כמעט את מלוא רוחב הפס, וכל stride גדול יותר קורס באותו אופן.

כיצד מודדים coalescing עם profiler. ה-Nsight Compute חושף את זה ישירות דרך היחס "sectors per request". כל transaction של 128B מורכב מ-4 sectors של 32B. קריאה מאוחדת של warp מביאה בדיוק 4 sectors לכל בקשה; קריאה מפוזרת לגמרי מביאה עד 32. המדד:

ncu --metrics \
  l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio \
  ./strided

ערך של ~4 פירושו coalescing מושלם; ערך שמטפס ל-16 או 32 הוא הדגל האדום של גישה לא מאוחדת. זהו המדד שנשתמש בו בתרגול כדי לאשר שהבעיה אכן ב-coalescing ולא במשהו אחר.

תיקון kernel לא מאוחד - transpose של דפוס הגישה

הדוגמה הקלאסית ל-kernel שנשבר בגלל coalescing היא עיבוד מטריצה שבו כל thread מטפל בשורה שלמה. נניח מטריצה M x N שמורה row-major (שורה-ראשית) בזיכרון, ואנחנו רוצים לסכם כל שורה. הגרסה הנאיבית נותנת ל-thread ה-row-י את כל השורה:

// not coalesced: thread row scans row row
__global__ void rowsum_bad(const float* __restrict__ A,
                           float* __restrict__ s, int M, int N) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row >= M) return;
    float acc = 0.f;
    for (int col = 0; col < N; ++col)
        acc += A[row * N + col];        // neighboring threads -> addresses N*4 bytes apart
    s[row] = acc;
}

הבעיה: בכל איטרציה של הלולאה, threads שכנים (row ו-row+1) ניגשים לכתובות row*N+col ו-(row+1)*N+col - מרוחקות N floats זו מזו. אם N = 1024, זה stride של 4096 בייט בין thread ל-thread. גישה קטסטרופלית: 32 transactions נפרדים לכל איטרציה.

התיקון הוא להחליף את התפקידים של threads ו-loop, כך שבכל צעד ה-warp יסרוק כתובות רציפות. במקום שכל thread ייקח שורה, ניתן ל-warp שלם לטפל בשורה, כאשר בכל צעד ה-32 threads קוראים 32 עמודות רצופות:

// coalesced: a whole warp scans a row, threads read consecutive columns
__global__ void rowsum_good(const float* __restrict__ A,
                            float* __restrict__ s, int M, int N) {
    int row  = blockIdx.x;                    // one block per row
    int lane = threadIdx.x;                   // 0..blockDim.x-1
    float acc = 0.f;
    for (int col = lane; col < N; col += blockDim.x)
        acc += A[row * N + col];              // neighboring threads -> consecutive addresses!
    // reduction within the block to sum the row (lesson 8 / chapter 9)
    acc = blockReduceSum(acc);
    if (lane == 0) s[row] = acc;
}

עכשיו, בכל איטרציה, ה-lane ה-i-י קורא את A[row*N + col_base + i] - כתובות רציפות ל-warp - וה-load מאוחד לחלוטין. אותה כמות עבודה בדיוק, אותם בייטים, אבל סדר גישה שונה. ההבדל במדידה דרמטי: על מטריצה גדולה, המעבר מ-rowsum_bad ל-rowsum_good נותן בקלות שיפור של פי 8 עד פי 20, בדיוק בטווח שראינו בטבלת ה-stride. המסקנה המעשית: ה-thread שרץ מהר ביותר הוא זה שה-index שלו רץ על הממד הרציף בזיכרון. למערך row-major, ה-threadIdx.x צריך לרוץ על העמודה (הממד הפנימי), לא על השורה.

בנקים בזיכרון המשותף - shared memory banks

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

המיפוי בין כתובת לבנק פשוט: מילים עוקבות של 32 סיביות ממופות לבנקים עוקבים. המילה במיקום w (כלומר בייט 4*w) שייכת לבנק w mod 32:

address (bytes):  0   4   8  12 ... 124 | 128 132 ... 252 | 256 ...
word (w):          0   1   2   3 ...  31 |  32  33 ...  63 |  64 ...
bank (w mod 32):   0   1   2   3 ...  31 |   0   1 ...  31 |   0 ...
                  |___ 32 banks ______| |___ 0..31 again __| 

מכאן נובעת עובדה קריטית: כתובות שמרוחקות זו מזו ב-128 בייט (32 מילים) נופלות על אותו בנק. מילה 0, מילה 32, מילה 64 - כולן בבנק 0. זהו בדיוק ה-stride שיפיל אותנו. שימו לב שרוחב הבנק הוא 4 בייט כי, כלשון הגלוסרי, "ה-GPU עוצבו עם floats ו-integers של 32 סיביות בראש", לא 64.

התרחיש הטוב הוא ה-default: כשכל thread ניגש למילה עוקבת, כל 32 ה-threads נופלים על 32 בנקים שונים, וכל הגישה מסתיימת במחזור אחד:

__shared__ float data[1024];
int tid = threadIdx.x;
float value = data[tid];   // thread 0->bank 0, thread 1->bank 1, ... thread 31->bank 31

כאן אין שום התנגשות: 32 בנקים, 32 threads, אחד לכל בנק. זהו המצב שאליו נשאף בכל גישה ל-shared memory.

התנגשות בנקים - bank conflict

bank conflict מתרחש, בהגדרת הגלוסרי, "כאשר כמה threads ב-warp מבקשים בו-זמנית זיכרון בתוך אותו בנק ב-shared memory, אך בכתובות שונות". שימו לב לשני התנאים, שניהם הכרחיים: אותו בנק וגם כתובות שונות. אם התנאים מתקיימים, החומרה אינה יכולה לשרת את הגישות במקביל - בנק אחד מגיש מילה אחת למחזור - ולכן היא מסדרת אותן בטור (serialization), מחזור אחר מחזור.

הקנס הוא פקטור שלם: התנגשות בדרגה N (N threads על אותו בנק) איטית פי ~N. המקרה הקיצוני הוא התנגשות 32-way, שבה כל 32 ה-threads נופלים על בנק יחיד. אז ה-latency, שבמצב תקין הוא בסדר גודל של ~10 מחזורים, מזנק למאות מחזורים - פי ~32. הנה איך זה קורה עם stride של 32:

__shared__ float data[1024];
int tid = threadIdx.x;
float value = data[tid * 32];  // thread i -> word 32*i -> bank (32*i) mod 32 = 0

כל thread ניגש למילה 32*i, וכל אחת מהן מתחלקת ב-32 ללא שארית, ולכן כולן בבנק 0. הכתובות שונות (מילים 0, 32, 64, ...) - שני התנאים מתקיימים - ולכן 32 גישות מסודרות בטור. פי 32 latency, מאותו קוד שנראה תמים.

32-way conflict (data[tid*32]):

thread:   0    1    2   ...  31
word:     0   32   64  ...  992
bank:     0    0    0   ...   0      <- all in bank 0!
          serialization: 32 cycles instead of 1

החריג: שידור (broadcast). יש מקרה אחד שבו "אותו בנק" אינו התנגשות - כאשר כל ה-threads ניגשים לאותה כתובת בדיוק. אז אין צורך בסריאליזציה, כי החומרה יכולה לשדר (broadcast/multicast) את המילה האחת לכל ה-threads בבת אחת. לכן קריאה אחידה של קבוע יחיד מ-shared memory (למשל data[0] על ידי כל ה-warp) חינמית לחלוטין. זכרו את שלוש האפשרויות: בנקים שונים = מקבילי (טוב); אותו בנק, כתובות שונות = התנגשות (רע); אותו בנק, אותה כתובת = שידור (חינם).

תיקון: padding ו-transpose - הריפוד [32][33]

הדוגמה הקלאסית של התנגשות בנקים היא transpose (שחלוף) של מטריצה דרך ה-shared memory. הנה ה-kernel הנאיבי: כל block טוען tile של 32x32 מ-global memory אל tile, ואז כותב אותו משוחלף ל-global memory:

#define TILE 32
// naive: column read from shared -> 32-way conflict
__global__ void transposeNaive(const float* __restrict__ in,
                               float* __restrict__ out,
                               int width, int height) {
    __shared__ float tile[TILE][TILE];               // 32x32
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;
    if (x < width && y < height)
        tile[threadIdx.y][threadIdx.x] = in[y * width + x];   // coalesced read from global
    __syncthreads();

    x = blockIdx.y * TILE + threadIdx.x;             // transposed coordinates
    y = blockIdx.x * TILE + threadIdx.y;
    if (x < height && y < width)
        out[y * height + x] = tile[threadIdx.x][threadIdx.y]; // <-- column access!
}

הקריאה מ-global (in[y*width+x]) והכתיבה ל-global (out[y*height+x]) שתיהן מאוחדות יפה, כי threadIdx.x רץ על הממד הרציף. הבעיה החבויה היא בקריאה מ-shared: tile[threadIdx.x][threadIdx.y]. נחשב את הבנק שכל thread ב-warp נוגע בו. ב-warp יחיד, threadIdx.y קבוע ו-threadIdx.x רץ 0..31. הכתובת (במילים) היא threadIdx.x * 32 + threadIdx.y, והבנק:

bank = (threadIdx.x * 32 + threadIdx.y) mod 32
     = (threadIdx.x * 32) mod 32  +  threadIdx.y
     = 0 + threadIdx.y
     = threadIdx.y   (constant for the whole warp!)

כל 32 ה-threads של ה-warp נופלים על אותו בנק (threadIdx.y), בכתובות שונות (שורות שונות של tile) - התנגשות 32-way מושלמת. פי 32 latency בכל קריאה מ-shared.

התיקון האלגנטי הוא ריפוד (padding): להוסיף עמודה מדומה אחת ולהצהיר tile[32][33] במקום tile[32][32]:

    __shared__ float tile[TILE][TILE + 1];   // 32x33 - the 33rd column is padding

השורה הזו לבדה מתקנת הכל. עכשיו כל שורה של tile תופסת 33 מילים ולא 32, ולכן stride השורה הוא 33 - זר ראשוני ל-32. נחשב שוב את הבנק לקריאה tile[threadIdx.x][threadIdx.y]:

address (words) = threadIdx.x * 33 + threadIdx.y
bank = (threadIdx.x * 33 + threadIdx.y) mod 32
     = (threadIdx.x * (32 + 1) + threadIdx.y) mod 32
     = (threadIdx.x + threadIdx.y) mod 32

עכשיו כש-threadIdx.x רץ 0..31 (ו-threadIdx.y קבוע), הביטוי (threadIdx.x + threadIdx.y) mod 32 עובר על כל 32 הבנקים בדיוק פעם אחת - אפס התנגשויות. שילמנו 33/32 מקום נוסף ב-shared (בזבוז של ~3%) וקנינו פי 32 מהירות בקריאה. זהו ה"טריק" הכי מפורסם באופטימיזציית CUDA, וההסבר שלו הוא בדיוק חשבון ה-mod שעשינו.

כיצד מודדים bank conflicts עם profiler. ה-Nsight Compute סופר את מספר ההתנגשויות ישירות:

ncu --metrics \
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,\
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum \
  ./transpose

עבור transposeNaive המונה הזה יהיה גדול (מיליוני התנגשויות); אחרי הריפוד ל-[32][33] הוא צונח כמעט לאפס. זהו האישור החד-משמעי שהתיקון עבד.

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

• בלבול coalescing עם bank conflicts - הטעות מספר אחת. coalescing הוא מושג של global memory; bank conflicts הם מושג של shared memory. ל-shared אין coalescing (אין bursts של DRAM), ול-global אין בנקים. אם שמעתם "בנק" - מדובר ב-shared; אם שמעתם "128B transaction" או "stride" - מדובר ב-global.
• "stride גדול תמיד מחצה את הביצועים לינארית" - שגוי. הדעיכה לא לינארית: חצייה בכל הכפלה עד stride ~8, ואז השטחה מ-stride 16 בגלל מעבר לצוואר בקבוק של TLB ולוקליות, לא coalescing.
• "התנגשות בנקים = תמיד פי 32" - שגוי. רק התנגשות 32-way היא פי 32. התנגשות בדרגה N איטית פי N; 2-way זה פי 2, 4-way פי 4. וגם: שידור (אותה כתובת) חינמי לגמרי, גם אם כל ה-warp נוגע באותו בנק.
• ריפוד עיוור מבזבז shared בלי צורך. הוסיפו +1 רק כשיש באמת גישה שיוצרת stride שמתחלק ב-32. ריפוד למערך שכבר נגיש בבנקים שונים סתם גוזל מקום ומוריד occupancy. תמיד אמתו עם ncu שהיתה התנגשות לפני שאתם מרפדים.
• "coalescing נשלט על ידי הcompiler" - שגוי. ב-CPU ה-cache שקוף ואוטומטי; ב-GPU אתם אחראים לסדר הגישה. אותו קוד עם threadIdx.x על הממד הלא-נכון יכול להיות פי 18 איטי יותר, והcompiler לא יציל אתכם.
• שחלוף (transpose) ל-global שונה מ-transpose ל-shared. בדוגמת ה-transpose, הקריאה והכתיבה ל-global מאוחדות (טוב), והבעיה כולה בקריאת ה-shared. אל תניחו שה-kernel איטי בגלל global רק כי הוא "משוחלף" - בדקו את שני המדדים בנפרד.

סיכום

• ה-coalescing (איחוד גישות) ו-bank conflicts (התנגשות בנקים) הן שתי אופטימיזציות שונות של דפוס גישה - הראשונה ל-global memory, השנייה ל-shared memory - ואסור לערבב ביניהן: אין בנקים ב-global ואין coalescing ב-shared.
• ה-coalescing ממזג את 32 הקריאות של warp ל-transaction פיזי אחד של 128 בייט (32 threads x 4B) כשהכתובות רציפות; זה מנוהל על ידי המתכנת ותקף רק ל-global memory שמגובה ב-DRAM (GDDR/HBM).
• מדד ה-stride על T4 מראה 206 GB/s ב-stride 1 שקורס ל-11 GB/s ב-stride 128 (פי ~18): הדעיכה לא לינארית - חצייה בכל הכפלה עד stride 8, ואז השטחה מ-stride 16 בגלל מעבר לצוואר בקבוק של TLB ולוקליות.
• התיקון ל-coalescing הוא לסדר את דפוס הגישה כך ש-threadIdx.x ירוץ על הממד הרציף בזיכרון (העמודה במערך row-major); מדידה עם l1tex__average_t_sectors_per_request... צריכה להראות ~4 sectors/request במצב מאוחד מול עד 32 במפוזר.
• ל-shared memory יש 32 בנקים של 4 בייט; מילים עוקבות ממופות לבנקים עוקבים (w mod 32), וכתובות שמרוחקות 128 בייט (32 מילים) נופלות על אותו בנק.
• ה-bank conflict הוא מצב שבו כמה threads על אותו בנק בכתובות שונות -> סריאליזציה, איטי פי N (עד פי 32 בהתנגשות 32-way, מ-~10 מחזורים למאות); שידור (אותה כתובת) חינמי.
• התיקון ל-bank conflict הוא ריפוד או transpose של המערך המשותף - למשל tile[32][33] במקום tile[32][32] - כך ש-stride השורה (33) זר ל-32 והבנק הופך ל-(threadIdx.x + threadIdx.y) mod 32 שעובר על כל 32 הבנקים.
• מדידה עם ncu מאשרת כל אחת בנפרד: sectors_per_request ל-coalescing ו-data_bank_conflicts... ל-bank conflicts - תמיד אמתו את הבעיה עם המדד הנכון לפני שמתקנים, ובמיוחד לפני שמרפדים.