CUDA Çekirdekleri ve Çok Boyutluluk

Öğrenim Hedefleri

  • Çok boyutlu CUDA çekirdekleri hakkında bilgi sahibi olmak

    • Çok boyutlu çekirdek yapısında blok ve iş parçacığı indisleri,

    • Blok ve iş parçacığı indislerinin data indisleriyle eşleştirilmesi.

Örnek: Çok Boyutlu CUDA Çekirdeği

Aşağıda görmüş olduğunuz görsel küçük bir çok boyutlu CUDA çekirdeğinin yapısını göstermektedir. Bu örnekte çekirdek konfigürasyonu 2 boyutlu şebeke içinde bulunan bloklar ve her blok içerisine 3 boyutlu şekilde dizilmiş iş parçacıkları gösterilmiştir.

../../../_images/013.png

Görmüş olduğunuz gibi her bir blok içerisinde (2 * 2 * 4 = 16 olacak şekilde) 16 adet iş parçacığı bulunmaktadır.

Örnek: Çok Boyutlu CUDA Çekirdeği ile 2 Boyutlu Resim İşleme

Çok boyutlu CUDA çekirdeğinin genel yapısını bir önceki bölümde bir örnek üzerinden gördük. Bu örnekle 2 boyutlu bir CUDA çekirdeğinin nasıl kullanıldığını inceleyeceğiz.

Aşağıda görmüş olduğunuz 62 piksel (boy) x 76 piksel (en) boyutlarındakı görseli 16x16 şeklinde konfigüre edilmiş iş parçacığı blokları kullanarak işleyeceğiz.

../../../_images/022.png

62 piksellik uzunluk için bize en az 4 adet alt alta dizilmiş blok gerekmektedir. 4 adet, bu şekilde konfigüre edilmiş blok bize 64 piksellik bir uzunluk sağlar bu yüzden resimde de gözüktüğü gibi 2 adet satır hesaplamaya katılmayacaktır. Ayrıca 76 piksel uzunluğu için bize en az 5 adet yan yana dizilmiş blok gerekmektedir. Bu şekilde konfigüre edilmiş 5 blok bize 80 piksellik bir uzunluk sağlar bu yüzden 4 adet sütun hesaplamaya katılmayacaktır.

Satır Bazlı Sıralama

C/C++ dillerinde oluşturduğumuz 2 boyutlu yapılar, bellek üzerinde lineer şekilde yerleştirilir ve bu yerleştirme işlemi yapılırken Satır Bazlı Sıralama (Row Major Access) kullanılır. 2 boyutlu yapılar lineerize edilirken satırların bölünmeden arka arkaya eklenerek tek boyutlu yapı oluşturmasına satır bazlı sıralama denir. Sütunların bölünmeden arka arkaya dizilmesine ise sütun bazlı sıralama denilmektedir. C/C++ dilleri 2 boyutlu yapıları lineerize ederken satır bazlı sıralama kullanır. Örnek için aşağıdaki görseli inceleyiniz.

../../../_images/032.png

Görüldüğü üzere 2 boyutlu yapıdaki satırlar bölünmeden arka arkaya eklenerek lineerize edilmiştir. Bu noktada satır ve sütun indisleri ile kaçıncı sırada bulunacağı hesaplanabilir. Örnek vermek gerekirse üçüncü satır ikinci eleman olan M2,1, 2 boyutlu yapı lineerize edildikten sonra 10. sırada yer almaktadır (indis = 9). Bu indisi hesaplamak için satır indisi ile bir satırda kaç adet girdi olduğu çarpılıp üzerine sütun indisi eklenebilir. Örnek üzerinden ilerlersek M2,1 girdisinde satır indisi 2, sütun indisi 1 ve her satırda 4 adet girdi bulunmaktadır. Lineerize edildikten sonraki indis (2 * 4) + 1 = 9 şeklinde bulunabilir.

2 Boyutlu Resim İşleme Çekirdek ve Sunucu Kodu

