cuda - 如何使用 OpenACC 优化矩阵乘法？

Question

我正在学习 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 库，结果如下:

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 获得更好的性能？
也许我选择的指令不对？

Answer 1

你得到了 14 倍的加速，根据我的经验，这对 PGI 的编译器来说非常好。

首先，您是否使用 -Minfo 进行编译？这会给你很多来自编译器的关于优化选择的反馈。

您使用的是 32x32 线程块，但根据我的经验，16x16 线程块往往会获得更好的性能。如果省略 vector(32) 子句，编译器会选择什么调度？

使用限制声明 a 和 b 可能会让编译器生成更好的代码。

仅通过查看您的代码，我不确定共享内存是否有助于提高性能。如果您的代码可以在其中存储和重用值而不是转到全局内存，则共享内存仅有助于提高性能。在这种情况下，您不会在阅读后重用 a 或 b 的任何部分。

还值得注意的是，在共享内存使用方面，我对 PGI 的编译器有过糟糕的经历。它有时会做一些有趣的事情并缓存错误的值(如果您向后迭代循环，似乎通常会发生)，产生错误的结果。我实际上必须使用未记录的 -ta=nvidia,nocache 选项编译我当前的应用程序，以使其正常工作，完全绕过共享内存的使用。