CEN310 Paralel Programlama
Hafta-11 (İleri GPU Programlama)
Bahar Dönemi, 2024-2025
Genel Bakış
Konular
- CUDA Bellek Modeli
- Paylaşımlı Bellek Optimizasyonu
- İş Parçacığı Senkronizasyonu
- Performans Optimizasyon Teknikleri
Hedefler
- CUDA bellek hiyerarşisini anlamak
- Paylaşımlı bellek kullanımını öğrenmek
- İş parçacığı senkronizasyonunda ustalaşmak
- Optimizasyon stratejilerini uygulamak
1. CUDA Bellek Modeli
Bellek Türleri
- Global Bellek
- Paylaşımlı Bellek
- Sabit Bellek
- Doku Belleği
- Yazmaçlar
Bellek Erişim Desenleri
// Birleştirilmiş bellek erişimi örneği
__global__ void birlesik_erisim(float* veri, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
// Birleştirilmiş erişim deseni
float deger = veri[idx];
// Değeri işle
veri[idx] = deger * 2.0f;
}
}
2. Paylaşımlı Bellek Optimizasyonu
Paylaşımlı Bellek Kullanımı
__global__ void matris_carpimi(float* A, float* B, float* C, int N) {
__shared__ float paylasimliA[BLOK_BOYUTU][BLOK_BOYUTU];
__shared__ float paylasimliB[BLOK_BOYUTU][BLOK_BOYUTU];
int satir = blockIdx.y * blockDim.y + threadIdx.y;
int sutun = blockIdx.x * blockDim.x + threadIdx.x;
float toplam = 0.0f;
for(int karo = 0; karo < N/BLOK_BOYUTU; karo++) {
// Veriyi paylaşımlı belleğe yükle
paylasimliA[threadIdx.y][threadIdx.x] =
A[satir * N + karo * BLOK_BOYUTU + threadIdx.x];
paylasimliB[threadIdx.y][threadIdx.x] =
B[(karo * BLOK_BOYUTU + threadIdx.y) * N + sutun];
__syncthreads();
// Paylaşımlı bellek kullanarak hesapla
for(int k = 0; k < BLOK_BOYUTU; k++) {
toplam += paylasimliA[threadIdx.y][k] * paylasimliB[k][threadIdx.x];
}
__syncthreads();
}
C[satir * N + sutun] = toplam;
}
3. İş Parçacığı Senkronizasyonu
Senkronizasyon Yöntemleri
- Blok seviyesi senkronizasyon
- Izgara seviyesi senkronizasyon
- Atomik işlemler
Örnek: Atomik İşlemler
__global__ void histogram(int* veri, int* hist, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
atomicAdd(&hist[veri[idx]], 1);
}
}
Optimizasyon Teknikleri
- Bellek Birleştirme
- Bank Çakışması Önleme
- Doluluk Optimizasyonu
- Döngü Açma
Örnek: Bank Çakışması Çözümü
// Kötü: Bank çakışmaları
__shared__ float paylasimli_veri[BLOK_BOYUTU][BLOK_BOYUTU];
// İyi: Bank çakışmalarını önlemek için dolgulu
__shared__ float paylasimli_veri[BLOK_BOYUTU][BLOK_BOYUTU + 1];
İleri Bellek Yönetimi
Birleşik Bellek
// Birleşik bellek ayır
float* birlesik_veri;
cudaMallocManaged(&birlesik_veri, boyut);
// Ana bilgisayar veya cihazdan erişim
// Açık transfer gerekmiyor
cekirdek<<<izgara, blok>>>(birlesik_veri);
// Birleşik belleği serbest bırak
cudaFree(birlesik_veri);
Akış İşleme
Eşzamanlı Yürütme
cudaStream_t akis1, akis2;
cudaStreamCreate(&akis1);
cudaStreamCreate(&akis2);
// Farklı akışlarda asenkron işlemler
cekirdek1<<<izgara, blok, 0, akis1>>>(veri1);
cekirdek2<<<izgara, blok, 0, akis2>>>(veri2);
cudaStreamSynchronize(akis1);
cudaStreamSynchronize(akis2);
cudaStreamDestroy(akis1);
cudaStreamDestroy(akis2);
Dinamik Paralellik
İç İçe Çekirdek Başlatma
__global__ void cocuk_cekirdek(float* veri) {
// Çocuk çekirdek kodu
}
__global__ void ebeveyn_cekirdek(float* veri) {
if(threadIdx.x == 0) {
cocuk_cekirdek<<<izgara, blok>>>(veri);
cudaDeviceSynchronize();
}
}
Laboratuvar Alıştırması
Görevler
- Paylaşımlı bellek ile matris çarpımı uygulama
- Global bellek versiyonu ile performans karşılaştırması
- Bellek erişim desenlerini analiz etme
- Farklı GPU mimarileri için optimize etme
- Çalışma süresi
- Bellek verimi
- Doluluk oranı
- Önbellek isabet oranı
Kaynaklar
Dokümantasyon
- CUDA C++ Programlama Kılavuzu
- CUDA En İyi Uygulamalar Kılavuzu
- GPU Hesaplama Webinarları
Araçlar
- Nsight Compute
- CUDA Profilleyici
- Visual Studio GPU Hata Ayıklayıcı
Sorular ve Tartışma