0

我正在使用 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 的一维块大小来保证这一点。

4

1 回答 1

2

不能保证整个块的所有线程都在同一时刻执行。这保证您只有一个经线（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;
}

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

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

于 2012-11-03T20:35:11.763 回答