单GPU验证通过双GPU失败的OpenMP多GPU域分解问题排查

在Intel开发者云的4x GPU Max 1100 + Xeon平台上，使用OpenMP、C++及icpx编译器开发程序。现有CPU并行代码及对应的GPU Offload代码，采用域分解拆分k循环到多GPU执行。单GPU验证结果正确，但双GPU执行时，虽能观测到双GPU均有计算活动，但最终计算结果错误。目前已完成以下操作：

• 确认数据映射逻辑正确
• 其他同架构多GPU代码可正常运行
• 设置了与GPU数量匹配的CPU线程数（OMP_NUM_THREADS=2）
• 添加了同步屏障

现寻求问题排查方向，请问可能遗漏了哪些关键点？

CPU参考代码 #

double rsd[ISIZ3][ISIZ2/2*2+1][ISIZ1/2*2+1][5];
double frct[ISIZ3][ISIZ2/2*2+1][ISIZ1/2*2+1][5];
double rho_i[ISIZ3][ISIZ2/2*2+1][ISIZ1/2*2+1];
double qs[ISIZ3][ISIZ2/2*2+1][ISIZ1/2*2+1];
double u[ISIZ3][ISIZ2/2*2+1][ISIZ1/2*2+1][5];

#pragma omp parallel
#pragma omp for
for(k=0; k<nz; k++){
    for(j=0; j<ny; j++){            
        for(i=0; i<nx; i++){
            for(m=0; m<5; m++){
                rsd[k][j][i][m]=-frct[k][j][i][m];
            }
            tmp=1.0/u[k][j][i][0];
            rho_i[k][j][i]=tmp;
            qs[k][j][i]=0.50*(u[k][j][i][1]*u[k][j][i][1]
                    +u[k][j][i][2]*u[k][j][i][2]
                    +u[k][j][i][3]*u[k][j][i][3])
                *tmp;
        }
    }
}

GPU Offload代码（多GPU方案） #

#pragma omp parallel
#pragma omp target teams distribute parallel for device(device_id) 
for (k = first_index; k < next_index; k++){
    for(j=0; j<ny; j++){            
        for(i=0; i<nx; i++){
            for(m=0; m<5; m++){
                rsd[k][j][i][m]=-frct[k][j][i][m];
            }
            tmp=1.0/u[k][j][i][0];
            rho_i[k][j][i]=tmp;
            qs[k][j][i]=0.50*(u[k][j][i][1]*u[k][j][i][1]
                    +u[k][j][i][2]*u[k][j][i][2]
                    +u[k][j][i][3]*u[k][j][i][3])
                *tmp;
        }
    }
}

数据映射示例 #

#pragma omp target enter data map(to: u[first_index:chunk_size_final][0:ISIZ2/2*2+1][0:ISIZ1/2*2+1][0:5]) device(device_id)

编译选项 #

icpx -std=c++14 -Ofast -qopenmp -fopenmp-targets=spir64 -xSAPPHIRERAPIDS -ipo -g -gline-tables-only -fdebug-info-for-profiling

环境变量 #

• OMP_TARGET_OFFLOAD=MANDATORY
• OMP_NUM_THREADS=2（与GPU数量匹配）

可能的排查方向 #

• 检查输出数组的映射完整性：确认rsd、rho_i、qs这些输出数组的map(from)或map(alloc)设置是否覆盖了各自的k循环分片区域。单GPU全量映射无问题，但多GPU分片时若输出区域映射不完整，会导致数据未正确回传或被覆盖。
• 验证k循环分片的边界计算：检查first_index和next_index的计算逻辑，确认两个GPU的k区间无重叠、无遗漏（比如最后一个分片的next_index是否等于nz）。边界错误会导致部分k层重复计算或未计算，结果自然出错。
• 检查临时变量tmp的作用域：GPU代码中tmp未显式声明，可能被默认视为共享变量。在OpenMP target区域中，未声明变量的默认数据共享属性可能引发竞态，建议显式声明double tmp;在循环内部，确保每个线程拥有独立的临时变量。
• 确认设备间数据同步时机：虽然添加了屏障，需检查屏障位置是否正确——必须在所有GPU完成计算并回传数据后，再执行后续CPU操作。比如是否在#pragma omp target区域结束后添加#pragma omp barrier，或因使用nowait导致同步失效。
• 检查数组维度的一致性：确认nz是否等于ISIZ3？如果nz小于ISIZ3，分片计算时可能越界访问未初始化的数组区域，导致结果异常。
• 关闭优化做最小化验证：暂时去掉-Ofast和-ipo优化选项，编译调试版本，对比单GPU和双GPU的中间结果，排查是否是优化导致分片计算逻辑被篡改。
• 检查设备ID的分配逻辑：确认device_id是每个CPU线程对应唯一的GPU ID（比如0和1），是否存在线程间设备ID冲突或未正确赋值的情况。

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