IF Notes

Hirarki Memori CUDA - Global, Shared, Register, Constant

5 min read

Back to IF3130 Sistem Paralel dan Terdistribusi

Hierarki Memori CUDA: Memanfaatkan Kecepatan untuk Performa Optimal

Questions/Cues

  • Mengapa memori jadi bottleneck?

  • Apa saja tipe memori di CUDA?

  • Apa itu Global Memory?

  • Apa itu Shared Memory?

  • Apa itu Registers?

  • Apa perbedaan Local vs. Shared Memory?

  • Apa itu Constant Memory?

  • Bagaimana scope & lifetime setiap memori?

  • Di mana variabel dideklarasikan?

Reference Points

  • 7 - IF-3230-07-GPU-03-2022.pdf

  • 6a - IF3230-06a-GPU-2022.pdf

Mengapa Memori Menjadi Bottleneck?

Banyak algoritma paralel, terutama pada GPU, bersifat memory-bound, bukan compute-bound. Artinya, performa mereka tidak dibatasi oleh seberapa cepat GPU bisa menghitung, melainkan oleh seberapa cepat data bisa disuplai ke unit komputasi.

Contoh: Perkalian Matriks Naif

Dalam perkalian matriks sederhana, untuk setiap elemen hasil, sebuah thread harus membaca satu baris penuh dari matriks pertama dan satu kolom penuh dari matriks kedua. Ini berarti setiap elemen data di matriks input dibaca berulang kali dari memori utama (Global Memory) yang lambat.

  • Rasio Komputasi-Memori: Perbandingannya sangat buruk. Untuk setiap operasi multiply-add, diperlukan dua akses memori. GPU dengan kekuatan 1500 GFLOPS mungkin memerlukan 6000 GB/s bandwidth memori, sementara perangkat kerasnya hanya menyediakan ~200 GB/s. Akibatnya, GPU menghabiskan sebagian besar waktunya menunggu data, dan hanya mencapai sebagian kecil dari performa puncaknya.

Solusinya adalah dengan cerdas memanfaatkan hierarki memori yang disediakan CUDA untuk meminimalkan akses ke memori yang lambat.

Tipe-Tipe Memori di CUDA

CUDA menyediakan beberapa jenis memori yang berbeda, masing-masing dengan karakteristik kecepatan, ukuran, scope, dan lifetime yang unik.

Global Memory

  • Lokasi: Off-chip (DRAM). Ini adalah memori VRAM utama GPU.

  • Ukuran: Paling besar (Gigabyte).

  • Kecepatan: Paling lambat, dengan latensi tinggi.

  • Scope: Grid. Semua thread dari semua block dalam satu pemanggilan kernel dapat mengaksesnya. Host (CPU) juga dapat membaca dan menulis ke memori ini.

  • Lifetime: Aplikasi. Data tetap ada selama Host tidak membebaskannya (cudaFree()).

  • Fungsi Utama: Komunikasi data antara Host dan Device, serta sebagai penyimpanan utama untuk dataset yang besar.

Shared Memory

  • Lokasi: On-chip, di dalam setiap Streaming Multiprocessor (SM).

  • Ukuran: Kecil (puluhan atau ratusan Kilobyte per SM).

  • Kecepatan: Sangat cepat, latensinya mendekati register.

  • Scope: Block. Data hanya dapat diakses oleh thread-thread di dalam block yang sama.

  • Lifetime: Block. Data hanya ada selama block tersebut dieksekusi. Begitu block selesai, isinya hilang.

  • Fungsi Utama: Kunci untuk performa tinggi. Digunakan sebagai cache yang dikelola secara manual oleh programmer untuk berbagi data antar thread dalam satu block dan secara drastis mengurangi lalu lintas ke Global Memory. Dideklarasikan dengan __shared__.

Registers

  • Lokasi: On-chip, di dalam setiap core komputasi.

  • Ukuran: Sangat kecil (ribuan register per SM, dibagi untuk semua thread).

  • Kecepatan: Paling cepat. Akses instan.

  • Scope: Thread. Setiap thread memiliki set registernya sendiri yang tidak bisa dilihat oleh thread lain.

  • Lifetime: Thread. Nilai ada selama thread dieksekusi.

  • Fungsi Utama: Menyimpan variabel lokal yang sering diakses (variabel otomatis dalam kernel). Compiler akan berusaha menempatkan variabel di register sebisa mungkin.

Local Memory

  • Lokasi: Off-chip (secara fisik berada di Global Memory).

  • Kecepatan: Lambat, sama seperti Global Memory.

  • Scope: Thread. Private untuk setiap thread.

  • Lifetime: Thread.

  • Fungsi Utama: Digunakan secara otomatis oleh compiler sebagai “tumpahan” (register spilling) ketika tidak ada cukup register yang tersedia, atau untuk variabel lokal berukuran besar seperti array. Penggunaannya harus dihindari karena dapat menurunkan performa secara signifikan.

Constant Memory

  • Lokasi: Off-chip, namun dengan cache on-chip.

  • Ukuran: Kecil (biasanya 64 KB).

  • Kecepatan: Cepat jika terjadi cache hit. Data di-broadcast ke semua thread dalam satu warp, sangat efisien untuk akses seragam.

  • Scope: Grid. Read-only untuk semua thread, ditulis oleh Host.

  • Lifetime: Aplikasi.

  • Fungsi Utama: Menyimpan data yang tidak berubah selama eksekusi kernel dan dibaca oleh banyak thread (misalnya, koefisien filter, konstanta fisika). Dideklarasikan dengan __constant__.

Ringkasan Scope dan Lifetime

Deklarasi VariabelMemoriScopeLifetime
int varLokal; (di dalam kernel)RegisterThreadThread
__device__ __shared__ int varBersama;SharedBlockBlock
__device__ int varGlobal;GlobalGridAplikasi
__device__ __constant__ int varKonstan;ConstantGridAplikasi

  • __device__ bersifat opsional jika digunakan bersama __shared__ atau __constant__.

  • Variabel yang dideklarasikan di luar fungsi apa pun secara default masuk ke Global Memory.

  • Variabel otomatis (tanpa qualifier) di dalam kernel secara default masuk ke Register.

Summary

CUDA menyediakan hierarki memori yang kaya, mulai dari Register on-chip yang super cepat namun terbatas hingga Global Memory off-chip yang besar namun lambat. Kunci untuk mencapai performa tinggi di GPU adalah menulis kernel yang memaksimalkan penggunaan memori on-chip (Register dan Shared Memory) untuk komputasi internal dan meminimalkan akses ke Global Memory. Shared Memory, khususnya, adalah alat vital yang memungkinkan thread dalam satu block untuk bekerja sama secara efisien dengan berbagi data dan mengurangi ketergantungan pada memori utama yang lambat.

Graph View

Backlinks