在 CUDA 内核中顺序递增计数器？

Question

我想不通这个逻辑，所以我有一个计数器 SOLUTIONS 每次我有每个线程的新解决方案时它都会自动递增，但在某些情况下，一个线程可以找到多个解决方案，但在这种情况下，我希望它再增加一次当前值。这个逻辑看似简单，由atomicAdd实现。但问题是，我需要将该计数器用作数组的索引。

示例：

SOME_ARRAY[tid] = STUFF; // puts stuff in individual indexes because of thread id.
atomicAdd(SOLUTIONS, 1); // increments the max solutions found.

但是现在单个线程已经找到了两个解，我希望它在当前的MAX解之后递增。

SOME_ARRAY[SOLUTIONS+1] = STUFF;

但据我了解，如果两个或更多线程尝试执行此操作，它不会起作用吗？我需要第二个线程位于 SOLUTIONS+2 而不是写入相同的位置。

我怎样才能做到这一点？感谢任何帮助。

编辑：我尝试实现，这仍然不起作用，因为当我打印偏移量时，两个线程打印相同的数字。当我访问 d_PROGRESS 数组时，2 元素总是数字 85，这是一个我从未放在那里的随机数。

 __device__
 bool isSafe (int row, int col, int cmp_row, int cmp_col) {
    if ((col - cmp_col) == 0) {return 0;}
    if (abs(row - cmp_row) == abs(col - cmp_col)) {return 0;}
    return 1;
 }

 __global__
  void nqueensKernel(int row, int n, bool one, bool all, int pitch,
               Solution * d_solution,
               Solution * d_PROGRESS,
               Solution *d_PROGRESS_OUT,
               unsigned long long * NUM_THREADS,
               unsigned long long * NUM_SOLUTIONS) {

int index = threadIdx.x + blockIdx.x * blockDim.x;
int isAlone = 0;
if (index == 0) {*NUM_THREADS = 0;}
__syncthreads();
if (row == 0) {
  d_PROGRESS_OUT[index * n + row] = index;
  atomicAdd(NUM_THREADS, 1);
}
else {
  int moresolutions = 0;
  for (int col = 0; col < n; col++) {
    for (int k = 0; k < row; k++) {
      int checkcol = d_PROGRESS[index * n + k];
      isAlone = isSafe(row, col, k, checkcol);
      if (!isAlone) { /*printf("Is Alone? %d\n", isAlone);*/ break; }
    }
    __syncthreads();
    if ( isAlone ) {
      moresolutions++;
      if (moresolutions == 1) {
          d_PROGRESS_OUT[(index * n) + row] = col;
          atomicAdd(NUM_THREADS, 1);
          for (int o = 0; o < row; o++) {
            d_PROGRESS_OUT[(index * n) + o] = d_PROGRESS[(index * n) + o];
          }
      } else if (moresolutions > 1) {
          int offset = atomicAdd(NUM_THREADS, 1);
          d_PROGRESS_OUT[((offset+1) * n) + row] = col;
          for (int m = 0; m < row; m++) { d_PROGRESS_OUT[((offset+1) * n) + m] = d_PROGRESS[(index * n) + m]; }
      }
      if (row == n-1) { atomicAdd(NUM_SOLUTIONS, 1); }
    }
  }
}

Answer 1

您想使用 atomicAdd 返回的值保留 SOME_ARRAY 中的存储空间，该值是存储在变量 old 中的值=25=]在添加之前，将原子变量递增要保留的插槽数。例如：

int offset = atomicAdd(SOLUTIONS, number_of_solutions_found_by_this_thread);
SOME_ARRAY[offset] = stuff;
SOME_ARRAY[offset + 1] = more stuff;
...
SOME_ARRAY[offset + number_of_solutions_found_by_this_thread - 1] = also more stuff;

假设 SOLUTIONS 最初是 0。

某个线程（我们将其亲切地称为#27）找到 4 个解决方案，并执行 atomicAdd。结果返回 0，SOLUTIONS 现在是 4.
在 #27 进入下一行之前，线程 #88 进入原子添加并将 SOLUTIONS 递增 9。结果返回 4 并且 SOLUTIONS 是现在 13.
#27 继续从索引 0 到 3 存储其解决方案，不干扰 #88。

这是许多需要跨线程聚合可变长度结果的并行算法中的常见模式。

在 CUDA 内核中顺序递增计数器？

Increment counter sequentially inside CUDA kernel?

counter

cuda