“CUDA’ya Kolay ve Güncel Bir Giriş”

Not: Bu blog yazısı ilk olarak 25 Ocak 2017’de yayımlandı, ancak yeni güncellemeleri yansıtacak şekilde düzenlendi.

Bu yazı, NVIDIA’nın popüler paralel hesaplama platformu ve programlama modeli olan CUDA hakkında süper basit bir giriş niteliğindedir. 2013 yılında yazdığım daha önceki bir yazı olan CUDA’ya Kolay Bir Giriş yıllar boyunca popüler oldu. Ancak CUDA programlama daha da kolaylaştı ve GPU’lar çok daha hızlı hale geldi, bu yüzden güncellenmiş (ve daha da kolay) bir giriş yapmak zamanıdır.

CUDA C++, CUDA ile çoklu paralel uygulamalar oluşturmanın yollarından sadece birisidir. Bu, güçlü C++ programlama dilini kullanarak yüksek performanslı algoritmalar geliştirmeyi sağlar ve GPU’lar üzerinde çalışan binlerce paralel iş parçacığı ile hızlandırılmış hesaplamalar gerçekleştirmenizi mümkün kılar. Birçok geliştirici, bu şekilde derin öğrenme olarak bilinen yapay zeka devriminin temelini oluşturan uygulamaları hızlandırmıştır.

Eğer CUDA hakkında bir şeyler duydunuz ve bunu kendi uygulamalarınızda kullanmak istiyorsanız, ve eğer bir C++ programcısıysanız, bu blog yazısı iyi bir başlangıç noktası olmalıdır. Devam edebilmek için, CUDA uyumlu bir GPU’ya sahip bir bilgisayara (Windows, WSL veya 64-bit Linux, herhangi bir NVIDIA GPU yeterlidir) ya da GPU’lar sunan bir bulut instance’ına (AWS, Azure, Google Colab ve diğer bulut hizmet sağlayıcılarını kullanabilirsiniz) ihtiyacınız olacak. Ayrıca, ücretsiz CUDA Toolkit’in yüklü olduğundan emin olun.

Basit Başlayalım

Bir milyon eleman içeren iki dizinin elemanlarını toplayan basit bir C++ programı ile başlayalım.

#include <iostream>
#include <math.h>

// İki dizinin elemanlarını toplama fonksiyonu
void add(int n, float *x, float *y)
{
 for (int i = 0; i < n; i++)
     y[i] = x[i] + y[i];
}

int main(void)
{
 int N = 1<<20; // 1M eleman

 float *x = new float[N];
 float *y = new float[N];

 // x ve y dizilerini ana bellek üzerinde başlat
 for (int i = 0; i < N; i++) {
   x[i] = 1.0f;
   y[i] = 2.0f;
 }

 // Ana bellek üzerinde 1M eleman için işlemi gerçekleştirin
 add(N, x, y);

 // Hataları kontrol et (tüm değerlerin 3.0f olması beklenir)
 float maxError = 0.0f;
 for (int i = 0; i < N; i++)
   maxError = fmax(maxError, fabs(y[i]-3.0f));
 std::cout << "Maksimum hata: " << maxError << std::endl;

 // Belleği serbest bırak
 delete [] x;
 delete [] y;

 return 0;
}
        

Öncelikle, bu C++ programını derleyin ve çalıştırın. Yukarıdaki kodu bir dosyaya yapıştırın ve add.cpp olarak kaydedin. Ardından, C++ derleyiciniz ile derleyin. Ben Linux’tayım, bu yüzden g++ kullanıyorum; ancak Windows’ta MSVC (veya WSL üzerinde g++) kullanabilirsiniz.

> g++ add.cpp -o add
        

Sonrasında çalıştırın:

> ./add
Maksimum hata: 0.000000
        

(Windows üzerinde, yürütülebilir dosyayı add.exe olarak adlandırmak ve .add ile çalıştırmak isteyebilirsiniz.)

Beklendiği gibi, toplama işlemi hata olmadan tamamlandı ve program sonlandı. Şimdi bu hesaplamayı GPU’nun çoklu çekirdeklerinde paralel olarak çalıştırmak istiyorum. Bunun ilk adımları oldukça basit.

