はじめに

プロメテックの南です。

今回はOpenACCのworker並列を取り上げます。
NVIDIA GPUをターゲットとして、worker並列を用いた高速化例を紹介します*1。

OpenACCの3階層の並列性

OpenACCでは、3階層の並列性が存在し、大きい順に gang, worker, vectorと呼ばれます。

プログラムをOpenACCでGPU化する時、通常は、これらの並列性を意識する必要はありません。
コンパイラがループに対して並列性を自動的に選択するためです。

OpenACCでは、明示的に、ループに対して並列性を指定することができます。
適切に指定すれば、並列性を活用して、さらなる高速化を図ることができるというわけです。

コンパイラによるループの半自動的な最適化

今回取り上げるworker並列は、実は、通常の利用ではあまり出てきません。
例えば、以下のような簡単な2次元ループをコンパイルすると、以下のようなコンパイラメッセージが確認できます。

$ cat example.c
#pragma acc parallel loop
for (int i = 0; i < N; i++) {
    for (int j = 0; j < M; j++) {
        // Processing...
    }
    ...
}
$ nvc -acc -Minfo=accel example.c
     21, #pragma acc loop gang /* blockIdx.x */            // 外側ループがgang並列
     22, #pragma acc loop vector(128) /* threadIdx.x */    // 内側ループがvector並列
...

このように並列化可能な(多次元)ループをOpenACCでGPU化する場合、典型的には、最外ループがgang, 最内ループがvector並列が選択されます。
コンパイラがworker並列を自動的に選択することは少なく、多くの場合、プログラマが明示的に指定（＝プログラマの協力）しないと登場しません。

worker並列が有効なケース：最内ループが短い場合の最適化

では、なぜworker並列は自動的に選択されないのでしょうか。
それはworker並列が有効な局面を、コンパイラが見抜くことが難しいためです。

worker並列が有効な典型的な例は、大きなvector長が不要な場合＝最内ループ長が小さい場合、です。
上記の例では、M が小さい場合です（典型的には32を下回る場合と考えられます）。
一般の数値計算プログラムは、実行時に入力配列の大きさからループ長が定まるため、
コンパイラが、コンパイル時にループ長を知り、それを活用することは難しいです。

小さなループ長であることを最適化に活用したい場合、プログラマの協力が必要です。
例えば、プログラマが入力配列はさほど大きくない（＝ループ長は大きくない）ことを知っているのであれば、
それをコンパイラにヒントとして知らせることで、最適化を促進することができます。
今回の場合は、コンパイラにヒントとして知らせる＝OpenACCの指示節を書き加える、です。

さて、最内ループ長が小さい場合、以下の2つの最適化が考えられます。
本記事では以下の2ケースを評価します。

1. ベクトル長を小さくする
   小さなベクトル長を指定し、無駄になるvectorを減らすことで、性能向上が狙える
2. 1. + worker並列を使う
   1つの大きなvectorを複数の小さなvectorに分けることで、性能向上が狙える

worker並列を使った性能最適化

最内ループ長が小さい例としてCSR形式を用いた疎行列ベクトル積を取り上げます*1。
ソースコード全文はBitbucketをご確認ください*2。

最適化前の、ナイーブなカーネルは以下です。
内側ループ長の row_end – row_start (疎行列のある行に含まれる非ゼロ成分の個数) は非常に小さいことが期待されます。
なお、説明のため、各ループには明示的にgang, vectorを指定しています（なくても並列化結果は同じになります）。

$ baseline.c
...
void sparse_matrix_vector_multiply(int num_rows,
                                   const int* row_offsets,
                                   const unsigned int* cols,
                                   const double* Acoefs,
                                   const double* xcoefs,
                                   double* ycoefs) {
    #pragma acc parallel loop gang
    for (int i = 0; i < num_rows; i++) {
        double sum = 0;
        int row_start = row_offsets[i];
        int row_end = row_offsets[i+1];

        #pragma acc loop vector reduction(+:sum)
        for (int j = row_start; j < row_end; j++) {
            unsigned int Acol = cols[j];
            sum += Acoefs[j] * xcoefs[Acol];
        }

        ycoefs[i] = sum;
    }
}
...

以下の2つの最適化を考えます。

1. ベクトル長を小さくする (opt1)

   
         #pragma acc parallel loop gang vector_length(32)
         for (int i = 0; i < num_rows; i++) {
           ...
           #pragma acc loop vector reduction(+:sum)
           for (int j = row_start; j < row_end; j++) {
        

2. opt1 + worker並列を使う (opt2)

   
         #pragma acc parallel loop gang worker vector_length(32) num_workers(4)
         for (int i = 0; i < num_rows; i++) {
           ...
           #pragma acc loop vector reduction(+:sum)
           for (int j = row_start; j < row_end; j++) {
        

さて、実行してみましょう。

$ make
...
// baseline
    21, #pragma acc loop gang /* blockIdx.x */
    27, #pragma acc loop vector(128) /* threadIdx.x */
...
// opt1
    22, #pragma acc loop gang /* blockIdx.x */
    28, #pragma acc loop vector(32) /* threadIdx.x */
...
// opt2
    22, #pragma acc loop gang, worker(4) /* blockIdx.x threadIdx.y */
    28, #pragma acc loop vector(32) /* threadIdx.x */
...
$ make run
...
./baseline
Execution time: 0.001513 seconds
...
./opt1
Execution time: 0.000734 seconds
...
./opt2
Execution time: 0.000501 seconds

baseline -> opt1 で 0.001513/0.000734 = 2.1倍程度
baseline -> opt2 で 0.001513/0.000501 = 3.0倍程度
の性能向上が得られているようです。

まとめ

本記事では、OpenACCのworker並列を紹介しました。
worker並列が有効な場面について説明し、worker並列による性能向上を実際のコードとその結果を交えて紹介しました。
今後のプログラム開発において、適切な並列化戦略を検討する際の参考にしていただければ幸いです。

出典

1. OpenACC Programming and Best Practices Guide
   疎行列ベクトル積の例は、5.6 Case Study – Optimize Loops に詳しい
2. worker並列 の実装例
   本コラムで紹介したプログラム