CUDAでこのカーネルを実行すると効果的です

Question

私のテスト機能はこのようなものです。

DIMENSION 20 
POPSIZE 5000 

    __global__ void repairT(int* H, int* diff){ 

     int tidx = blockDim.x * blockIdx.x + threadIdx.x; 
     int ii = tidx * DIMENSION; 

     //if (ii < DIMENSION * POPSIZE) 
     //{ 
      int Hdiff[DIMENSION] = { 0 }; 
      int diffcount = 0; 
      bool isInIndiv = false; 

      //complement set H 
      for (int i = 1; i <= DIMENSION; i++) 
      { 
       for (int j = ii; j < ii + DIMENSION; j++) //H for 
       { 
        if (i == H[j]) 
        { 
         isInIndiv = isInIndiv || true; 
        } 
       } 
       if (isInIndiv == false) 
       { 
        Hdiff[diffcount] = i; 
        diffcount++; 
       } 
       else 
        isInIndiv = false; 
      } 
      // diff to array 
      int diffc = ii * DIMENSION; 
      for (int i = 0; i < DIMENSION; i++) 
      { 
       diff[diffc] = Hdiff[i]; 
       diffc++; 
      } 
     //} 
    } 

私はH（POPSIZE * DIMENSION）と呼ばれる大きな1D配列を持っています。そして、新しい配列diffを作成して、間隔0-19,20-39などで不足要素を節約したいと考えています。

そしてこのコードを効果的にparralel 5000timesで実行する必要があります。これを試しましたが、 0-12の間隔でH

dim3 nbThreadsR1(128); 
dim3 nbBlocksR1((POPSIZE/nbThreadsR1.x) + 1); 
repairT << <nbBlocksR1, nbThreadsR1 >> >(d_H, d_diff); 

Pls私にいくつかのアドバイスをお願いします。

Answer 1

あなたはHにアクセスし、diffはcoalescedではありません。つまり、メモリユニットの効率はあまり高くありません。データの順序を変更するか、アクセスを統合するようにコードを変更する必要があります。

また、H [j]をかなりの回数読み取っているようです。あなたはそれが過剰を避けるためにプリロード別の小さな配列Hcacheを定義することもできます読み取ります。最後に

 int Hcache[DIMENSION]; 
     for (int j = 0; j < DIMENSION; j++) //H for 
     { 
      Hcache[j] = H[j+ii]; 
     } 
     for (int i = 1; i <= DIMENSION; i++) 
     { 
      for (int j = 0; j < ii; j++) //H for 
      { 
       if (i == Hcache[j]) 
       { 
        isInIndiv = isInIndiv || true; 
       } 
      } 
      if (isInIndiv == false) 
      { 
       Hdiff[diffcount] = i; 
       diffcount++; 
      } 
      else 
       isInIndiv = false; 
     } 

、あなたはコンパイラがレジスタに十分な自由を取得し、お使いのデバイスがHcacheとHdiffであることを、多くのように扱うことができることを確認しますレジスタファイルに格納されます（maxrregcountオプションhereを参照）。