Nvidia CUDA Programming untuk Developer - Belajar GPU Computing dari Nol

Laptop saya panas. Kipasnya bunyi kayak pesawat lepas landas. Dan yang paling bikin frustrasi: saya cuma jalanin satu proses training model kecil, tapi CPU usage mentok 100% selama 40 menit sementara GPU Nvidia di laptop itu nganggur, diam, nggak ngapa-ngapain.

Kalau kamu pernah ngalamin hal yang sama, artikel ini buat kamu. Saya akan tunjukin cara mulai pakai CUDA, dari cek apakah GPU kamu support sampai nulis kernel pertama. Semua pakai contoh nyata yang bisa langsung kamu copy-paste.

CUDA itu bukan sihir. Dia cuma cara buat bilang ke GPU: "eh, kerjaan ini bisa kamu pecah ke ribuan core, kerjain paralel aja." Bedanya, cara mikirnya beda total sama CPU. Dan begitu kamu paham bedanya, banyak hal yang tadinya kelihatan lambat jadi masuk akal.

Kenapa CPU Nggak Cocok Buat Semua Hal

CPU kamu punya 4 sampai 16 core. Masing-masing core-nya pintar, bisa branch prediction, out-of-order execution, punya cache besar. Dia dirancang buat ngerjain satu tugas rumit dengan cepat.

GPU punya ribuan core. Tapi setiap core-nya bodoh. Nggak ada branch prediction, cache kecil, dan semua core dalam satu warp harus jalanin instruksi yang sama persis di waktu yang sama. Kalau salah satu core ambil jalur berbeda, semua core lain harus nunggu.

Jadi aturan mainnya sederhana:

• CPU logika bercabang, I/O, parsing, hal-hal yang urutannya nggak bisa diprediksi. Contoh: membaca file JSON, menangani request HTTP, mengurai konfigurasi.
• GPU operasi matematika identik yang diulang jutaan kali. Contoh: perkalian matriks, konvolusi gambar, hashing password, simulasi partikel.

Perkalian matriks 5000x5000 itu sekitar 125 miliar operasi floating point. CPU dengan 8 core mungkin butuh puluhan detik. GPU yang bisa ngerjain 10.000 operasi sekaligus? Hitungannya jadi di bawah satu detik. Tapi kalau kamu mau sorting array 10 elemen, GPU malah lebih lambat karena overhead transfer data dan launching kernel lebih besar dari kerjaannya sendiri.

Cek Dulu: GPU Kamu Support CUDA atau Nggak

Sebelum install apa-apa, pastikan GPU kamu punya. CUDA cuma jalan di GPU Nvidia. AMD pakai ROCm atau HIP, Intel pakai oneAPI. Kalau kamu punya AMD, artikel ini tetap berguna buat konsepnya, tapi tooling-nya beda.

# Cek GPU yang terdeteksi
lspci | grep -i nvidia

# Cek driver dan versi CUDA runtime yang terpasang
nvidia-smi

Output nvidia-smi bakal nunjukin sesuatu kayak gini:

+-----------------------------------------------------------------------------+
| NVIDIA-SMI 550.54.14 Driver Version: 550.54.14 CUDA Version: 12.4 |
|-------------------------------+----------------------+----------------------+
| GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. |
|===============================+======================+======================|
| 0 NVIDIA GeForce RTX 3060 On | 00000000:01:00.0 On | N/A |
| 45% 52C P8 18W / 170W | 412MiB / 12288MiB | 0% Default |
+-------------------------------+----------------------+----------------------+

Yang penting di cek: Driver Version dan CUDA Version. Angka CUDA di situ adalah versi maksimum yang bisa dipakai driver kamu. Kalau mau pakai toolkit yang lebih baru, driver harus di-update dulu.

Kalau nvidia-smi bilang "command not found", berarti driver belum terpasang. Di Ubuntu, cara paling aman:

# Cek driver yang direkomendasikan
ubuntu-drivers devices

