Fortran OpenACC中transfer$r意外数据传输问题排查

一、意外transfer$r数据拷贝的排查 #

transfer$r代表设备到主机的数据回拷，这类隐式拷贝通常源于编译器隐式规则、变量作用域或库调用冲突，可按以下步骤定位：

1. 强化编译输出分析
   编译时添加更详细的诊断选项，查看编译器对每个变量的内存管理决策：

   nvfortran -acc -Minfo=accel,all -acc=verbose your_code.f90
   

   重点关注工作区数组的present状态、拷贝方向，若编译器未显示相关拷贝，说明是运行时触发的隐式操作（如库内部逻辑、变量作用域溢出）。

2. 检查工作区的设备端生命周期

   • 若工作区在主程序分配后传入子例程，必须在子例程开头显式声明其设备端存在性：

     subroutine batch_dgemm(workspace, ...)
       real(C_DOUBLE), intent(inout) :: workspace(:)
       !$acc declare present(workspace)
       ...
     end subroutine
     

   • 避免在并行区域内对工作区指针做重定向操作，哪怕是条件分支中的临时赋值，都会触发编译器重新评估内存状态。

3. 排查BLAS库的隐式拷贝
   若调用的是OpenACC封装的cuBLAS，需确认工作区是否为库预期的设备端内存：

   • 改用acc_malloc直接在设备端分配工作区，跳过主机端初始化拷贝：

     use iso_c_binding
     real(C_DOUBLE), pointer :: workspace(:)
     integer(c_size_t) :: ws_size = ... ! 计算工作区大小
     call acc_malloc(ws_size, c_loc(workspace))
     !$acc enter data create(workspace)
     ! 批量DGEMM逻辑
     !$acc exit data delete(workspace)
     call acc_free(c_loc(workspace))
     

4. 运行时跟踪拷贝触发点
   用cuda-gdb调试，结合NV_ACC_NOTIFY=2的输出，在拷贝发生时查看调用栈：

   NV_ACC_NOTIFY=2 cuda-gdb ./your_executable
   

   断点设置在CUDA的cudaMemcpy函数，可直接定位到触发拷贝的代码行。

二、移除vector_length(1) num_gangs(1)后代码失效的原因 #

该指令强制单线程执行，会掩盖并行化中的数据竞争、变量作用域错误，需从以下角度排查：

1. 检查循环变量的私有化
   批量DGEMM的外层循环若未显式私有化循环变量，多gangs执行时会出现竞争：

   !$acc parallel loop private(i, mat_a, mat_b, mat_c)
   do i = 1, batch_count
     call dgemm(..., mat_a(:,:,i), mat_b(:,:,i), mat_c(:,:,i), workspace)
   end do
   

   确保循环内的临时矩阵、索引变量被标记为private。

2. 验证工作区的访问安全性
   若多个gangs共用同一工作区，需添加同步机制或为每个gang分配独立工作区：

   • 若工作区无法拆分，在DGEMM调用前后添加!$acc atomic或!$acc critical；
   • 若允许，按gang数量拆分工作区，避免并发访问冲突。

3. 检查编译器自动并行化策略
   移除显式gang/vector设置后，编译器会自动划分线程块，若代码存在循环间依赖（如工作区未正确重置），会导致计算错误。可先改用-acc=multicore在CPU上测试，排除GPU硬件特有的问题。

4. 确认DGEMM实现的线程安全性
   若为手动实现的DGEMM，需确保其并行逻辑支持多gangs/vectors执行；若调用cuBLAS，需显式设置CUDA流，避免多线程调用时的流冲突：

   integer(c_int) :: stream
   call acc_get_cuda_stream(stream)
   call cublasSetStream(cublas_handle, stream)
   

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