Loop führte Abhängigkeit von `->` verhindert Parallelisierung

Question

Ich habe eine Model-Klasse, die Daten für ein Modell enthält und mehrere Funktionen auf diesen Daten ausführt. Die Details sind wahrscheinlich nicht so wichtig, außer dass es das folgende Design hat:

• Variablen werden in dem Klassennamespace gespeichert.
• Variablen werden von einer der Klassenmethoden initialisiert und freigegeben.
• Variablen werden von mehreren anderen Methoden verwendet.

A MWE der Klasse sieht wie folgt aus:

#include <cstdlib> 


class Model { 
private: 
    int width; 
    int height; 
    int size; 

    int nshift[8];  //Offset from a focal cell's index to its neighbours 
    double *restrict h; //Digital elevation model (height) 
    int *restrict rec; //Index of receiving cell 

    const int NO_FLOW = -1; 
    const double SQRT2 = 1.414213562373095048801688724209698078569671875376948; 
    const double dr[8] = {1,SQRT2,1,SQRT2,1,SQRT2,1,SQRT2}; 

private: 
    void GenerateRandomTerrain(){ 
    //srand(std::random_device()()); 
    for(int y=0;y<height;y++) 
    for(int x=0;x<width;x++){ 
     const int c = y*width+x; 
     h[c] = rand()/(double)RAND_MAX; 
    } 
    } 


public: 
    Model(const int width0, const int height0) 
    : nshift{-1,-width0-1,-width0,-width0+1,1,width0+1,width0,width0-1} 
    { 
    width = width0; 
    height = height0; 
    size = width*height; 

    h  = new double[size]; 

    GenerateRandomTerrain(); 
    } 

    ~Model(){ 
    delete[] h; 
    } 

private: 
    void FindDownstream(){ 
    //! computing receiver array 
    #pragma acc parallel loop collapse(2) independent present(h,rec,width,height) 
    for(int y=2;y<height-2;y++) 
    for(int x=2;x<width-2;x++){ 
     const int c  = y*width+x; 

     //The slope must be greater than zero for there to be downhill flow; 
     //otherwise, the cell is marekd NO_FLOW 
     double max_slope = 0; 
     int max_n  = NO_FLOW; 

     #pragma acc loop seq 
     for(int n=0;n<8;n++){ 
     double slope = (h[c] - h[c+nshift[n]])/dr[n]; 
     if(slope>max_slope){ 
      max_slope = slope; 
      max_n  = n; 
     } 
     } 
     rec[c] = max_n; 
    }  
    } 

public: 
    void run(const int nstep){ 
    rec = new int[size]; 

    #pragma acc enter data copyin(h[0:size],nshift[0:8],height,width,this) create(rec[0:size]) 

    for(int step=0;step<=nstep;step++) 
     FindDownstream(); 

    #pragma acc exit data copyout(h[0:size]) delete(this,rec) 

    delete[] rec; 
    } 

}; 

int main(int argc, char **argv){ 
    Model model(300,300); 
    model.run(100); 

    return 0; 
} 

Wenn ich mit kompilieren:

pgc++ -acc -ta=tesla,pinned,cc60 -Minfo=accel -fast test.cpp -std=c++11 

ich die folgende Warnung:

51, Loop without integer trip count will be executed in sequential mode 
    Complex loop carried dependence of rec->,nshift prevents parallelization 
    Loop carried dependence of rec-> prevents parallelization 
    Loop carried backward dependence of rec-> prevents vectorization 

Einige Graben auf Das Internet zeigt, dass eine typische Ursache dafür das Pointe-Potenzial ist r Aliasing, um Abhängigkeiten zu verursachen.

Ich habe versucht, *restrict und independent zu verwenden (wie gezeigt), um dem Compiler zu sagen, alles ist in Ordnung, aber es ignoriert mich und parallelisiert die Schleife nicht.

Die Übergabe von Zeigern als Argumente für die Funktion mit der entsprechenden Verwendung von restrict beseitigt den Fehler, aber ich habe eine ästhetische Vorliebe gegen dieses Design. Alternativ könnten alle Verfahren, die jeweils im Wesentlichen ein Kern sind, in der Funktion aneinander gereiht werden, aber dies ist wiederum nicht wünschenswert.

Wenn ich independent auf der inneren Schleife verwenden, erhalte ich:

PGCC-W-0155-Innenschleife geflieste/zusammengeführten Schleife Nest sollte nicht eine weitere Schleife Richtlinie haben (actual_code.cpp: 227)