# Install otomatis driver yang direkomendasikan
sudo ubuntu-drivers autoinstall

# Reboot wajib
sudo reboot

Jangan pernah download driver .run langsung dari situs Nvidia kalau kamu pakai Ubuntu dengan secure boot. Nanti kejadiannya kernel module ditolak, dan kamu bakal stuck di layar hitam. Pakai repository distro, hidup jadi lebih mudah.

Install CUDA Toolkit

Setelah driver beres, baru install toolkit. Toolkit ini isinya compiler nvcc, library (cuBLAS, cuFFT, cuRAND), dan profiler.

# Tambah repository Nvidia
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2404/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt update

# Install toolkit (sesuaikan versi dengan yang ada di repo)
sudo apt install -y cuda-toolkit-12-4

# Tambahkan ke PATH
echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc

# Verifikasi
nvcc --version

Kalau nvcc --version keluar outputnya, berarti compiler siap. Ukuran install-nya besar, sekitar 4-6 GB, jadi pastikan disk kamu cukup.

Program CUDA Pertama: Vector Addition

Contoh klasik yang wajib kamu tulis sekali seumur hidup. Ini program CUDA paling sederhana yang lengkap: alokasi memori GPU, transfer data, launch kernel, ambil hasil.

// add.cu
#include <stdio.h>
#include <cuda_runtime.h>

// Kernel: jalan di GPU, dipanggil dari host (CPU)
__global__ void vectorAdd(const float *A, const float *B, float *C, int n) {
 int i = blockIdx.x * blockDim.x + threadIdx.x;
 if (i < n) {
 C[i] = A[i] + B[i];
 }
}

int main() {
 const int N = 1 << 20; // sekitar 1 juta elemen
 const size_t bytes = N * sizeof(float);

 // 1. Alokasi memori di host
 float *h_A = (float*)malloc(bytes);
 float *h_B = (float*)malloc(bytes);
 float *h_C = (float*)malloc(bytes);

 for (int i = 0; i < N; i++) {
 h_A[i] = 1.0f;
 h_B[i] = 2.0f;
 }

 // 2. Alokasi memori di device (GPU)
 float *d_A, *d_B, *d_C;
 cudaMalloc(&d_A, bytes);
 cudaMalloc(&d_B, bytes);
 cudaMalloc(&d_C, bytes);

 // 3. Copy host -> device
 cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
 cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

 // 4. Launch kernel
 int threadsPerBlock = 256;
 int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
 vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

 // 5. Copy hasil device -> host
 cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);

 // Cek error
 cudaError_t err = cudaGetLastError();
 if (err != cudaSuccess) {
 printf("CUDA error: %s\n", cudaGetErrorString(err));
 return 1;
 }

 printf("C[0]=%f C[last]=%f\n", h_C[0], h_C[N-1]);

 // 6. Bersihkan
 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
 free(h_A); free(h_B); free(h_C);
 return 0;
}

Compile dan jalankan:

nvcc -o add add.cu
./add
# Output: C[0]=3.000000 C[last]=3.000000

Perhatikan tiga baris yang bikin CUDA beda dari C biasa:

• __global__ fungsi ini jalan di GPU, tapi dipanggil dari CPU. Ini yang disebut kernel.
• <<<blocksPerGrid, threadsPerBlock>>> konfigurasi eksekusi. 4096 block, 256 thread per block berarti 1.048.576 thread jalan sekaligus.
• blockIdx.x * blockDim.x + threadIdx.x cara setiap thread menghitung "saya nomor berapa". Ini pola paling fundamental di CUDA.

Paham Model Eksekusi: Grid, Block, Thread

Ini bagian yang paling sering bikin orang bingung. Analoginya begini: bayangkan kamu punya perusahaan dengan beberapa divisi (block), setiap divisi punya beberapa karyawan (thread), dan semua divisi diatur oleh satu manajer besar (grid).

// Organisasi 1D
dim3 block(256); // 256 thread per block
dim3 grid((N + 255) / 256); // jumlah block dibulatkan ke atas

