如何使用acc_set_cuda_stream？CUDA与OpenACC流同步问题排查

背景 #

我通过以下方式创建CUDA流，并将其绑定到cuFFT计划：

创建CUDA流 #

integer(kind=cuda_stream_kind) :: stream1
istat = cudaStreamCreate(stream1)

绑定到cuFFT计划 #

err_dir = err_dir + cufftPlan2D(plan_dir1,NY,NY,CUFFT_D2Z)
err_dir = err_dir + cufftSetStream(plan_dir1,stream1)

cuFFT执行及后续OpenACC循环 #

在执行cuFFT的子程序中，传入plan_dir1的代码如下：

subroutine new_fft_dir(z,plan)

  !$acc host_data use_device(z)
  ierr = ierr + cufftExecD2Z(plan,z,z)
  !$acc end host_data

  !$acc parallel loop collapse(2) present(z)
  do i=1,NXP2
  do j=1,NY
  z(i,j) = z(i,j)/NY**2
  enddo
  enddo
  !$acc end parallel loop

问题重现 #

我尝试将OpenACC流与CUDA流stream1绑定，使用以下代码时触发语法错误：

integer(kind=cuda_stream_kind) :: stream1
istat = cudaStreamCreate(stream1)
integer :: stream
istat = cudaStreamCreate(stream1)
acc_set_cuda_stream(stream,stream1)

错误提示：

NVFORTRAN-S-0034-Syntax error at or near end of line (main.f90: 48)

我的目标是给后续OpenACC循环添加async子句，让它和FFT在同一个CUDA流中执行：

!$acc parallel loop collapse(2) present(z) async(stream)
  do i=1,NXP2
  do j=1,NY
  z(i,j) = z(i,j)/NY**2
  enddo
  enddo
  !$acc end parallel loop

请问问题是否出在使用integer(kind=cuda_stream_kind)而非cudaStream_t stream？

问题分析与解决 #

1. 语法错误的直接原因 #

你重复调用cudaStreamCreate(stream1)属于冗余代码，但并非语法错误根源。真正的问题是**acc_set_cuda_stream的参数类型与使用方式不匹配**：

• OpenACC的acc_set_cuda_stream第一个参数需要的是OpenACC异步队列的整数标识符（需通过acc_create_stream创建并返回），而你直接声明了普通integer :: stream却未初始化。
• 第二个参数是CUDA流句柄，用integer(kind=cuda_stream_kind)声明是正确的，它和cudaStream_t在Fortran中等价（cudaStream_t本质是指针类型，Fortran中用对应kind的整数存储）。

2. 正确的绑定流程 #

要让OpenACC循环与cuFFT在同一CUDA流执行，正确步骤如下：

步骤1：创建CUDA流

integer(kind=cuda_stream_kind) :: stream1
istat = cudaStreamCreate(stream1)

步骤2：创建OpenACC异步队列并绑定到CUDA流

integer :: acc_stream_id
! 创建OpenACC异步队列
istat = acc_create_stream(acc_stream_id)
! 将队列绑定到已有的CUDA流stream1
istat = acc_set_cuda_stream(acc_stream_id, stream1)

步骤3：在OpenACC循环中使用该异步队列

!$acc parallel loop collapse(2) present(z) async(acc_stream_id)
do i=1,NXP2
do j=1,NY
z(i,j) = z(i,j)/NY**2
enddo
enddo
!$acc end parallel loop

3. 额外注意事项 #

• 你已经通过cufftSetStream将plan_dir1绑定到stream1，这部分是正确的，确保了cuFFT执行在目标流中。
• 若需要同步流，可使用cudaStreamSynchronize(stream1)或acc_wait(acc_stream_id)。

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