1

我正在研究 Cuda Parallel reduction Whitepaper，但不幸的是我的算法似乎反复产生不正确的结果，我似乎无法弄清楚为什么（当然教科书示例必须有效？当然我只是在做一些非常明显的错误？） . 这是我的内核函数：

我的定义：

 #define BLOCK_SIZE 512

我的内核功能：

 __global__ void total(float * inputList, float * outputList, int len) {
      __shared__ float sdata[2*BLOCK_SIZE];
      unsigned int tid = threadIdx.x;
      unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
      sdata[t] = inputList[i]+inputList[i+blockDim.x];
      __syncthreads();
      for (unsigned int s=blockDim.x/2; s>0; s>>=1) {
        if (tid < s) {
          sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
      }
      if (tid == 0) 
        outputList[blockIdx.x] = sdata[0];
}

我的内存分配：

  outputSize = inputSize / (BLOCK_SIZE<<1);
  cudaMalloc((void**) &deviceInput, inputSize*sizeof(float));
  cudaMalloc((void**) &deviceOutput, outputSize*sizeof(float));
  cudaMemcpy(deviceInput, hostInput, inputSize*sizeof(float), cudaMemcpyHostToDevice);

我的设备呼叫：

 dim3 dimGrid((inputSize-1)/BLOCK_SIZE +1, 1, 1);
 dim3 dimBlock(BLOCK_SIZE,1,1);

 total<<<dimBlock, dimGrid>>>(deviceInput, deviceOutput, outputSize);
 cudaDeviceSynchronize();

我的记忆提取：

 cudaMemcpy(hostOutput, deviceOutput, outputSize*sizeof(float), cudaMemcpyDeviceToHost);

最后是我的最终计算：

 for (int counter = 1; counter < outputSize; counter++) {
    hostOutput[0] += hostOutput[counter];
 }

任何帮助，将不胜感激。

4

2 回答 2

5

以下代码行中的内核启动配置不正确。

total<<<dimBlock, dimGrid>>>(deviceInput, deviceOutput, outputSize);

内核配置的第一个参数是网格大小，第二个参数是块大小。

你应该这样做：

total<<<dimGrid, dimBlock>>>(deviceInput, deviceOutput, outputSize);

请始终对 CUDA Runtime 函数调用执行错误检查，并检查返回的错误代码以获取程序失败的原因。

您的内核启动应该在您当前的代码中失败。对cudaDeviceSynchronize调用进行错误检查会导致您找到不正确结果的原因。

于 2013-01-14T11:06:13.390 回答
3

该代码假定输入大小是块大小的倍数。如果 inputSize 不是块大小的倍数，它将读取 inputList 数组的末尾。

于 2013-01-14T10:29:51.113 回答