OpenACC und objektorientierte C++

Question

Ich versuche, einen objektorientierten C++ - Code zu schreiben, der mit OpenACC parallelisiert ist. Ich konnte einige Stackoverflow-Fragen und GTC-Gespräche auf OpenACC finden, aber ich konnte keine realen Beispiele für objektorientierten Code finden.

In this question wurde ein Beispiel für eine OpenACCArray gezeigt, die einige Speicherverwaltung im Hintergrund (Code verfügbar unter http://www.pgroup.com/lit/samples/gtc15_S5233.tar). Ich frage mich jedoch, ob es möglich ist, eine Klasse zu erstellen, die die Arrays auf einer höheren Ebene verwaltet. Z.B.

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

Interessanter dieses Programm funktioniert, aber sobald ich OpenACCArray<float> a; Kommentar-, das heißt ein anderes Mitglied, dass die Daten-Struktur hinzufügen, ich Speicherfehler bekommen. FATAL ERROR: variable in data clause is partially present on the device.

Da die Struktur OpenACCArray eine flache Struktur ist, die die Zeiger-Umleitungen selbst behandelt, sollte es funktionieren, um es als Mitglied zu kopieren? Oder muss ein Zeiger auf die Struktur sein und die Zeiger müssen mit Direktiven fest verdrahtet sein? Dann fürchte ich das Problem, dass ich Alias-Zeiger wie von Jeff Larkin bei the above mentioned question vorgeschlagen verwenden müssen. Es macht mir nichts aus, die Arbeit zu machen, um das in Gang zu bringen, aber ich kann keinen Hinweis finden, wie das geht. Mithilfe von Compiler-Direktiven keepgpu,keepptx hilft ein wenig zu verstehen, was der Compiler tut, aber ich würde eine Alternative zum Reverse Engineering generierten Ptx-Code bevorzugen.

Alle Hinweise auf hilfreiche Referenzprojekte oder Dokumente werden sehr geschätzt.

Answer 1

Entfernen Sie im Header OpenACCArray1.h die beiden Pragma "#pragma acc enter data create (this)". Was passiert, ist, dass der Konstruktor "Data" die Objekte "a" und "a3" auf dem Gerät erzeugt. Wenn der zweite Dateneingabedatenbereich im OpenACCArray-Konstruktor gefunden wird, ist dieser Zeiger daher bereits vorhanden.

Es funktioniert, wenn es nur ein Datenelement gibt, da "a3" und "Data" die gleiche Adresse für den This-Zeiger teilen. Wenn das zweite Eingabe-Daten-Pragma gefunden wird, sieht die aktuelle Überprüfung daher, dass es sich bereits auf dem Gerät befindet, so dass es nicht erneut erstellt wird. Wenn "a" hinzugefügt wird, ist die Größe von "Data" doppelt so groß wie die von "a". Daher sieht die vorliegende Überprüfung, dass der this-Zeiger bereits vorhanden ist, aber eine andere Größe als zuvor hat. Das ist der Fehler "teilweise vorhanden". Die Daten sind da, haben aber eine andere Größe als erwartet.

Nur die Elternklasse/Struktur sollte den This-Zeiger auf dem Gerät erstellen.

Hope this, Mat