Memory Coalescing

Penjelasan paling gampang

Bayangin GPU punya truk pengantar barang yang aneh. Truk ini cuma bisa jalan kalau ngangkut satu peti penuh isi 8 kotak. Mau kamu butuh 1 kotak doang atau 8 kotak sekaligus, truknya tetap bawa satu peti penuh tiap jalan. Nggak bisa cuma ambil 1 kotak.

Sekarang misalkan ada 8 anak yang masing-masing butuh 1 kotak.

Kalau 8 anak ini tinggal sebelah-sebelahan (kamar 0, 1, 2, ... 7), kotak mereka kebetulan ada di satu peti yang sama. Truk cukup jalan 1 kali, satu peti, semua kebagian. Hemat!
Kalau 8 anak ini tinggal kepencar (kamar 0, 4, 8, 12, ...), tiap kotak ada di peti yang beda-beda. Truk harus jalan 4 kali, dan tiap peti cuma kepake 2 kotak dari 8 — sisanya keangkut percuma. Boros!

Hasil yang dibawa pulang sama persis (8 kotak). Tapi yang pertama jauh lebih cepat, cuma karena posisi kotaknya nempel.

Itulah coalescing. Kalau data yang diminta nempel-nempelan di memory → cepat (coalesced). Kalau kepencar → lambat (strided).

Terjemahan istilahnya:

Truk sekali jalan = satu transaksi memory.
Peti isi 8 kotak = satu cache line / segmen (di GPU asli biasanya 32 float = 128 byte, bukan 8 — 8 cuma biar gampang dibayangin).
Anak = satu work-item (atau lane di dalam warp).

Apa yang sebenarnya terjadi

Warp / wavefront

Work-item di GPU nggak dijalanin satu-satu. GPU ngumpulin mereka jadi grup — 32 di NVIDIA (disebut warp) atau 64 di AMD (wavefront) — dan satu grup ngejalanin instruksi yang sama di waktu yang sama, cuma beda data (model SIMT). Jadi pas baris baca-memory dijalanin, semua lane minta alamatnya barengan.

get_global_id(0) menentukan lane

Pas work-item disusun jadi lane, dimensi 0 yang paling cepat berubah. Artinya lane 0, 1, 2, 3, ... punya get_global_id(0) = 0, 1, 2, 3, ... berurutan. Ini bukan soal nama variabel, tapi soal indeks mana yang nempel ke get_global_id(0).

Transaksi memory

GPU nggak bisa ambil 1 float doang. Tiap akses ke global memory ngangkut satu segmen utuh (mis. 128 byte = 32 float). Jadi pertanyaan kuncinya cuma satu:

"32 alamat yang diminta satu warp itu jatuh ke berapa segmen?"

Semua masuk 1 segmen → 1 transaksi (coalesced).
Kepencar ke 32 segmen → 32 transaksi (strided), dan tiap transaksi tetap ngangkut 32 float padahal cuma 1 kepake.

Contoh: kernel matmul

Lihat baris yang nulis output di kernel C = A × B.

Versi strided (lambat)

int x = get_global_id(0);   // x = lane index (baris)
int y = get_global_id(1);   // y = kolom

// ...
C[x * colsB + y] = acc;     // y sama untuk semua lane, x naik 1 tiap lane

Dalam satu warp, y sama untuk semua lane dan x jalan 0, 1, 2, ... Jadi alamatnya:

lane 0 -> C[0*colsB + y]
lane 1 -> C[1*colsB + y]   <- lompat colsB float
lane 2 -> C[2*colsB + y]   <- lompat colsB float lagi
...

Tiap lane lompat colsB float dari tetangganya → strided. Kalau colsB gede, tiap lane jatuh di segmen beda → banyak transaksi, bandwidth kebuang.

Versi coalesced (cepat)

Cukup tuker dimensi mana yang ke get_global_id(0):

int col = get_global_id(0); // col = lane index
int row = get_global_id(1);

// ...
C[row * colsB + col] = acc; // row sama untuk semua lane, col naik 1 tiap lane

Sekarang row sama untuk semua lane dan col jalan 0, 1, 2, ... → alamat ...+0, ...+1, ...+2 = berurutan → semua jatuh ke 1 segmen → 1 transaksi, ~100% kepake.

Logika dot product-nya nol perubahan. Hasilnya sama persis. Yang berubah cuma pola akses memory-nya.

Strided vs coalesced
Aspek	Strided (`id(0)` = baris)	Coalesced (`id(0)` = kolom)
Alamat tetangga (per lane)	lompat `colsB`	lompat 1
Transaksi per warp	banyak (sampai 32)	1
Bandwidth kepakai	~ 1 / `colsB` (boros)	~ 100%

Tip: {{{1}}}

Kasus khusus: broadcast

Akses ke B di kernel itu, B[k * colsB + y], nggak gantung ke x sama sekali. Jadi semua lane minta alamat yang sama persis → GPU cukup baca sekali terus broadcast ke semua lane. Ini gratis dan udah efisien — jadi yang perlu dibenerin cuma akses A dan C.

Aturan cepat ngecek

Nggak usah mikirin warp tiap kali. Cukup:

Lihat indeks mana yang nempel ke get_global_id(0).
Tanya: kalau indeks itu naik 1, alamat memory-nya naik 1 juga nggak?
- Naik 1 → coalesced (bagus).
- Naik colsB / colsA / angka lain → strided (mapping kebalik, perlu dituker).

Pegangan: dimensi yang contiguous di memory harus dipetakan ke get_global_id(0).

Kenapa penting

Naive matmul itu memory-bound — rasio kerjanya cuma 1 multiply-add per 2 load global. Karena bottleneck-nya di memory, efisiensi tiap transaksi yang nentuin segalanya. Coalescing jelek bukan cuma rugi beberapa persen; dia ngaliin beban memory lipat ganda sesuai stride. Tuker dua baris mapping aja udah bisa ngasih speedup gede tanpa ngubah hasil.

Langkah optimasi berikutnya setelah coalescing biasanya tiling pakai __local memory: load satu blok A dan B ke local memory dengan pola coalesced sekali, terus reuse berkali-kali dari local (yang nggak kena aturan coalescing global).

Lihat juga