CUDA Çekirdek ve SPMD Paralelleşmesi

Öğrenim Hedefleri

Basit CUDA çekirdek işlemleri hakkında bilgi sahibi olmak
- Çekirdek deklarasyonu
- Yerleşik değişkenler (ing., built-in variables)
- İş parçacığı indisleri ile veri indisleri arasında ilişki kurmak

Örnek: Vektör Toplama İşlemi

CUDA Paralelleşme modeli ve çekirdek kullanımını daha yakından incelemek için vektör toplama örneğinin cihaz kodu üzerinden aşama aşama ilerleyelim.

__global__ void sumVector(int *vec1, int* vec2, int* outVec, int size)
 {
     int index = threadIdx.x + blockIdx.x * blockDim.x;
     if(index < size)
     {
         outVec[index] = vec1[index] + vec2[index];
     }
 }

__global__ etiketi cihaz ve sunucu (GPU ve CPU) tarafından çalıştırılabilen fonksiyonları işaretlemek için kullanılır. CUDA çekirdekleri yazarken __global__ etiketi gibi etiket(ler) kullanmamız gerekir. Bu etiketlerin neler olduğunu ve ne işe yaradıklarını bir sonraki bölümde görebilirsiniz.
void sumVector(int *vec1, int* vec2, int* outVec, int size) kısmından görebileceğiniz üzere çekirdek fonksiyonunun geri döndürdüğü bir değer yoktur. CUDA, çekirdek fonksiyonlarının geri değer döndürmesini desteklemez. Bu yüzden CUDA çekirdeklerinizi yazarken void fonskiyonlar yazmanız gerekmektedir. Bunun dışında örnekte kullanılan parametreler toplanacak olan vektörler, sonuç vektörü ve vektörlerin büyüklüğü olarak sıralanabilir. Daha önceki bölümlerde CUDA bellek ayırma ve kopyalama işlemlerinden bahsetmiştik. Bu noktada gerekli ayırma ve kopyalama işlemlerinin yapıldığını varsayıyoruz. Yani çekirdek fonksiyonuna parametre olarak gelen adresler grafik işlem biriminin belleğinde bir adres olup toplanacak vektörlerin elemanları bu adrese halihazırda kopyalanmış olmalıdır.
Aynı işlemi yapan çok sayıda iş parçacığını kullanarak elde ettiğimiz paralel hesaplama modelimiz için hangi iş parçacığının hangi veri üzerinde çalışacağını belirlememiz gerekir. int index = .. kısmındaki hesaplama ile her bir iş parçacığı için evrensel indis bulunur ve toplama işleminin gerçekleştirileceği vektörlerin index. elemanı bu iş parçacığı tarafından gerçekleştirilir.
İndis hesabından sonra görmüş olduğumuz if(index<size) kontrolü ise vektörlerin eleman sayısıdan fazla sayıda iş parçacığı bulunduğu koşulda yanlış/gereksiz bir toplama işlemi yapmamak için gereklidir.

CUDA çekirdeğinin nasıl başlatıldığını görmek için şimdi de sunucu tarafındaki kodu inceleyelim.

...
sumVector<<<ceil(size/256),256>>>(d_input1,d_input2,d_output,inputLength);
...

Bu kod bloğunda bellek ayırma ve kopyalama işlemleri atlanmıştır.

Görüldüğü üzere bir CUDA çekirdeği başlatmak için normal bir C fonksiyonundan farklı olarak <<<…>>> işaretleri arasına alınmış konfigürasyon parametreleri bulunmaktadır. Bu parametreler şebeke (grid) ve blok boyutlarını kontrol etmemizi sağlar. Bu da, iş parçacıklarını içinde bulunduran blokların kaç boyutlu olduğu (bizim örneğimizde 1 boyutlu), bir blokta kaç adet iş parçacığı olduğu (256) ve bu blokların oluşturduğu şebekenin kaç boyutlu olduğu (bizim örneğimizde 1 boyutlu), kaç adet blok içerdiği (vektör uzunluğu / 256) gibi değerleri ayarlamamızı sağlar. Fark ettiğiniz üzere bu örnekte boyut anlamında bir bilgi girilmeden sadece kaç adet blok ve her blokta kaç adet iş parçacığı olacağı girilmiştir. Tek boyutlu blok ve şebeke yapılarında, özel veri yapıları kullanmanıza gerek yoktur.

