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 row = get_global_id(0);   // row = lane index
int col = get_global_id(1);

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

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

lane 0 -> C[0*colsB + col]
lane 1 -> C[1*colsB + col]   <- lompat colsB float
lane 2 -> C[2*colsB + col]   <- 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%

Catatan: Di host, kalau col dipetakan ke get_global_id(0), maka global[0] = colsB dan global[1] = rowsA.

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