我是 openacc 的新手,只有高级知识,所以任何帮助和解释我做错的事情都将不胜感激。
我正在尝试加速(并行化)一个不太直接的嵌套循环,该循环使用 openacc 指令更新扁平化(3D 到 1D)数组。当使用
编译时,我在下面发布了一个简化的示例代码pgcc -acc -Minfo=accel test.c
出现以下错误:
调用 cuStreamSynchronize 返回错误 700:内核执行期间地址非法
代码:
#include <stdio.h>
#include <stdlib.h>
#define min(a,b) (a > b) ? b : a
#define max(a,b) (a < b) ? b : a
#define NX 10
#define NY 10
#define NZ 10
struct phiType {
double dx, dy, dz;
double * distance;
};
typedef struct phiType Phi;
#pragma acc routine seq
double solve(Phi *p, int index) {
// for simplicity just returning a value
return 2;
}
void fast_sweep(Phi *p) {
// removing boundaries
int x = NX - 2;
int y = NY - 2;
int z = NZ - 2;
int startLevel = 3;
int endLevel = x + y + z;
#pragma acc data copy(p->distance[0:NX*NY*NZ])
for(int level = startLevel; level <= endLevel; level++){
int ks = max(1, level-(y + z));
int ke = min(x, level-2);
int js = max(1, level-(x + z));
int je = min(y, level-2);
#pragma acc region
{
#pragma acc loop independent
for(int k = ks; k <= ke; k++){
#pragma acc loop independent
for(int j = js; j <= je; j++){
int i = level - (k + j);
if(i > 0 && i <= z){
int index = i * NX * NY + j * NX + k;
p->distance[index] = solve(p, index);
}
}
}
}
}
}
void create_phi(Phi *p){
p->dx = 1;
p->dy = 1;
p->dz = 1;
p->distance = (double *) malloc(sizeof(double) * NX * NY * NZ);
for(int i = 0; i < NZ; i++){
for(int j = 0; j < NY; j++){
for(int k = 0; k < NX; k++){
int index = i * NX * NY + j * NX + k;
p->distance[index] = (i*j*k == 0) ? 0 : 1;
}
}
}
}
int main()
{
printf("start \n");
Phi *p = (Phi *) malloc(sizeof(Phi));
create_phi(p);
printf("calling fast sweep \n");
fast_sweep(p);
printf(" print the results \n");
for(int i = 0; i < NZ; i++){
for(int j = 0; j < NY; j++){
for(int k = 0; k < NX; k++){
int index = i * NX * NY + j * NX + k;
printf("%f ", p->distance[index]);
}
printf("\n");
}
printf("\n");
}
return 0;
}
不使用 region
和 loop
指令,而是使用
#pragma acc kernels
产生以下错误:
solve:
19, Generating acc routine seq
fast_sweep:
34, Generating copy(p->distance[:1000])
42, Generating copy(p[:1])
45, Loop carried dependence due to exposed use of p[:1] prevents parallelization
Accelerator scalar kernel generated
47, Loop carried dependence due to exposed use of p[:i1+1] prevents parallelization
我正在运行这段代码
GNU/Linux
CentOS release 6.7 (Final)
GeForce GTX Titan
pgcc 15.7-0 64-bit target on x86-64 Linux -tp sandybridge
最佳答案
错误来自 GPU 上的计算内核取消引用 CPU 指针。这是一个非常普遍的问题,OpenACC 委员会正在努力解决这个问题。像这样的动态数据结构确实会导致很多问题,所以我们想修复它。这里有两种可能的解决方法。
1) 在编译器安装期间通过 PGI“统一内存评估包”选项使用“托管内存”。这是一个 beta 功能,但它会将您的所有数据放入一种对 CPU 和 GPU 都可见的特殊类型的内存中。您应该在文档中阅读很多注意事项,最重要的是,您受限于 GPU 上可用的内存量,并且在 GPU 上使用内存时无法从 CPU 访问内存,但它是一种可能的解决方法。假设您在安装期间启用了该选项,只需将 -ta=tesla:managed
添加到您的编译器标志中即可将其打开。我用你的代码试过了,它成功了。
2) 添加指向代码的指针,这样您就不会通过 p
访问 distance
,而是直接访问它,如下所示:
double *distance = p->distance;
#pragma acc data copy(p[0:1],distance[0:NX*NY*NZ])
for(int level = startLevel; level <= endLevel; level++){
int ks = max(1, level-(y + z));
int ke = min(x, level-2);
int js = max(1, level-(x + z));
int je = min(y, level-2);
#pragma acc parallel
{
#pragma acc loop independent
for(int k = ks; k <= ke; k++){
#pragma acc loop independent
for(int j = js; j <= je; j++){
int i = level - (k + j);
if(i > 0 && i <= z){
int index = i * NX * NY + j * NX + k;
distance[index] = solve(p, index);
}
}
}
}
我知道当有很多数据数组要执行此操作时,这可能会很痛苦,但这是我在很多代码中成功使用的解决方法。不幸的是,这是必要的,这就是为什么我们希望在未来版本的 OpenACC 中提供更好的解决方案。
希望对您有所帮助!如果我能想出一个不需要额外指针的解决方案,我会更新这个答案。
关于c - 使用 OpenACC 并行化嵌套循环,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/32190696/