CUDA Rehberi: NVIDIA GPU Programlamaya Giriş

GPU’lar artık yalnızca oyun dünyasının değil, bilimsel hesaplama, makine öğrenmesi ve yüksek başarımlı hesaplama (HPC) alanlarının da vazgeçilmez bileşeni haline geldi. NVIDIA’nın CUDA (Compute Unified Device Architecture) platformu, GPU’ların paralel işlem gücünü genel amaçlı hesaplamalar için kullanmayı mümkün kılan en yaygın çerçevedir. Bu rehberde CUDA’nın temellerini, GPU mimarisini ve pratik kod örnekleriyle başlangıç seviyesinden orta seviyeye bir yol haritası sunuyoruz.

GPU Mimarisi: CPU’dan Farkı Ne?

Modern bir CPU, birkaç güçlü çekirdekten oluşur ve karmaşık, sıralı görevler için optimize edilmiştir. Buna karşın bir NVIDIA GPU, binlerce daha küçük çekirdeği paralel olarak çalıştırabilir. Bu yapısal fark, GPU’ları aynı türden işlemin büyük veri kümeleri üzerinde tekrarlanması gereken hesaplamalarda (matris çarpımı, vektör işlemleri, konvolüsyon) son derece verimli kılar.

CUDA mimarisinde temel hiyerarşi şu şekilde işler:

Bir GPU, bellekte veri bekleme sürelerini binlerce thread arasında geçiş yaparak gizler. Bu nedenle GPU programlamasında en kritik beceri, işi yeterince ince parçalara bölüp binlerce thread’e dağıtmaktır.

CUDA Kurulumu ve İlk Adımlar

CUDA geliştirme ortamı kurmak için önce sisteminizde uyumlu bir NVIDIA sürücüsü ve CUDA Toolkit’in yüklü olduğundan emin olmanız gerekir. Kurulum sonrası nvcc derleyicisi ve çeşitli profil araçları (nvprof, Nsight) kullanılabilir hale gelir.

Derleme için temel komut:

nvcc -o cikti program.cu

CUDA kaynak dosyaları .cu uzantısını taşır ve hem standart C/C++ kodunu hem de GPU çekirdeği tanımlamalarını içerir.

İlk Kernel: Vektör Toplama

CUDA’da GPU üzerinde çalışan bir fonksiyona kernel denir. Kernel, __global__ anahtar kelimesiyle tanımlanır ve CPU tarafından çağrılır, GPU üzerinde çalışır.

Aşağıdaki örnek, iki dizinin elemanlarını toplar:

#include <stdio.h>
#include <cuda_runtime.h>

// GPU üzerinde çalışacak kernel fonksiyonu
__global__ void vektorTopla(float *a, float *b, float *c, int n) {
    // Her thread'in global indeksini hesapla
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    int n = 1024;
    size_t boyut = n * sizeof(float);

    // CPU (host) bellekleri
    float *h_a = (float*)malloc(boyut);
    float *h_b = (float*)malloc(boyut);
    float *h_c = (float*)malloc(boyut);

    // Başlangıç değerlerini ayarla
    for (int i = 0; i < n; i++) {
        h_a[i] = (float)i;
        h_b[i] = (float)(i * 2);
    }

    // GPU (device) bellekleri
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, boyut);
    cudaMalloc(&d_b, boyut);
    cudaMalloc(&d_c, boyut);

    // Veriyi CPU'dan GPU'ya kopyala
    cudaMemcpy(d_a, h_a, boyut, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, boyut, cudaMemcpyHostToDevice);

    // Kernel'i çalıştır: 4 blok, her blokta 256 thread
    int bloklardaThread = 256;
    int blokSayisi = (n + bloklardaThread - 1) / bloklardaThread;
    vektorTopla<<<blokSayisi, bloklardaThread>>>(d_a, d_b, d_c, n);

    // Sonucu GPU'dan CPU'ya kopyala
    cudaMemcpy(h_c, d_c, boyut, cudaMemcpyDeviceToHost);

    printf("c[0] = %.1f, c[1023] = %.1f\n", h_c[0], h_c[1023]);

    // Bellek temizleme
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

