基于OpenMP在GPU实现矩阵求逆时行归一化数据竞争问题求助

你遇到的数据竞争本质是同一行归一化操作中，对角线元素matrix[i*dim + i]的读（作为归一化因子）和写（除以自身得到1.0）发生在并行迭代的不同线程/工作组中，没有保证读操作全部完成后再执行写操作。你当前通过host中转factor的方案虽然能临时解决问题，但会引入host-device数据传输开销，性能损失极大。

• 方案1：在设备侧先完成归一化因子的读取，再执行并行更新
  完全去掉host和device之间的factor来回传输逻辑，直接在target区域内先让单个线程读取对角线值存入共享变量，同步后再执行并行计算，代码修改参考：

#pragma omp target data map(tofrom: matrix[0:dim*dim], iden[0:dim*dim])
for (int i = 0; i < dim; i++) {
    #pragma omp target teams
    {
        float factor;
        // 仅由一个线程读取对角线元素作为归一化因子
        #pragma omp master
        factor = matrix[i * dim + i];
        // 保证所有线程都能拿到正确的factor
        #pragma omp barrier
        
        #pragma omp distribute parallel for
        for (int x = i; x < dim + i + 1; x++) {
            if (x < dim) {
                matrix[i * dim + x] /= factor;
            } else {
                iden[i * dim + x - dim] /= factor;
            }
        }
    }
    // 后续消元逻辑
    ...
}

该方案不需要调整迭代逻辑，也没有额外的host-device传输开销。

• 方案2：拆分迭代逻辑，最后更新对角线元素
  先并行处理除了对角线之外的所有元素，最后单独更新对角线元素为1.0，从根源上消除读写冲突：

#pragma omp target data map(tofrom: matrix[0:dim*dim], iden[0:dim*dim])
for (int i = 0; i < dim; i++) {
    float factor;
    // 先在设备侧读取一次对角线元素
    #pragma omp target map(from: factor)
    factor = matrix[i * dim + i];

    #pragma omp target teams distribute parallel for shared(factor)
    for (int x = i+1; x < dim + i + 1; x++) { // 跳过x=i的对角线位置
        if (x < dim) {
            matrix[i * dim + x] /= factor;
        } else {
            iden[i * dim + x - dim] /= factor;
        }
    }
    // 单独设置对角线为1，无任何读写冲突
    #pragma omp target
    matrix[i * dim + i] = 1.0f;
    // 后续消元逻辑
    ...
}

• 你当前的写法每次循环都启动一次target kernel，建议把外层的i循环也放到target区域内，减少kernel启动开销，该优化仅当dim较小时收益明显。
• 注意teams distribute指令后需要搭配parallel for才能让每个team内的线程并行执行，部分OpenMP编译器如果缺失该指令会默认仅用team主线程执行迭代，导致性能远低于预期。

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