A primer on GPUs and kernel optimization

Pendahuluan: Ketika Otak Komputer Mulai Ngopi dan Ngerokok

Halo, Bosque sejagad raya! Kembali lagi dengan sang baginda teknologinya kaum pelarian rig karatan, si Wong Edan yang budiman tapi rada sengklek kalau udah pegang terminal Linux. Pernah nggak lu, pas lagi asyik-asyiknya training model LLM (Large Language Models) atau nambang aset digital, tiba-tiba GPU lu bunyi kaya baling-baling helikopter Chinook mau evakuasi tapi FPS-nya seret kayak kuah bakso sisa kemarin? Nah, itu tandanya GPU lu lagi ngambek, bukan karena kurang doa, tapi karena Kernel Optimization lu acakadut, Bos!

Banyak dari kita mikir kalau beli GPU seri sultan dengan VRAM seluas lapangan bola bakal otomatis bikin kode jalan secepat kilat. Padahal, tanpa pemahaman mendalam tentang arsitektur perangkat keras dan bagaimana kernel dieksekusi, lu cuma buang-buang daya listrik PLN. Di artikel super panjang, super detail, dan tentunya rada edan ini, kita bakal bedah tuntas apa itu GPU, anatomi jeroannya, kenapa kernel optimization itu harga mati, dan bagaimana teknik-teknik mutakhir seperti automated discovery dan operator fusion mengubah peta jalan komputasi paralel modern. Singsingkan lengan baju, seduh kopi hitam lu, dan mari kita mulai pembantaian baris kode ini!

1. Anatomi Dasar GPU: Kenapa Dia Bukan CPU yang Disuntik Steroid?

Kesalahan fatal pemula (dan kadang senior yang lagi khilaf) adalah memperlakukan GPU (Graphics Processing Unit) seperti CPU (Central Processing Unit) tapi dengan jumlah inti yang dibanyakin. Padahal secara filosofis, mereka ini beda aliran. CPU ituibarat seorang manajer proyek jenius yang bisa ngerjain tugas-tugas rumit secara sekuensial dengan latensi super rendah—dia punya cache gede dan logika branch prediction yang kompleks.

Sebaliknya, GPU itu ibarat pasukan semut sebanyak 10.000 ekor yang disuruh mindahin gula pasir satu-satu. Tiap inti (core) di GPU itu relatif “bodoh” dan lambat kalau dibandingin sama core CPU secara individual. Tapi, kekuatan mereka ada di Massive Parallelism (paralelisme masif). Menurut literatur dari Aman.ai Primer on GPU Architecture, menjalankan model AI pada GPU NVIDIA melibatkan lapisan-lapisan terkoordinasi, mulai dari definisi model tingkat tinggi hingga pengiriman operator ke CUDA kernels yang sangat dioptimalkan.

Mari kita lihat perbandingan struktur dasarnya:

CPU: Didesain untuk meminimalkan latensi eksekusi utas tunggal (single-thread latency). Punya sedikit core (misal 8 hingga 64 core), tapi masing-masing core sangat kuat dan dilengkapi cache L1/L2/L3 berukuran jumbo.
GPU: Didesain untuk memaksimalkan throughput (kapasitas total). Terdiri dari ribuan ALU (Arithmetic Logic Units) yang dikelompokkan dalam Streaming Multiprocessors (SM) pada arsitektur NVIDIA.

Ketika lu nge-dispatch pekerjaan ke GPU, lu nggak ngirim satu instruksi buat dikerjakan sendiri. Lu ngirim grid ribuan threads yang diatur dalam blok-blok. Kalau struktur data lu salah naruh di memori, siap-siap aja nunggu antrean sembako digital yang nggak ada ujungnya.

2. Memahami Kernel: Apaan Sih Sebenarnya ‘GPU Kernel’?

Banyak orang salah kapok ngira kernel di sini adalah Linux Kernel yang ngurusin memori sistem operasi lu. Bukan, Bos! Dalam konteks komputasi paralel, GPU Kernel adalah sebuah fungsi yang ditulis dalam bahasa pemrograman khusus (seperti CUDA untuk NVIDIA, HIP untuk AMD ROCm, atau bahasa tingkat menengah seperti Triton) yang dieksekusi secara serentak (simultan) oleh ribuan thread di atas hardware GPU.

Bayangin lu mau ngaliin dua matrik raksasa berukuran 10000×10000. Kalau pakai CPU konvensional, lu bakal bikin perulangan for bersarang (nested loop) yang bikin CPU ngos-ngosan ngecek satu per satu indeksnya. Di GPU, lu nulis satu buah fungsi kernel yang cuma ngitung satu elemen tunggal dari matriks hasil perkalian tersebut. Nah, fungsi tunggal itu nantinya dicopy paste dan dijalankan secara paralel oleh jutaan thread sekaligus!

