還元によってCUDAの配列の総和を見つける方法

Question

私は、配列を32 * 32の要素を持ち、値が0〜1023の縮小を使って配列の和を求める関数を実装しています。 私の予想される合計値は523776ですが、私の再考は15872です、それは間違っています。ここ は私のコードです：

#include <stdio.h> 
#include <cuda.h> 

#define w 32 
#define h 32 
#define N w*h 

__global__ void reduce(int *g_idata, int *g_odata); 
void fill_array (int *a, int n); 

int main(void) { 
    int a[N], b[N]; // copies of a, b, c 
    int *dev_a, *dev_b; // device copies of a, b, c 
    int size = N * sizeof(int); // we need space for 512 integers 

    // allocate device copies of a, b, c 
    cudaMalloc((void**)&dev_a, size); 
    cudaMalloc((void**)&dev_b, size); 

    fill_array(a, N); 

    // copy inputs to device 
    cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_b, b, size, cudaMemcpyHostToDevice); 

    dim3 blocksize(16,16); 
    dim3 gridsize; 

    gridsize.x=(w+blocksize.x-1)/blocksize.x; 
    gridsize.y=(h+blocksize.y-1)/blocksize.y; 

    reduce<<<gridsize, blocksize>>>(dev_a, dev_b); 

    // copy device result back to host copy of c 
    cudaMemcpy(b, dev_b, sizeof(int) , cudaMemcpyDeviceToHost); 

    printf("Reduced sum of Array elements = %d \n", b[0]); 

    cudaFree(dev_a); 
    cudaFree(dev_b); 

    return 0; 
} 

__global__ void reduce(int *g_idata, int *g_odata) { 

    __shared__ int sdata[256]; 

    // each thread loads one element from global to shared mem 
    int i = blockIdx.x*blockDim.x + threadIdx.x; 

    sdata[threadIdx.x] = g_idata[i]; 

    __syncthreads(); 
    // do reduction in shared mem 
    for (int s=1; s < blockDim.x; s *=2) 
    { 
     int index = 2 * s * threadIdx.x;; 

     if (index < blockDim.x) 
     { 
      sdata[index] += sdata[index + s]; 
     } 
     __syncthreads(); 
    } 

    // write result for this block to global mem 
    if (threadIdx.x == 0) 
     atomicAdd(g_odata,sdata[0]); 
} 

// CPU function to generate a vector of random integers 
void fill_array (int *a, int n) 
{ 
    for (int i = 0; i < n; i++) 
     a[i] = i; 
} 

Answer 1

あなたのコード内の少なくとも2つの問題

1. あなたのdev_b配列の最初の要素にatomicAddを行っているが、あなたがその要素を初期化されていませんがあります既知の値（すなわち0）。確かに、カーネルを実行する前にbをdev_bにコピーしていますが、bを既知の値に初期化していないので、それは役に立ちません。アレイbは、CまたはC++で自動的にゼロに初期化されません。 をdev_bにコピーする前に、b[0]をゼロに設定してこの問題を解決できます。

2. リダクションカーネルは1Dのケースを扱うように書かれています（すなわち、使用されるスレッドインデックスは.xに基づく1Dスレッドインデックスのみです）が、2Dスレッドブロックとグリッドを持つカーネルを起動しています。この不一致は正しく動作せず、1Dのスレッドブロックとグリッドを起動するか、カーネルを2Dインデックス（つまり.xと.y）で動作させる必要があります。私は前者（1D）を選んだ。

   $ cat t1218.cu 
   #include <stdio.h> 
   
   #define w 32 
   #define h 32 
   #define N w*h 
   
   __global__ void reduce(int *g_idata, int *g_odata); 
   void fill_array (int *a, int n); 
   
   int main(void) { 
       int a[N], b[N]; // copies of a, b, c 
       int *dev_a, *dev_b; // device copies of a, b, c 
       int size = N * sizeof(int); // we need space for 512 integers 
   
       // allocate device copies of a, b, c 
       cudaMalloc((void**)&dev_a, size); 
       cudaMalloc((void**)&dev_b, size); 
   
       fill_array(a, N); 
       b[0] = 0; //initialize the first value of b to zero 
       // copy inputs to device 
       cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice); 
       cudaMemcpy(dev_b, b, size, cudaMemcpyHostToDevice); 
   
       dim3 blocksize(256); // create 1D threadblock 
       dim3 gridsize(N/blocksize.x); //create 1D grid 
   
       reduce<<<gridsize, blocksize>>>(dev_a, dev_b); 
   
       // copy device result back to host copy of c 
       cudaMemcpy(b, dev_b, sizeof(int) , cudaMemcpyDeviceToHost); 
   
       printf("Reduced sum of Array elements = %d \n", b[0]); 
       printf("Value should be: %d \n", ((N-1)*(N/2))); 
       cudaFree(dev_a); 
       cudaFree(dev_b); 
   
       return 0; 
   } 
   
   __global__ void reduce(int *g_idata, int *g_odata) { 
   
       __shared__ int sdata[256]; 
   
       // each thread loads one element from global to shared mem 
       // note use of 1D thread indices (only) in this kernel 
       int i = blockIdx.x*blockDim.x + threadIdx.x; 
   
       sdata[threadIdx.x] = g_idata[i]; 
   
       __syncthreads(); 
       // do reduction in shared mem 
       for (int s=1; s < blockDim.x; s *=2) 
       { 
        int index = 2 * s * threadIdx.x;; 
   
        if (index < blockDim.x) 
        { 
         sdata[index] += sdata[index + s]; 
        } 
        __syncthreads(); 
       } 
   
       // write result for this block to global mem 
       if (threadIdx.x == 0) 
        atomicAdd(g_odata,sdata[0]); 
   } 
   
   // CPU function to generate a vector of random integers 
   void fill_array (int *a, int n) 
   { 
       for (int i = 0; i < n; i++) 
        a[i] = i; 
   } 
   $ nvcc -o t1218 t1218.cu 
   $ cuda-memcheck ./t1218 
   ========= CUDA-MEMCHECK 
   Reduced sum of Array elements = 523776 
   Value should be: 523776 
   ========= ERROR SUMMARY: 0 errors 
   $ 
   

   注：ここでは

があなたのコードにこれらの変更で働いた例であり、正しい結果を生成するようだ

1. カーネルと書かれたとして、あなたのコードスレッドブロックサイズ（256）の正確な倍数であるNに依存します。それはこのケースでは満足ですが、そうでないと物事が壊れます。

2. 私はproper cuda error checkingというエビデンスはありません。ここには何も出てこないだろうが、その良い習慣だ。クイックテストとして、ここで行ったようにcuda-memcheckでコードを実行してください。