OpenACC et orienté objet C++

Question

J'essaye d'écrire un code C++ orienté objet parallélisé avec OpenACC. J'ai été capable de trouver des questions de stackoverflow et des discussions GTC sur OpenACC, mais je n'ai pas trouvé d'exemples concrets de code orienté objet.

Dans this question a été montré un exemple pour un OpenACCArray qui fait de la gestion de la mémoire en arrière-plan (code disponible à http://www.pgroup.com/lit/samples/gtc15_S5233.tar). Cependant, je me demande s'il est possible de créer une classe qui gère les tableaux à un niveau supérieur. Par exemple.

struct Data 
{ 

// OpenACCArray<float> a; 

    OpenACCArray<Vector3<float>> a3; 

    Data(size_t len) { 
#pragma acc enter data copyin(this) 
//  a.resize(len); 
     a3.resize(len); 
    } 
    ~Data() { 
#pragma acc exit data delete(this) 
    } 
    void update_device() { 
//  a.update_device(); 
     a3.update_device(); 
    } 
    void update_host() { 
//  a.update_host(); 
     a3.update_host(); 
    } 
}; 

int main(int argc, char *argv[]) 
{ 
    const size_t len = 32*128; 
    Data d(len); 

    d.update_device(); 
#pragma acc kernels loop independent present(d) 
    for (int i=0; i < len; ++i) { 
    float val = (float)i/(float)len; 

    d.a3[i].x = val; 
    d.a3[i].y = i; 
    d.a3[i].z = d.a3[i].x/d.a3[i].y; 
    } 
    d.update_host(); 
    for (int i=0; i < len/128; ++i) { 
     cout << i << ": " << d.a3[i].x << "," << d.a3[i].y << "," << d.a3[i].z << endl; 
    } 
    cout << endl; 
    return 0; 
} 

Il est intéressant de ce programme fonctionne, mais dès que je décommenter OpenACCArray<float> a;, à savoir ajouter un autre membre à cette struct de données, je reçois des erreurs de mémoire. FATAL ERROR: variable in data clause is partially present on the device.

Étant donné que la struct OpenACCArray est une structure plate qui gère les indirections de pointeur seul, il devrait fonctionner pour le copier en tant que membre? Ou doit être un pointeur vers la structure et les pointeurs doivent être câblés avec des directives? Ensuite, je crains le problème que je dois utiliser des pointeurs d'alias comme suggéré par Jeff Larkin au the above mentioned question. Cela ne me dérange pas de faire le travail pour que cela fonctionne, mais je ne trouve aucune référence sur la façon de le faire. L'utilisation des directives du compilateur keepgpu,keepptx aide un peu à comprendre ce que fait le compilateur, mais je préférerais une alternative au reverse engineering du code ptx.

Tous les pointeurs vers un projet de référence ou des documents utiles sont très appréciés.

Answer 1

Dans l'en-tête OpenACCArray1.h, supprimez les deux pragmas "#pragma acc enter data create (this)". Ce qui se passe, c'est que le constructeur "Data" crée les objets "a" et "a3" sur l'appareil. Par conséquent, lorsque la seconde région de données d'entrée est rencontrée dans le constructeur OpenACCArray, le périphérique ce pointeur est déjà là.

Cela fonctionne lorsqu'il n'y a qu'un seul membre de données puisque "a3" et "Data" partagent la même adresse pour ce pointeur. Ainsi, lorsque le second pragma de saisie de données est rencontré, le contrôle actuel voit qu'il est déjà sur le périphérique et ne le crée donc pas à nouveau. Quand "a" est ajouté, la taille de "Data" est deux fois celle de "a", donc le contrôle actuel voit que le pointeur this est déjà là mais a une taille différente qu'avant. C'est ce que signifie l'erreur "partiellement présent". Les données sont là mais ont une taille différente de celle attendue.

Seul le parent/struct doit créer le pointeur this sur le périphérique.

Hope this helps, Mat