Herhangi bir döngü yazarken genelde (int i =0; i<N; i++) gibi şeyler kullanarak içinde bulunduğumuz indexi ayarlarız. Aynı olayı GPU da yapamıyoruz. Bize sunulan threadIdx, blockIdx gibi ifadelerle içinde bulunduğumuz indexi hesaplayabiliriz. Bu kavramlara daha önceki yazıda değinmiştik. Şimdi biraz daha detaylı bakalım
- threadIdx : İçinde bulunduğumuz thread’in ID’sini verir.
- blockIdx : İçinde bulunduğumuz block’un ID’sini verir.
- blockDim : Bir block’ta kaç tane thread bulunduğu.
- gridDim : Bir grid’de kaç tane block olduğu
#include <cstdio>
__global__ void hello()
{
// Device code
printf(
"GridDim=>(%d,%d,%d)\
\tBlockDim=>(%d,%d,%d)\
\tblockIdx=>(%d,%d,%d)\
\tthreadIdx=>(%d,%d,%d)\
\n",
gridDim.x, gridDim.y, gridDim.z,
blockDim.x, blockDim.y, blockDim.z,
blockIdx.x, blockIdx.y, blockIdx.z,
threadIdx.x, threadIdx.y, threadIdx.z);
}
int main()
{
// Host code
dim3 grid_dim = dim3(2, 1, 1);
dim3 block_dim = dim3(2, 4, 1);
hello<<<grid_dim, block_dim>>>();
cudaError_t err = cudaDeviceSynchronize();
if (err != cudaSuccess)
{
printf("CUDA error: %s\n", cudaGetErrorString(err));
return 1;
}
return 0;
}esat@fedora:~/dev/CUDA/build$ ./hello
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(0,0,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(1,0,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(0,1,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(1,1,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(0,2,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(1,2,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(0,3,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(1,0,0) threadIdx=>(1,3,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(0,0,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(1,0,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(0,1,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(1,1,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(0,2,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(1,2,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(0,3,0)
GridDim=>(2,1,1) BlockDim=>(2,4,1) blockIdx=>(0,0,0) threadIdx=>(1,3,0) Indexing
Eğer çok boyutlu arraylerin indexlerinin nasıl hesaplandığının mantığını biliyorsanız bunun mantığını da aynı.
- block boyutu * block Index + thread Index
#include <cstdio>
__global__ void hello()
{
int myId = blockDim.x * blockIdx.x + threadIdx.x;
printf("CurentID -> %d * %d + %d = %d \n",
blockDim.x , blockIdx.x , threadIdx.x, myId);
}
int main()
{
dim3 grid_dim = dim3(2, 1, 1);
dim3 block_dim = dim3(3, 1, 1);
hello<<<grid_dim, block_dim>>>();
cudaError_t err = cudaDeviceSynchronize();
// ...
}CurentID -> 3*1 + 0 = 3
CurentID -> 3*1 + 1 = 4
CurentID -> 3*1 + 2 = 5
CurentID -> 3*0 + 0 = 0
CurentID -> 3*0 + 1 = 1
CurentID -> 3*0 + 2 = 2 index = blockDim.x * blockIdx.x + threadIdx.x
2D Indexing
Aynı olay, değişen pek bir şey yok. 2D array diyince aklıma ilk gelen şey resimler. Bir problem üretelim ve çözleim.
- 1920×1080 boyutunda bir resmimiz olsun. Bu resmin her pikselinde işlem yapacağım fakat öncelikle bana piksellerin IDsi lazım
- 1920×1080 tane thread i tek bir block’a koyamam çünkü 1 block ta maks 1024 thread olabiliyordu. Bu durumda resmi 32×32 lik parçalara bölüp her parçayı 1 block a atayabilirim.
- 1 parçanın boyutunu biliyorsam yatayda ve dikeyde kaç parça gerektiğini bulabilirim ki bu da gridDim olur.
- Fakat şöyle bir problemim var 1920 sayısı 32 ye tam bölünür fakat 1080 sayısı 32 ye tam bölünmez. Bu durumu
- Ya fazladan thread oluşturup if() ile index kontrolü yapacağım
- Ya da parça boyutunu değiştireceğim. Ben şuanlık bunu seçiyorum
__global__ void hello()
{
int xID = blockDim.x * blockIdx.x + threadIdx.x;
int yID = blockDim.y * blockIdx.y + threadIdx.y;
//printf("x:%d\ty: %d\n", xID, yID);
}
int main()
{
int nX = 1920, nY = 1080;
int X = 32, Y = 30;
dim3 block_dim = dim3(X, Y, 1);
dim3 grid_dim = dim3(nX/X, nY/Y, 1);
hello<<<grid_dim, block_dim>>>();
cudaError_t err = cudaDeviceSynchronize();printf biraz yavaş o sebeple comment in yaptım. 3D’de buna benzer şekilde yapılıyor. Tabi bu günün sonunda size kalmış bir şey.
Bounds checking
1920×1080 piksellik resmi anlamlı parçalara tam olarak bölebildik ama her zaman karşımıza güzel sayılar gelmeyecek. Böyle durumlarda fazladan thread çalıştırıp gereksizleri elememiz lazım.
- 15 tane sayıyı işlemem gereksin
- her blokta 4 thread çalıştırsam en az 4 tane blok çalıştırmam lazım
- yani 16 thread çalıştırmamız lazım
- 1 thread boşa gidecek
- device code’da id yi hesaplattıktan sonra if kontrolü eklersem 16. thread i pass geçebilirim.
__global__ void hello()
{
int xID = blockDim.x * blockIdx.x + threadIdx.x;
if(xID<15)
printf("x:%d\n", xID);
}
int main()
{
int nX = 15;
int X = 4;
dim3 block_dim = dim3(X, 1, 1);
dim3 grid_dim = dim3(nX/X+1, 1, 1);
hello<<<grid_dim, block_dim>>>();
cudaError_t err = cudaDeviceSynchronize();- nX/X+1 : integer division floor verir -> 15/4 => 3 bu sonuca 1 ekliyoruz
- fakat nX =16 => 16/4 + 1 => 5 eder ki 1 blok fazlalık demek bunu da istemiyoruz. Daha hoş bir formül var:
(nX + X - 1) / X
Bu formülle beraber gerekli olan block sayısını bulabiliriz.
Leave a Reply