c++ - CUDA 缩减 - 竞争条件？

Question

请考虑我从教程中获得的以下代码和随附的解释性图像。其目的是演示 CUDA 的并行缩减。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <numeric>
using namespace std;

__global__ void sumSingleBlock(int* d)
{
  int tid = threadIdx.x;

  // Number of participating threads (tc) halves on each iteration
  for (int tc = blockDim.x, stepSize = 1; tc > 0; tc >>= 1, stepSize <<= 1)
  {
    // Thread must be allowed to write
    if (tid < tc)
    {
      // We need to do A + B, where B is the element following A, so first we 
      // need to find the position of element A and of element B      
      int posA = tid * stepSize * 2;
      int posB = posA + stepSize;

      // Update the value at posA by adding the value at posB to it
      d[posA] += d[posB];
    }
  }
}

int main()
{
  cudaError_t status;

  const int count = 8;
  const int size = count * sizeof(int);
  int* h = new int[count];
  for (int i = 0; i < count; ++i)
    h[i] = i+1;

  int* d;
  status = cudaMalloc(&d, size);

  status = cudaMemcpy(d,h,size, cudaMemcpyHostToDevice);

  sumSingleBlock<<<1,count/2>>>(d);

  int result;
  status = cudaMemcpy(&result,d,sizeof(int),cudaMemcpyDeviceToHost);

  cout << "Sum is " << result << endl;

  getchar();

  cudaFree(d);
  delete [] h;

  return 0;
}

现在，我可以理解图中概述的一般归约原理。我不明白的是添加的内容中没有竞争条件 (*):

很明显，所有四个线程将运行相同次数的循环；仅当tid < tc他们会做一些有用的事情吗？线程 #0 将 1 和 2 相加并将结果存储在元素 0 中。它的第二次迭代然后访问元素 2。同时，线程 #1 的第一次迭代将 3 和 4 相加并将结果存储在元素 2 中。

如果线程 #0 在线程 #1 完成迭代 1 之前开始迭代 2 怎么办？这意味着线程 #0 可以读取 3 而不是 7，或者可能是一个撕裂的值(？)这里没有任何同步，所以代码是错误的吗？

(*) 注意:我不确定是否存在竞争条件，我完全相信教程中的安全代码是正确的。

Answer 1

代码有误，需要调用__syncthreads()，如下所示。

__global__ void sumSingleBlock(int* d)
{
  int tid = threadIdx.x;

  // Number of participating threads (tc) halves on each iteration
  for (int tc = blockDim.x, stepSize = 1; tc > 0; tc >>= 1, stepSize <<= 1)
  {
    // Thread must be allowed to write
    if (tid < tc)
    {
      // We need to do A + B, where B is the element following A, so first we 
      // need to find the position of element A and of element B      
      int posA = tid * stepSize * 2;
      int posB = posA + stepSize;

      // Update the value at posA by adding the value at posB to it
      d[posA] += d[posB];
    }
     __syncthreads();
  }
}