openacc 和缓存分拣
openacc and cache tilling
-----示例代码------------
for (body1 = 0; body1 < NBODIES; body1 ++) {
for (body2=0; body2 < NBODIES; body2++) {
OUT[body1] += compute(body1, body2);
}
}
-----阻塞代码-----
for (body2 = 0; body2 < NBODIES; body2 += BLOCK) {
for (body1=0; body1 < NBODIES; body1 ++) {
for (body22=0; body22 < BLOCK; body22 ++) {
OUT[body1] += compute(body1, body2 + body22);
}
}
}
我插入 OpenACC 指令以将代码卸载到 GPU。
但是性能却在下降。
我搜索了一些论文,他们得出的结论是 OpenACC 无法利用 GPU 中的共享内存。但我认为主要原因是 tilling/blocking 阻止了并行。因为耕作带来了数据依赖。
如果 OpenACC 不提供或不鼓励代码挖掘?
如果有耕作技术改进OpenACC代码的解决方案或例子。
OpenACC 可以进行自动和显式平铺(通过 tile 子句),但是,我认为这不是您的问题。我看到的问题是,由于对 "OUT[body1]" 的依赖,body2 循环不可并行化。 OpenACC 可以并行执行标量缩减,因此您可以尝试以下操作:
#pragma acc parallel loop
for (body1 = 0; body1 < NBODIES; body1 ++) {
sum = 0.0;
#pragma acc loop reduction(+:sum)
for (body2=0; body2 < NBODIES; body2++) {
sum += compute(body1, body2);
}
OUT[body1] += sum;
}
当然,我是在猜测,所以如果这没有帮助,请 post 一个符合问题的示例。如果您正在使用 PGI,请 post 编译器反馈消息 (-Minfo=accel)。
-----示例代码------------
for (body1 = 0; body1 < NBODIES; body1 ++) {
for (body2=0; body2 < NBODIES; body2++) {
OUT[body1] += compute(body1, body2);
}
}
-----阻塞代码-----
for (body2 = 0; body2 < NBODIES; body2 += BLOCK) {
for (body1=0; body1 < NBODIES; body1 ++) {
for (body22=0; body22 < BLOCK; body22 ++) {
OUT[body1] += compute(body1, body2 + body22);
}
}
}
我插入 OpenACC 指令以将代码卸载到 GPU。 但是性能却在下降。 我搜索了一些论文,他们得出的结论是 OpenACC 无法利用 GPU 中的共享内存。但我认为主要原因是 tilling/blocking 阻止了并行。因为耕作带来了数据依赖。 如果 OpenACC 不提供或不鼓励代码挖掘? 如果有耕作技术改进OpenACC代码的解决方案或例子。
OpenACC 可以进行自动和显式平铺(通过 tile 子句),但是,我认为这不是您的问题。我看到的问题是,由于对 "OUT[body1]" 的依赖,body2 循环不可并行化。 OpenACC 可以并行执行标量缩减,因此您可以尝试以下操作:
#pragma acc parallel loop
for (body1 = 0; body1 < NBODIES; body1 ++) {
sum = 0.0;
#pragma acc loop reduction(+:sum)
for (body2=0; body2 < NBODIES; body2++) {
sum += compute(body1, body2);
}
OUT[body1] += sum;
}
当然,我是在猜测,所以如果这没有帮助,请 post 一个符合问题的示例。如果您正在使用 PGI,请 post 编译器反馈消息 (-Minfo=accel)。