Зависимость переноса цикла `-> предотвращает распараллеливание

У меня есть класс Model, который содержит данные для модели и выполняет несколько функций на этих данных. Детали, вероятно, не слишком важны, за исключением того, что они имеют следующий дизайн:

  • Переменные хранятся в пространстве имен класса.
  • Переменные инициализируются и освобождаются одним из методов класса.
  • Переменные используются несколькими другими методами.

MWE этого класса выглядит следующим образом:

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

Когда я компилирую с:

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

Я получаю следующее предупреждение:

 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

Некоторые исследования в Интернете показывают, что типичной причиной этого является возможность наложения указателей на зависимости.

Я пытался использовать *restrict а также independent (как показано), чтобы сказать компилятору все в порядке, но он игнорирует меня и не распараллеливает цикл.

Передача указателей в качестве аргументов функции с соответствующим использованием restrict устраняет ошибку, но у меня есть эстетическое предпочтение этому дизайну. Альтернативно, все методы, каждый из которых, по сути, является ядром, могут быть связаны в run() функция, но опять же, это не желательно.

Если я использую independent по внутренней петле я получаю:

PGCC-W-0155-внутренний цикл в гнезде циклических / свернутых циклов не должен иметь другой директивы цикла (actual_code.cpp: 227)

Но цикл действительно распараллеливается.

Я компилирую с PGI 17.9.

3

Решение

Проблема здесь в том, что «высота» и «ширина» являются членами данных класса. Следовательно, компилятор должен предполагать, что они могут иметь внешние ссылки на них и, следовательно, могут изменять значения во время выполнения этих циклов.

Решение состоит в том, чтобы скопировать значения в локальные переменные, а затем использовать локальные переменные в качестве границ цикла.

Обратите внимание, что поскольку у вас есть «collapse (2)» во внешнем цикле, условие «независимо» уже применяется к обоим циклам. (Хотя «независимый» является значением по умолчанию для «параллельных» вычислительных областей, поэтому он не нужен.) Вторая конструкция «цикла» не допускается при свертывании нескольких циклов.

% 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[:])
1

Другие решения

Других решений пока нет …

По вопросам рекламы [email protected]