Насколько эффективно выполнить это ядро ​​в 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++; 
      } 
     //} 
    } 

У меня есть большой массив 1D под названием H (POPSIZE * DIMENSION). И я хочу создать новый массив diff, который сохраняет недостающие элементы в интервале 0-19, 20-39 и так далее ...

И мне нужно выполнить этот код эффектно в parralel 5000times Я пробовал это, но выполнял его только для интервала 0-19 в 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).