Pour les boucles nestedes avec CUDA

J’ai un problème avec certaines boucles nestedes que je dois convertir de C / C ++ en CUDA. Fondamentalement, j’ai 4 boucles nestedes qui partagent le même tableau et effectuent des opérations de décalage de bits.

#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 

C’est le code que je devrais convertir en CUDA. J’ai essayé de différentes manières, mais à chaque fois, j’ai la mauvaise sortie. Ici, je poste ma version de la conversion CUDA (la partie de la partie du kernel)

 #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++; } } 

Je suis nouveau dans CUDA et j’apprends toujours, mais je ne trouve vraiment pas de solution pour les boucles nestedes. Merci d’avance.

Comme nous l’avons déjà indiqué, il existe un problème d’initialisation. Ce que je recommanderais est que vous réécriviez le programme comme suit

 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; } 

Je ne l'ai pas testé, je ne peux donc pas garantir l'exactitude des indices. Je vous laisse ça.

Un peu plus d'explication: dans le code ci-dessus, seules les boucles sur i1 et i2 sont parallélisées. Cela suppose que N ** 2 est suffisamment grand par rapport au nombre de cœurs que vous avez sur votre GPU. Si ce n'est pas le cas. Les quatre boucles doivent être parallélisées pour obtenir un programme efficace. L'approche serait alors un peu différente.