OpenACC：带有私有变量的嵌套循环

Question

我想加速这些嵌套循环。由于 v (NMAX=MAX(NX1, NX2, NX3)) 的维度，我知道这可能是两个外部循环并行化的冲突。我尝试使用 private 子句：

  static double **v;

  if (v == NULL) {
    v = ARRAY_2D(NMAX_POINT, NVAR, double);
  }
  
  #pragma acc parallel loop present(V, U) private(v[:NMAX_POINT][:NVAR])
  for (k = kbeg; k <= kend; k++){ g_k = k;
  #pragma acc loop
  for (j = jbeg; j <= jend; j++){ g_j = j;

    #pragma acc loop collapse(2)
    for (i = ibeg; i <= iend; i++) {  
    for (nv = 0; nv < NVAR; nv++){ 
      v[i][nv] = V[nv][k][j][i];
    }}

    #pragma acc routine(PrimToCons) seq
    PrimToCons (v, U[k][j], ibeg, iend);
  }}

我收到以下错误：

Generating present(V[:][:][:][:],U[:][:][:][:])

     Generating Tesla code

    144, #pragma acc loop seq

    146, #pragma acc loop seq

    151, #pragma acc loop gang, vector(128) collapse(2) /* blockIdx.x threadIdx.x */

    154,   /* blockIdx.x threadIdx.x collapsed */

144, Accelerator restriction: induction variable live-out from loop: g_k

     Complex loop carried dependence of v->-> prevents parallelization

146, Accelerator restriction: induction variable live-out from loop: g_j

     Loop carried dependence due to exposed use of v prevents parallelization

     Complex loop carried dependence of V->->->->,v->-> prevents parallelization

g_k 和 g_j 是 extern int。我以前从未见过消息“induction variable live-out from loop”。

编辑： 我按照建议修改了循环，但它仍然不起作用

#pragma acc parallel loop collapse(2) present(U, V) private(v[:NMAX_POINT][:NVAR])
  for (k = kbeg; k <= kend; k++){
  for (j = jbeg; j <= jend; j++){

    #pragma acc loop collapse(2)
    for (i = ibeg; i <= iend; i++) {  
    for (nv = 0; nv < NVAR; nv++){ 
      v[i][nv] = V[nv][k][j][i];
    }}
    PrimToCons (v, U[k][j], ibeg, iend, g_gamma);
  }}

我收到此错误：

Failing in Thread:1
call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

就好像编译器找不到 v、U 或 V 但在主函数中我使用了这些指令：

#pragma acc enter data copyin(data)
 #pragma acc enter data copyin(data.Vc[:NVAR][:NX3_TOT][:NX2_TOT][NX1_TOT], data.Uc[:NX3_TOT][:NX2_TOT][NX1_TOT][:NVAR])

data.Vc 和 data.Uc 是我要并行化的这个例程中的 V 和 U。

Answer 1

g_k 和 g_j 是外部整数。我从未见过“感应”的消息 变量活出循环”之前。

并行运行时，循环迭代的执行顺序是不确定的。因此，一旦退出循环，g_k 和 g_j 的值将是最后一次迭代。这会产生依赖关系，因为为了获得正确的答案（即与串行运行时一致的答案），“k”和“j”循环必须按顺序运行。

如果“g_k”和“g_j”是局部变量，那么编译器将隐式私有化它们以消除这种依赖关系。然而，由于它们是全局变量，它必须假设代码的其他部分使用结果，因此不能假设它们可以设为私有。如果您知道变量没有在其他地方使用，那么您可以通过将它们添加到“私有”子句来解决此问题。请注意，这些变量似乎并未在循环本身中使用，因此可以将其删除并在循环外分配值“kend”和“jend”。

除非在“PrimToCons”子程序中使用“g_k”和“g_j”？在这种情况下，您会遇到更大的问题，因为这会导致竞争条件，因为变量值可能会被其他线程更新，并且不再是子例程所期望的值。在这种情况下，解决方法是将“k”和“j”作为参数传递给“PrimToCons”，而不是使用“g_k”和“g_j”。

至于“v”，它也应该是“j”循环私有的，而不仅仅是“k”循环。为了解决这个问题，我建议在“k”循环的编译指示中添加一个“collapse(2)”子句，并删除关于“j”循环的循环指令。