// Organisasi 2D cocok untuk gambar atau matriks
dim3 block(16, 16); // 256 thread, 16x16
dim3 grid((width + 15) / 16, (height + 15) / 16);

// Di dalam kernel, akses koordinat
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;

Kenapa harus dibulatkan ke atas? Karena jumlah thread harus kelipatan ukuran block. Kalau N = 1000 dan block = 256, kamu butuh 4 block (1024 thread). 24 thread terakhir nggak punya kerjaan. Makanya di kernel ada pengecekan if (i < n). Kalau lupa, kamu bakal nulis di luar batas array dan hasilnya sampah atau crash.

Kesalahan yang Pasti Kamu Lakukan Pertama Kali

Saya sudah melakukannya semua. Ini daftar lengkapnya supaya kamu nggak perlu buang waktu berjam-jam.

1. Lupa cek error launch

Kernel launch itu asynchronous. Artinya vectorAdd<<<...>>>() langsung balik ke baris berikutnya sebelum kernel selesai. Kalau ada error, kamu nggak tahu sampai nanti.

// WRONG error silently ignored
myKernel<<<grid, block>>>(data);
processResult(data); // pakai data yang belum siap!

// CORRECT cek error, lalu sinkronkan
myKernel<<<grid, block>>>(data);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
 printf("Launch failed: %s\n", cudaGetErrorString(err));
 return 1;
}
cudaDeviceSynchronize(); // tunggu sampai kernel selesai

Aturan praktis: selalu pasang cudaGetLastError() setelah setiap launch. Tanpa itu, kamu debugging dalam kegelapan total.

2. Transfer memori bolak-balik di dalam loop

Ini penyebab paling umum kenapa "CUDA saya lebih lambat dari CPU". Transfer RAM ke VRAM lewat PCIe cuma sekitar 12 GB/detik. Bandwidth VRAM sendiri bisa 400+ GB/detik.

// WRONG transfer 1000 kali
for (int step = 0; step < 1000; step++) {
 cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice);
 kernel<<<grid, block>>>(d_data);
 cudaMemcpy(h_data, d_data, bytes, cudaMemcpyDeviceToHost);
}

// CORRECT transfer sekali, kerja di GPU sampai selesai
cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice);
for (int step = 0; step < 1000; step++) {
 kernel<<<grid, block>>>(d_data);
}
cudaMemcpy(h_data, d_data, bytes, cudaMemcpyDeviceToHost);

Prinsipnya: data datang sekali, diolah berkali-kali, baru pulang sekali. Kalau kamu masih transfer di dalam loop, semua keuntungan GPU hilang dimakan overhead.

3. Ukuran block yang asal-asalan

Ukuran block 256 atau 512 hampir selalu jadi pilihan aman. Kenapa? Karena GPU memproses thread dalam grup 32 yang disebut warp. Ukuran block yang kelipatan 32 berarti nggak ada warp yang setengah kosong.

// Buruk 100 bukan kelipatan 32, ada warp yang cuma terisi 4 thread
dim3 block(100);

// Bagus kelipatan 32, semua warp terisi penuh
dim3 block(256); // 8 warp penuh
dim3 block(512); // 16 warp penuh

Batasan keras: maksimal 1024 thread per block. Lebih dari itu, kernel gagal launch dengan error invalid configuration argument.

4. Akses memori yang nggak berurutan

Ini soal coalescing. Kalau thread 0 baca elemen 0, thread 1 baca elemen 1, dan seterusnya GPU bisa ambil semua data itu dalam satu transaksi. Kalau aksesnya melompat-lompat, jadi banyak transaksi.

// Bagus coalesced, akses berurutan
int i = blockIdx.x * blockDim.x + threadIdx.x;
float val = input[i];

// Buruk strided, akses melompat
int i = (blockIdx.x * blockDim.x + threadIdx.x) * 4;
float val = input[i];

