• トップ
  • G-DEPについて
  • ご購入ガイド
  • サポート
  • お問い合わせ

G-DEPトップ  >  第12回 OpenACCとPGIコンパイラについて③ -Tips-

第12回 OpenACCとPGIコンパイラについて③ -Tips-

第12回 「OpenACCとPGIコンパイラについて③」

<< 第11回   |   目次に戻る   

 前回、OpenACCの指示文について踏み込んで見ていきましたが、お試しいただけたでしょうか?

今回は、OpenACCのオプションであるClause(節)および、実際にPGIコンパイラでコンパイルする際のオプションについても触れていきます。

PGIコンパイラの詳細仕様については(株)ソフテック様のHPおよびマニュアルに記載されていますので、それも合わせてご確認頂けたらと思います。

 
 
12.1 OpenACCのClause;クラウス=「節」

 繰り返しになりますが、OpenACCで最低限覚えておくべきことは

①Data 構文(データ移動指示) acc data
②Accelerate Compute 構文(offload領域指示) acc parallel
③Loop 構文(並列化のためのマッピング)  acc loop

という3つの構文です。

 これらの構文を挿入しOpenACC対応のコンパイラでコンパイルすることで、指定した計算領域が自動的に並列化され、GPUにて高速に計算をさせることができるようになります。
ここで、それぞれの構文の後に続けてClauseを書くことで、より細かい処理を指定することができるようになります。
 
