1

OpenACC のネストされた機能を使用して、GPU カードの動的並列処理をアクティブにしようとしています。私は Tesla 40c を使用しており、OpenACC コンパイラは PGI バージョン 15.7 です。

私のコードはとてもシンプルです。次のコードをコンパイルしようとすると、コンパイラからこれらのメッセージが返されます

PGCC-S-0155-Illegal context for pragma: acc  parallel loop (test.cpp: 158)
PGCC/x86 Linux 15.7-0: compilation completed with severe errors

私のコード構造:

#pragma acc parallel loop
for( i = 0; i < N; i++ )
{
  // << computation >>

  int ss = A[tid].start;
  int ee = A[tid].end;

  #pragma acc parallel loop
  for(j = ss; j< ( ee + ss); j++)
  {
    // << computation >>
  }

また、ルーチン ディレクティブを使用するようにコードを変更しようとしました。しかし、私は再びコンパイルできませんでした

#pragma acc routine workers
foo(...)
{

  #pragma acc parallel loop
  for(j = ss; j< ( ee + ss); j++)
  {
    // << computation >>
  }
}

#pragma acc parallel loop
for( i = 0; i < N; i++ )
{
  // << computation >>

  int ss = A[tid].start;
  int ee = A[tid].end;

  foo(...);

}

もちろん、内部の並列ループディレクティブを使用せずにルーチン (seq、worker、gang) でのみ試しました。コンパイラはされていますが、動的並列処理は有効化されていません。

    37, Generating acc routine worker
         Generating Tesla code
         42, #pragma acc loop vector, worker /* threadIdx.x threadIdx.y */
         Loop is parallelizable

OpenACC で動的並列処理を使用するにはどうすればよいですか?

4

2 に答える 2

3

OpenACC で動的並列処理を使用するにはどうすればよいですか?

ネストされた領域 (おそらく動的並列処理を使用する) はOpenACC 2.0 仕様の新機能ですが、PGI 15.7 ではまだ実装されていないと思います。PGI 15.7 は、OpenACC 2.0 仕様の部分的な実装を表しています。

この制限については、PGI 15.7 コンパイラに同梱される PGI 15.7 リリース ノート (pgirn157.pdf) のセクション 2.7 に記載されています (これらのリリース ノートは現在、こちらから入手できます)。

OpenACC 2.0 にない機能

‣ グローバル データの宣言リンク ディレクティブは実装されていません。

‣ ネストされた並列処理 (並列またはカーネル領域内の並列およびカーネル構造) は実装されていません。

コメントに基づいて、 についていくつかの懸念がある#pragma acc routine workerため、その PGI 15.7 を使用した完全に機能する例を次に示します。

$ cat t1.c
#include <stdio.h>
#include <stdlib.h>
#define D1 4096
#define D2 4096
#define OFFS 2

#pragma acc routine worker
void my_set(int *d, int len, int val){
  int i;
  for (i = 0; i < len; i++) d[i] += val+OFFS;
}

int main(){


  int i,*data;
  data = (int *)malloc(D1*D2*sizeof(int));
  for (i = 0; i < D1*D2; i++) data[i] = 1;

#pragma acc kernels copy(data[0:D1*D2])
  for (i = 0; i < D1; i++)
    my_set(data+(i*D2), D2, 1);

  printf("%d\n", data[0]);

  return 0;
}
$ pgcc -acc -ta=tesla -Minfo=accel t1.c -o t1
my_set:
      8, Generating acc routine worker
         Generating Tesla code
         10, #pragma acc loop vector, worker /* threadIdx.x threadIdx.y */
         Loop is parallelizable
main:
     20, Generating copy(data[:16777216])
     21, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         21, #pragma acc loop gang /* blockIdx.x */
$ ./t1
4
$

ギャング並列処理は外側のループで実行され、ワー​​カー並列処理は内側の (ルーチン) ループで実行されていることに注意してください。

このメソッドは、動的並列処理に依存せず (代わりに、ルーチン レベルのワーカーと呼び出し元レベルのギャング間の並列処理の分割に依存します)、動的並列処理を呼び出しません。

動的並列処理 (CDP) のネイティブ使用は、現在 PGI 15.7ではサポートされていません。OpenACC コードから CDP を利用する他の関数 (CUDA やライブラリなど)を呼び出す (相互運用する)ことができるはずですが、現在、PGI 15.7 ではネイティブに使用されていません (サポートされていません)。

于 2015-08-12T15:11:14.300 に答える
0

「#pragma acc parallel loop」を「#pragma acc loop」に置き換えてみてください

于 2018-03-04T03:52:48.077 に答える