1

書こうとしている CUDA プログラムに問題があります。約 524k の浮動小数点値 (1.0) の配列があり、リダクション手法を使用してすべての値を追加しています。1 回だけ実行したい場合は問題なく動作しますが、実際にはカーネルを数回実行して、最終的に 10 億を超える値を合計できるようにしたいと考えています。

これを 524k のチャンクで行っている理由は、GPU で約 100 万を超えると常にゼロが返されるためです。カードのメモリを超えてはいけませんが、その時点で常に失敗します。

とにかく、カーネルを 1 回だけループすると、すべて正常に動作します。つまり、ループしなくても問題ありません。ループを実行すると、ゼロが返されます。どこかで限界を超えているのではないかと思いますが、それがわかりません。それは私を夢中にさせています。

どんな助けでも大歓迎です、

ありがとう、

アル

コードは次のとおりです。

#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

First, thanks to everyone who gave this a look and any help.

Second, I finally figured out what I was doing wrong. BLOCKS should have been defined as MAX_WIDTH/TILE_WIDTH, not WIDTH/TILE_WIDTH. Stupid, silly mistake from my part.

Thanks again.

于 2011-11-30T21:03:29.397 に答える