Öncelikle, add fonksiyonumuzu GPU’nun çalıştırabileceği bir fonksiyon olan kernel haline dönüştürmem gerekiyor. Bunu yapmak için, fonksiyonun başına __global__ belirtecini eklemem yeterlidir. Bu, CUDA C++ derleyicisine bu fonksiyonun GPU’da çalışacağını ve CPU kodundan çağrılabileceğini belirtir.

// Dizilerin elemanlarını toplama kernel fonksiyonu
__global__
void add(int n, float *sum, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    sum[i] = x[i] + y[i];
}
        

Bu __global__ fonksiyonu anlık bir CUDA kernel‘idir ve GPU’da çalışır. GPU üzerinde çalışan kodlara genellikle cihaz kodu, CPU üzerinde çalışan kodlara ise ana kod denir.

CUDA’da Bellek Tahsisi

GPU üzerinde hesaplama yapmak için GPU’nun erişebileceği belleği tahsis etmem gerekiyor. CUDA’da Birleşik Bellek bunu kolaylaştırıyor ve tüm sistemdeki tüm GPU’lar ve CPU’lar tarafından erişilebilen tek bir bellek alanı sağlıyor. Birleşik bellek tahsis etmek için cudaMallocManaged() fonksiyonunu kullanmalısınız, bu da ana (CPU) kodu veya cihaz (GPU) kodu tarafından erişilebilen bir gösterici döndürür. Verileri serbest bırakmak için yalnızca cudaFree() fonksiyonuna bu göstericiyi iletmeniz yeterlidir.

Bu yüzden önceki kodda new çağrılarını cudaMallocManaged() çağrıları ile değiştirmem gerekiyor ve delete [] çağrılarını cudaFree() çağrıları ile değiştireceğim.

   // Birleşik Bellek Tahsisi – CPU veya GPU'dan erişilebilir
   float *x, *y, *sum;
   cudaMallocManaged(&x, N*sizeof(float));
   cudaMallocManaged(&y, N*sizeof(float));

   ...
   // Belleği serbest bırak
   cudaFree(x);
   cudaFree(y);
        

Son olarak, add() kernel’ını başlatmalıyım ki bu da onu GPU’da çalıştırır. CUDA kernel başlatmaları, üçlü açı parantez sözdizimi <<< >>> ile belirtilir. Tek yapmam gereken bu ifadeyi add çağrısının önüne eklemek.

  add<<<1, 1>>>(N, sum, x, y);
        

Kolay! Üzerinde hangi açılı parantezlerin bulunduğunun içeriği hakkında yakında daha fazla bilgi vereceğim; şu anda yapmanız gereken tek şey, bu satırın add() fonksiyonunu çalıştırmak için bir GPU iş parçacığı başlattığıdır.

Bir şey daha var: Ana kodun, kernel tamamlanmadan sonuçlara erişmeden önce beklemesi gerekiyor (çünkü CUDA kernel başlatmaları, çağrılan CPU iş parçacığını engellemez). Bunu yapmak için, sonuçları kontrol etmeden önce cudaDeviceSynchronize() çağrısını ekliyorum.

İşte tam kod:

#include <iostream>
#include <math.h>

// Dizilerin elemanlarını toplama kernel fonksiyonu
__global__
void add(int n, float *x, float *y)
{
 for (int i = 0; i < n; i++)
   y[i] = x[i] + y[i];
}

