OpenACC - 嵌套循环奇怪的行为

Question

我正在使用 OpenACC 处理 LU decomposition 中的 block diagonal matrices。
当我按顺序运行我的代码时，我得到了正确的分解，而在 OpecACC 指令下执行它时，我在进行分解时得到了错误的结果。

LU 分解涉及该类型的嵌套循环（参见hereLUPSolve 函数）：

for (unsigned int i = 0; i < N; i++)
   for (unsigned int k = 0; k < i; k++)

看来，当在并行区域内的 routine seq 指令中使用这种类型的嵌套循环时，设备总是设法进入嵌套循环 即使 i=0（这也是由于k<i 条件，这是不可能的）。

我做了一个简单的代码来检查它：

#pragma acc routine seq
void test ( int* x, int const n ) {
   for (unsigned int i = 0; i < n; i++) {
      x[i] = -1;
      for (unsigned int k = 0; k < i; k++)
         x[i] = k < i;
   }
}

int main ( ) {
   unsigned const n(4);
   unsigned const nb(3);
   int x[nb*n];
   #pragma acc parallel loop copyout(x[:nb*n])
   for (unsigned int b = 0; b < nb; b++)
      test(x+b*n,n);
   // display x
}

我得到的结果是这个：

x = 0, 1, 1, 1, 0, 1, 1, 1, 0, 1, 1, 1,

但是正确的（当我在没有 OpenACC 的情况下运行代码时得到的）应该是：

x = -1, 1, 1, 1, -1, 1, 1, 1, -1, 1, 1, 1,

我一定是做错了什么，因为它不应该在i=0时进入嵌套循环...
另外，当我将循环直接放在并行区域中（不使用函数调用）时，它确实可以正常工作。

Answer 1

看起来像编译器代码生成器问题，即使 k 和 i 都为零，它也总是执行内部循环。我已提交问题报告 (TPR#24317) 并将其发送给我们的编译器工程师进行进一步评估。作为一种变通方法，在内部循环中添加一个“if”检查。

% cat test.cpp
#include <stdio.h>
#include <stdlib.h>

#pragma acc routine seq
void test ( int* x, int const n ) {
   for (unsigned int i = 0; i < n; i++) {
      x[i] = -1;
      for (unsigned int k = 0; k < i; k++) {
         if (k < i)
            x[i] = (k<i);
      }
   }
}

int main ( ) {
   unsigned const n(4);
   unsigned const nb(3);
   int x[nb*n];
   #pragma acc parallel loop copyout(x[:nb*n])
   for (unsigned int b = 0; b < nb; b++)
      test(x+b*n,n);

   for (int i=0; i <nb; ++i) {
   for (int j=0; j <n; ++j) {
     printf("%d:%d %d\n", i,j, x[i*n+j]);
  } }
   exit(0);
}
% pgc++ -acc -Minfo=acc -ta=tesla:cc60 test.cpp; a.out
test(int *, int):
      5, Generating acc routine seq
         Generating Tesla code
main:
     18, Generating copyout(x[:])
         Accelerator kernel generated
         Generating Tesla code
         20, #pragma acc loop gang, vector(3) /* blockIdx.x threadIdx.x */
0:0 -1
0:1 1
0:2 1
0:3 1
1:0 -1
1:1 1
1:2 1
1:3 1
2:0 -1
2:1 1
2:2 1
2:3 1