CUDA - 多次调用内核答案

【问题标题】：CUDA - Calling a Kernel Multiple TimesCUDA - 多次调用内核
【发布时间】：2011-11-29 21:20:07
【问题描述】：

我在尝试编写 CUDA 程序时遇到了困难。我有一个大约 524k 浮点值（1.0）的数组，我正在使用归约技术来添加所有值。如果我只想运行一次，问题就可以解决，但我真的想多次运行内核，以便最终汇总超过 10 亿个值。

我在 524k 的块中这样做的原因是，当我在 gpu 上超过大约 100 万时，我总是得到零。这不应该超过卡上的内存，但它总是在那个时候失败。

无论如何，当我只循环一次内核时，一切正常。也就是说，没有循环是好的。当我使用循环运行时，它会返回零。我怀疑我会在某个地方越界，但我无法弄清楚。快把我逼疯了。

感谢您的帮助，

谢谢，

阿尔

代码如下：

#include <stdio.h>
#include <stdlib.h>
#include "cutil.h"

#define TILE_WIDTH     512
#define WIDTH          524288 
//#define WIDTH          1048576
#define MAX_WIDTH      524288

#define BLOCKS         WIDTH/TILE_WIDTH

__global__ void PartSum(float * V_d)
{
   int tx = threadIdx.x;
   int bx = blockIdx.x;

   __shared__ float partialSum[TILE_WIDTH];

   for(int i = 0; i < WIDTH/TILE_WIDTH; ++i)
   {
      partialSum[tx] = V_d[bx * TILE_WIDTH + tx];
      __syncthreads();


      for(unsigned int stride = 1; stride < blockDim.x; stride *= 2)
      {
         __syncthreads();
         if(tx % (2 * stride) == 0)
            partialSum[tx] += partialSum[tx + stride];
      }
   }

   if(tx % TILE_WIDTH == 0)
      V_d[bx * TILE_WIDTH + tx] = partialSum[tx];
}

int main(int argc, char * argv[])
{
   float * V_d;
   float * V_h;
   float * R_h;
   float * Result;
   float * ptr;

   dim3 dimBlock(TILE_WIDTH,1,1);
   dim3 dimGrid(BLOCKS,1,1);

   // Allocate memory on Host
   if((V_h = (float *)malloc(sizeof(float) * WIDTH)) == NULL)
   {
      printf("Error allocating memory on host\n");
      exit(-1);
   }

   if((R_h = (float *)malloc(sizeof(float) * MAX_WIDTH)) == NULL)
   {
      printf("Error allocating memory on host\n");
      exit(-1);
   }

   // If MAX_WIDTH is not a multiple of WIDTH, this won't work
   if(WIDTH % MAX_WIDTH != 0)
   {
      printf("The width of the vector must be a multiple of the maximum width\n");
      exit(-3);
   }

   // Initialize memory on host with 1.0f
   ptr = V_h;
   for(long long i = 0; i < WIDTH; ++i)
   {
      *ptr = 1.0f;
      ptr = &ptr[1];
   }

   ptr = V_h;

   // Allocate memory on device in global memory
   cudaMalloc((void**) &V_d, MAX_WIDTH*(sizeof(float)));
   float Pvalue = 0.0f;
   for(int i = 0; i < WIDTH/MAX_WIDTH; ++i)
   {


   if((Result = (float *) malloc(sizeof(float) * WIDTH)) == NULL)
   {
      printf("Error allocating memory on host\n");
      exit(-4);
   }

   for(int j = 0; j < MAX_WIDTH; ++j)
   {
      Result[j] = *ptr;
      ptr = &ptr[1];
   }

      ptr = &V_h[i*MAX_WIDTH];
      // Copy portion of data to device
      cudaMemcpy(V_d, Result, MAX_WIDTH*(sizeof(float)), cudaMemcpyHostToDevice);

      // Execute Kernel
      PartSum<<<dimGrid, dimBlock>>>(V_d);

      // Copy data back down to host
      cudaMemcpy(R_h, V_d, MAX_WIDTH*(sizeof(float)), cudaMemcpyDeviceToHost);

      for(int i = 0; i < MAX_WIDTH; i += TILE_WIDTH)
      {
         Pvalue += R_h[i];
      }
printf("Pvalue == %f\n", Pvalue);

  free(Result);


   }

//   printf("WIDTH == %d items\n", WIDTH);
//   printf("Value: %f\n", Pvalue);

   cudaFree(V_d);
   free(V_h);
   free(R_h);
   return(1);
}

好的，我想我已经将问题范围缩小到设备上的 V_d。我怀疑我以某种方式超出了数组的范围。如果我分配 2 倍于我实际需要的内存量，程序将以预期的结果结束。问题是，我无法弄清楚是什么导致了问题。

阿尔

【问题讨论】：

您使用 cudaMemcpyDeviceToHost 是否有特殊原因？由于您想对内核进行更多迭代，因此您可以考虑改用cudaMemcpyDeviceToDevice。
我不敢相信你真的需要 CUDA 来做这件事——这将是 I/O 主导的，因为你每个点只有一个添加操作——你还不如只使用 CPU。您是否真的对这个实现的 CPU 实现进行了基准测试？考虑到所有数据移动且几乎没有计算，您认为 CUDA 实现可能会快多少？
这对我来说是一个学习实验。我意识到这没有效率。
@karlphillip：我不确定你的意思。
您正在将内存来回传输到 GPU。我认为没有必要在每次迭代时将数据传输到 CPU。

标签： c cuda reduction

【解决方案1】：

我想我在这里发现了第一个错误：

  if(tx % TILE_WIDTH == 0)
      V_d[bx * TILE_WIDTH + tx] = partialSum[tx];

tx 的范围是 0-511，从未达到 512。所以 if 条件永远不会为真。您可以将其写为 if(tx % (TILE_WIDTH-1) == 0)。

【讨论】：

啊！我没看到！谢谢。

【解决方案2】：

首先，感谢所有给予此关注和帮助的人。

其次，我终于弄清楚我做错了什么。 BLOCKS 应该被定义为 MAX_WIDTH/TILE_WIDTH，而不是 WIDTH/TILE_WIDTH。我犯了一个愚蠢而愚蠢的错误。

再次感谢。

【讨论】：