int main(void)
{
 int N = 1<<20;
 float *x, *y;

 // Birleşik Bellek Tahsisi – CPU veya GPU'dan erişilebilir
 cudaMallocManaged(&x, N*sizeof(float));
 cudaMallocManaged(&y, N*sizeof(float));

 // x ve y dizilerini ana bellek üzerinde başlat
 for (int i = 0; i < N; i++) {
   x[i] = 1.0f;
   y[i] = 2.0f;
 }

 // GPU üzerinde 1M eleman için kernel'ı çalıştır
 add<<<1, 1>>>(N, x, y);

 // Ana kodun GPU'nun tamamlanmasını beklemesi için
 cudaDeviceSynchronize();

 // Hataları kontrol et (tüm değerlerin 3.0f olması beklenir)
 float maxError = 0.0f;
 for (int i = 0; i < N; i++) {
   maxError = fmax(maxError, fabs(y[i]-3.0f));
 }
 std::cout << "Maksimum hata: " << maxError << std::endl;

 // Belleği serbest bırak
 cudaFree(x);
 cudaFree(y);
  return 0;
}
        

CUDA dosyaları .cu uzantısına sahiptir. Yani bu kodu add.cu adlı bir dosyaya kaydedin ve nvcc, CUDA C++ derleyicisi ile derleyin.

> nvcc add.cu -o add_cuda
> ./add_cuda
Maksimum hata: 0.000000
        

Bu sadece ilk adımdır çünkü bu haliyle, bu kernel yalnızca bir iş parçacığı için doğrudur; çünkü çalıştıran her bir iş parçacığı tüm dizi üzerindeki toplamayı gerçekleştirecektir. Ayrıca, birden fazla paralel iş parçacığının aynı yerleri okuyup yazması nedeniyle bir yarış durumu da vardır.

Not: Windows üzerinde Microsoft Visual Studio’daki Proje Ayarlarında Platform’u x64 olduğundan emin olmanız gerekir.

Profilini Al!

Kernel’in ne kadar sürdüğünü öğrenmenin iyi bir yolu, onu NSight Systems CLI aracı olan nsys ile çalıştırmaktır. Ancak, yalnızca kernelin çalışma süresini öğrenmek istediğimiz için komut satırında nsys profile -t cuda --stats=true ./add_cuda yazmak yeterlidir. Ancak bu, ayrıntılı istatistikler üretir. Bu yazıda yalnızca istediğimiz çıktıyı üreten bir nsys_easy adında bir sarmalayıcı betik yazdım. Betik GitHub’da bulunmaktadır. Sadece nsys_easy‘yı indirin ve PATH’ınıza (veya mevcut dizine) bir yere koyun.

> nsys_easy ./add_cuda
Maksimum hata: 0
'/tmp/nsys-report-bb25.qdstrm' oluşturuluyor
[1/1] [========================100%] nsys_easy.nsys-rep
Oluşturuldu:
   /home/nfs/mharris/src/even_easier/nsys_easy.nsys-rep
SQLite dosyası nsys_easy.sqlite oluşturuluyor ve işleniyor...
** CUDA GPU Özeti (Kernels/MemOps) (cuda_gpu_sum):
Zaman (%)  Toplam Zaman (ns)  Sayı  Kategori   İşlem
--------  ---------------  ----- ----------- --------------------------
    98.5       75,403,544      1 CUDA_KERNEL add(int, float *, float *)
     1.0          768,480     48 MEMORY_OPER [memcpy Unified H2D]
     0.5          352,787     24 MEMORY_OPER [memcpy Unified D2D]
        

CUDA GPU Özeti tablosu, add işlevine yapılan tek bir çağrıyı gösteriyor. Bu işlem, bir NVIDIA T4 GPU üzerinde yaklaşık 75ms sürüyor. Şimdi paralellik ile bunu daha hızlı hale getirelim.

İş Parçacıklarını Al!

Artık bir kernel ile hesaplama yaptıklarına göre, bu işlemi paralel hale getirmenin yollarını keşfedelim. Anahtar, CUDA’nın <<<1, 1>>> sözdizimindedir. Bu, yürütme konfigürasyonu olarak adlandırılır ve CUDA çalışma zamanına GPU üzerinde kaç paralel iş parçacığı kullanılacağını söyler. Buradaki iki parametre vardır, ancak öncelikle ikinciyi değiştirelim: bloktaki iş parçacığı sayısını. CUDA GPU’ları kernel’ları, boyutları 32 katlarının bir katı olan iş parçacığı blokları kullanarak çalıştırır; 256 iş parçacığı seçmek makul bir boyut olacaktır.

