optimisation CUDA: boucles imbriquées

Question

Je suis en train d'importer CUDA dans ce code:

double square=0; 
for(int j=0;j<width; j++) { 
    double Up=0,Down=0; 
    for(int i=0;i<height; i++) { 
    if(array1[i]>0 && array2[i]>0){ 
     square = source[i*width+j]; 
     square = square*square; 
     Up += square*array2[i]/array1[i]; 
     Down += square; 
    } 
    } 
    if(Down>0){ 
    out[j] *= (1.+(Up/Down-1.)); 
    } 
} 

Dans la première tentative, je réduit la première boucle. (fonctionne bien)

int j = blockDim.x * blockIdx.x + threadIdx.x; 

double Up=0, Down=0, square=0; 
if (j<width) { 
    for(int i=0;i<height;i++) { 
    if(array1[i]>0 && array2[i]>0){ 
     square = source[i*width+j]; 
     square = square*square; 
     Up += square*array2[i]/array1[i]; 
     Down += square; 
    } 
    } 
    if(Down>0){ 
    out[j] *= (1.+(Up/Down-1.)); 
    } 
} 

Je voudrais également réduire la seconde pour la boucle, je l'ai essayé avec une grille 2D ne fonctionne pas. Ce est le noyau:

int j = blockDim.x * blockIdx.x + threadIdx.x; 
int i = blockDim.y * blockIdx.y + threadIdx.y; 
int offset = j + i * blockDim.x * gridDim.x; 

double Up[width],Down[width], square[height]; 
if (j>=width && i>=height) return; 

if(array1[i]>0 && array2[i]>0){ 
    square[i] = source[offset]*source[offset]; 
    Up[j] += square[i]*array2[i]/array1[i]; 
    Down[j] += square[i]; 
} 
if(Down[j]>0){ 
    out[j] *= (1.+(Up[j]/Down[j]-1.)); 
} 

ce qui est l'appel du noyau:

dim3 blocks(32,32); 
dim3 grid(width/32,height/32); 
kernel <<< grid, blocks >>> (...); 
cudaDeviceSynchronize(); 

... Quelle est l'erreur? il y a des solutions plus efficaces? (Je pourrais utiliser le parallélisme dynamique?)

Merci beaucoup!

Answer 1

Dans votre dernier noyau, il semble que vous vouliez le tableau de Up, Down et square persister entre les fils, mais ces réseaux sont fil local, afin que les données qu'ils contiennent n'est pas partagé entre les threads. Malheureusement, votre approche ne fonctionnerait pas même si elles étaient partagées entre les threads.

Dans votre boucle interne, le cycle en cours de la boucle utilise les données calculées lors du cycle précédent. Il n'est pas entièrement trivial de paralléliser de telles boucles, et parfois cela ne peut pas être fait du tout. Dans votre cas, une solution simple consisterait à utiliser des opérateurs atomiques pour augmenter les compteurs Up et Down, mais cela ne serait pas efficace car les opérateurs atomiques provoquent une sérialisation implicite des opérations.

Vous devriez probablement chercher à résoudre ce problème avec les primitives parallèles existantes, telles que les préfixes-sommes, qui ont déjà été optimisées. Par exemple, ceux en CUB ou Thrust.