CUDA - 并行归约和

Question

我正在尝试在 CUDA 7.5 中实现并行归约和。我一直在尝试遵循NVIDIA PDF，它会引导您完成初始算法，然后稳步优化版本。我目前正在制作一个用 1 作为每个数组位置的值的数组，以便我可以检查输出是否正确，但我得到的值为 -842159451 的大小为 64 的数组。我期待内核代码是正确的，因为我遵循了 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];

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

选项 2

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

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

这适用于任何值numThreadsPerBlock（当然，前提是它在允许的范围内）。缺点是如果有多个extern共享数组，就需要自己想办法把then放入内存中，这样一个不会覆盖另一个。

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