Satır bazlı sıralama hakında bilgi sahibi olduktan sonra 2 boyutlu resim işleme örneğinin çekirdek ve sunucu kodunu inceleyelim.

CUDA Çekirdek Kodu

__global__ void PictureKernel(float* d_Pin, float* d_Pout, int height, int width)
{
   //d_Pin ve d_Pout icin satir indisini hesaplama
   int Row = blockIdx.y*blockDim.y + threadIdx.y;

   //d_Pin ve d_Pout icin sutun indisini hesaplama
   int Col = blockIdx.x*blockDim.x + threadIdx.x;

   // Eger satir indisi toplam gereken boy uzunlugundan kucuk ise ve
   // Eger sutun indisi toplam gereken en uzunlugundan kucuk ise hesaplama yapılacak
   if ((Row < height) && (Col < width)) {
      // 2 boyutlu resim tek boyutlu halde geldiginden ve
      // satır bazlı sıralama ile lineerize oldugundan oturu
      // indis hesaplaması yapılması gerekmektedir.
      d_Pout[Row*width+Col] = 2.0*d_Pin[Row*width+Col];
   }
}

CUDA Host Kodu

// islenecek resimin boyutları = m X n,
// y ekseninde m piksel ve x ekseninde n piksel
// d_Pin icin gerekli yer cihaz uzerinde ayrilmis ve cihaza kopyalama islemi tamanlanmis oldugu varsayilmakta
// d_Pout icin gerekli yer cihaz uzerinde ayrilmis oldugu varsayilmakta
...
dim3 DimGrid((n-1)/16 + 1, (m-1)/16+1, 1);
dim3 DimBlock(16, 16, 1); PictureKernel<<<DimGrid,DimBlock>>>(d_Pin, d_Pout, m, n);
...

Şebeke boyutu daha önceki bölümlerde gördüğümüz gibi en az hesaplamayı başarıyla tamamlayacak kadar blok içermektedir. Örnek resim boyuları ile düşünürsek x ekseninde 5, y ekseninde 4 adet blok bulunduracak bir 2 boyutlu şebeke yapısı konfigüre edilmektedir. Ayrıca örneğin başında bahsettiğimiz gibi her bir blok 16x16 iş parçacığı içerecek şekilde 2 boyutlu bir yapıya sahiptir.

62x76 Piksellik Resimi 16x16 Bloklar ile Kaplamak

Blok içerisindeki program kontrol akışı düşünüldüğünde işlemiş olduğumuz örnekte 4 farklı bölge bulunmaktadır. Aşağıdaki resimde 1 ile işaretlenmiş bölgedeki blokların içindeki her iş parçacığı resimdeki bir piksele denk gelmektedir ve bu yüzden bu bloklardaki her bir iş parçacığı bir piksel işleyecektir. 2. bölgedeki bloklar ise 4 sütunda bulunan iş parçacıklarına resim üzerindeki herhangi bir piksel denk gelmemektedir. Bu iş parçacıkları hesaplama yapmayacaklardır. 3. bölgedeki bloklardaki 2 satırda bulunan iş parçacıklarına resim üzerinde herhangi bir piksel denk gelmemektedir. Bu iş parçacıkları hesaplama yapmayacaklardır. Son olarak 4. bölgedeki bloklardaki 4 sütun ve 2 satırda bulunan iş parçacıklarına resim üzerindeki herhangi bir piksel denk gelmemektedir. Bu yüzden bu iş parçacıkları hesaplama yapmayacaktır.

../../../_images/041.png

Görüldüğü üzere 4 farklı bölgedeki program kontrol akışı farklılaşmaktadır. 2. ve 3. bölgelerdeki fark şu anda çok belirgin olmasa da ilerleyen bölümlerde Kontrol Akış Dağılması ve Warp kavramını yakından inceledğimizde 2. ve 3. bölgelerdeki performans farkı hakkında fikir sahibi olacağız.