在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）。我想創建新的數組差異，它保存間隔0-19,20-39等缺失的元素...

我需要在parralel中有效執行此代碼5000次 我試過這個，但它只執行對於區間0-19在H

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

請給我一些建議。

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）。