openacc verschachtelte Schleifen mit dynamischen Arrays

Question

Ich versuche, openacc anwenden, um Multicore und GPU beschleunigte Binärdateien zu entwickeln. Ich habe das Farber Buch gelesen und erfolgreich Testprogramme von dort und durch einige Online-Kurse von NVIDIA angeboten. Dann habe ich versucht, unseren Legacy-Code zu parallelisieren.

Das Unterprogramm, das ich schreibe, hat eine Abhängigkeit von einem dreidimensionalen Array, das innerhalb von drei verschachtelten Schleifen verwendet wird. Dies ist ein typisches N (N + 1)/2-paarweises Entfernungsproblem für eine Trajektorie von n Schritten. Für unser wissenschaftliches Problem ist nsteps typischerweise 1E5 bis 1E7, n Partikel sind 1E4 bis 5E5.

for(i=0 ; i < nsteps ; i++){ 
    pair = 0 ; 
    for(j=0 ; j < nparticles-1 ; j++){ 
     x1 = position[i][j][0] ; 
     y1 = position[i][j][1] ; 
     z1 = position[i][j][2] ; 
     for(k=j+1 ; k < nparticles ; k++){ 
      x2 = position[i][k][0] ; 
      y2 = position[i][k][1] ; 
      z2 = position[i][k][2] ; 
      dx2 = (x1 - x2) * (x1 - x2) ; 
      dy2 = (y1 - y2) * (y1 - y2) ; 
      dz2 = (z1 - z2) * (z1 - z2) ; 
      sdist = sqrt(dx2 + dy2 + dz2) ; 
      dist[pair] += sdist ; 
      pair++ ; 
     } 
    } 
} 

Ich habe keine Kontrolle über die Eingabeposition-Array (nsteps, nParticles) während der Kompilierung als der Code durch Python zu einer C++ Erweiterung ausgeführt wird, der den Python numpy Array zu einem Typ C-Array-Daten umwandelt. Der openacc-Code soll als Quellobjektbibliothek kompiliert werden. Die C++ Erweiterung ruft den Openacc-Code auf. Dies ist ein Legacy-Code, der diese Anordnung erfordert. Die Schritte, um das System in Python zu definieren, die C++ Erweiterung aufzurufen und auf die Datei openacc * .so zuzugreifen, funktionieren einwandfrei.

Das Problem ist, dass dieses Problem nicht leicht in dem Buch oder den Übungen aus dem Kurs offensichtlich ist und daher die Lösung und/oder Kommentare zu dem Problem anderen helfen können.

Ich habe versucht, die parallelen, Kernel und Schleife & Daten Anweisungen auf dem Code-Schnipsel mit einigen Beispielen auf Stack Overflow und anderen Quellen, die ich gefunden habe. Soweit ich weiß, befassen sich die im Faber-Buch und anderen Quellen verwendeten Beispiele nicht mit dem in dieser Frage gestellten Anwendungsfall. Im besten Fall kann ich die ersten beiden Schleifen parallelisieren, aber die innerste Schleife wird nicht parallelisiert (Schleife nicht vektorisiert: Datenabhängigkeit). Da ich nach einer allgemeinen Anleitung suche, schreibe ich nicht die Bits, die gescheitert sind, so dass eine pädagogischere Diskussion/Tipps zur Verfügung gestellt werden.

Okay, jetzt zu meinen Fragen.

1. Wie handhabe ich die unbekannten Dimensionen der Eingabe Position Arrays mit Pragma-Direktiven?
2. Wie verwalte ich die Akkumulation des Arrays dist [] (das selbst eine Länge hat, die zur Kompilierzeit aufgrund der Anzahl der Partikelpaare unbekannt ist)?
3. Welche Pragma-Direktiven werden im Allgemeinen für dieses Problem empfohlen?
4. Wie behandelt man die Abhängigkeit der "k-Schleife" von der "j-Schleife"?
5. Sollte man das Problem reduzieren, um die zu verwendenden Anweisungen zu definieren?

