CUDA线程同步

220 阅读1分钟

考虑如下程序:

`__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int index = threadIdx.x;
  int reindex = n-index-1;
  s[index] = d[index];
  __syncthreads();
  d[index] = s[reindex];
}`

采用了线程同步,因此可以保证所有线程都会执行完 s[index] = d[index]; 再继续执行。

该程序的作用是数组倒序排列。

但是假如我们只想倒序前一半的话:

`__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];

  int index = threadIdx.x;

  int reindex = n/2-index-1;

    if(reindex >= 0){

        s[index] = d[index];

        __syncthreads();

        d[index] = s[reindex];
    }  
}`

这样可能就卡死了,因为同步需要多有线程都调用,才能继续执行,所以不要定义在分支结构中,否则就有可能有GPU线程调用不到,导致线程阻塞。