1

我在尝试编写 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 倍于我实际需要的内存量，程序将以预期的结果结束。问题是，我无法弄清楚是什么导致了这些问题。

铝

4

2 回答 2

2

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

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

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

于 2011-11-30T09:20:22.180 回答
2

首先，感谢所有看过这个并提供任何帮助的人。

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

再次感谢。

于 2011-11-30T21:03:29.397 回答