0

本 GPU Gems 3、Chapter 39: Parallel Prefix Sum (Scan) with CUDAでカーネルを呼び出すコードを書きました。

ただし、得られる結果は、プレフィックススキャンではなく、一連の負の数です。

カーネル呼び出しが間違っているのでしょうか、それとも GPU Gems 3 ブックのコードに何か問題がありますか?

これが私のコードです：

#include <stdio.h>
#include <sys/time.h>
#include <cuda.h>

__global__ void kernel(int *g_odata, int  *g_idata, int n, int dim)
{
    extern __shared__ int temp[];// allocated on invocation
    int thid = threadIdx.x;
    int offset = 1;

    temp[2*thid] = g_idata[2*thid]; // load input into shared memory
    temp[2*thid+1] = g_idata[2*thid+1];
    for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree
    {
    __syncthreads();
    if (thid < d)
    {
    int ai = offset*(2*thid+1)-1;
    int bi = offset*(2*thid+2)-1;
    temp[bi] += g_idata[ai];
    }
    offset *= 2;
    }
    if (thid == 0) { temp[n - 1] = 0; } // clear the last element
    for (int d = 1; d < n; d *= 2) // traverse down tree & build scan
    {
    offset >>= 1;
    __syncthreads();
    if (thid < d)
    {
    int ai = offset*(2*thid+1)-1;
    int bi = offset*(2*thid+2)-1;
    int t = temp[ai];
    temp[ai] = temp[bi];
    temp[bi] += t;
    }
    }
    __syncthreads();
    g_odata[2*thid] = temp[2*thid]; // write results to device memory
    g_odata[2*thid+1] = temp[2*thid+1];
}

void Initialize(int  *h_in,int num_items)
{
    int j;
    for(j=0;j<num_items;j++)

        h_in[j]=j;
        printf(" input: ");
            printf("\n\n");
}

int main(int argc, char** argv)
{
    int num_items = 512;

    int*  h_in = new int[num_items];

    // Initialize problem 
    Initialize(h_in, num_items);

    int *d_in = NULL;
    cudaMalloc((void**)&d_in, sizeof(int) * num_items);

    if(cudaSuccess != cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice)) fprintf(stderr,"could not copy to gpu");

    // Allocate device output array
    int *d_out = NULL;
    cudaMalloc((void**)&d_out, sizeof(int) * (num_items+1));

    kernel<<<1,256,num_items*sizeof(int)>>>(d_out, d_in,num_items, 2);

    int* h_out= new int[num_items+1];
    if(cudaSuccess != cudaMemcpy(h_out,d_out,sizeof(int)*(num_items+1),cudaMemcpyDeviceToHost))fprintf(stderr,"could not copy back");
    int i;
    printf(" \n");
    for(i=0;i<num_items;i++)
    printf(" ,%d ",h_out[i]);
    // Cleanup
    if (h_in) delete[] h_in;
    if (h_out) delete[] h_out;
    if (d_in) cudaFree(d_in);
    if (d_out) cudaFree(d_out);

    printf("\n\n");

    return 0;
}
4

1 に答える 1

6

GPU Gems 3 の章のコードをカーネルに書き写す際に、少なくとも 1 つのエラーが発生したようです。この行は正しくありません:

temp[bi] += g_idata[ai];

そのはず：

temp[bi] += temp[ai];

あなたが今投稿したコードにその1つの変更を加えると、正しい（排他的スキャン）プレフィックスの合計が出力されるようです。私が言及する他のいくつかのことがあります：

  1. その変更がなくても、ほぼ正しい結果が得られます。したがって、大幅に異なるもの (負の数など) を取得している場合は、マシンのセットアップまたは CUDA のインストールに問題がある可能性があります。現在使用しているエラー チェックよりも厳密なcuda エラー チェックを使用することをお勧めします (ただし、チェックの 1 つにマシンのセットアップの問題が示されているはずです)。

  2. 作成されたルーチンにはいくつかの制限があります。これは単一のスレッドブロックでのみ使用でき、共有メモリアクセスでバンク競合が発生し、データセットのサイズが単一のスレッドブロックで処理できるサイズに制限されます (このルーチンはスレッドごとに 2 つの出力要素を生成するため、データセットのサイズは、スレッド数の 2 倍に等しいと予想されます)。すでに説明したように、動的な共有メモリの割り当ては、データ セットのサイズと同じ大きさにする必要があります (つまり、要素数でスレッド サイズの 2 倍)。

  3. これは学習に役立つかもしれませんが、堅牢で高速なプレフィックス スキャンが必要な場合は、この (古い) 記事から派生したものであっても、独自のコードではなく、thrustまたはcubのルーチンを使用することをお勧めします。

次のコードはあなたのものに似ていますが、上記の問題が修正されており、さまざまなデータ型で使用できるようにカーネルをテンプレート化しました。

#include <stdio.h>
#define DSIZE 512
#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


typedef int mytype;

template <typename T>
__global__ void prescan(T *g_odata, T *g_idata, int n)
{
  extern __shared__ T temp[];  // allocated on invocation
  int thid = threadIdx.x;
  int offset = 1;
  temp[2*thid] = g_idata[2*thid]; // load input into shared memory
  temp[2*thid+1] = g_idata[2*thid+1];
  for (int d = n>>1; d > 0; d >>= 1)                    // build sum in place up the tree
  {
    __syncthreads();
    if (thid < d)
    {
      int ai = offset*(2*thid+1)-1;
      int bi = offset*(2*thid+2)-1;
      temp[bi] += temp[ai];
    }
    offset *= 2;
  }
  if (thid == 0) { temp[n - 1] = 0; } // clear the last element
  for (int d = 1; d < n; d *= 2) // traverse down tree & build scan
    {
      offset >>= 1;
      __syncthreads();
      if (thid < d)
      {
         int ai = offset*(2*thid+1)-1;
         int bi = offset*(2*thid+2)-1;
         T t = temp[ai];
         temp[ai] = temp[bi];
         temp[bi] += t;
      }
    }
  __syncthreads();
  g_odata[2*thid] = temp[2*thid]; // write results to device memory
  g_odata[2*thid+1] = temp[2*thid+1];
}

int main(){

  mytype *h_i, *d_i, *h_o, *d_o;
  int dszp = (DSIZE)*sizeof(mytype);

  h_i = (mytype *)malloc(dszp);
  h_o = (mytype *)malloc(dszp);
  if ((h_i == NULL) || (h_o == NULL)) {printf("malloc fail\n"); return 1;}
  cudaMalloc(&d_i, dszp);
  cudaMalloc(&d_o, dszp);
  cudaCheckErrors("cudaMalloc fail");
  for (int i = 0 ; i < DSIZE; i++){
    h_i[i] = i;
    h_o[i] = 0;}
  cudaMemset(d_o, 0, dszp);
  cudaCheckErrors("cudaMemset fail");
  cudaMemcpy(d_i, h_i, dszp, cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy 1 fail");
  prescan<<<1,DSIZE/2, dszp>>>(d_o, d_i, DSIZE);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail");
  cudaMemcpy(h_o, d_o, dszp, cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy 2 fail");
  mytype psum = 0;
  for (int i =1; i < DSIZE; i++){
    psum += h_i[i-1];
    if (psum != h_o[i]) {printf("mismatch at %d, was: %d, should be: %d\n", i, h_o[i], psum); return 1;}
    }
  return 0;
}
于 2015-06-14T22:05:42.263 に答える