cuda - 在 OpenACC 中使用共享内存

Question

我正在尝试使用共享内存来缓存 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 上我会知道我的 block 的确切大小，并且能够:

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

我知道我可以使用

#pragma acc cached

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

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

Answer 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 。