我正在尝试使用共享内存来缓存 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
指令,并且我可以使用vector和gang选项指定块大小,但是我在理解如何将其映射到 CUDA 架构时遇到了一些麻烦。
有没有办法用 OpenACC 实现类似的功能?是否有关于使用缓存指令或如何将共享内存的一些功能从 CUDA 映射到 OpenACC 的良好教程/资源?