thx,

SB

UPDATE:

Um Ergebnisse pro Anregung von @jefflarkin liefere ich den Code für die Indizierung der dist-Array und hinzugefügt, um die vorgeschlagenen Pragmas zu berücksichtigen geändert . Das Unterprogramm kompiliert fein und läuft parallel. Ich werde jetzt mit dem Profiling beginnen, um zu sehen, wie ich die Routine optimieren kann, um die Ressourcennutzung zu maximieren. Hier ist eine Kopie des Arbeitscodes:

#pragma acc data copyin(position[nsteps][nparticles][3]) copy(dist[npairs]) 
for(i=0 ; i < nsteps ; i++){ 
    #pragma acc parallel loop 
    for(j=0 ; j < nparticles-1 ; j++){ 
     x1 = position[i][j][0] ; 
     y1 = position[i][j][1] ; 
     z1 = position[i][j][2] ; 
     #pragma acc loop 
     for(k=j+1 ; k < nparticles ; k++){ 
      x2 = position[i][k][0] ; 
      y2 = position[i][k][1] ; 
      z2 = position[i][k][2] ; 
      dx2 = (x1 - x2) * (x1 - x2) ; 
      dy2 = (y1 - y2) * (y1 - y2) ; 
      dz2 = (z1 - z2) * (z1 - z2) ; 
      sdist = sqrt(dx2 + dy2 + dz2) ; 
      local_count = ((j*nparticles)-((j*(j+1))/2))+k-(j+1) ; 
      dist[local_count] += sdist ; 
     } 
    } 
} 

Ich bekomme dieses Compiler - Ergebnis (pgC++ (16.10): CFLAGS = -fPIC -c -Fast -ACC -Minfo = accel -ta = mehradrige -O3 -shared)

38, Generating Multicore code 
39, #pragma acc loop gang 
45, Loop is parallelizable 

und GPU (CFLAGS = -v -fPIC -c -Fast -ACC - MINFO accel = -ta = tesla: cuda8, ​​FastMath -O3 -shared)

35, Generating copyin(coor[:nframes][:nparticles][:3]) 
    Generating copy(dist[:npairs]) 
38, Accelerator kernel generated 
    Generating Tesla code 
    39, #pragma acc loop gang /* blockIdx.x */ 
    45, #pragma acc loop vector(128) /* threadIdx.x */ 
45, Loop is parallelizable 

, wo die Leitung 35 i-Schleife (nsteps), 38-Linie ist die j-Schleife und die Leitung 45 ist die k-Schleife.

Answer 1

Sie sagten, dass die Größe der Arrays unbekannt ist, aber Sie können sie für Ihre Datenklauseln basierend auf nParticles berechnen. Das Positionsarray ist die Größe [nPartikel] [nPartikel] [3]. Das dist-Array ist ein wenig komplizierter, weil k für jedes j schrumpft, aber es ist die Summe von (npartikel - 1) + (npartikel - 2) + ... + 1, was ich glaube (npartikel * (npartikel - 1))/2. Obwohl die Anzahl der Partikel zur Laufzeit nicht bestimmt werden kann, muss der Compiler etwas über das Array wissen, sonst würde [] nicht funktionieren.

Ich würde vorschlagen, die Indizierung in das dist-Array zu ändern, so dass Sie den Index in diesem Array basierend auf j und k berechnen können, andernfalls benötigen Sie eine atomare Operation, um das Inkrement des Paares zu schützen. t garantieren, dass für jeden Schritt das gleiche Punktepaar an der gleichen Stelle im Entfernungsfeld steht.

Damit fixierte, würde ich denke, ein acc data auf der Stufe Schleife, acc parallel loop auf der j-Schleife, und eine einfaches acc loop auf der k-Schleife (zu helfen, den Compiler mit Abhängigkeitsanalyse out) sollten Ihnen und Laufen bringen .