Burada dikkat edilmesi gereken temel noktalar şunlardır:

• blockIdx.x, blockDim.x ve threadIdx.x her thread’e özgü yerleşik değişkenlerdir.
• Veri önce cudaMalloc ile GPU belleğine ayrılır, ardından cudaMemcpy ile aktarılır.
• Kernel çağrısındaki <<<blokSayisi, bloklardaThread>>> sözdizimi, NVIDIA’ya özgüdür ve standart C++ derleyicileriyle derlenemez.

Bellek Hiyerarşisi ve Yönetimi

CUDA’da performansın büyük bölümü doğru bellek stratejisinden gelir. GPU’nun birden fazla bellek katmanı vardır:

Global Bellek

En geniş kapasiteye sahip bellek türüdür (tipik olarak birkaç GB). cudaMalloc ile ayrılır. Erişim gecikmesi yüksektir; bu nedenle mümkün olduğunca birleşik (coalesced) erişim kalıpları tercih edilmelidir. Bitişik thread’lerin bitişik bellek adreslerine erişmesi durumunda, donanım bu erişimleri tek bir işlemde birleştirir ve bant genişliği kullanımı dramatik biçimde artar.

Paylaşılan Bellek

Aynı thread bloğundaki tüm thread’lerin erişebildiği, chip üzerinde konumlandırılmış hızlı bellektir. __shared__ anahtar kelimesiyle tanımlanır. Büyüklüğü SM başına genellikle 48-96 KB ile sınırlıdır; ancak global belleğe kıyasla onlarca kat daha hızlıdır. Matris çarpımı gibi algoritmaların blok tabanlı uygulamalarında paylaşılan belleği etkin kullanmak, performansı 5-10 kat artırabilir.

Sabit ve Doku Belleği

Sabit bellek (__constant__), salt okunur ve önbelleklenmiş küçük veri yapıları için uygundur. Tüm thread’ler aynı adrese erişiyorsa broadcast mekanizması devreye girer ve tek bir erişim tüm warp’a yayılır. Doku belleği ise 2B uzamsal lokaliteyi kullanan algoritmalar için optimize edilmiştir.

Thread Organizasyonu ve Boyut Seçimi

Kernel çağrısında blok sayısı ve blok başına thread sayısını dikkatli belirlemek gerekir:

• Thread sayısı per blok genellikle 128, 256 veya 512 olarak seçilir. 32’nin katı olması warp uyumu açısından önemlidir.
• Blok sayısı, işlenecek veri boyutuna ve SM sayısına göre belirlenir. GPU’nun tüm SM’lerini dolu tutmak için yeterli blok sayısı sağlanmalıdır.
• 2B ve 3B grid/block yapıları, matris ve görüntü işleme gibi çok boyutlu problemler için dim3 türüyle tanımlanabilir:

dim3 blokBoyutu(16, 16);
dim3 gridBoyutu((genislik + 15) / 16, (yukseklik + 15) / 16);
matrisKerneli<<<gridBoyutu, blokBoyutu>>>(d_giris, d_cikis, genislik, yukseklik);

Eşzamanlılık ve Akışlar (Streams)

CUDA’da akışlar (streams), farklı GPU işlemlerinin örtüşmesine olanak tanır. Varsayılan akışta tüm işlemler sırayla yürütülür. Birden fazla akış oluşturarak bellek kopyaları ile kernel çalışmasını paralel hale getirmek mümkündür:

cudaStream_t akis1, akis2;
cudaStreamCreate(&akis1);
cudaStreamCreate(&akis2);

// Birinci parça
cudaMemcpyAsync(d_a1, h_a1, boyut/2, cudaMemcpyHostToDevice, akis1);
kernel<<<grid, blok, 0, akis1>>>(d_a1, d_c1, n/2);

