使用 OpenACC 共享内存答案

【问题标题】：Use of shared memory with OpenACC使用 OpenACC 共享内存
【发布时间】：2012-10-07 07:05:56
【问题描述】：

我正在尝试使用共享内存来缓存 OpenACC 的内容。

基本上我正在研究的是矩阵乘法，而我所拥有的是：

typedef float ff; 

// Multiplies two square row-major matrices a and b, puts the result in c. 
void mmul(const restrict ff* a, 
          const restrict ff* b, 
          restrict ff* c, 
          const int n) { 
#pragma acc data copyin(a[0:n*n], b[0:n*n]) copy(c[0:n*n]) 
{ 

#pragma acc region 
{ 

#pragma acc loop independent vector(16) 
  for (int i = 0; i < n; ++i) { 
#pragma acc loop independent vector(16) 
    for (int j = 0; j < n; ++j) { 
      ff sum = 0; 
      for (int k = 0; k < n; ++k) { 
        sum += a[i + n * k] * b[k + n * j]; 
      } 
      c[i + n * j] = sum; 
    } 
  } 

} 
}
}

我想做的是使用 共享内存 来缓存矩阵 'a' 和 'b' 的切片，以用于计算 'c'，其方式类似于CUDA mmul 算法可以。

基本上在 CUDA 上，我会知道我的块的确切大小，并且能够：

用块的大小声明一个共享内存
将数据的“相关”部分复制到块中
使用这些数据

我知道我可以使用

#pragma acc cached

directive，并且我可以使用 vector 和 gang 选项指定块大小，但我在理解如何将其映射到 CUDA 架构时遇到了一些麻烦.

有没有办法用 OpenACC 实现类似的功能？是否有关于使用缓存指令或如何将共享内存的一些功能从 CUDA 映射到 OpenACC 的良好教程/资源？

【问题讨论】：

PGI 加速器编译器可能已经在使用共享内存。您是否使用 -Minfo 开关检查了输出？这个tutorial 可能很有趣。
是的，但是 Minfo 开关只告诉我我的实现使用了多少共享内存。虽然这很有用，但我更想知道是否有办法显式操纵这种记忆。不过，能够看到生成的高级 cuda 非常有帮助。
@leo 您找到问题的答案了吗？您是否能够在 OpenACC 中显式定义共享内存？

标签： cuda openacc

【解决方案1】：

如果您使用的是 PGI Accelerator Compiler，您可以转储生成的 PTX 文件并查看执行的底层发生了什么：

pgcc -acc -fast -Minfo -ta=nvidia,cc13,keepptx matrixMult.c -o matrixMult

生成的 PTX 将保存在当前目录中。

编辑：您可能更喜欢查看高级代码（C 或 Fortran 的 CUDA）。所以请使用-ta=nvidia,cc13,keepptx,keepgpu。

【讨论】：