Dikkat edilmesi gereken nokta ceil(size/256) adet blok oluşturuluyor olmasıdır. İşlemimizi gerçekleştirmeye yetecek kadar iş parçacığına sahip olmak amacıyla toplam kullanılacak iş parçacığı sayısı (her iş parçacığı 1 eleman toplayacağından vektörlerin kaç eleman içerdiği), her blokta 256 adet iş parçacığı bulunduğundan 256 ile bölünmektedir. Sonuç olarak size/256 adet blok gereklidir. Ancak 256 sayısı size sayısını tam bölemediği durumlarda çıkan sayısının bir üst tam sayıya yuvarlanması sağlanmalıdır. Blok sayısı tam sayı olmak durumunda olduğundan, yukarı yuvarlama işlemi yapılmazsa virgülden sonrası atılacağından eksik sayıda blok oluşabilir. Bu durumda ise sonuç yanlış/eksik çıkacaktır. Bunu engellemek adına bir üst tam sayıya yuvarlanarak ihtiyacımız kadar veya daha fazla sayıda iş parçacığının vektör toplama işleminde çalışması sağlanır. Fazla kalan iş parçacıkları çekirdek kodundaki if ifadesi ile durdurulmaktadır.

Çok boyutlu şebeke ve blok kullanılmak istediğinizde bunu aşağıda örneği verilen dim3 değişkenleri ile yapabilirsiniz. Bu veri yapısı 3 tamsayıdan oluşur. Bunlar şebeke ve blok**ların **x, y, ve z boyutlarının uzunluklarını göstermektedir.

dim3 DimGrid((n-1)/256 + 1, 1, 1);
dim3 DimBlock(256, 1, 1);
sumVector<<<DimGrid,DimBlock>>>(d_input1,d_input2,d_output,inputLength);

Yukarda görmüş olduğunuz örnek üstteki çekirdek başlatımı ile eşdeğerdir. Ayrıca (n-1)/256 + 1 ifadesinin yukarı yuvarlamak ile aynı işleve sahip olduğunu da dikkat ediniz. Bir boyutlu şebeke ve blok yapısına sahip olduğumuzu DimGrid ve DimBlock değişkenlerinin y ve z eksenlerinde 1 adet elemana sahip olduğundan anlayabilirsiniz. Daha sönraki bölümlerde 2 boyutlu çekirdek oluşturmayı uygun bir örnek üzerinden inceleyceğiz.

Özet: CUDA Çekirdek Başlatma

Yukardaki örnek üzerinden gösterdiğimiz CUDA çekirdek başlatma işlemini kısaca özetleyelim.

Sunucu tarafından konfigüre edilen çekirdeğin başlamasıyla beraber aşağıdaki görselle modellenmiş bir yapı ortaya çıkar. Bu yapının fiziksel bir karşılığı olmadığını sadece CUDA Çekirdek ve İş Parçacığı mantığının böyle kurgulandığını unutmayınız.

Görselde görmüş olduğunuz şebeke içerisindeki her blok, birden fazla sayıda iş parçacığı içermekte ve her iş parçacığı ise cihaz kodu olan çekirdeği çalıştırmaktadır.
Her iş parçacığı çekirdek kodunda bulunan evrensel indisiyle çekirdekte gerçekleştirilen işleme katılıp katılmayacağına karar verir. Bu gerekenden az sayıda iş parçası oluşturmayı engellemek için kullandığımız blok sayısını yukarı yuvarlama tekniği nedeniyle gereklidir. Fazladan oluşan iş parçacıkları çekirdek kodunun devamını çalıştırmaz.

Çeşitli Çekirdek Deklerasyonları

__global__ ile işaretlenmiş CUDA çekirdekleri (fonksiyonları) sunucu tarafından veya diğer __global__ çekirdeklerden çağrılabilir (compute capability 3.5 ve üstü için geçerlidir) ve bu fonksiyonlar cihaz üzerinde çalışır.
__device__ ile işaretlenmiş CUDA fonksiyonları cihaz tarafından çağrılabilir ve bu fonksiyonlar cihaz üzerinde çalışır. __device__ ile işaretli fonksiyonlar void dışında bir değer döndürebilir.
__host__ ile işaretlenmiş CUDA fonksiyonları sunucu tarafından çağrılabilir ve bu fonksiyonlar sunucu üzerinde çalışır. Bir fonksiyonu sadece __host__ ile işaretlemek opsiyoneldir, herhangi başka bir etiket ile etiketlenmemiş fonksiyonlar derleyici tarafından __host__ ile işaretlenmiş sayılır.
__device__ ve __host_ birlikte kullanılabilir. Birlikte kullanıldıklarında derleyici sunucu da ve cihaz da çalışacak iki farklı fonksiyon yaratır.