CUDA-并行归约和

CUDA - Parallel Reduction Sum

本文关键字：归约 并行 CUDA-      更新时间：2023-10-16

我正在尝试在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: %dn", sumGPU);
std::string wait;
std::cin >> wait;
return 0;
}

我还尝试过输入的数组大小越来越大，无论数组大小如何，我都会得到非常大的负值的相同结果。

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

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的任何值(当然，前提是它在允许的范围内)。缺点是，如果你有多个外部共享数组，你需要自己想办法把它们放在内存中，这样一个就不会覆盖另一个。

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

相关文章：