例えば、前回例で用いたC言語プログラムの場合
int main()
{
   int ...
  ...
  #pragma data copy (a[0:n], copyout(r)); // data移動指示
  #pragma acc kernels // 並列実行kernell領域の指定
  #pragma acc loop gang, vector(64) // mapping
        for( i = 0; i < n; ++i ){
          s = sinf(a[i]);
          c = cosf(a[i]);
          r[i] = s*s + c*c;
        }...
太字で書かれた部分がClauseです。

 

それでは、各構文に対するClauseをまとめていきます。
 
12.2 Data構文のClause

 copy(list)  copyin + copyout
 copyin(list)  ホスト⇒デバイスへのコピー
 copyout(list)  デバイス⇒ホストへのコピー
 create(list)  デバイスのみで使用するローカル変数の割り当て指示
 present(list)  すでにデバイス側に存在していることを指示
 present_or_copy(list)  デバイス側にデータの存在のテストを行う
 present_or_copyin(list)  デバイス側にデータの存在のテストを行う
 present_or_copyout(list)  デバイス側にデータの存在のテストを行う
 present_or_create(list)  デバイス側にデータの存在のテストを行う
 deviceptr(list) l istがデバイス側のポインタであることを宣言

Data構文を使用する際の最大のポイントは、挿入する場所です。

 例えば以下のような複数の階層で構成されたループ構文を考えてみます。この場合、もしループ内にdata構文を入れてしまうと繰り返しごとにデータ転送が行われるため、転送時間がかかり性能が出なくなってしまいます。そこでこの場合は、以下のようにループの前にdata構文を入れてでデータ転送を行い、ループ計算中は配列をGPU上に常駐させてあげるべきです。

for  ( i=0;  i<N; i++) {
    for  ( j= 0;  j<N; j++) {
        for  ( k=0;  k<N; k++) {
            c[ i*N + j ] += a[ i*N + k ] + b[ k*N + j ];
        }
    }
}

 ↓

#pragma acc data copyin( a, b ),  copyout(c)) // a,bをデバイスに渡し、cをホストに返す。
for  ( i = 0;  i < N; i++) {
#pragma acc kernels
    for  ( j 0;  j<N; j++) {
        for  ( k = 0;  k<N; k++) {
            c[ i*N + j ] += a[ i*N + k ] + b[ k*N + j ];
        }
    }
}

 その他にも、main()からサブルーチンを呼び出す際、データ転送した変数をサブルーチン側でpresent()のClauseを用いることで、すでにその変数がGPU上に存在していることをコンパイラに伝え無駄なデータ転送を省くことなどができます。

 

12.3 Accelerate Compute構文とloop構文


 前回も触れましたが、Accelerate Compute構文は並列化したい領域を指定するための指示文です。また、実行する処理によって acc parallel および acc kernels の2つがあり、それぞれでloop構文のClauseが若干異なってきます

 ここでまた、OpenACCの説明の上で重要な用語が出てきます。それは、GangWorkerVectorという3つの階層概念です。どこかで同じようなものを見かけましたね。そう、CUDAにおけるGridBlockThread(1Warp=32 threads)という階層概念を覚えているでしょうか。CUDAの階層概念はNVIDIA GPUのアーキテクチャに絡んだものであるため、全く同じ意味合いというわけではないですがOpenACCもCUDA同様に処理の分割を階層的に分けることができるというのが、ポイントです。

  • Gang = PEs:プロセッサ要素の基本集合体 → SM:ストリーミングマルチプロセッサ
  • Worker = each PE:マルチスレッディング処理 → warps within SM
  • Vector = Each thread of the PE:ベクトル命令処理 → warp内のスレッド数(=32)

ただし、NVIDIA GPUを使用する際にはハードウェアとのマッピングが難しいため、通常はGangとVectorの2つについて留意して指定すれば良さそうです。

 12.3.1 "parallel" 構文

 
 parallel構文におけるClause(acc parallel [clause])は以下のように指定。
 num_gangs (expression)
 生成するgangの数
 num_workers (expression)
 各gang内で生成するworkerの数
 vector_length (expression)
 各woker内のvector長(SIMD実行)
 private (list)
 各gang上でプライベート変数の生成
 firstprivate (list)
 ホスト側からの初期値をプライベート変数にセット
 reduction (list)
 gang間でのプライベートコピーを使用し、リダクション処理を行う
 
 parallell構文におけるloop構文のClause(acc loop [clause])は以下のように指定。
 collapse(n)
 下に続くn段数のnested(入れ子)ループに指示文を適用する
 gang
 並列領域の各gangでiterationを共有(gangの数でiterationを分割)。
 引数はなく、parallel構文のClauseで指定する。
 worker
 gang内のworkersでiterationを共有(wokerの数でiterationを分割)。
 引数はなく、parallel構文のClauseで指定する。
 vector
 SIMDモードでiterationを実行する。上のgang、workerどちらかと共に使用する。
 seq
 GPU内でloopをシーケンシャルに実行する。
 private (list)
 loopの各iterationに対してlistの変数のコピーを生成し、プライベート変数を生成
 reduction (operator: list)
 iteration間でプライベートコピーを使ってリダクション処理を行う。
 

 

 12.3.2 "kernels" 構文

 kernels構文におけるgang数やvector数の指定はloop構文のClauseで指定する。
 
kernels構文におけるloop構文のClause(acc loop [clause])は以下のように指定。
 collapse(n)
 下に続くn段数のnested(入れ子)ループに指示文を適用する
 gang (expression)
 並列領域の各gangでiterationを共有(gangの数でiterationを分割)。
 worker (expression)
 gang内のworkersでiterationを共有(wokerの数でiterationを分割)。
 vector (expression)
 iterationをlengthの数で分割してSIMDモードで実行する。
 seq
 GPU内でloopをシーケンシャルに実行する。
 Independent  直下のloopは依存性が無いことを指示
 private (list)
 loopの各iterationに対してlistの変数のコピーを生成し、プライベート変数を生成
 reduction (operator: list)
 iteration間でプライベートコピーを使ってリダクション処理を行う。
 

 

 12.3.3 Combined Directives(複合ディレクティブ)

 ちなみに、Accelerate Computing構文とloop構文は以下のように複合ディレクティブとして使用することも出来る。(直下のループを対象)

 "parallel"
  C: #pragma acc parallel loop
  FORTRAN: !$acc parallel loop ~ !$acc end parallel loop
 
 "kernels"
  C: #pragma acc kernels loop
  FORTRAN: !$acc kernels loop ~ !$acc end kernelsl loop

 

 

12.4 PGIコンパイラ側のオプション


 OpenACCのClauseをまとめてみましたが、もちろんコンパイラ側にもコンパイルオプションがあり、コンパイルの仕方を指定することが出来ます。

以下のように指定することで、OpenACCに対応したコンパイルを行い、そのコンパイルに関した情報を返します。詳細については製品のマニュアル等もご参照ください。

 FORTRAN
  $ pgfortran -acc -Minfo=accel program.f90

 
 C (C99)  ※C++は未実装
  $ pgcc -acc -Minfo=accel program.c
  • -acc :OpenACCを用いたプログラムであることを指定する。
  • -ta=nvidia :アクセラレータをNVIDIA GPUとして指定する。-accと違いはほとんどない。
  • -Minfo :コンパイルする際の情報を表示する。-Minfo=accelとするとデータの転送状況やどこが並列化されたかなどの情報が分かる。

 

12.5 実際に試してみる

 以上のClauseおよびコンパイラオプションを用いて実際にプログラムを動かして見ましょう。

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

#define N 1024 // 正方行列のサイズを指定(N×N)

void multiplication(int * restrict matA, int * restrict matB, int * restrict matC)
{
    int i,j,k,sum;

    // 行列の計算C = A * B
#pragma acc data copyin(matA[0:N*N], matB[0:N*N]), copyout(matC[0:N*N])
#pragma acc kernels
#pragma acc loop independent gang

    for (i = 0; i < N; i++) {
#pragma acc loop independent gang(32), vector(64)
        for (j = 0; j < N; j++) {
            sum = 0;
#pragma acc loop seq
            for (k = 0; k < N; k++) {
                sum += matA[i * N + k] * matB[k * N + j];
                matC[i * N + j] = sum;
            }
        }
    }
}


int main( int argc, char** argv)
{
    int i, j, k, t;
    int time;

    int* matA;
    int* matB;
    int* matC;
    int sum;

    // 時間計測
    clock_t start, stop;

    // 行列A~Cのメモリ確保
    matA = (int*)malloc(sizeof(int) * N * N);
    matB = (int*)malloc(sizeof(int) * N * N);
    matC = (int*)malloc(sizeof(int) * N * N);
    
    // 行列A、Bの各要素に乱数を入れる
    // 行列Cは初期値0
    for (i = 0; i < N; i++){
        for (j = 0; j < N; j++){
            matA[i * N + j] = rand() % (N * N);
            matB[i * N + j] = rand() % (N * N);
            matC[i * N + j] = 0;
        }
    }

    // 時間計測開始
    start = clock();

    multiplication(matA, matB, matC);

    // 時間計測終了
    stop = clock();
    
    // 計算時間結果の表示
    printf("It takes %dmsec\n", stop - start);
    printf("matC =\n");
    for(i=0; i<3; i++){
        for(j=0; j<3; j++){
            printf("%d ", matC[i*N+j]);
        }printf("\n");
    }

    // 各行列のメモリ開放
    free(matA);
    free(matB);
    free(matC);

    return 0;
}

 今回は応用として、行列積の計算を取り扱ってみます。

 ちなみにC言語の場合は、matA, matB, matC配列の各要素のポインタレベルの重複性(エイリアス)をコンパイラが判断することができず、必ずデータ依存性のコンパイル警告「~, prevents parallelization」が出ます。C99言語からrestrict keywordが追加され、その機能を使ってポインタレベルでの重複がないことを明言すると、コンパイラは最適化を行うようになります。青文字で示した部分がそうです。このrestrict keywordについては下に参考ページを載せておきます。

 まずは、-accを付けずに並列化させない場合を見てみましょう。計算結果が等しいことを示すため、配列Cの始めから3×3の部分を抜き出しています。

1024×1024の行列積計算で、約9秒かかりました。

 それでは次に「-acc -Minfo=accel」を付け足してコンパイルし、変換の情報を見てみましょう。上記のプログラムではfor (j = 0; j < N; j++) のループをgang(32), loop(64)としています。

 loop is parallelizableとなっている様に、for文が並列化可能と認識され、GPUカーネルが生成されます。ただし、matCの各要素内の計算(matAの行およびmatBの列)において連続性があるため、並列化させることはできません。

結果は約0.19秒と大幅に短縮されました。

gangやloopの( )内の数字を変えてみると計算時間はどのようになるでしょう?上手くチューニングすることで計算時間をより短縮させることができます。

 

 今回はよりOpenACCのClauseについて掘り下げてみました。ディレクティブの挿入だけでポータビリティ性を失わず、簡易に高速化させることができるのは大変便利です。

 OpenACCに関するコラムは今回が最後となります。次回からはCUDA本題に戻り、CUBLAS、CUFFT、サードパーティライブラリなどについて解説していきたいと思います。

 

<< 第11回   |   目次に戻る   


参考資料

・「プログラミング言語Cの新機能 7.6 restrict ポインタ」(http://seclan.dll.jp/c99d/c99d07.htm#dt19991018
 
・「yohhoyの日記 restrictキーワード」(http://d.hatena.ne.jp/yohhoy/20120223/p1
 
謝辞
本稿の作成にあたり、(株)ソフテック HPCエグゼクティブ・コンサルタントの加藤努 様のご協力をいただきました。
ここに記して謝意を表します。

 

(執筆 G-DEP Associate Research Engineer 東京大学大学院工学系研究科 岡安優)