openacc 在内核中运行时创建数据

Question

我的任务是由 OpenACC 加速。我需要在内核计算中进行动态内存分配。我已经为它构建了一个更简单的演示，如下所示。

#include <iostream>

using namespace std;

#pragma acc routine seq
int *routine(int init) {
    int *ptr;
    #pragma acc data create(ptr[:10])
    for (int i = 0; i < 10; ++i) {
        ptr[i] = init + i;
    }
    return ptr;
}

void print_array(int *arr) {
    for (int i = 0; i < 10; ++i) {
        cout << arr[i] << " ";
    }
    cout << endl;
}

int main(void) {
    int *arrs[5];

#pragma acc kernels
    for (int i = 0; i < 5; ++i) {
        arrs[i] = routine(i);
    }

    for (int i = 0; i < 5; ++i) {
        print_array(arrs[i]);
    }
    return 0;
}

在这个演示中，我试图在内核结构中运行时调用例程。例程过程想要在 GPU 中创建一些数据，然后将一些值放入其中。

虽然我可以编译代码，但它报告运行时问题如下。

lisanhu@lisanhu-XPS-15-9550:create_and_copyout$ pgc++ -o test main.cc -acc -Minfo=accel
routine(int):
      6, Generating acc routine seq
main:
     23, Generating implicit copyout(arrs[:])
     26, Accelerator restriction: size of the GPU copy of arrs is unknown
         Loop is parallelizable
         Generating implicit copy(arrs[:][:])
         Accelerator kernel generated
         Generating Tesla code
         26, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
lisanhu@lisanhu-XPS-15-9550:create_and_copyout$ ./test 
call to cuStreamSynchronize returned error 715: Illegal instruction

我想知道我应该做些什么来完成这项任务（在内核构造的处理中动态分配内存）。如果您能提供帮助，我们将不胜感激。

Answer 1

这是未经测试的，可能很慢，但这可能会满足您的需要。

int main() {
    const int num = 20;
    int a[x] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 0};
    int* sizes = (int *)malloc(num * sizeof(int));
    int *ptrs[num];
    int* temp, *temp2;
    int sum;
    int* finished = (int *)malloc(num * sizeof(int));
    for (int x = 0; x < num; ++x){
        finished[x] = 0;
    }
    #pragma acc kernels copyin(a[0:10]) copyout(ptrs[:num][:1]) async(num*2+1)
    {
        #pragma acc loop private(temp)
        for (int i = 0; i < num; ++i){
            #pragma acc loop seq async(i)
            for (int j = 0; j < 1; ++j){
                temp = ptrs[x];
                sizes[i] = ...
            }
            while (ptrs[x] != x);
            ptrs[x] = routine(a, sizes[i]);
        }
    }

    while (true){
        sum = 0;
        for (int x = 0; x < num; ++x){
            sum += finished[x];
        }
        if (sum == num){
            break;
        }
        for (int x = 0; x < num; ++x){
            if (acc_async_test(x) != 0 && finished[x] == 0){
                finished[x] = 1;
                #pragma acc update host(sizes[x:1])
                temp = (int *)malloc(size[x] * sizeof(int));
                #pragma acc enter data copyin(temp[0:x])
                temp2 = acc_deviceptr(temp);
                ptrs[x] = temp2;
                #pragma acc update device(ptrs[x:1][0:1])
            }
        }
    }
}