OpenACC：'*(*(b))‘的复杂循环携带依赖项阻止并行化

Question

我使用的是带有动态数组分配的OpenACC。下面是我如何分配：

float **a;
float **b;
float **c;
float **seq;
a=(float**)malloc(SIZE*sizeof(float*));
b=(float**)malloc(SIZE*sizeof(float*));
c=(float**)malloc(SIZE*sizeof(float*));
seq=(float**)malloc(SIZE*sizeof(float*));
for(i=0; i<SIZE; i++){
        a[i]=(float*)malloc(SIZE*sizeof(float));
        b[i]=(float*)malloc(SIZE*sizeof(float));
        c[i]=(float*)malloc(SIZE*sizeof(float));
        seq[i]=(float*)malloc(SIZE*sizeof(float));
}

下面是我是如何并行矩阵add的：

#pragma acc kernels copyin(a[0:SIZE][0:SIZE],b[0:SIZE][0:SIZE]) copy(c[0:SIZE][0:SIZE])
        for (i = 0; i < SIZE; ++i) {
                for (j = 0; j < SIZE; ++j) {
                        c[i][j] = a[i][j] + b[i][j];
                }
        }

当我用pgcc编译这段代码时，它会在循环迭代中检测对float**指针的依赖，并生成所有标量内核(1个块，1个线程/块)，这与预期的性能很差：

 40, Complex loop carried dependence of '*(*(b))' prevents parallelization
     Complex loop carried dependence of '*(*(a))' prevents parallelization
     Complex loop carried dependence of '*(*(c))' prevents parallelization
     Accelerator scalar kernel generated
     CC 1.0 : 11 registers; 40 shared, 4 constant, 0 local memory bytes
     CC 2.0 : 22 registers; 0 shared, 56 constant, 0 local memory bytes

循环显然是并行的，我认为编译器也可以检测到这一点。我很好奇如何向pgcc解释？

提前谢谢。

Answer 1

我想我找到答案了。关键是使用independent子句：

    #pragma acc data copyin(a[0:SIZE][0:SIZE],b[0:SIZE][0:SIZE]) copy(c[0:SIZE][0:SIZE])
    {
             # pragma acc region 
             {
                    #pragma acc loop independent vector(16)
                    for (i = 0; i < SIZE; ++i) {
                            #pragma acc loop independent vector(16)
                            for (j = 0; j < SIZE; ++j) {
                                   c[i][j] = a[i][j] + b[i][j];
                            }
                    }
             }
    }