add<<<1, 256>>>(N, x, y);
        

Eğer sadece bu değişiklikle kodu çalıştırırsam, hesaplamayı her bir iş parçacığı için bir kez yapar ve paralel iş parçacıklarına yaymak yerine tüm işlemi tek bir iş parçacığı gerçekleştirmiş olur. Bunu düzgün bir şekilde yapmak için, kernel’i değiştirmem gerekiyor. CUDA C++’ta, çalışan iş parçacıklarının dizinlerini almak için anahtar kelimeler vardır. threadIdx.x, mevcut iş parçacığının kendi bloğundaki indeksini içerir ve blockDim.x ise bloğun içindeki iş parçacığı sayısını içerir. Döngüyü, paralel iş parçacıkları ile diziyi taramak için değiştireceğim.

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}
        

add işlevi çok fazla değişmedi. Aslında, index‘i 0 yapmak ve stride‘ı 1 yapmak, onu anlam bakımından ilk versiyona eşdeğer yapıyor.

Dosyayı add_block.cu olarak kaydedin ve nvprof ile derleyip çalıştırın. Yazının geri kalanında yalnızca çıktının ilgili satırını göstereceğim.

Zaman (%) Zaman (ns) Sayı Kategori   İşlem             
-------- --------- ----- ----------- --------------------------
79.0     4,221,011     1 CUDA_KERNEL add(int, float *, float *)
        

Bu büyük bir hızlanma (75ms’den 4ms’ye), ama şaşırtıcı değil çünkü yürütme bir iş parçacığından 256 iş parçacığına geçti. Devam edip daha fazla performans elde edelim.

Blokların Dışına Çıkalım

CUDA GPU’ları, Streaming Multiprocessor veya SM olarak gruplandırılan birçok paralel işleyiciye sahiptir. Her SM, birden fazla eşzamanlı iş parçacığı bloğu çalıştırabilir; ancak her iş parçacığı bloğu tek bir SM üzerinde çalışır. Örneğin, bir NVIDIA T4 GPU’su, Turing GPU Mimarisi’na dayanmaktadır ve 40 SM ve 2560 CUDA çekirdeğine sahiptir; ve her SM, en fazla 1024 aktif iş parçacığını destekleyebilir. Tüm bu iş parçacıklarını tam olarak kullanmak için, kernel’i birden fazla iş parçacığı bloğu ile başlatmalıyım.

Artık tahmin ettiğiniz üzere, yürütme konfigürasyonunda ilk parametre blok sayısını belirtir. Paralel iş parçacıkları grupları, grid olarak adlandırılır. İşleyeceğim N elemanı var ve 256 iş parçacığı kullanacak olduğum için bloğun sayısını hesaplayarak en az N iş parçacığı elde etmeliyim. Yalnızca N’yi blockSize ile bölerim (N’nin blockSize’a bölünmediği durumda yukarıyu yuvarladığıma dikkat ederek).

int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);
        

Ayrıca kernel kodunu, iş parçacığı bloklarının tamamını hesaba katacak şekilde güncellemeli olduğum. CUDA, gridDim.x‘i sağlıyor; bu, grid’deki blok sayısını içerir ve blockIdx.x ise grid’deki mevcut iş parçacığı bloğunun indeksini içerir. Şekil 1, CUDA kernel’larında dizinin (tek boyutlu) dizinlemesini blockDim.x, gridDim.x ve threadIdx.x kullanarak çizimle gösterir. Fikir, her bir iş parçacığın indeksini, kendi bloğunun başındaki offset’i (blok indeksinin blok boyutu ile çarpılması: blockIdx.x * blockDim.x) ve iş parçacığının bloğundaki indeksini (threadIdx.x) toplamak. blockIdx.x * blockDim.x + threadIdx.x kodu, CUDA’ya özgü bir ifadedir.

__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}
        

Güncellenen kernel, stride değerini grid’deki toplam iş parçacığı sayısı olarak da belirler (blockDim.x * gridDim.x). Bu tür bir kernel döngüsü CUDA’da sıkça grid-stride döngüsü olarak adlandırılır.

