c++ - 循环携带 `->` 的依赖性阻止并行化

Question

我有一个模型类，它保存模型的数据并在该数据上运行多个函数。除了具有以下设计之外，细节可能不太重要:

• 变量存储在类命名空间中。
• 变量由类的方法之一初始化和释放。
• 变量被其他几种方法使用。

该类的 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-inner loop of tiled/collapsed loop nest should not have another loop directive (actual_code.cpp: 227)

但是循环似乎是并行化的。

我正在使用 PGI 17.9 进行编译。

Answer 1

这里的问题是“height”和“width”是类数据成员。因此，编译器必须假定它们可能具有对它们的外部引用，因此可以在这些循环的执行期间更改值。

解决方案是将值复制到局部变量，然后使用局部变量作为循环边界。

请注意，由于您在外层循环中使用了“collapse(2)”，因此“independent”子句已经适用于两个循环。 (不过，“独立”是“并行”计算区域的默认设置，因此不需要。)折叠多个循环时不允许使用第二个“循环”构造。

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