cuda - 在 Cuda Fortran 中使用其他文件中的托管模块变量的问题
问题描述
我在文件p2.f95
中有一个托管变量module a1
。我将此托管变量传递给文件中的内核p6.f95
。我能够访问文件中的变量,p6.f95
但无法从内核访问相同的变量,即;我无法从设备打印变量值。有人可以解释为什么......?
文件1:p2.f95
module a1
implicit none
integer(kind=4),managed::lnode
end module a1
文件2:p6.f95
module kernelsubroutine
contains
attributes(global) subroutine kernel(lnode)
implicit none
integer(kind=4)::lnode
print*,"on device",lnode
end subroutine kernel
end module kernelsubroutine
program main
use a1
use cudafor
use kernelsubroutine
implicit none
integer(kind=4)::n
lnode = 1
print*,"on host",lnode
call kernel<<<1,1>>>(lnode)
n = cudaDeviceSynchronize()
end program main
编译命令1:$pgf95 -Mcuda=rdc p2.f95 p6.f95
编译命令2:
$pgf95 -Mcuda=rdc -c p2.f95
$pgf95 -Mcuda=rdc p6.f95 p2.o
p6.f95:
使用两组编译命令,输出保持不变
输出命令:$ cuda-memcheck ./a.out
输出:
========= CUDA-MEMCHECK
on host 1
========= Invalid __global__ read of size 4
========= at 0x000010c8 in /home/vsriram/Documents/fortran_programs/p6.f95:11:kernelsubroutine_kernel_
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x0060cf40 is out of bounds
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so [0x25428a]
========= Host Frame:/opt/nvidia/hpc_sdk/Linux_x86_64/21.7/cuda/11.4/lib64/libcudart.so.11.0 [0x1402c]
========= Host Frame:/opt/nvidia/hpc_sdk/Linux_x86_64/21.7/cuda/11.4/lib64/libcudart.so.11.0 (cudaLaunchKernel + 0x1d8) [0x67e58]
========= Host Frame:/opt/nvidia/hpc_sdk/Linux_x86_64/21.7/compilers/lib/libcudafor.so (__pgiLaunchKernel + 0x1a6) [0x11194]
========= Host Frame:./a.out [0x14f4]
========= Host Frame:./a.out [0x1393]
=========
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
========= Saved host backtrace up to driver entry point at error
========= Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so [0x355b43]
========= Host Frame:/opt/nvidia/hpc_sdk/Linux_x86_64/21.7/cuda/11.4/lib64/libcudart.so.11.0 (cudaDeviceSynchronize + 0x127) [0x43217]
========= Host Frame:/opt/nvidia/hpc_sdk/Linux_x86_64/21.7/compilers/lib/libcudafor_114.so (cudadevicesynchronize_ + 0x11) [0x7aca1]
========= Host Frame:./a.out [0x14fb]
========= Host Frame:./a.out [0x1393]
=========
========= ERROR SUMMARY: 2 errors```
解决方案
使用“托管”属性时启用的 CUDA 统一内存目前仅可用于动态分配的内存。虽然我们希望将来支持静态数据,但您将无法在像“lnode”这样的静态变量上使用“托管”。
最简单(也可能是最好)的解决方案是按值传递 lnode。尽管您也可以使“lnode”可分配,或使主机lnode 和“设备”lnode。
% cat sol1.cuf
module a1
implicit none
integer(kind=4)::lnode
end module a1
module kernelsubroutine
contains
attributes(global) subroutine kernel(lnode)
implicit none
integer(kind=4),value::lnode
print*,"on device",lnode
end subroutine kernel
end module kernelsubroutine
program main
use a1
use cudafor
use kernelsubroutine
implicit none
integer(kind=4)::n
lnode = 1
print*,"on host",lnode
call kernel<<<1,1>>>(lnode)
n = cudaDeviceSynchronize()
end program main
% nvfortran sol1.cuf; a.out
on host 1
on device 1
% cat sol2.cuf
module a1
implicit none
integer(kind=4),allocatable,managed::lnode
end module a1
module kernelsubroutine
contains
attributes(global) subroutine kernel(lnode)
implicit none
integer(kind=4)::lnode
print*,"on device",lnode
end subroutine kernel
end module kernelsubroutine
program main
use a1
use cudafor
use kernelsubroutine
implicit none
integer(kind=4)::n
allocate(lnode)
lnode = 1
print*,"on host",lnode
call kernel<<<1,1>>>(lnode)
n = cudaDeviceSynchronize()
deallocate(lnode)
end program main
% nvfortran sol2.cuf ; a.out
on host 1
on device 1
% cat sol3.cuf
module a1
implicit none
integer(kind=4)::lnode_h
integer(kind=4),device::lnode_d
end module a1
module kernelsubroutine
contains
attributes(global) subroutine kernel(lnode)
implicit none
integer(kind=4)::lnode
print*,"on device",lnode
end subroutine kernel
end module kernelsubroutine
program main
use a1
use cudafor
use kernelsubroutine
implicit none
integer(kind=4)::n
lnode_h = 1
print*,"on host",lnode_h
lnode_d=lnode_h
call kernel<<<1,1>>>(lnode_d)
n = cudaDeviceSynchronize()
end program main
% nvfortran sol3.cuf ; a.out
on host 1
on device 1
推荐阅读
- android - 可访问性 - 将焦点设置到列表视图后向左或向右滑动选择视图层次结构的顶部
- ios - watchOS vs iOS 和 Delegate
- java - EntityStoreException:不支持的属性类型 id 9
- sas - SAS如何通过xy等格式输入整数
- c++ - 如何为这个函数给出有意义的Ω和Big-O界限?
- firebase - ERROR_INVALID_EMAIL,电子邮件地址格式错误。,null
- android - 房间中的多对多关系 CRUD
- python - 事件正在无缘无故地删除机器人嵌入
- html - 滑动时的布局问题(增加空间)
- excel - 如果找不到图像,则显示默认图像