Dosyayı add_grid.cu olarak kaydedin ve nvprof ile derleyip çalıştırın.

Zaman (%) Zaman (ns) Sayı Kategori    İşlem             
-------- --------- ----- ----------- ----------------------------
79.6     4,514,384     1 CUDA_KERNEL add(int, float *, float *)
        

İlginçtir ki, bu değişiklikten bir hız artışı sağlamadık; muhtemelen hafif bir yavaşlama bile olmuş olabilir. Neden? Eğer hesaplamayı 40 kat (SM sayısı) artıramıyorsak, toplam süreyi düşürmeyen neden bu. Çünkü hesaplama, dar boğaz değil.

Birleşik Bellek Ön Bellek Alma

Profil oluşturucudan gelen bu darboğaz hakkında bir ipucu var:

Zaman (%) Zaman (ns) Sayı Kategori    İşlem
-------- --------- ------ ----------- --------------------------
   79.6 4,514,384      1 CUDA_KERNEL add(int, float *, float *)         
   14.2   807,245     64 MEMORY_OPER [CUDA memcpy Unified H2D]
    6.2   353,201     24 MEMORY_OPER [CUDA memcpy Unified D2H]
        

Burada, 64 ana bellekten cihaza (H2D) ve 24 cihazdan ana belleğe (D2H) “birleşik” memcpy işlemlerinin olduğunu görebiliriz. Ancak, kodda hiçbir memcpy çağrısı yok. CUDA’da Birleşik Bellek sanal bellek olarak çalışır. Bireysel sanal bellek sayfaları, sistemdeki herhangi bir cihazin (GPU veya CPU) belleğinde bulunabilmektedir ve bu sayfalar talep üzerine taşınmaktadır. Bu program önce CPU üzerinde dizileri bir döngü ile başlatır ve daha sonra kernel başlatıldığında diziler GPU tarafından okunur ve yazılır. Fakat bellekteki sayfa hatalarının çoğu CPU’da bulunduğundan, birden fazla sayfa hatası oluşur ve GPU belleğindeki sayfalar bu hataların meydana geldiği zamanlarda taşınır. Bu durum, bir bellek darboğazına neden olur, bu nedenle hız artışı görmüyoruz.

Taşımalar, maliyetlidir çünkü sayfa hataları tekil olarak gerçekleşir ve GPU iş parçacıkları, sayfa taşınmasını beklerken duraklatılır. Kernel’in şeylerine ihtiyaç duyduğumdan, x ve y dizilerinin bellekte olmasını sağlamak için ön belleğe alma kullanabilirim. Bunu, kernel’i başlatmadan önce cudaMemPrefetchAsync() fonksiyonunu kullanarak yapacağım:

// x ve y dizilerini GPU'ya ön belleğe al
cudaMemPrefetchAsync(x, N*sizeof(float), 0, 0);
cudaMemPrefetchAsync(y, N*sizeof(float), 0, 0);
        

Bunu profilde çalıştırdığımda, kernel süresi artık 50 mikro saniyenin altında!

Zaman (%) Zaman (ns) Sayı Kategori    İşlem
-------- --------- ----- -----------  --------------------------
   63.2   690,043     4 MEMORY_OPER  [CUDA memcpy Unified H2D]
   32.4   353,647    24 MEMORY_OPER  [CUDA memcpy Unified D2H]
    4.4    47,520     1 CUDA_KERNEL  add(int, float *, float *)
        

Sonuçları Toplamak

Dizilerin tüm sayfalarını aynı anda ön belleğe almak, bireysel sayfa hatalarından çok daha hızlıdır. Bu değişikliği, tüm sürümlere ekleyip tekrar profilleyelim. İşte bir özet tablosu.

Bir kez veriler bellekte olduğunda, tek bir bloktan birden fazla bloğa geçmenin hız artışı, GPU üzerindeki SM sayısına orantılıdır (40).

