Memory Coalescing - Revision history

SanzSkuy at 04:46, 23 June 2026

2026-06-23T04:46:11Z

← Older revision		Revision as of 04:46, 23 June 2026
Line 88:		Line 88:
	\|}		\|}

	{{Catatan\|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}		{{Catatan\|1=Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}

	== Kasus khusus: broadcast ==		== Kasus khusus: broadcast ==

SanzSkuy at 04:43, 23 June 2026

2026-06-23T04:43:54Z

← Older revision		Revision as of 04:43, 23 June 2026
Line 6:		Line 6:

	* '''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 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!		* '''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.		Hasil yang dibawa pulang '''sama persis''' (8 kotak). Tapi yang pertama jauh lebih cepat, cuma karena posisi kotaknya nempel.
Line 14:		Line 14:
	Terjemahan istilahnya:		Terjemahan istilahnya:
	* '''Truk sekali jalan''' = satu '''transaksi memory'''.		* '''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).		* '''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).		* '''Anak''' = satu ''work-item'' (atau ''lane'' di dalam warp).

Line 21:		Line 21:
	=== Warp / wavefront ===		=== 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'''.		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 ===		=== get_global_id(0) menentukan lane ===
Line 43:		Line 43:

	<syntaxhighlight lang="c">		<syntaxhighlight lang="c">
	int x = get_global_id(0); // x = lane index ~~(baris)~~		int row = get_global_id(0); // row = lane index
	int y = get_global_id(1); ~~// y = kolom~~		int col = get_global_id(1);

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

	Dalam satu warp, <code>y</code> sama untuk semua lane dan <code>x</code> jalan 0, 1, 2, ... Jadi alamatnya:		Dalam satu warp, <code>col</code> sama untuk semua lane dan <code>row</code> jalan 0, 1, 2, ... Jadi alamatnya:

	<pre>		<pre>
	lane 0 -> C[0*colsB + y]		lane 0 -> C[0*colsB + col]
	lane 1 -> C[1*colsB + y] <- lompat colsB float		lane 1 -> C[1*colsB + col] <- lompat colsB float
	lane 2 -> C[2*colsB + y] <- lompat colsB float lagi		lane 2 -> C[2*colsB + col] <- lompat colsB float lagi
	...		...
	</pre>		</pre>
Line 92:		Line 92:
	== Kasus khusus: broadcast ==		== Kasus khusus: broadcast ==

	Akses ke B di kernel itu, <code>B[k * colsB + y]</code>, '''nggak''' gantung ke <code>x</code> 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.		Akses ke B di kernel itu, <code>B[k * colsB + y]</code>, '''nggak''' gantung ke <code>x</code> 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 ==		== Aturan cepat ngecek ==
Line 107:		Line 107:
	== Kenapa penting ==		== 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.		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 <code>__local</code> 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).		Langkah optimasi berikutnya setelah coalescing biasanya '''tiling pakai <code>__local</code> 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).

SanzSkuy at 03:49, 23 June 2026

2026-06-23T03:49:38Z

← Older revision		Revision as of 03:49, 23 June 2026
Line 88:		Line 88:
	\|}		\|}

	{{~~Tip~~\|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}		{{Catatan\|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}

	== Kasus khusus: broadcast ==		== Kasus khusus: broadcast ==

SanzSkuy at 03:47, 23 June 2026

2026-06-23T03:47:39Z

← Older revision		Revision as of 03:47, 23 June 2026
Line 88:		Line 88:
	\|}		\|}

	{{~~Note~~\|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}		{{Tip\|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}

	== Kasus khusus: broadcast ==		== Kasus khusus: broadcast ==

SanzSkuy at 03:46, 23 June 2026

2026-06-23T03:46:31Z

← Older revision		Revision as of 03:46, 23 June 2026
Line 88:		Line 88:
	\|}		\|}

	{{~~Catatan~~\|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}		{{Note\|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}

	== Kasus khusus: broadcast ==		== Kasus khusus: broadcast ==

SanzSkuy: Created page with "== 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..."

2026-06-23T03:44:51Z

Created page with "== 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..."

New page

== 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 <code>get_global_id(0)</code> = 0, 1, 2, 3, ... berurutan. Ini bukan soal nama variabel, tapi soal '''indeks mana yang nempel ke <code>get_global_id(0)</code>'''.

=== 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 <code>C = A × B</code>.

=== Versi strided (lambat) ===

<syntaxhighlight lang="c">
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
</syntaxhighlight>

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

<pre>
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
...
</pre>

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

=== Versi coalesced (cepat) ===

Cukup '''tuker''' dimensi mana yang ke <code>get_global_id(0)</code>:

<syntaxhighlight lang="c">
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
</syntaxhighlight>

Sekarang <code>row</code> sama untuk semua lane dan <code>col</code> jalan 0, 1, 2, ... → alamat <code>...+0, ...+1, ...+2</code> = '''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.

{| class="wikitable"
|+ Strided vs coalesced
! Aspek !! Strided (<code>id(0)</code> = baris) !! Coalesced (<code>id(0)</code> = kolom)
|-
| Alamat tetangga (per lane) || lompat <code>colsB</code> || lompat 1
|-
| Transaksi per warp || banyak (sampai 32) || 1
|-
| Bandwidth kepakai || ~ 1 / <code>colsB</code> (boros) || ~ 100%
|}

{{Catatan|Di host, kalau <code>col</code> dipetakan ke <code>get_global_id(0)</code>, maka <code>global[0] = colsB</code> dan <code>global[1] = rowsA</code>.}}

== Kasus khusus: broadcast ==

Akses ke B di kernel itu, <code>B[k * colsB + y]</code>, '''nggak''' gantung ke <code>x</code> 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 <code>get_global_id(0)</code>.
# Tanya: '''kalau indeks itu naik 1, alamat memory-nya naik 1 juga nggak?'''
#* Naik 1 → '''coalesced''' (bagus).
#* Naik <code>colsB</code> / <code>colsA</code> / angka lain → '''strided''' (mapping kebalik, perlu dituker).

Pegangan: '''dimensi yang contiguous di memory harus dipetakan ke <code>get_global_id(0)</code>.'''

== 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 <code>__local</code> 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 ==

* [[Matmul]]
* [[OpenCL indexing]]
* [[Local memory tiling]]

[[Category:GPU]]
[[Category:OpenCL]]
[[Category:Performance]]