并行处理-用计数器代替blockId

Question

最初我使用的是blockIdx。x在我的代码中，但我想删除它，而是有一个全局值，并在我的块中使用它，而不是blockidx.x。因为我的代码太大了，当我用大的输入大小运行它时，它会挂起，我认为think可以帮助。我自动增加计数器，但当我运行代码时，它挂起。谁能看看我的代码，看看我做错了什么?

__device__ int counter = 0;
__global__ void kernel(int * ginput, int * goutput)
{
  const int tid = threadIdx.x;
  const int id = threadIdx.x + blockIdx.x * blockDim.x;
  in myval = ginput[id];  
  if (tid == 0) {
    atomicAdd(&counter, 1);
  }
  __syncthreads();
  if (counter == 0) {
    goutput[tid] = ...;
  }
  if (counter > 0) {
   ...
  }
}

如果我使用blockIdx。x而不是计数器在我的代码中它可以工作，但我只想用计数器

Answer 1

如果您希望counter取代blockIdx.x的使用(即您希望每个块都有一个从counter读取的唯一值)，那么像这样的东西应该工作:

__device__ int counter = 0;
__global__ void kernel(int * ginput, int * goutput)
{
  const int tid = threadIdx.x;
  const int id = threadIdx.x + blockIdx.x * blockDim.x;
  __shared__ int my_block_id;

  if (tid == 0) {
    my_block_id = atomicAdd(&counter, 1);
  }
  __syncthreads();
  if (my_block_id == 0) {
    goutput[tid] = ...;
  }
  if (my_block_id > 0) {
   ...
  }
}

你的方法会很麻烦，因为如果你这样做:

if (counter > 5) ....

您可能正在从全局内存中读取counter的新更新值，并且任何数量的块都可能已经更新了该值，因此行为将是不可预测的。