Görüldüğü üzere, GPU’lar çok yüksek bant genişliğine ulaşabilir. Add kernel’i çok bant genişliği sınırlıdır (265 GB/s, T4’ün 320GB/s’lik zirve bant genişliğinin %80’inden fazladır); ancak GPU’lar yoğun hesaplama gerektiren hesaplamalarda da mükemmeldir. Örneğin, yoğun matris lineer cebir işlemleri, derin öğrenme, görüntü ve sinyal işleme, fiziksel simülasyonlar gibi.

Uygulamalar

Devam etmenizi sağlamak için kendi başınıza denemeniz gereken bazı şeyler burada.

1. CUDA Toolkit dokümantasyonunu inceleyin. Eğer hala CUDA kurmadıysanız, Hızlı Başlangıç Rehberi ve kurulum kılavuzlarına göz atın. Ardından Programlama Kılavuzu ve En İyi Uygulamalar Kılavuzu‘nu inceleyin. Her mimari için ayar rehberleri de mevcuttur.
2. Kernel içinde printf() ile denemeler yapın.threadIdx.x ve blockIdx.x değerlerini bazı veya tüm iş parçacıkları için yazdırarak kontrol edin. Acaba sıralı bir şekilde mi yazdırdınız? Neden böyle oldu?
3. threadIdx.y veya threadIdx.z (veya blockIdx.y) değerlerini yazdırın. (Aynı şekilde blockDim ve gridDim için de). Neden bu işlevler var? Bu değerleri 0’dan (1) farklı değerlere nasıl alabilirsiniz?

İleride Ne Olacak?

Umarım bu yazı, CUDA hakkında merakınızı uyandırmıştır ve CUDA C++ kullanarak kendi hesaplamalarınızı yapma isteğiyle dolmuşsunuzdur. Sorularınız veya yorumlarınız varsa, lütfen aşağıdaki yorum bölümünü kullanarak benimle iletişime geçin.

Devam edebileceğiniz bir dizi eski tanıtım yazısı var:

• CUDA C++ ile Performans Ölçümleri Uygulama
• Aygıt Özelliklerini Sorgulamak ve Hataları Yönetmek
• CUDA C++ ile Veri Aktarımlarını Optimize Etme
• CUDA C++ ile Veri Aktarımlarını Üst Üste Getirme
• CUDA C++ ile Global Belleğe Etkili Erişim
• CUDA C++ ile Paylaşımlı Belleği Kullanma
• CUDA C++ ile Verimli Matris Transpoz
• CUDA C++ ile Sonlu Fark Yöntemleri, Bölüm 1
• CUDA C++ ile Sonlu Fark Yöntemleri, Bölüm 2
• CUDA ile Bir Hafta Sonunda Hızlandırılmış Işın İzleme

Bunun dışında CUDA Fortran ile paralel bir seri oluşturan yazılar da mevcut; bunlara, CUDA Fortran’a Kolay Bir Giriş ile başlayarak inceleyebilirsiniz.

CUDA C++ ve diğer GPU hesaplama konuları hakkında daha fazla içerik bulmak için NVIDIA Geliştirici Blogu’nda göz atmanızı öneririm!

Bu yazıyı beğendiyseniz ve daha fazlasını öğrenmek istiyorsanız, NVIDIA DLI birkaç derinlemesine CUDA programlama kursu sunmaktadır.

• Yeni başlayanlar için,Modern CUDA C++ ile Hızlandırılmış Hesaplamaya Giriş sunulmaktadır. Bu kurs, özel GPU kaynakları, daha sofistike bir programlama ortamı, NVIDIA Nsight Systems görsel profillerini kullanma, düzinelerce etkileşimli egzersiz, ayrıntılı sunumlar, 8 saatten fazla içerik ve DLI Yeterlilik Sertifikası kazanma imkanı sunmaktadır.
• Python programcıları için,CUDA Python ile Hızlandırılmış Hesaplamanın Temelleri mevcuttur.
• Daha ileri düzey bazı CUDA programlama materyalleri için, NVIDIA DLI’nin Hızlandırılmış Hesaplama bölümüne göz atabilirsiniz, özgür olarak sunulan hazırlık kılavuzlarına ulaşabilirsiniz.