Aber die Schleife scheint zu parallelisieren.

Ich kompiliere mit PGI 17.9.

Answer 1

Das Problem hier ist, dass "height" und "width" Klassendatenelemente sind. Daher muss der Compiler annehmen, dass er externe Referenzen auf sie haben kann und daher Werte während der Ausführung dieser Schleifen ändern könnte.

Die Lösung besteht darin, die Werte in lokale Variablen zu kopieren und dann die lokalen Variablen als Schleifengrenzen zu verwenden.

Beachten Sie, dass, da Sie auf der äußeren Schleife "collapse (2)" haben, die "unabhängige" Klausel bereits für beide Schleifen gilt. (Obwohl "unabhängig" für "parallele" Rechenbereiche standardmäßig ist, ist es nicht erforderlich.) Ein zweites "Schleifen" -Konstrukt ist nicht erlaubt, wenn mehrere Schleifen kollabiert werden.

% cat test.cpp 
#include <cstdlib> 


class Model { 
private: 
    int width; 
    int height; 
    int size; 

    int nshift[8];  //Offset from a focal cell's index to its neighbours 
    double *restrict h; //Digital elevation model (height) 
    int *restrict rec; //Index of receiving cell 

    const int NO_FLOW = -1; 
    const double SQRT2 = 1.414213562373095048801688724209698078569671875376948; 
    const double dr[8] = {1,SQRT2,1,SQRT2,1,SQRT2,1,SQRT2}; 

private: 
    void GenerateRandomTerrain(){ 
    //srand(std::random_device()()); 
    for(int y=0;y<height;y++) 
    for(int x=0;x<width;x++){ 
     const int c = y*width+x; 
     h[c] = rand()/(double)RAND_MAX; 
    } 
    } 


public: 
    Model(const int width0, const int height0) : nshift{-1,-width0-1,-width0,-width0+1,1,width0+1,width0,width0-1} 
    { 
    width = width0; 
    height = height0; 
    size = width*height; 

    h  = new double[size]; 

    GenerateRandomTerrain(); 
    } 

    ~Model(){ 
    delete[] h; 
    } 

private: 
    void FindDownstream(){ 
    //! computing receiver array 
    int hgt = height; 
    int wdt = width; 
    #pragma acc parallel loop collapse(2) present(h,rec) 
    for(int y=2;y<hgt-2;y++) 
    for(int x=2;x<wdt-2;x++){ 
     const int c  = y*wdt+x; 

     //The slope must be greater than zero for there to be downhill flow; 
     //otherwise, the cell is marekd NO_FLOW 
     double max_slope = 0; 
     int max_n  = NO_FLOW; 

     #pragma acc loop seq 
     for(int n=0;n<8;n++){ 
     double slope = (h[c] - h[c+nshift[n]])/dr[n]; 
     if(slope>max_slope){ 
      max_slope = slope; 
      max_n  = n; 
     } 
     } 
     rec[c] = max_n; 
    } 
    } 

public: 
    void run(const int nstep){ 
    rec = new int[size]; 

    #pragma acc enter data copyin(this,h[0:size],nshift[0:8]) create(rec[0:size]) 

    for(int step=0;step<=nstep;step++) 
     FindDownstream(); 

    #pragma acc exit data copyout(h[0:size]) delete(rec,nshift,this) 

    delete[] rec; 
    } 

}; 

int main(int argc, char **argv){ 
    Model model(300,300); 
    model.run(100); 

    return 0; 
} 
% pgc++ test.cpp -w --c++11 -Minfo=accel -ta=tesla:cc60 -V17.10; a.out 
Model::FindDownstream(): 
    49, Generating present(h[:]) 
     Accelerator kernel generated 
     Generating Tesla code 
     51, #pragma acc loop gang, vector(128) collapse(2) /* blockIdx.x threadIdx.x */ 
     52, /* blockIdx.x threadIdx.x collapsed */ 
     61, #pragma acc loop seq 
    49, Generating implicit copy(this[:]) 
     Generating present(rec[:]) 
    61, Loop carried scalar dependence for max_slope at line 63 
Model::run(int): 
    74, Generating enter data copyin(nshift[:],h[:size]) 
     Generating enter data create(rec[:size]) 
     Generating enter data copyin(this[:1]) 
    83, Generating exit data delete(this[:1],rec[:1]) 
     Generating exit data copyout(h[:size]) 
     Generating exit data delete(nshift[:])