2

私はopenaccに非常に慣れておらず、高レベルの知識しか持っていないので、私が間違っていることについての助けと説明をいただければ幸いです。

openacc ディレクティブを使用してフラット化された (3D から 1D へ) 配列を更新する、それほど単純ではないネストされたループを高速化 (並列化) しようとしています。を使用してコンパイルしたときの簡略化されたサンプルコードを以下に投稿しました

pgcc -acc -Minfo=accel test.c

次のエラーが発生します。

call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

コード：

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

regionandloopディレクティブを使用する代わりに、

#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
4

2 に答える 2

5

このエラーは、GPU 上のコンピューティング カーネルが CPU ポインターを逆参照しているために発生しています。これはかなり一般的な問題であり、OpenACC 委員会が解決に取り組んでいます。このような動的なデータ構造は、実際に多くの問題を引き起こす可能性があるため、修正したいと考えています。考えられる回避策は 2 つあります。

1) コンパイラのインストール時に、PGI の「統合メモリ評価パッケージ」オプションを介して「マネージド メモリ」を使用します。これはベータ版の機能ですが、CPU と GPU の両方から見える特別なタイプのメモリにすべてのデータが格納されます。ドキュメントには、読むべき注意事項がたくさんあります。そのほとんどは、GPU で使用できるメモリの量に制限されていることと、GPU でメモリが使用されている間は CPU からメモリにアクセスできないことです。考えられる回避策の 1 つ。インストール中にそのオプションを有効にしたと仮定すると-ta=tesla:managed、コンパイラフラグに追加してオンにするだけです。私はあなたのコードでこれを試しましたが、うまくいきました。

2) コードにポインタを追加して、次のように をdistance介してアクセスするpのではなく、直接アクセスするようにします。

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 の将来のバージョンでより良いソリューションを提供したいと考えています。

これが役立つことを願っています！追加のポインターを必要としない解決策を思いつくことができれば、この回答を更新します。

于 2015-08-24T21:31:49.713 に答える