Untuk matriks yang disimpan row-major, ini berarti thread berurutan harus mengakses kolom yang sama saat melakukan iterasi baris. Terdengar aneh, tapi ini yang bikin perbedaan performa bisa 10x.

5. Nggak cek kapasitas VRAM

Kalau cudaMalloc gagal karena VRAM habis, dia balik dengan pointer NULL. Kalau kamu langsung pakai, hasilnya segfault yang susah dilacak.

float *d_big;
cudaError_t err = cudaMalloc(&d_big, huge_bytes);
if (err != cudaSuccess) {
 printf("OOM: butuh %zu MB, cuma ada %zu MB free\n",
 huge_bytes / (1024*1024), getFreeVRAM() / (1024*1024));
 return 1;
}

Dari CUDA ke PyTorch: Jalan Pintas yang Realistis

Jujur saja: kalau tujuan kamu adalah machine learning, kamu nggak perlu nulis kernel CUDA manual. PyTorch sudah punya ribuan kernel yang sudah dioptimasi bertahun-tahun oleh orang yang jauh lebih paham hardware dari kita.

import torch

# Cek GPU tersedia
print(torch.cuda.is_available())
print(torch.cuda.get_device_name(0))
print(f"VRAM: {torch.cuda.get_device_properties(0).total_memory / 1e9:.1f} GB")

# Pindahkan tensor ke GPU
device = torch.device("cuda" if torch.cuda.is_available() else "cpu")
a = torch.randn(5000, 5000, device=device)
b = torch.randn(5000, 5000, device=device)

import time

# Di CPU
a_cpu = a.cpu(); b_cpu = b.cpu()
t0 = time.time()
c_cpu = a_cpu @ b_cpu
print(f"CPU: {time.time() - t0:.3f}s")

# Di GPU (sinkronkan dulu supaya pengukuran akurat)
torch.cuda.synchronize()
t0 = time.time()
c_gpu = a @ b
torch.cuda.synchronize()
print(f"GPU: {time.time() - t0:.3f}s")

Di RTX 3060, matriks 5000x5000 biasanya selesai sekitar 0.05 detik di GPU versus 3-8 detik di CPU. Perbedaan 60-150x. Itu bukan angka marketing, itu memang begitu hasilnya.

Yang penting diingat: torch.cuda.synchronize() wajib dipanggil sebelum mengukur waktu. Tanpa itu, kamu mengukur waktu antrean perintah, bukan waktu eksekusi sebenarnya. Saya pernah buang satu jam gara-gara ini hasil benchmark-nya kelihatan 1000x lebih cepat, ternyata cuma karena pengukurannya salah.

Kapan CUDA Benar-Benar Layak Dipakai

Nggak semua pekerjaan cocok dipindah ke GPU. Ini panduan praktis berdasarkan pengalaman saya.

• Layak: perkalian matriks skala besar, training model deep learning, rendering, simulasi fisika, pemrosesan gambar batch, hashing password massal (dengan catatan etis), analisis data genomik.
• Nggak layak: query database, parsing JSON, logika bisnis dengan banyak percabangan, pekerjaan yang datanya kecil (di bawah 100 ribu elemen), proses yang butuh latensi rendah satu per satu.
• Abu-abu: sorting data besar. Bisa di-GPU dan memang cepat, tapi implementasinya ribet dan sering nggak sepadan usahanya.

Ukuran cepat: kalau data kamu nggak muat di VRAM, atau kamu butuh transfer bolak-balik lebih dari beberapa kali, GPU kemungkinan besar bukan jawabannya. Perbaiki dulu algoritmanya di CPU sering kali masalahnya bukan di hardware.

Profiling: Cara Tahu Kernel Kamu Cepat atau Cuma Kelihatan Cepat

Jangan pernah percaya intuisi soal performa. Ukur. Nvidia punya dua tool bawaan yang wajib kamu kenal.

# nsys timeline level sistem, lihat kapan GPU idle
nsys profile --stats=true ./my_program

# ncu analisis detail per kernel, termasuk memory throughput
ncu --set full ./my_program

