cuda by example

int offset= x+y*dim

 

x 線程塊內的線程索引

y 線程塊索引

dim 線程塊的維度

 

tid = threadIdx.x+blockIdx.x*blockDim.x

　　

計算大於或等於128的最小倍數（127+x）/128

 

kernel<<<(x+127)/128,128>>>(a,b,c)

　　

 

規約求和

 

int i= blockDim.x/2;
while(i != 0){
    if(cacheIndex < i)
        cache[cacheIndex] += cache[cacheIndex + i];
    __synthreads();
    i /= 2;
}

　　

 

 

const int N = 33*1024
const int threadsperblock = 256;
const int blockpergrid = imin(32,(N+threadperblock-1)/threadsperblock);
 
kernel<<<blockpergrid,threadsperblock>>>(a,b,c);
 
__global__ static void kenel(int *a,int *b,int *c){
    ...
    int tid = threadIdx.x+blockIdx.x*blockDim.x;
    ...
    while(tid<N){
        ...
        tid += blockDim.x*gridDim.x;
        ...
    }
}

　　

 

if(threadIdx.x % 2){
    ...
    __synthreads();
}

　　

這會形成 線程發散：

    當某些線程須要執行一條指令，而其餘線程不須要執行時，這種狀況成爲線程發散。

 

__synthreads會當全部的線程都執行後才釋放，而有些線程若是不執行，那麼kernel函數會無止境的等待。