I have a managed variable in file p2.f95
in module a1
. I’m passing this managed variable to kernel in file p6.f95
.I am able to access the variable in file p6.f95
but unable to access the same from the kernel i.e; I am not able to print variable value from the device. Can someone please explain why…?
file1: p2.f95
module a1
implicit none
integer(kind=4),managed::lnode
end module a1
file2: 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
compilation command 1: $pgf95 -Mcuda=rdc p2.f95 p6.f95
compilation command 2:
$pgf95 -Mcuda=rdc -c p2.f95
$pgf95 -Mcuda=rdc p6.f95 p2.o
p6.f95:
with both sets of compilation commands the output remains same
output command: $ cuda-memcheck ./a.out
output:
========= 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