Namun, di sinilah letak masalah utamanya. Berdasarkan diskusi hangat di Hacker News mengenai A primer on GPUs and kernel optimization, disebutkan bahwa pada masa kini di mana banyak kernel GPU performan tinggi digenerate atau di-hill-climb secara otomatis, kesalahan konfigurasi kernel menjadi sumber error terbesar yang sering ditemui. Kalau thread-thread ini berebutan akses memori tanpa koordinasi yang bener, terjadilah yang namanya Memory Bottleneck.

3. Hierarki Memori GPU: Jangan Biarkan ALU Nganggur Kelaparan

Satu hal yang paling sering bikin programmer pemula nangis darah di pojokan ruangan adalah mengabaikan hierarki memori GPU. ALU di GPU itu cepet banget kerjanya, bagaikan kesetanan nyiapin perhitungan matematika. Masalahnya, mereka sering kelaparan data karena sibuk nunggu data dikirim dari RAM utama (VRAM global memory).

Mari kita bedah hierarki memori GPU dari yang paling lambat sampai yang secepat kilat:

Global Memory (VRAM): Kapasitasnya paling gede (bisa 24GB, 48GB, atau lebih di kartu enterprise), tapi aksesnya paling lambat. Berdasarkan The Ultra-Scale Playbook dari Hugging Face, inisialisasi CUDA kernel standar saja biasanya langsung memakan sekitar 1-2 GB memori GPU sekadar untuk overhead pustaka dasar (misalnya saat lu manggil import torch; torch.ones((1, 1)).to("cuda")).
L2 Cache: Tempat singgah sementara sebelum data masuk ke SM. Ukurannya lumayan tapi tetap terbatas.
Shared Memory / Scratchpad: Memori super cepat yang dikontrol secara manual oleh programmer (atau compiler kernel). Memori ini dibagi-bagi di dalam satu Streaming Multiprocessor (SM). Kalau data yang sering diakses bisa ditaruh di sini, performa kernel bisa melonjak drastis.
Registers: Memori paling cepat yang melekat langsung di ALU. Variabel lokal di dalam kernel lu biasanya tinggal di sini. Kalau register lu habis (register spilling), compiler bakal terpaksa mindahin datanya ke memori lokal yang jauh lebih lambat.

Berikut adalah contoh kerangka kode sederhana dalam menulis kernel CUDA dasar (menggunakan sintaks C++/CUDA) untuk operasi penjumlahan vektor sederhana:


__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements) {
C[i] = A[i] + B[i];
}
}

Perhatikan baris int i = blockDim.x * blockIdx.x + threadIdx.x;. Ini adalah formula sakral untuk ngitung indeks global thread. Kalau formula ini salah hitung barang satu angka saja, lu bakal kena Segmentation Fault atau yang lebih serem: Silent Data Corruption (data korup tanpa pesan error, tahu-tahu hasil training AI lu jadi ngawur nyembah berhala).

4. Teknik Optimasi Mutakhir: Dari Operator Fusion sampai Automated Discovery

Zaman sekarang, nulis kernel CUDA pakai tangan kosong dari baris pertama sampai akhir itu udah mulai ditinggalkan para dewa koding, kecuali buat mereka yang hobi masokis atau bikin riset底层 (low-level research). Dunia perkopian teknologi sekarang udah bergeser ke arah otomatisasi dan kompilasi cerdas.

Operator Fusion di PyTorch 2.0

Kenapa sih model AI zaman now bisa jauh lebih ngebut dibanding beberapa tahun lalu? Jawabannya ada pada Operator Fusion. Seperti yang dibahas dalam artikel How PyTorch 2.0 Accelerates Deep Learning with Operator Fusion and Code Generation, PyTorch memperkenalkan fungsi torch.compile() yang mampu men-generate kernel teroptimasi secara dinamis untuk GPU menggunakan OpenAI Triton.

Secara tradisional, kalau lu punya operasi seperti $ReLU(Activation(A + B))$, GPU harus melakukan hal berikut secara terpisah:

Ambil matriks A dan B dari VRAM, jumlahkan, lalu tulis hasilnya kembali ke VRAM.
Baca hasil penjumlahan itu dari VRAM, lakukan fungsi aktivasi, lalu tulis lagi ke VRAM.
Baca lagi dari VRAM untuk fungsi ReLU, lalu tulis lagi ke VRAM.

Kebiasaan bolak-balik nulis dan baca VRAM ini disebut Memory Bandwidth Bound. Dengan Operator Fusion, compiler menggabungkan beberapa operasi matematika tersebut menjadi satu kernel tunggal. Data dijumlahkan, langsung dikenai fungsi aktivasi, langsung dikenai ReLU di dalam register/shared memory, baru kemudian hasil akhirnya ditulis sekali saja ke VRAM. Kecepatan naiknya bisa berlipat-lipat ganda, Bos!

Automated GPU Kernel Discovery (OpenEvolve & Hill-Climbing)

