Для вложенных циклов с CUDA

Question

У меня проблема с некоторыми для вложенных циклов, которые мне нужно преобразовать из C/C++ в CUDA. В основном у меня есть 4 для вложенных циклов, которые используют один и тот же массив и выполняют операции сдвига бит.

#define N 65536 

// ---------------------------------------------------------------------------------- 

int a1,a2,a3,a4, i1,i2,i3,i4; 

int Bit4CBitmapLookUp[16] = {0, 1, 3, 3, 7, 7, 7, 7, 15, 15, 15, 15, 15, 15, 15, 15}; 

int _cBitmapLookupTable[N]; 

int s = 0; // index into the cBitmapLookupTable 

for (i1 = 0; i1 < 16; i1++) 
{ 
    // first customer 
    a1 = Bit4CBitmapLookUp[i1] << 12; 

    for (i2 = 0; i2 < 16; i2++) 
    { 
     // second customer 
     a2 = Bit4CBitmapLookUp[i2] << 8; 

     for (i3 = 0; i3 < 16; i3++) 
     { 
      // third customer 
      a3 = Bit4CBitmapLookUp[i3] << 4; 

      for (i4 = 0;i4 < 16;i4++) 
      { 
       // fourth customer 
       a4 = Bit4CBitmapLookUp[i4]; 

       // now actually set the sBitmapLookupTable value 
       _cBitmapLookupTable[s] = a1 | a2 | a3 | a4; 

       s++; 

      } // for i4 
     } // for i3 
    } // for i2 
} // for i1 

Это код, который я должен преобразовать в CUDA. Я пробовал разные способы, но каждый раз, когда у меня был неправильный вывод. Здесь я отправляю свою версию преобразования CUDA (штучный из части ядра)

#define N 16 

//---------------------------------------------------------------------------------- 

// index for the GPU 
int i1 = blockDim.x * blockIdx.x + threadIdx.x; 
int i2 = blockDim.y * blockIdx.y + threadIdx.y; 
int i3 = i1; 
int i4 = i2; 

__syncthreads(); 
for(i1 = i2 = 0; i1 < N, i2 < N; i1++, i2++) 
{ 
    // first customer 
    a1 = Bit4CBitmapLookUp_device[i1] << 12; 

    // second customer 
    a2 = Bit4CBitmapLookUp_device[i2] << 8; 

    for(i3 = i4 = 0; i3 < N, i4 < N; i3++, i4++){ 
     // third customer 
     a3 = Bit4CBitmapLookUp_device[i3] << 4; 

     // fourth customer 
     a4 = Bit4CBitmapLookUp_device[i4]; 

     // now actually set the sBitmapLookupTable value 
     _cBitmapLookupTable[s] = a1 | a2 | a3 | a4; 
     s++; 
    } 
} 

Я совершенно новое в CUDA и я все еще учусь, но на самом деле я не могу найти решение для тех, для вложенных циклов , Спасибо заранее.

Answer 1

Как уже отмечалось, проблема с инициализацией. То, что я рекомендовал бы, что вы переписать программу следующим образом

int i1 = blockDim.x * blockIdx.x + threadIdx.x; 
int i2 = blockDim.y * blockIdx.y + threadIdx.y; 
int i3; 
int i4; 

while(i1 < N && i2 < N){ 
    a1 = ..; 
    a2 = ..; 
    for(i3 = i4 = 0; i3 < N, i4 < N; i3++, i4++){ 
    // third customer 
    a3 = Bit4CBitmapLookUp_device[i3] << 4; 

    // fourth customer 
    a4 = Bit4CBitmapLookUp_device[i4]; 

    // now actually set the sBitmapLookupTable value 
    _cBitmapLookupTable[s] = a1 | a2 | a3 | a4; 
    s ++; 
    } 
    s += blockDim.x*gridDim.x*blockDim.y*gridDim.y; 
    i1 += blockDim.x*gridDim.x; 
    i2 += blockDim.y*gridDim.y; 
} 

Я не проверял, так что я не могу гарантировать, что индексы являются правильными. Я оставлю это вам.

Немного больше объяснений: В приведенном выше коде распараллеливаются только циклы над i1 и i2. Это предполагает, что N ** 2 достаточно велико по сравнению с количеством ядер, которые у вас есть на вашем GPU. Если это не так. Все четыре петли необходимо распараллелить, чтобы получить эффективную программу. Тогда подход будет немного иным.