c++ - 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;
}

我还为输入尝试了更大和更小的数组大小，无论数组大小如何，我都得到了非常大的负值的相同结果。

Answer 1

似乎您正在使用动态分配的共享数组:

extern __shared__ int sdata[];

但您没有在内核调用中分配它:

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

你有两个选择:

选项 1

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

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

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

选项 2

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

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

这将适用于 numThreadsPerBlock 的任何值(当然前提是它在允许的范围内)。缺点是如果你有多个外部共享数组，你需要自己弄清楚如何将 then 放入内存中，这样一个就不会覆盖另一个。

---

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