Beralih ke ranah riset terdepan, bagaimana cara industri mencari kernel tercepat di muka bumi? Topik ini ramai didiskusikan di komunitas Machine Learning terkait OpenEvolve dan riset automated GPU kernel discovery. Alih-alih manusia yang pusing tujuh keliling mikirin unrolling loop atau sinkronisasi warp, sistem otomatis menggunakan teknik evolutionary search atau hill-climbing algorithms untuk men-generate ribuan variasi kernel, mengujinya secara real-time di atas hardware fisik, lalu bermutasi mencari kode tercepat.

Meskipun demikian, seperti catatan di Hacker News tadi, tantangan terbesar dari kernel yang di-generate otomatis ini adalah potensi error logika halus yang sulit dideteksi jika tidak divalidasi dengan unit testing yang ketat.

5. Profiling Tools: Jangan Jadi Dukun, Gunakan Alat Ukur!

Sering banget Wong Edan nemuin junior yang sok tahu pas ngadepin masalah performa GPU. “Wah, ini pasti gara-gara thread block-nya kurang gede!” teriaknya sambil ngasal ubah angka. Padahal dia nggak pernah nyolok alat ukur sama sekali.

Buat jadi optimizer handal, lu wajib berteman akrab dengan tool profiling resmi. Jangan pakai tebak-tebakan ala dukun beranak. Gunakan perangkat tempur berikut:

NVIDIA Nsight Systems: Digunakan untuk melihat timeline aplikasi secara makro. Kapan CPU ngirim kerjaan ke GPU, di mana terjadi jeda (idle time), dan bagaimana transfer data PCIe berjalan.
NVIDIA Nsight Compute: Senjata pamungkas untuk melihat jeroan kernel secara mikro. Tool ini bakal ngasih tahu lu apakah kernel lu Compute Bound (ALU sibuk mentok) atau Memory Bound (ALU kelaparan nunggu data VRAM), persentase penggunaan shared memory, hingga masalah warp divergence.
Triton Profiler / PyTorch Profiler: Buat anak-anak AI yang pakai ekosistem Python, tool ini ngebantu banget melacak operator mana yang makan waktu eksekusi paling lama.

Kalau lu mengabaikan profiling, optimasi yang lu lakukan itu ibarat menembak dalam gelap: peluru habis, target nggak kena, tetangga malah ngamuk karena rumahnya bolong.

6. Jebakan Batman dalam Kernel Optimization

Sebelum kita tutup kitab suci ini, mari kita bahas beberapa dosa besar yang sering bikin kernel optimization gagal total:

Warp Divergence: Ingat bahwa GPU mengeksekusi instruksi dalam kelompok yang disebut warp (biasanya 32 thread sekaligus). Kalau di dalam warp itu ada percabangan kondisi (misal if-else) di mana separuh thread masuk ke if dan separuh lagi ke else, GPU terpaksa mengeksekusi kedua jalur tersebut secara berurutan (serial), bukan paralel! Kinerja langsung terjun bebas.
Bank Conflicts pada Shared Memory: Shared memory dibagi ke dalam beberapa bank memori terpisah yang bisa diakses secara simultan. Kalau ada beberapa thread dalam satu warp yang mencoba mengakses alamat memori di bank yang sama secara bersamaan, terjadi antrean (bank conflict) yang bikin akses melambat drastis.
Coalesced Memory Access: Saat thread membaca data dari Global Memory, pastikan akses memorinya berurutan (contiguous). Kalau thread 0 baca alamat 0, thread 1 baca alamat 1, dan seterusnya, hardware bisa menggabungkannya dalam satu transaksi memori tunggal (coalesced). Kalau mencar-mencar tidak jelas, bandwidth VRAM bakal terbuang sia-sia.

Kesimpulan: Menjadi Master Kernel Tanpa Harus Masuk Rumah Sakit Jiwa

Wahai para pejuang bit dan bait, menguasai arsitektur GPU dan teknik kernel optimization bukanlah sesuatu yang bisa dipelajari dalam semalam sambil leyeh-leyeh nonton drakor. Ini butuh kombinasi pemahaman perangkat keras tingkat rendah, penguasaan tool profiling yang tepat, serta kesabaran ekstra menghadapi compiler yang kadang hobi ngasih pesan error penuh teka-teki.

Dengan memanfaatkan teknologi modern seperti torch.compile() untuk operator fusion, memahami hierarki memori dengan benar, serta memanfaatkan otomatisasi kernel saat dibutuhkan, rig kentang lu yang tadinya megap-megap bakal berubah jadi monster komputasi yang siap melibas beban kerja seberat apapun. Ingat pesan Wong Edan: Kode yang efisien bukan cuma soal jalan, tapi soal bagaimana hardware lu nggak perlu nangis darah tiap kali tombol ‘Execute’ ditekan!

Sampai jumpa di artikel teknis edan berikutnya, tetap semangat, jangan lupa stretching, dan salam CUDA buat kita semua!