使用OpenMP将C语言三重指针热传导代码卸载到NVIDIA GPU的可行性咨询

核心结论 #

你当前使用的多层离散分配的三重指针无法直接通过OpenMP卸载到NVIDIA GPU运行，强烈建议替换为扁平化连续数组格式，原因如下：

• 三重指针的每一层内存都是独立调用malloc分配的，地址不连续，OpenMP的target内存映射机制无法自动递归映射所有层级的指针指向的内存空间，手动逐层映射不仅代码冗余度极高，还容易出现地址访问错误。
• GPU的内存访问优化依赖连续地址的合并访存特性，离散三重指针的访问会产生大量随机访存请求，即使勉强完成映射成功运行，性能也会比连续数组低5~10倍，完全无法发挥GPU的算力优势。

修改方案 #

1. 内存分配替换为连续分配 #

把原来的三层malloc替换为单块连续内存分配，用索引计算替代三维指针访问，索引映射规则为：array[i][j][k] 等价于 flat_array[i*N*N + j*N + k]，如果想保留三维写法可以用宏定义封装，示例如下：

// 宏定义封装三维索引，保留原有写法习惯
#define IDX(i,j,k) ((i)*N*N + (j)*N + (k))

// 分配代码修改示例
double *arrayOld = (double*)malloc(N*N*N*sizeof(double));
double *arrayNew = (double*)malloc(N*N*N*sizeof(double));

边界初始化、计算核函数的访问方式只需要把原来的arrayOld[i][j][k]改成arrayOld[IDX(i,j,k)]即可，原有业务逻辑不需要改动。

2. OpenMP卸载指令修改 #

原有CPU并行的omp parallel for指令需要替换为适配GPU的target系列指令，核心计算部分修改示例：

void compute_heat_transfer(double *arrayOld, double *arrayNew){
    int i,j,k;
    // 把两个数组映射到GPU设备，NUM_STEPS次循环全部在设备端执行，避免反复拷贝数据
    #pragma omp target data map(tofrom: arrayOld[0:N*N*N]) map(alloc: arrayNew[0:N*N*N])
    for(int nsteps=0; nsteps < NUM_STEPS; nsteps++){
        #pragma omp target teams distribute parallel for collapse(3) private(i,j,k)
        for(i=1; i<N-1; i++){
            for(j=1; j<N-1; j++){
                for(k=1;k<N-1;k++){
                    arrayNew[IDX(i,j,k)] = (arrayOld[IDX(i-1,j,k)] + arrayOld[IDX(i+1,j,k)] + arrayOld[IDX(i,j-1,k)] + arrayOld[IDX(i,j+1,k)] +
                             arrayOld[IDX(i,j,k-1)] + arrayOld[IDX(i,j,k+1)])/6.0;
                }
            }
        }
        #pragma omp target teams distribute parallel for collapse(3) private(i,j,k)
        for(i=1; i<N-1; i++){
            for(j=1; j<N-1; j++){
                for(k=1; k<N-1; k++){
                    arrayOld[IDX(i,j,k)] = arrayNew[IDX(i,j,k)];
                }   
            }
        }
    }
}

3. 编译指令 #

使用NVIDIA的nvc编译器编译时的指令示例：
nvc -fast -mp=gpu -o heat_sim heat_sim.c

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