【发布时间】：2015-11-18 08:55:57
【问题描述】：
我是 openacc 的新手，并且只有高级知识，因此对于我做错的任何帮助和解释将不胜感激。
我正在尝试加速（并行化）一个不那么简单的嵌套循环，该循环使用 openacc 指令更新扁平（3D 到 1D）数组。我在下面发布了一个简化的示例代码，当使用
编译时pgcc -acc -Minfo=accel test.c
给出以下错误：
call to cuStreamSynchronize returned error 700: Illegal address during kernel execution
代码：
#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
【问题讨论】：
-
我正在努力为您提供完整的修复，但您应该立即更改的一件事是将
region替换为parallel。前者来自旧的 PGI 加速器模型（OpenACC 的前身），后者是 OpenACC 语法。就个人而言，我还将copyin(p[0:1])添加到数据子句中，以明确p只是指向Phi的单个实例而不是数组的指针（编译器无法确定这一点）。即使进行了这些更改，您也可能会遇到错误，这就是为什么我只是发表评论而不发布答案的原因。 -
感谢@jefflarkin 的提醒，我将立即更换。
-
几乎可以肯定的是编译器没有正确地将数据结构的某些部分映射到设备，因此内核正在从设备中取消引用主机内存。如果您在编译器安装期间碰巧启用了“统一内存评估包”，那么有一个非常简单的解决方案。您可以使用 '-ta=tesla:managed" 将所有指向的内存放入托管内存中，这可以从主机和 GPU 中看到。您应该查看文档以了解随之而来的限制，但它应该在至少让你再次动起来。