为什么我的 CUDA 程序的输出不可靠？

Question

我正在使用 CUDA 提供的编程指南自学 CUDA。为了练习，我做了一个简单的内核，它确定数组的最大值并将其返回给 CPU：

  __global__ void getTheMaximum(float* myArrayFromCPU, float* returnedMaximum) {
    // Store my current value in shared memory.
    extern __shared__ float sharedData[];
    sharedData[threadIdx.x] = myArrayFromCPU[threadIdx.x];

    // Iteratively calculate the maximum.
    int halfScan = blockDim.x / 2;
    while (halfScan > 0 && threadIdx.x < halfScan) {
      if (sharedData[threadIdx.x] < sharedData[threadIdx.x + halfScan]) {
        sharedData[threadIdx.x] = sharedData[threadIdx.x + halfScan];
      }
      halfScan = halfScan / 2;
    }

    // Put maximum value in global memory for later return to CPU.
    returnedMaximum[0] = sharedData[0];
  }

myArrayFromCPU 是一个大小为 1024 的浮点值数组。returnedMaximum 是一个包含单个项目的普通数组：计算的最大值。

我对这个算法的想法是，它会迭代地确定最大值，因为它会从当前值的一半块大小中缩减值。

但是，当我运行此代码时，我得到了不可靠的输出。返回的最大值会有所不同。这是为什么？单个算法如何每次产生不同的值？

更新：

我也只是在一个街区上跑步。我通过设置 X=1024 的一维块大小来保证这一点。

Answer 1

不能保证整个块的所有线程都在同一时刻执行。这保证您只有一个经线（32 个线程组）。

为了避免块内的并发风险 - 您可以使用 __syncthreads() 内在函数，它会停止线程到达它，直到所有线程都到达该点。 请注意，您不应该将__syncthreads() 放在无法保证所有线程一致到达该点的分支代码中。

尝试以下循环：

__syncthreads();
while (halfScan > 0) {
  if (threadIdx.x < halfScan) {
    if (sharedData[threadIdx.x] < sharedData[threadIdx.x + halfScan]) {
      sharedData[threadIdx.x] = sharedData[threadIdx.x + halfScan];
    }
  }
  __syncthreads();
  halfScan = halfScan / 2;
}

请注意，我从 while 循环中删除了条件 threadIdx.x < halfScan，因为我希望 所有 线程在同一位置以相同的次数执行 __syncthreads()。

此外，循环之前的__syncthreads() 可能有助于确保循环开始之前来自myArrayFromCPU 的加载完成（对于所有线程）。