CUDA’ya Giriş

GPU’lar başlangıçta grafik çizmek için geliştirildi. Bir görüntü çizmek için milyonlarca pikselin rengini hesaplamaları gerekir ve bu küçük hesaplamaların çoğunu aynı anda yapabilirler. CUDA, bu gücü grafiklerin yanı sıra kendi programlarımızda da kullanmamızı sağlayan bir NVIDIA platformudur. NVIDIA GPU’larıyla çalışır. Bu yazıda hiçbir şey kurmadan temel fikri göreceğiz: büyük bir işin çok sayıda küçük parçaya nasıl bölündüğünü.

CPU’da az sayıda güçlü çekirdek, GPU’da ise binlerce daha basit çekirdek bulunur. Bir görüntünün tüm piksellerini boyamamız gerektiğini düşünün. CPU aynı anda az sayıda pikseli çok hızlı boyar. GPU’nun her çekirdeği daha yavaş olsa da aynı anda çok sayıda piksel boyar. İkisini de çalıştırın, ardından en küçük boyutu deneyin.

İş küçük olduğunda CPU kazanır. GPU’nun başlamasından önce verilerin genellikle kendi belleğine kopyalanması ve işin başlatılması gerekir. Yalnızca birkaç piksel olduğunda bu hazırlık, işin kendisinden uzun sürer. İş büyüdükçe GPU, CPU’yu geride bırakır. Bu yüzden GPU’lar, görüntüler ve sinir ağlarının matematiği gibi birbirinden bağımsız küçük parçalara bölünebilen işler için uygundur.

CUDA’daki en küçük çalışan birim bir threaddır. İş parçacıkları blok halinde gruplanır; tüm blokların oluşturduğu yapıya ise grid denir. GPU’da bir iş başlatırken kaç blok olacağını ve her blokta kaç thread bulunacağını seçeriz.

Her thread aynı kodu çalıştırır. Peki, hangi iş parçasının kendisine ait olduğunu nasıl bilir? CUDA her threade birkaç yerleşik değer verir. blockIdx.x blok numarasıdır, blockDim.x bloktaki thread sayısıdır, threadIdx.x ise threadin blok içindeki numarasıdır. Küçük bir formülle bunları her thread için benzersiz bir numaraya dönüştürürüz. Görmek için aşağıdaki threadlerin üzerine gelin. Anladığınızda Beni test et düğmesine basıp doğru threadi kendiniz bulmayı deneyin.

Her adın neden .x ile bittiğini merak edebilirsiniz. Bloklar ve threadler, .y ve .z ile iki veya üç boyutlu da düzenlenebilir. Bu, görüntüler ve üç boyutlu veriler için kullanışlıdır. Burada basit tutmak için tek boyut kullanıyoruz.

GPU’da bu şekilde başlattığımız fonksiyona kernel denir. __global__ ile işaretlenir ve özel <<<blocks, threads>>> söz dizimiyle başlatılır. Aşağıda klasik bir ilk CUDA programı var: iki diziyi toplamak. Her thread yalnızca bir sayı çiftini toplar ve threadler paralel çalışır.

Dikkat edilmesi gereken bir nokta var. İş parçacıkları tam bloklar halinde geldiği için çoğu zaman ihtiyacımızdan biraz fazlasını başlatırız. n değeri 10 ve bir blok 4 thread içeriyorsa 3 blok gerekir; bu da 12 thread demektir. Bu yüzden her thread yazmadan önce i < n koşulunu kontrol eder. Gerçek programlar genellikle 256 thread gibi daha büyük bloklar kullanır; ancak küçük sayıları görmek daha kolaydır. Eleman sayısını değiştirip sınır kontrolünü kapatmayı deneyin.

__global__ void add(float *a, float *b, float *c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    c[i] = a[i] + b[i];
}
 
// launch
add<<<3, 4>>>(a, b, c, 10);

Çoğu bilgisayarda CPU ve GPU’nun bellekleri ayrıdır. Bir kernel çalışmadan önce veriler, örneğin cudaMemcpy ile GPU’ya kopyalanır; işlem sonunda da sonuç geri kopyalanır. Bu kopyalamalar işin kendisinden uzun sürebilir. Her kernel öncesi ve sonrası verileri taşırsak GPU zamanının büyük kısmını bekleyerek geçirebilir. Kernel çalıştırmaları arasında verileri GPU’da tutarsak kopyalama maliyetini yalnızca bir kez öderiz.

Kendiniz hiç CUDA yazmasanız da bu durumla karşılaşırsınız. PyTorch’ta .to("cuda"), bir tensörü GPU belleğine kopyalar; .cpu() veya .item() ise verileri CPU’ya geri getirir. .item(), CPU’nun GPU’nun işini bitirmesini beklemesine de neden olur. Bunları eğitim döngüsünde tekrar tekrar çağırmak bu yüzden yavaşlığa yol açabilir: veriler sürekli gidip gelir ve CPU bekler.

Artık tüm parçaları biliyoruz; şimdi birleştirelim. Küçük bir CUDA programı beş adımdan oluşur: GPU’da bellek ayır, girdileri kopyala, kerneli çalıştır, sonucu geri kopyala ve belleği serbest bırak. Aşağıda adım adım ilerleyip iki belleği de izleyin.

const int n = 10;
size_t size = n * sizeof(float);
float a[n], b[n], c[n]; // on the CPU, a and b are filled
float *d_a, *d_b, *d_c;          // will point to GPU memory
 
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
int threads = 256;
int blocks = (n + threads - 1) / threads;
add<<<blok, thread>>>(d_a, d_b, d_c, n);
cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

CUDA’nın temel fikri budur. İşi küçük parçalara bölün, her threadin basit bir formülle kendi parçasını bulmasını sağlayın ve verileri mümkün olduğunca az taşıyın. Görüntü filtrelerinden sinir ağlarına kadar birçok GPU programı bu fikirlerle oluşturulur.

@aykutkardas