用计数器替换 blockId

【问题标题】：replacing blockId with a counter用计数器替换 blockId
【发布时间】：2016-07-05 20:44:11
【问题描述】：

最初我在我的代码中使用了 blockIdx.x，但我想删除它，取而代之的是一个全局值并在我的块中使用它而不是 blockidx.x。由于我的代码太大，并且当我以大输入大小运行它时它会挂起，我认为这会有所帮助。我以原子方式递增计数器，但是当我运行代码时它挂起。谁能看看我的代码，看看我是不是做错了什么？

__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 而不是计数器，它可以工作，但我只想用计数器替换它

【问题讨论】：

标签： parallel-processing cuda atomic nvidia

【解决方案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 的新更新值，并且任何数量的块都可能更新了该值，因此行为将是不可预测的。

【讨论】：