6

私は OpenACC (PGI のコンパイラを使用) を学習しており、行列乗算の例を最適化しようとしています。これまでに思いついた最速の実装は次のとおりです。

void matrix_mul(float *restrict r, float *a, float *b, int N, int accelerate){

#pragma acc data copyin (a[0: N * N ], b[0: N * N]) copyout (r [0: N * N ]) if(accelerate)
{
# pragma acc region if(accelerate)
{
# pragma acc loop independent vector(32) 
for (int j = 0; j < N; j ++)
{    
   # pragma acc loop independent vector(32) 
   for (int i = 0; i < N ; i ++ )
   {
      float sum = 0;
      for (int k = 0; k < N ; k ++ ) {
         sum += a [ i + k*N ] * b [ k + j * N ];
      }
      r[i + j * N ] = sum ;
   }
}
}
}

これにより、サイズ 32x32 スレッドのスレッド ブロックが生成され、これまでのところ最高のパフォーマンスが得られます。ベンチマークは次のとおりです。

Matrix multiplication (1500x1500): 
GPU: Geforce GT650 M, 64-bit Linux 

Data sz             : 1500     
Unaccelerated:
     matrix_mul() time    : 5873.255333 msec
Accelerated:
     matrix_mul() time    : 420.414700 msec

Data size             : 1750 x 1750     
    matrix_mul() time    : 876.271200 msec
Data size             : 2000 x 2000     
    matrix_mul() time    : 1147.783400 msec
Data size             : 2250 x 2250     
    matrix_mul() time    : 1863.458100 msec
Data size             : 2500 x 2500     
    matrix_mul() time    : 2516.493200 msec

残念ながら、生成された CUDA コードは非常に原始的 (共有メモリを使用していないなど) であるため、手動で最適化された CUDA プログラムと競合できないことに気付きました。リファレンス実装として、Arrayfire lib を使用して次の結果を得ました。

Arrayfire 1500 x 1500 matrix mul
CUDA toolkit 4.2, driver 295.59
GPU0 GeForce GT 650M, 2048 MB, Compute 3.0 (single,double)
Memory Usage: 1932 MB free (2048 MB total)
af:  0.03166 seconds

Arrayfire 1750 x 1750 matrix mul
 af:  0.05042 seconds
Arrayfire 2000 x 2000 matrix mul
 af:  0.07493 seconds
Arrayfire 2250 x 2250 matrix mul
 af:  0.10786 seconds
Arrayfire 2500 x 2500 matrix mul
 af:  0.14795 seconds

OpenACC からより良いパフォーマンスを得る方法について何か提案があるのだろうか? おそらく、ディレクティブの選択が正しくないのでしょうか?

4

1 に答える 1

5

あなたは14倍のスピードアップで正しくやっています。これは私の経験ではPGIのコンパイラにとってかなり良いことです。

まず、-Minfoでコンパイルしていますか？これにより、最適化の選択に関してコンパイラから多くのフィードバックが得られます。

32x32スレッドブロックを使用していますが、私の経験では、16x16スレッドブロックの方がパフォーマンスが向上する傾向があります。vector（32）句を省略した場合、コンパイラはどのスケジューリングを選択しますか？

制限付きでaとbを宣言すると、コンパイラーがより適切なコードを生成できるようになる可能性があります。

コードを見ただけでは、共有メモリがパフォーマンスに役立つかどうかはわかりません。共有メモリは、コードがグローバルメモリに移動する代わりに、そこに値を格納して再利用できる場合にのみ、パフォーマンスの向上に役立ちます。この場合、読んだ後にaまたはbのどの部分も再利用していません。

共有メモリの使用に関しては、PGIのコンパイラで悪い経験をしたことも注目に値します。それは時々面白いことをし、間違った値をキャッシュし（ループを逆方向に繰り返すとほとんど起こるようです）、間違った結果を生成します。共有メモリの使用を完全にバイパスして、ドキュメント化されていない-ta = nvidia、nocacheオプションを使用して現在のアプリケーションをコンパイルし、正しく動作させる必要があります。

于 2012-08-13T11:54:09.627 に答える