OpenACC - Comportement étrange de la boucle imbriquée

Question

Je travaille sur LU decomposition de block diagonal matrices en utilisant OpenACC.
Je reçois la décomposition correcte lorsque je cours mon code de manière séquentielle, alors que lorsque je l'exécute sous les directives OpecACC, j'obtiens un mauvais résultat lors de la décomposition.

décomposition LU implique boucle imbriquée de ce type (voir la fonction hereLUPSolve):

for (unsigned int i = 0; i < N; i++) 
    for (unsigned int k = 0; k < i; k++) 

Il semble que lorsque ce type de boucle imbriquée est utilisée dans une directive routine seq dans une région parallèle, le dispositif gère toujours pour entrer dans la boucle imbriquée même lorsquei=0 (ce qui ne devrait pas être possible en raison de la condition k<i).

J'ai fait un code simple à vérifier:

#pragma acc routine seq 
void test (int* x, int const n) { 
    for (unsigned int i = 0; i < n; i++) { 
     x[i] = -1; 
     for (unsigned int k = 0; k < i; k++) 
     x[i] = k < i; 
    } 
} 

int main () { 
    unsigned const n(4); 
    unsigned const nb(3); 
    int x[nb*n]; 
    #pragma acc parallel loop copyout(x[:nb*n]) 
    for (unsigned int b = 0; b < nb; b++) 
     test(x+b*n,n); 
    // display x 
} 

Le résultat que je reçois est celui-ci:

x = 0, 1, 1, 1, 0, 1, 1, 1, 0, 1, 1, 1, 

Mais le bon (que je reçois quand je lance le code sans OpenACC) devrait être:

x = -1, 1, 1, 1, -1, 1, 1, 1, -1, 1, 1, 1, 

Je dois faire quelque chose de mal parce qu'il ne doit pas entrer dans la boucle imbriquée lorsque i=0 ...
Plus quand je mets les boucles directement dans la région parallèle (sans utiliser un appel de fonction) cela fonctionne très bien.

Answer 1

Cela ressemble à un problème de générateur de code de compilation où il exécute toujours la boucle interne même lorsque k et i sont tous deux à zéro. J'ai soumis un rapport de problème (TPR # 24317) et l'ai envoyé à nos ingénieurs de compilation pour une évaluation plus poussée. En guise de solution, ajoutez un contrôle "if" dans la boucle interne.

% cat test.cpp 
#include <stdio.h> 
#include <stdlib.h> 

#pragma acc routine seq 
void test (int* x, int const n) { 
    for (unsigned int i = 0; i < n; i++) { 
     x[i] = -1; 
     for (unsigned int k = 0; k < i; k++) { 
     if (k < i) 
      x[i] = (k<i); 
     } 
    } 
} 

int main () { 
    unsigned const n(4); 
    unsigned const nb(3); 
    int x[nb*n]; 
    #pragma acc parallel loop copyout(x[:nb*n]) 
    for (unsigned int b = 0; b < nb; b++) 
     test(x+b*n,n); 

    for (int i=0; i <nb; ++i) { 
    for (int j=0; j <n; ++j) { 
    printf("%d:%d %d\n", i,j, x[i*n+j]); 
    } } 
    exit(0); 
} 
% pgc++ -acc -Minfo=acc -ta=tesla:cc60 test.cpp; a.out 
test(int *, int): 
     5, Generating acc routine seq 
     Generating Tesla code 
main: 
    18, Generating copyout(x[:]) 
     Accelerator kernel generated 
     Generating Tesla code 
     20, #pragma acc loop gang, vector(3) /* blockIdx.x threadIdx.x */ 
0:0 -1 
0:1 1 
0:2 1 
0:3 1 
1:0 -1 
1:1 1 
1:2 1 
1:3 1 
2:0 -1 
2:1 1 
2:2 1 
2:3 1