CUDA - 并行归约求和

Question

我正在尝试在CUDA 7.5中实现并行归约求和。我一直在尝试遵循 NVIDIA 的PDF，其中逐步介绍了最初的算法，然后是逐渐优化的版本。我当前正在创建一个数组，该数组中每个位置的值都为1，以便我可以检查输出是否正确，但是对于大小为64的数组，我得到了-842159451的值。我期望内核代码是正确的，因为我已经按照NVIDIA的完全相同的代码进行了操作，以下是我的内核：

__global__ void reduce0(int *input, int *output) {
    extern __shared__ int sdata[];

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = input[i];

    __syncthreads();

    for (unsigned int s = 1; s < blockDim.x; s *= 2) {
        if (tid % (2 * s) == 0) {
            sdata[tid] += sdata[tid + s];
        }

        __syncthreads();
    }

    if (tid == 0) output[blockIdx.x] = sdata[0];
}

这是我的调用内核的代码，我认为我的问题就在这里：

int main()
{
    int numThreadsPerBlock = 1024;

    int *hostInput;
    int *hostOutput; 
    int *deviceInput; 
    int *deviceOutput; 

    int numInputElements = 64;
    int numOutputElements; // number of elements in the output list, initialised below

    numOutputElements = numInputElements / (numThreadsPerBlock / 2);
    if (numInputElements % (numThreadsPerBlock / 2)) {
        numOutputElements++;
    }

    hostInput = (int *)malloc(numInputElements * sizeof(int));
    hostOutput = (int *)malloc(numOutputElements * sizeof(int));


    for (int i = 0; i < numInputElements; ++i) {
        hostInput[i] = 1;
    }

    const dim3 blockSize(numThreadsPerBlock, 1, 1);
    const dim3 gridSize(numOutputElements, 1, 1);

    cudaMalloc((void **)&deviceInput, numInputElements * sizeof(int));
    cudaMalloc((void **)&deviceOutput, numOutputElements * sizeof(int));

    cudaMemcpy(deviceInput, hostInput, numInputElements * sizeof(int), cudaMemcpyHostToDevice);

    reduce0 << <gridSize, blockSize >> >(deviceInput, deviceOutput);

    cudaMemcpy(hostOutput, deviceOutput, numOutputElements * sizeof(int), cudaMemcpyDeviceToHost);

    for (int ii = 1; ii < numOutputElements; ii++) {
        hostOutput[0] += hostOutput[ii]; //accumulates the sum in the first element
    }

    int sumGPU = hostOutput[0];

    printf("GPU Result: %d\n", sumGPU);

    std::string wait;
    std::cin >> wait;

    return 0;
}

我还尝试了更大和更小的数据输入数组，但无论数组大小如何，都得到一个非常大的负值。

- Terramet

Answer 1

看起来您正在使用动态分配的共享数组：

extern __shared__ int sdata[];

但你在内核调用中没有分配它：

reduce0 <<<gridSize, blockSize >>>(deviceInput, deviceOutput);

你有两个选择：

选择1

在内核中静态地分配共享内存，例如：

constexpr int threadsPerBlock = 1024;
__shared__ int sdata[threadsPerBlock];

通常情况下，我发现这是最干净的方法，因为当您在共享内存中有多个数组时，它可以无问题地工作。缺点是，虽然大小通常取决于块中的线程数，但需要在编译时知道大小。

选项2

在内核调用中指定动态分配的共享内存量。

reduce0 <<<gridSize, blockSize, numThreadsPerBlock*sizeof(int) >>>(deviceInput, deviceOutput);

这将适用于numThreadsPerBlock的任何值（当然，前提是它在允许的范围内）。缺点是，如果您有多个外部共享数组，则需要自己找出如何将它们放入内存中，以确保不会互相覆盖。

---

请注意，您的代码可能存在其他问题。我没有测试过它。这只是我浏览您的代码时立即发现的一个问题。

- CygnusX1