使用OpenMP将double**双指针卸载到GPU运行报错如何解决？

你这个崩溃本质是double**二级指针的内存结构和你用的map语法不匹配。double**实现的二维数组本质是指针数组：先存n个指向每一行首地址的指针，每一行的double数据是单独分配的，行和行之间内存不连续。
你写的map(to: a[0:n][0:n],b[0:n][0:n])是OpenMP给连续存储二维数组（比如double a[n][n]这种栈上连续二维数组）设计的映射语法，套在不连续的二级指针上根本没法正确拷贝数据：OpenMP会默认先找a指向的连续n个指针，但这些行指针本身存在主机内存，GPU端拿不到有效地址，一访问就触发非法内存访问，直接coredump。

你要求必须保留double**类型不能改，有两种可落地的实现方式：

方案1：分层映射内存（兼容性最好，无需修改现有数据结构） #

先把主机上存的行指针数组本身映射到设备，再逐行映射每一行的实际数据，不要用collapse(2)——不连续内存下collapse带来的地址转换开销远大于并行收益，直接对外层i循环做任务划分即可。
参考代码：

// 第一步：先映射行指针数组本身到设备
#pragma omp target enter data map(to: a[0:n], b[0:n], c[0:n])
// 第二步：逐行映射每行的实际数据，c只需要在设备分配空间不需要提前拷入
for (int i = 0; i < n; i++) {
    #pragma omp target enter data map(to: a[i][0:n], b[i][0:n])
    #pragma omp target enter data map(alloc: c[i][0:n])
}

// 执行GPU并行计算
#pragma omp target teams distribute parallel for
for (int i = 0; i < n; i++) {
    for (int j = 0; j < n; j++) {
        c[i][j] = a[i][j] + b[i][j];
    }
}

// 计算完成后拷回c的结果，释放设备端内存
for (int i = 0; i < n; i++) {
    #pragma omp target exit data map(from: c[i][0:n])
    #pragma omp target exit data map(delete: a[i][0:n], b[i][0:n])
}
#pragma omp target exit data map(delete: a[0:n], b[0:n], c[0:n])

方案2：手动管理设备端指针（性能最优，适合n较大的性能敏感场景） #

如果你的编译器支持OpenMP 4.5以上的设备内存管理接口（GCC9+、Clang12+、NVHPC等主流编译器都支持），可以手动在设备端构造和主机端结构一致的二级指针数组，减少隐式映射的开销，这种写法可以正常用collapse(2)。
参考代码：

double **a_dev, **b_dev, **c_dev;
// 先在设备端分配行指针数组的空间
#pragma omp target map(from:a_dev, b_dev, c_dev)
{
    a_dev = (double**)omp_target_alloc(n*sizeof(double*), omp_get_default_device());
    b_dev = (double**)omp_target_alloc(n*sizeof(double*), omp_get_default_device());
    c_dev = (double**)omp_target_alloc(n*sizeof(double*), omp_get_default_device());
}
// 逐行分配设备端数据空间，拷贝a、b的主机数据到设备
for (int i=0; i<n; i++) {
    double *a_row, *b_row, *c_row;
    #pragma omp target map(from:a_row, b_row, c_row)
    {
        a_row = (double*)omp_target_alloc(n*sizeof(double), omp_get_default_device());
        b_row = (double*)omp_target_alloc(n*sizeof(double), omp_get_default_device());
        c_row = (double*)omp_target_alloc(n*sizeof(double), omp_get_default_device());
    }
    // 主机到设备的数据拷贝
    omp_target_memcpy(a_row, a[i], n*sizeof(double), 0, 0, omp_get_default_device(), omp_get_initial_device());
    omp_target_memcpy(b_row, b[i], n*sizeof(double), 0, 0, omp_get_default_device(), omp_get_initial_device());
    // 把设备端的行地址写入设备上的行指针数组
    #pragma omp target map(to:a_row, b_row, c_row)
    {
        a_dev[i] = a_row;
        b_dev[i] = b_row;
        c_dev[i] = c_row;
    }
}

// 执行计算，用is_device_ptr告诉编译器传入的是设备端指针，不需要做地址转换
#pragma omp target teams distribute parallel for collapse(2) is_device_ptr(a_dev, b_dev, c_dev)
for (int i=0; i<n; i++){
    for (int j=0; j<n; j++){
        c_dev[i][j] = a_dev[i][j] + b_dev[i][j];
    }
}

// 拷回c的计算结果，释放所有设备端内存
for (int i=0; i<n; i++) {
    double *c_row = c_dev[i];
    omp_target_memcpy(c[i], c_row, n*sizeof(double), 0, 0, omp_get_initial_device(), omp_get_default_device());
    #pragma omp target is_device_ptr(c_row, a_dev[i], b_dev[i])
    {
        omp_target_free(c_row, omp_get_default_device());
        omp_target_free(a_dev[i], omp_get_default_device());
        omp_target_free(b_dev[i], omp_get_default_device());
    }
}
#pragma omp target is_device_ptr(a_dev, b_dev, c_dev)
{
    omp_target_free(a_dev, omp_get_default_device());
    omp_target_free(b_dev, omp_get_default_device());
    omp_target_free(c_dev, omp_get_default_device());
}

• 不要对double**实现的不连续二维数组直接用单条map+collapse(2)的写法，OpenMP的数组切片映射默认要求底层内存连续，不连续结构必须逐块处理
• 编译时需要指定对应GPU架构的参数，比如NVIDIA GPU编译加-fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_XX，XX替换成你的GPU算力版本（比如75对应16系显卡、86对应30系显卡）；AMD GPU对应加-fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target -march=gfxXXX，XXX替换为对应AMD GPU的架构号
• 如果n的规模比较小，逐行映射的开销占比会很高，性能敏感场景优先选方案2

内容的提问来源于stack exchange，提问作者supe345