使用OpenACC并行化嵌套循环
我是openacc的新手，并且只具备高级知识，所以对我所做错的任何帮助和解释都将不胜感激。
我试图加速（并行化）一个不那么简单的嵌套循环，它使用openacc指令更新扁平（3D到1D）数组。 我在下面发布了一个简化的示例代码，使用时编译
pgcc -acc -Minfo=accel test.c
给出以下错误：
call to cuStreamSynchronize returned error 700: Illegal address during kernel execution
码：
#include #include #define min(a,b) (a > b) ? b : a #define max(a,b) (a 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 0 && i 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 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 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“统一内存评估包”选项使用“托管内存”。 这是一个测试版function，但它会将您的所有数据放入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); } } } }
我知道当有很多数据arrays可以做到这一点时会很痛苦，但这是我在很多代码中成功使用的一种解决方法。 遗憾的是，这是必要的，这就是我们希望在未来版本的OpenACC中提供更好的解决方案的原因。
我希望这有帮助！ 如果我能想出一个不需要额外指针的解决方案，我会更新这个答案。
Jeff是正确的，OpenACC委员会仍在研究如何使用动态数据成员标准化对聚合数据类型的支持。 但是对于PGI版本14.9或更高版本，我们已经为结构体和C ++类添加了更好的支持，因此在这种情况下，您可以通过添加create(p[0:1])来简化代码。 会发生什么是编译器将创建p的设备副本，其中仅为数据成员分配内存。 然后当你执行p->distance的复制时，将为“距离”分配内存，然后将其附加到p 。 （即运行时将填充结构中的设备指针）。
有警告。 首先，这种行为尚未标准化，因此其他编译器（如Cray，Pathscale，GCC和其他编译器）可能具有不同的行为。 第二，订单很重要。 需要在distance附近之前创建p 。 第三，更复杂的数据结构变得非常难以管理。 正如Jeff所说，使用CUDA统一内存是管理复杂数据结构的一个很好的选择。
如果您有兴趣，我的GTC2015大部分演讲都会讨论这个主题（ 链接 ）。 本演讲的重点是C ++类数据管理，但也适用于C结构。
希望这有帮助，Mat
% cat test1.c #include #include #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 create(p[0:1]) 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; } % pgcc -acc -ta=tesla:cc35 -Minfo=accel test1.c -V15.7 ; a.out solve: 19, Generating acc routine seq fast_sweep: 34, Generating create(p[:1]) Generating copy(p->distance[:1000]) 45, Loop is parallelizable 47, Loop is parallelizable Accelerator kernel generated Generating Tesla code 45, #pragma acc loop gang /* blockIdx.y */ 47, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */ start calling fast sweep print the results 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000