// İkinci parça eş zamanlı
cudaMemcpyAsync(d_a2, h_a2, boyut/2, cudaMemcpyHostToDevice, akis2);
kernel<<<grid, blok, 0, akis2>>>(d_a2, d_c2, n/2);

cudaStreamSynchronize(akis1);
cudaStreamSynchronize(akis2);

Bu yaklaşım özellikle veri transferinin darboğaz oluşturduğu büyük hesaplamalarda verimlilik artışı sağlar.

Hata Ayıklama ve Profilleme

Her CUDA API çağrısı bir hata kodu döndürür; üretim kodunda bunları kontrol etmek zorunludur:

cudaError_t hata = cudaMalloc(&d_a, boyut);
if (hata != cudaSuccess) {
    fprintf(stderr, "cudaMalloc hatası: %s\n", cudaGetErrorString(hata));
    exit(EXIT_FAILURE);
}

Profilleme araçları arasında NVIDIA Nsight Systems ve Nsight Compute öne çıkar. Nsight Systems timeline görünümüyle kernel sürelerini, bellek kopyalarını ve CPU-GPU etkileşimini görselleştirir. Nsight Compute ise tek bir kernel’in roofline analizi, bellek erişim kalıpları ve warp verimliliği gibi düşük düzeyli metriklerini sunar.

Yaygın Performans Tuzakları

CUDA’ya yeni başlayan geliştiricilerin en sık düştüğü hatalar şunlardır:

Aşırı bellek kopyası: Her küçük hesaplama için CPU-GPU arası veri transferi yapmak, PCI-Express bant genişliğini tıkar ve GPU’nun sunduğu hız avantajını ortadan kaldırır. Mümkünse veriyi GPU’da tutun ve birden fazla kernel’i arka arkaya çalıştırın.

Düşük doluluk oranı (occupancy): Her SM’nin belirli sayıda aktif warp barındırabilmesi gerekir. Blok başına çok az thread veya çok fazla register/paylaşılan bellek kullanımı doluluk oranını düşürür. NVIDIA’nın Occupancy Calculator aracı bu dengeyi bulmaya yardımcı olur.

Dallanma ıraksaması (warp divergence): Aynı warp içindeki thread’ler farklı if-else dallarına girerse, GPU bu dalları sırayla işler ve paralellik bozulur. Mümkün olduğunca warp içi thread’lerin aynı kodu çalıştırması sağlanmalıdır.

Hizasız bellek erişimi: Thread’lerin rastgele bellek adreslerine erişmesi birleşik okuma/yazma işlemlerini engeller. Veri yapılarını ve erişim kalıplarını bu kısıtı göz önünde bulundurarak tasarlamak önemlidir.

Sonraki Adımlar

CUDA temellerini kavradıktan sonra pratik HPC projelerinde sıkça başvurulan kütüphaneleri incelemek mantıklıdır: cuBLAS (doğrusal cebir), cuFFT (hızlı Fourier dönüşümü), cuDNN (derin öğrenme primitifler) ve Thrust (paralel STL benzeri algoritmalar). Bu kütüphaneler, NVIDIA’nın kendi mühendislerinin optimizasyon tecrübelerini hazır API’ler aracılığıyla sunar ve çoğu durumda sıfırdan yazılan kernel’lerden daha iyi performans gösterir.

CUDA ekosistemi, multi-GPU programlama (NCCL, MPI+CUDA), birleşik bellek (Unified Memory) ve en güncel GPU mimarilerine (Hopper, Blackwell) özgü özelliklerle sürekli gelişmektedir. Bu alanlara hâkim olmak, modern HPC ve yapay zeka altyapılarında rekabet avantajı sağlar.

Mevasis olarak CUDA programlama, GPU küme kurulumu ve yüksek başarımlı hesaplama altyapısı konusunda size destek olmaktan memnuniyet duyarız. İletişim için formu doldurun.