c - 使用 OpenACC 并行化嵌套循环

Question

我是 openacc 的新手，只有高级知识，所以任何帮助和解释我做错的事情都将不胜感激。

我正在尝试加速(并行化)一个不太直接的嵌套循环，该循环使用 openacc 指令更新扁平化(3D 到 1D)数组。当使用

编译时，我在下面发布了一个简化的示例代码

pgcc -acc -Minfo=accel test.c

出现以下错误:

调用 cuStreamSynchronize 返回错误 700:内核执行期间地址非法

代码:

#include <stdio.h>
#include <stdlib.h>

#define min(a,b) (a > b) ? b : a
#define max(a,b) (a < b) ? b : a

#define NX 10
#define NY 10
#define NZ 10

struct phiType {
  double dx, dy, dz;
  double * distance;
};

typedef struct phiType Phi;

#pragma acc routine seq
double solve(Phi *p, int index) {
  // for simplicity just returning a value
  return 2;
}

void fast_sweep(Phi *p) {

  // removing boundaries
  int x = NX - 2; 
  int y = NY - 2;
  int z = NZ - 2;

  int startLevel = 3;
  int endLevel   = x + y + z;

  #pragma acc data copy(p->distance[0:NX*NY*NZ])
  for(int level = startLevel; level <= endLevel; level++){
    int ks = max(1, level-(y + z));
    int ke = min(x, level-2);

    int js = max(1, level-(x + z));
    int je = min(y, level-2);

    #pragma acc region
    {
      #pragma acc loop independent
      for(int k = ks; k <= ke; k++){
        #pragma acc loop independent
        for(int j = js; j <= je; j++){
          int i = level - (k + j);
          if(i > 0 && i <= z){
            int index = i * NX * NY + j * NX + k;
            p->distance[index] = solve(p, index);
          }
        }
      }
    }
  }
}


void create_phi(Phi *p){

  p->dx = 1;
  p->dy = 1;
  p->dz = 1;

  p->distance = (double *) malloc(sizeof(double) * NX * NY * NZ);
  for(int i = 0; i < NZ; i++){
    for(int j = 0; j < NY; j++){
      for(int k = 0; k < NX; k++){
        int index = i * NX * NY + j * NX + k;
        p->distance[index] = (i*j*k == 0) ? 0 : 1;
      }
    }
  }

}


int main()
{
  printf("start \n");

  Phi *p = (Phi *) malloc(sizeof(Phi));
  create_phi(p);

  printf("calling fast sweep \n");
  fast_sweep(p);

  printf(" print the results \n");
  for(int i = 0; i < NZ; i++){
    for(int j = 0; j < NY; j++){
      for(int k = 0; k < NX; k++){
        int index = i * NX * NY + j * NX + k;
        printf("%f ", p->distance[index]);
      }
      printf("\n");
    }
    printf("\n");
  }

  return 0;
}

不使用 region 和 loop 指令，而是使用

#pragma acc kernels

产生以下错误:

solve:
     19, Generating acc routine seq
fast_sweep:
     34, Generating copy(p->distance[:1000])
     42, Generating copy(p[:1])
     45, Loop carried dependence due to exposed use of p[:1] prevents parallelization
         Accelerator scalar kernel generated
     47, Loop carried dependence due to exposed use of p[:i1+1] prevents parallelization

我正在运行这段代码

GNU/Linux
CentOS release 6.7 (Final)
GeForce GTX Titan
pgcc 15.7-0 64-bit target on x86-64 Linux -tp sandybridge 

Answer 1

错误来自 GPU 上的计算内核取消引用 CPU 指针。这是一个非常普遍的问题，OpenACC 委员会正在努力解决这个问题。像这样的动态数据结构确实会导致很多问题，所以我们想修复它。这里有两种可能的解决方法。

1) 在编译器安装期间通过 PGI“统一内存评估包”选项使用“托管内存”。这是一个 beta 功能，但它会将您的所有数据放入一种对 CPU 和 GPU 都可见的特殊类型的内存中。您应该在文档中阅读很多注意事项，最重要的是，您受限于 GPU 上可用的内存量，并且在 GPU 上使用内存时无法从 CPU 访问内存，但它是一种可能的解决方法。假设您在安装期间启用了该选项，只需将 -ta=tesla:managed 添加到您的编译器标志中即可将其打开。我用你的代码试过了，它成功了。

2) 添加指向代码的指针，这样您就不会通过 p 访问 distance，而是直接访问它，如下所示:

double *distance = p->distance;
#pragma acc data copy(p[0:1],distance[0:NX*NY*NZ])
  for(int level = startLevel; level <= endLevel; level++){
    int ks = max(1, level-(y + z));
    int ke = min(x, level-2);

    int js = max(1, level-(x + z));
    int je = min(y, level-2);

    #pragma acc parallel
    {
      #pragma acc loop independent
      for(int k = ks; k <= ke; k++){
        #pragma acc loop independent
        for(int j = js; j <= je; j++){
          int i = level - (k + j);
          if(i > 0 && i <= z){
            int index = i * NX * NY + j * NX + k;
            distance[index] = solve(p, index);
          }
        }
      }
    }

我知道当有很多数据数组要执行此操作时，这可能会很痛苦，但这是我在很多代码中成功使用的解决方法。不幸的是，这是必要的，这就是为什么我们希望在未来版本的 OpenACC 中提供更好的解决方案。

希望对您有所帮助!如果我能想出一个不需要额外指针的解决方案，我会更新这个答案。