使用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中，我会知道我的块的确切大小，并且能够：

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

我了解我可以使用

#pragma acc cached

我知道在CUDA中可以使用向量(vector)和块(gang)选项来指定块大小，但我有些困惑它们将如何映射到CUDA架构。

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

- leo

---

PGI加速器编译器可能已经在使用共享内存。您是否使用-Minfo开关检查了输出？这个教程可能会引起您的兴趣。 - Robert Crovella

是的，但Minfo开关只告诉我我的实现正在使用多少共享内存。虽然这很有用，但我更想知道是否有一种显式操作这种内存的方法。能够看到高级别的CUDA生成代码确实非常有帮助。 - leo

@leo，你找到问题的答案了吗？你能否在OpenACC中明确定义共享内存？ - mgNobody

Answer 1

如果您正在使用PGI加速编译器，您可以导出生成的PTX文件并查看执行底层的情况：

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

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

编辑：您可能更喜欢查看高级代码（CUDA for C 或 Fortran）。因此，请使用以下参数 -ta=nvidia,cc13,keepptx,keepgpu。

- lashgar