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函数会无止境的等待。