Я пытаюсь использовать общую память для кэширования вещей с помощью 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' аналогично алгоритму mmul CUDA. делает.
В основном на CUDA я бы знал точный размер своих блоков и мог бы:
- объявить разделяемую память размером блока
- скопируйте "соответствующую" часть данных в блок
- использовать эти данные
Я понимаю, что могу использовать
#pragma acc cached
и что я могу указать размеры блоков с помощью параметров vector и gang, но у меня возникли некоторые проблемы с пониманием того, как это будет отображаться в архитектуре CUDA.
Есть ли способ добиться чего-то подобного с OpenACC? Есть ли хороший учебник / ресурс по использованию директивы cached или о том, как сопоставить некоторые возможности общей памяти от CUDA до OpenACC?
