OpenACC - 私有二维数组

OpenACC - Private 2D array

我正在处理块对角矩阵(每个块具有相同的大小)并且当我使用 private 动态分配的二维数组时我有一个 illegal address error...

// NB is the number of block
// N is the block size
// A is the main matrix (block diagonal)

double** B; // a block
B = new double*[N];
for (unsigned int i = 0; i < N; i++)
    B[i] = new double[N];

#pragma acc parallel loop private(B[:N][:N]) copyin(A[:NB*N][:NB*N])
for (unsigned int b = 0; b < NB; b++) {
    #pragma acc loop
    for (unsigned int i = 0; i < N; i++) {
        #pragma acc loop
        for (unsigned int j = 0; j < N; j++) {
            B[i][j] = A[b*N+i][b*N+j];
        }
    }
    // process B
}

for (unsigned int i = 0; i < N; i++)
    delete[] B[i];
delete[] B;

我得到的错误是:

call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

如果我将数组展平为一维数组并使用词典索引或静态二维数组,效果很好,但我使用的函数需要 double** 作为参数,所以我更愿意坚持使用动态二维数组...

我已经阅读了规范中的 private 条款,但它并没有说不支持动态二维数组,所以我想我做错了什么...

抱歉,不支持在私有子句中使用指针数组。问题在于编译器运行时必须为每个 gang、worker 或 vector(取决于带有 private 子句的循环)动态创建一个 private 并填充所有设备指针。这将带来极高的管理费用。

如果"B"是一个固定大小的静态数组,"double B[N][N]",那么你可以在私有子句中使用它。

否则,我建议通过添加第三维来手动私有化数组。

 // NB is the number of block
// N is the block size
// A is the main matrix (block diagonal)

double*** B; // a block
B = new double**[NB];
for (unsigned int i = 0; i < NB; i++) {
  B[i] = new double*[N];
  for (unsigned int j = 0; j < N; j++) {
    B[i][j] = new double[N];
}}


#pragma acc parallel loop create(B[:NB][:N][:N]) copyin(A[:NB*N][:NB*N])
for (unsigned int b = 0; b < NB; b++) {
    #pragma acc loop
    for (unsigned int i = 0; i < N; i++) {
        #pragma acc loop
        for (unsigned int j = 0; j < N; j++) {
            B[b][i][j] = A[b*N+i][b*N+j];
        }
    }
    // process B
}

for (unsigned int i = 0; i < NB; i++) {
  for (unsigned int j = 0; j < N; j++) {
     delete[] B[i][j];
  }
  delete[] B[i];
}
delete[] B;