CUDA 阵列缩减答案

【问题标题】：CUDA Array ReductionCUDA 阵列缩减
【发布时间】：2013-06-06 21:06:11
【问题描述】：

我知道已经回答了多个与此类似的问题，但除了我可能错误地索引某些内容之外，我无法从他们那里拼凑出任何非常有用的东西。

我正在尝试将输入向量 A 的顺序寻址缩减为输出向量 B。

完整的代码在这里http://pastebin.com/7UGadgjX，但这是内核：

__global__ void vectorSum(int *A, int *B, int numElements) {
  extern __shared__ int S[];
  // Each thread loads one element from global to shared memory
  int tid = threadIdx.x;
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  if (i < numElements) {
    S[tid] = A[i];
    __syncthreads();
    // Reduce in shared memory
    for (int t = blockDim.x/2; t > 0; t>>=1) {
      if (tid < t) {
        S[tid] += S[tid + t];
      }
      __syncthreads();
    }
    if (tid == 0) B[blockIdx.x] = S[0];
  }
}

这些是内核启动语句：

// Launch the Vector Summation CUDA Kernel
  int threadsPerBlock = 256;
  int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
  vectorSum<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, numElements);

我收到了一个未指定的启动错误，我读到该错误类似于段错误。我一直在密切关注 nvidia 缩减文档，并试图将我的内核保持在 numElements 的范围内，但考虑到代码的简单程度，我似乎遗漏了一些关键。

【问题讨论】：

请不要在 pastebin 中发布代码。将相关代码放入问题中。
另外，您可能对stackoverflow.com/q/14038589/681865 上的代码感兴趣 - 在此示例中它将大大简化您的主机代码。

标签： arrays vector cuda reduction

【解决方案1】：

您的问题是缩减内核需要动态分配的共享内存才能正确运行，但您的内核启动没有指定任何内容。结果是越界/非法共享内存访问会中止内核。

在 CUDA 运行时 API 语法中，内核启动语句有四个参数。前两个是启动的网格和块尺寸。后两者是可选的，默认值为零，但指定动态分配的共享内存大小和流。

要解决此问题，请按如下方式更改启动代码：

// Launch the Vector Summation CUDA Kernel
  int threadsPerBlock = 256;
  int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
  size_t shmsz = (size_t)threadsPerBlock * sizeof(int);
  vectorSum<<<blocksPerGrid, threadsPerBlock, shmsz>>>(d_A, d_B, numElements);

[免责声明：在浏览器中编写的代码，未经编译或测试，使用风险自负]

这至少应该可以解决您的代码中最明显的问题。

【讨论】：