У меня есть класс 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.
Проблема здесь в том, что «высота» и «ширина» являются членами данных класса. Следовательно, компилятор должен предполагать, что они могут иметь внешние ссылки на них и, следовательно, могут изменять значения во время выполнения этих циклов.
Решение состоит в том, чтобы скопировать значения в локальные переменные, а затем использовать локальные переменные в качестве границ цикла.
Обратите внимание, что поскольку у вас есть «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[:])
Других решений пока нет …