我尝试自己使用 CUDA 实现向量和缩减，但遇到了一个我可以修复但不了解实际问题的错误。
我在下面实现了内核，它与 NVIDIA 的示例中使用的内核几乎相同。
__global__
void reduce0(int *input, int *output)
{
extern __shared__ int s_data[];
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x + threadIdx.x;
s_data[tid] = input[i];
__syncthreads();
for( int s=1; s < blockDim.x; s *= 2) {
if((tid % 2*s) == 0) {
s_data[tid] += s_data[tid + s];
}
__syncthreads();
}
if(tid == 0) {
output[blockIdx.x] = s_data[0];
}
}
此外，我在主机端计算共享内存空间如下
int sharedMemSize = numberOfValues * sizeof(int);
如果使用了超过 1 个线程块，则代码运行良好。仅使用 1 个块会导致上述索引超出范围错误。通过将我的主机代码与我发现以下行的示例之一进行比较来查找我的错误：
int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T);
稍微玩一下我的块/网格设置，我得到了以下结果：
- 块，任意数量的线程 => 代码崩溃
- >2 块，任意数量的线程 => 代码运行良好
- 1 块，任意数量的线程，共享内存大小 2*#threads => 代码运行良好
尽管考虑了几个小时，但我不明白为什么在使用太少数量的线程或块时会出现越界错误。
更新：主机代码按要求调用内核
int numberOfValues = 1024 ;
int numberOfThreadsPerBlock = 32;
int numberOfBlocks = numberOfValues / numberOfThreadsPerBlock;
int memSize = sizeof(int) * numberOfValues;
int *values = (int *) malloc(memSize);
int *result = (int *) malloc(memSize);
int *values_device, *result_device;
cudaMalloc((void **) &values_device, memSize);
cudaMalloc((void **) &result_device, memSize);
for(int i=0; i < numberOfValues ; i++) {
values[i] = i+1;
}
cudaMemcpy(values_device, values, memSize, cudaMemcpyHostToDevice);
dim3 dimGrid(numberOfBlocks,1);
dim3 dimBlock(numberOfThreadsPerBlock,1);
int sharedMemSize = numberOfThreadsPerBlock * sizeof(int);
reduce0 <<< dimGrid, dimBlock, sharedMemSize >>>(values_device, result_device);
if (cudaSuccess != cudaGetLastError())
printf( "Error!\n" );
cudaMemcpy(result, result_device, memSize, cudaMemcpyDeviceToHost);