如何使用 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) 子句，编译器会选择什么调度？

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

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

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