Problem with using managed module variables from other files in Cuda Fortran

user_15925363 0
Viewed 203

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```
cuda fortran
1 Answers

CUDA Unified Memory, enabled when using the "managed" attribute, is currently only available for dynamically allocated memory. While we hope to support static data in the future, you wont be able to use 'managed' on a static variable like 'lnode'.

The simplest (and probably best) solution is to pass lnode in by value. Though you could also make 'lnode' allocatable, or make a host lnode and 'device' 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
user_32044840
Related