8.4 Coalescing ו bank conflicts פתרון
פתרון - Coalescing ו-bank conflicts¶
נעבור תרגיל אחרי תרגיל עם הקוד, הפקודות, החישובים ביד והפלט הצפוי. כל המספרים כאן הם מ-H100 (SXM) עם CUDA 12.4; רוחבי הפס והזמנים המושגים משתנים מעט בין הרצות, דרייברים ודורות, והם ערכים אופייניים - העיקר הוא היחסים, הצורה של העקומות, וספירות ה-sectors וההתנגשויות שהן דטרמיניסטיות. מי שעובד על T4 ב-Colab יראה מספרים מוחלטים אחרים לגמרי (למשל stride 1 סביב 206 GB/s כמו בטבלת הגלוסרי), אבל אותה תמונה בדיוק.
פתרון תרגיל 1 - בניית מדד ה-stride ושחזור קריסת ה-coalescing¶
// strided.cu
#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)
__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];
acc = acc * 1.000000119f + v;
}
if (t < N) out[t] = acc;
}
int main(int argc, char** argv) {
const size_t N = 64ull << 20; // 64M floats = 256 MB
const size_t bytes = N * sizeof(float);
float *d_in, *d_out;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMemset(d_in, 1, bytes));
const int grid = 4096, block = 256;
int strides[] = {1, 2, 4, 8, 16, 32, 64, 128};
cudaEvent_t a, b;
CUDA_CHECK(cudaEventCreate(&a));
CUDA_CHECK(cudaEventCreate(&b));
const int ITERS = 30;
for (int s : strides) {
strided_read_kernel<<<grid, block>>>(d_in, d_out, N, s); // warmup
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaEventRecord(a));
for (int k = 0; k < ITERS; ++k)
strided_read_kernel<<<grid, block>>>(d_in, d_out, N, s);
CUDA_CHECK(cudaEventRecord(b));
CUDA_CHECK(cudaEventSynchronize(b));
float ms = 0.f;
CUDA_CHECK(cudaEventElapsedTime(&ms, a, b));
double sec = (ms / 1e3) / ITERS;
double usefulBytes = ((double)N / s) * sizeof(float); // N/stride elements
double gbps = usefulBytes / sec / 1e9;
printf("stride=%3d : %8.1f GB/s\n", s, gbps);
}
CUDA_CHECK(cudaFree(d_in)); CUDA_CHECK(cudaFree(d_out));
return 0;
}
$ nvcc -arch=sm_90 -O3 strided.cu -o strided
$ ./strided
stride= 1 : 2870.4 GB/s
stride= 2 : 1735.9 GB/s
stride= 4 : 980.2 GB/s
stride= 8 : 512.6 GB/s
stride= 16 : 270.1 GB/s
stride= 32 : 233.7 GB/s
stride= 64 : 205.4 GB/s
stride=128 : 171.8 GB/s
הצורה מדויקת: בין stride 1 ל-8 כל הכפלה כמעט חוצה את רוחב הפס (2870 -> 1736 -> 980 -> 513), ואז מ-stride 16 הדעיכה מתמתנת (270 -> 234 -> 205 -> 172). המספרים המוחלטים גבוהים מ-T4 כי ה-HBM3 של H100 (peak 3350 GB/s) מהיר בהרבה מ-GDDR6 של T4, אבל היחס stride1/stride128 (~פי 17) זהה כמעט למדידת הגלוסרי (206/11.2 ~= פי 18).
למה זה עבד: ב-stride 1 כל warp גורף 128B רציפים ל-transaction אחד - ניצול מלא. כל הכפלת stride מכניסה חצי מהבייטים השימושיים לכל burst של 128B, ולכן צריך כפול transactions - ומכאן החצייה. מ-stride 16 כל thread כבר צורך burst נפרד ממילא, וצוואר הבקבוק עובר ל-TLB וללוקליות, ולכן העקומה משתטחת. איך להכליל: זהו המדד הקנוני ל-coalescing. השתמשו בו לכל דפוס גישה חשוד: אם רוחב הפס המושג נופל בהרבה מ-peak וגדל כשמסדרים מחדש את הגישה - הבעיה ב-coalescing.
פתרון תרגיל 2 - אישור ה-coalescing עם ncu¶
הפכנו את ה-stride לארגומנט (atoi(argv[1])) והרצנו פעמיים:
$ ncu --metrics \
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio \
./strided 1
strided_read_kernel(const float *, float *, unsigned long, int)
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio 4.0
$ ncu --metrics \
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio \
./strided 32
strided_read_kernel(const float *, float *, unsigned long, int)
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio 32.0
בגישה מאוחדת (stride 1): 4 sectors לבקשה בדיוק - warp שקורא 32 floats = 128B = 4 sectors של 32B, כולם שימושיים. בגישה מפוזרת (stride 32): 32 sectors לבקשה - כל thread נופל ב-sector אחר, וכל 32 ה-threads מכריחים 32 sectors נפרדים, שמכל אחד רק 4B שימושיים. יעילות של 4/32 = 12.5%, בדיוק הצניחה שמדדנו בתרגיל 1.
למה זה עבד: 128 בייט ה-transaction מחולקים ל-4 sectors של 32B. המדד sectors_per_request מודד כמה sectors פיזיים כל בקשה לוגית של warp מפעילה. 4 = אידיאלי; 32 = הגרוע ביותר. איך להכליל: ה-sectors_per_request הוא מדד ה-coalescing הישיר ביותר - הוא מודד את המנגנון עצמו (transactions לבקשה), בעוד ה-GB/s מתרגיל 1 מודד את התוצאה (רוחב פס שנגזר גם מ-TLB, מ-occupancy ומדברים אחרים). כשרוצים לבודד את בעיית ה-coalescing, מסתכלים על ה-sectors.
פתרון תרגיל 3 - תיקון kernel לא מאוחד ומדידת הרווח¶
// rowsum.cu (main part)
__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[(size_t)row * N + col]; // neighboring threads N*4 bytes apart
s[row] = acc;
}
// one block per row; threadIdx.x runs over the columns -> coalesced read
__global__ void rowsum_good(const float* __restrict__ A,
float* __restrict__ s, int M, int N) {
int row = blockIdx.x;
__shared__ float red[256];
float acc = 0.f;
for (int col = threadIdx.x; col < N; col += blockDim.x)
acc += A[(size_t)row * N + col]; // neighboring threads -> consecutive addresses
red[threadIdx.x] = acc;
__syncthreads();
for (int off = blockDim.x / 2; off > 0; off >>= 1) {
if (threadIdx.x < off) red[threadIdx.x] += red[threadIdx.x + off];
__syncthreads();
}
if (threadIdx.x == 0) s[row] = red[0];
}
הפעלה: rowsum_bad<<<(M+255)/256, 256>>> מול rowsum_good<<<M, 256>>>.
$ nvcc -arch=sm_90 -O3 rowsum.cu -o rowsum
$ ./rowsum
M=8192 N=8192
rowsum_bad : 4.812 ms
rowsum_good: 0.361 ms
speedup : x13.3
max |diff| : 0.0 (both versions agree)
אישור ב-ncu:
$ ncu --metrics \
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio \
./rowsum
rowsum_bad(...) l1tex__..._op_ld.ratio 32.0
rowsum_good(...) l1tex__..._op_ld.ratio 4.0
הגרסה הרעה מראה 32 sectors לבקשה, הטובה 4 - וזה בדיוק ה-x13.3 שמדדנו. אותה כמות עבודה, אותם בייטים לוגיים, רק סדר הגישה השתנה.
למה זה עבד: ב-rowsum_bad, בכל איטרציה threads שכנים (row, row+1) קוראים כתובות במרחק N*4 = 32KB - 32 transactions נפרדים. ב-rowsum_good, ה-threadIdx.x רץ על העמודות, כך שבכל צעד ה-warp קורא 32 עמודות רצופות - transaction אחד מאוחד. איך להכליל: הכלל היחיד שצריך לזכור הוא ש-threadIdx.x ירוץ על הממד הרציף בזיכרון. למערך row-major זה הממד הפנימי (העמודה); ל-column-major (למשל פלט של cuBLAS, שיעור 6) זה השורה. כשקוד מעבד מטריצה לאורך הממד הלא-נכון, היפוך תפקידי thread ו-loop הוא התיקון.
פתרון תרגיל 4 - transpose נאיבי והתנגשות בנקים 32-way¶
// transpose.cu
#define TILE 32
__global__ void transposeNaive(const float* __restrict__ in,
float* __restrict__ out,
int width, int height) {
__shared__ float tile[TILE][TILE];
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[(size_t)y * width + x]; // coalesced from global
__syncthreads();
x = blockIdx.y * TILE + threadIdx.x;
y = blockIdx.x * TILE + threadIdx.y;
if (x < height && y < width)
out[(size_t)y * height + x] = tile[threadIdx.x][threadIdx.y]; // column -> conflict
}
הפעלה: dim3 block(32, 8); dim3 grid(width/32, height/32); (עם לולאה פנימית או block(32,32); כאן נשתמש ב-block(32,32) לפשטות הניתוח).
החישוב ביד (סעיף 3). בקריאה tile[threadIdx.x][threadIdx.y], ב-warp יחיד threadIdx.y קבוע ו-threadIdx.x רץ 0..31. הכתובת במילים:
address = 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 all 32 threads of the warp!
כל 32 ה-threads על אותו בנק (threadIdx.y), בכתובות שונות (שורות שונות של tile) - התנגשות 32-way.
$ nvcc -arch=sm_90 -O3 transpose.cu -o transpose
$ ./transpose
4096 x 4096 transpose
transposeNaive : 0.298 ms (correct: yes)
$ ncu --metrics \
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum \
./transpose
transposeNaive(const float *, float *, int, int)
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum 16,252,928
מעל 16 מיליון התנגשויות - החתימה המובהקת של הבעיה.
למה זה עבד: הקריאה והכתיבה ל-global מאוחדות (ה-threadIdx.x על הממד הרציף), אז ה-global נקי. אבל קריאת ה-shared בעמודה מכריחה את כל ה-warp על בנק יחיד. הגורם threadIdx.x * 32 תמיד מתחלק ב-32, ולכן "נעלם" ב-mod ומשאיר את כולם על אותו בנק. איך להכליל: כל פעם שמערך shared דו-ממדי עם ממד פנימי בגודל 32 (או כפולה של 32) נקרא בעמודה (הממד החיצוני משתנה בין threads של warp), מובטחת התנגשות 32-way. זהו הדפוס לחפש.
פתרון תרגיל 5 - ריפוד ל-[32][33] ומדידת השיפור¶
השינוי הוא תו אחד:
החישוב ביד (סעיף 2). עכשיו stride השורה הוא 33 מילים:
address = 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.x + threadIdx.y) mod 32 עובר על כל 32 הבנקים בדיוק פעם אחת - אפס התנגשויות.
$ nvcc -arch=sm_90 -O3 transpose.cu -o transpose
$ ./transpose
4096 x 4096 transpose
transposeNaive : 0.298 ms (correct: yes)
transposePadded : 0.121 ms (correct: yes)
speedup : x2.46
$ ncu --metrics \
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum \
./transpose
transposeNaive(...) l1tex__...op_ld.sum 16,252,928
transposePadded(...) l1tex__...op_ld.sum 0
ההתנגשויות צנחו מ-16 מיליון ל-0, והזמן השתפר פי ~2.5. (הרווח כאן פחות מפי 32 כי ה-kernel נשלט גם על ידי תעבורת ה-global; הקריאה מ-shared כבר לא צוואר הבקבוק.) עלות הריפוד: 33/32 = 1.031, כלומר ~3% shared memory נוסף - 32*33*4 = 4224 בייט במקום 4096.
למה זה עבד: 33 = 32 + 1 זר ראשוני ל-32, ולכן threadIdx.x * 33 mod 32 = threadIdx.x - כל thread נופל על בנק אחר. שילמנו עמודה מדומה אחת (~3% מקום) וקנינו גישה מקבילית מלאה. איך להכליל: כשמערך shared נקרא בעמודה עם ממד פנימי שכפולה של 32, הוסיפו +1 (או כל מספר זר ל-32) לממד הפנימי. אבל אמתו קודם עם ncu שיש התנגשות - ריפוד עיוור למערך שכבר נגיש בבנקים שונים סתם מבזבז shared ומוריד occupancy.
פתרון תרגיל 6 - שידור מול התנגשות (broadcast)¶
__global__ void bcast(const float* g, float* out) {
__shared__ float data[1024];
data[threadIdx.x] = g[threadIdx.x];
__syncthreads();
float v = data[0]; // the whole warp -> the exact same address
out[blockIdx.x * blockDim.x + threadIdx.x] = v;
}
$ ncu --metrics \
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum \
./bcast
bcast(const float *, float *)
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum 0
אפס התנגשויות, למרות שכל 32 ה-threads קוראים מבנק 0 - כי הם קוראים את אותה כתובת (data[0]), וזה שידור (broadcast): החומרה קוראת פעם אחת ומשדרת ל-32.
עכשיו data[(threadIdx.x % 2) * 32] - חצי מה-threads על מילה 0, חצי על מילה 32, שתיהן בבנק 0 אך שתי כתובות מובחנות:
זו התנגשות 2-way: שתי כתובות מובחנות בבנק אחד -> 2 מחזורים (בתוך כל קבוצת-אותה-כתובת החומרה עדיין משדרת). הקנס הוא כמספר הכתובות המובחנות בבנק, לא כמספר ה-threads.
טבלת שלושת המצבים:
pattern | result | penalty
-----------------------------------+-----------------+--------------
different banks (data[tid]) | parallel | 1x (none)
same bank, same address (data[0]) | broadcast | 1x (none)
same bank, different addresses | N-way conflict | Nx (# addresses)
למה זה עבד: התנאי להתנגשות דורש שני מרכיבים - אותו בנק וגם כתובות שונות. שידור מפר את התנאי השני (אותה כתובת), ולכן חינמי. הקנס נקבע לפי מספר הכתובות המובחנות בבנק, כי בתוך כל כתובת החומרה משדרת. איך להכליל: קריאה אחידה של קבוע מ-shared (כל ה-warp קורא אותו אלמנט) תמיד חינמית - נצלו זאת. חשש להתנגשות מוצדק רק כשה-threads קוראים כתובות שונות שנופלות על אותו בנק.
פתרון תרגיל 7 (בונוס) - למה אחת ל-global והשנייה ל-shared¶
coalescing לא רלוונטי ל-shared (סעיף 1): ה-128B transaction הוא תכונה של ה-DRAM (GDDR/HBM) שמגבה את ה-global memory - burst שמביא שורת זיכרון שלמה. ל-shared memory אין DRAM ואין bursts; הוא SRAM על השבב שנגיש ברזולוציה של מילה בודדת דרך 32 הבנקים. לכן אין מה "לאחד" - אין transaction גדול שאפשר לבזבז.
bank conflicts לא רלוונטיים ל-global (סעיף 2): 32 הבנקים הם ארגון פיזי של ה-SRAM המשותף בתוך ה-SM. ל-global memory (ב-HBM, מחוץ ל-SM) אין את המבנה הזה; היחידה שלו היא ה-cache line / ה-sector של 32B וה-transaction של 128B. לכן "בנק" הוא מושג חסר משמעות ל-global.
טבלת 2x2 (סעיף 3):
| global memory | shared memory
--------------+--------------------------+---------------------
coalescing | lives here (128B burst) | meaningless
bank conflict | meaningless | lives here (32 banks)
התרחיש (סעיף 4): ב-kernel שמעתיק global -> shared, מעבד ב-shared, וכותב shared -> global:
- בשלב הקריאה מ-global ובשלב הכתיבה ל-global דואגים ל-coalescing: ש-
threadIdx.xירוץ על הממד הרציף כדי שכל warp יגרוף 128B מאוחדים. - בשלב העיבוד ב-shared דואגים ל-bank conflicts: שדפוס הגישה לא יפיל כמה threads על אותו בנק, ובמידת הצורך מרפדים (
[32][33]).
זהו בדיוק המבנה של transpose מתרגילים 4-5: קריאה/כתיבה מאוחדות ל-global, ועיבוד נקי-בנקים ב-shared.
למה זה עבד: שתי האופטימיזציות עונות על אותה שאלה - "האם דפוס הכתובות של ה-warp ידידותי?" - אבל על שני זיכרונות עם שני מנגנוני חומרה שונים לגמרי (bursts של DRAM מול בנקי SRAM). לכל זיכרון החוק שלו. איך להכליל: בכל kernel זהו את שלבי ה-global ואת שלבי ה-shared בנפרד, והפעילו את הכלי הנכון בכל שלב - sectors_per_request ל-global, bank_conflicts ל-shared. אל תחפשו בנקים ב-global ואל תחפשו coalescing ב-shared.