# Kalau ncu nggak ada, pakai nvprof (versi lama)
nvprof ./my_program

Yang paling sering saya temukan dari profiling: waktu habis bukan di kernel, tapi di transfer memori. Kernel-nya cuma 5ms, tapi transfernya 200ms. Kalau kamu nggak profiling, kamu akan mengoptimasi kernel yang salah.

Metrik yang perlu kamu lihat di ncu:

• Memory Throughput % kalau di atas 80%, kernel kamu sudah mendekati batas bandwidth. Nggak banyak lagi yang bisa diperbaiki.
• Compute Throughput % kalau rendah sementara memory tinggi, kernel kamu memory-bound. Optimasi akses memori, bukan aritmatika.
• Occupancy persentase slot thread yang terisi. Di bawah 30% berarti GPU kamu banyak menganggur.
• Warp Stall Reasons alasan kenapa warp menunggu. Ini petunjuk paling langsung untuk perbaikan.

Kalau GPU Kamu Nggak Support CUDA

Jangan berkecil hati. Ada beberapa jalan alternatif yang realistis:

• Google Colab gratis, dapat GPU T4. Cukup buat belajar dan eksperimen. Batas waktunya sekitar 12 jam per sesi.
• Kaggle Notebooks gratis, GPU P100 atau T4, 30 jam per minggu. Bagus buat eksperimen lebih serius.
• Cloud GPU sewaan RunPod, Vast.ai, atau Lambda Labs. Mulai dari sekitar $0.20 per jam untuk GPU kelas menengah. Jauh lebih murah daripada beli GPU sendiri kalau cuma butuh sesekali.
• AMD dengan ROCm sudah cukup matang di 2026 untuk PyTorch. Kodenya nyaris identik, cuma perlu install ROCm versi PyTorch.
• Apple Silicon dengan MPS PyTorch support Metal Performance Shaders. Nggak secepat Nvidia, tapi cukup untuk prototyping.

Untuk belajar konsep CUDA, Colab sudah lebih dari cukup. Kamu bisa nulis dan compile kode .cu langsung di sana.

Checklist Praktis Sebelum Mulai

# 1. GPU terdeteksi?
nvidia-smi

# 2. Driver versi berapa?
nvidia-smi --query-gpu=driver_version --format=csv,noheader

# 3. Toolkit terpasang?
nvcc --version

# 4. PyTorch lihat GPU?
python3 -c "import torch; print(torch.cuda.is_available(), torch.cuda.device_count())"

# 5. Berapa VRAM bebas?
nvidia-smi --query-gpu=memory.free --format=csv,noheader

# 6. Ada proses lain yang pakai GPU?
nvidia-smi --query-compute-apps=pid,process_name,used_memory --format=csv

Kalau semua langkah di atas lolos, kamu siap. Kalau ada yang gagal, balik ke bagian install masalahnya 90% ada di driver, bukan di kode kamu.

Kesimpulan

CUDA itu bukan hal mistis. Intinya cuma tiga: pindahkan data ke VRAM sekali, pecah kerja ke ribuan thread, ambil hasilnya sekali. Yang bikin orang menyerah biasanya bukan konsepnya, tapi error yang nggak jelas asalnya dan itu semua bisa dihindari dengan cek error yang rajin dan profiling yang jujur.

Buat kamu yang baru mulai, urutannya begini: pastikan driver dan toolkit jalan, tulis vector addition sampai berhasil, pahami model grid-block-thread, baru pindah ke PyTorch untuk kerjaan nyata. Jangan langsung loncat ke custom kernel 95% kebutuhan real-world sudah tercakup library yang ada.

Kalau kamu sudah pernah coba CUDA, apa masalah pertama yang bikin kamu stuck? Driver yang nggak mau ke-load, atau kernel yang hasilnya salah terus tanpa pesan error? Ceritakan di kolom komentar mungkin pengalaman kamu bisa bantu pembaca lain yang lagi di posisi yang sama.