How to pass a device allocated array into a function properly in Fortran?

[novus]

Hi I am new to Cuda but I am trying to allocate an array in my main program, and pass it to another function that
combines a few related Cuda/fft calls but it seems like the device array is being read as not associated when passed into the function, I understand the function itself is on the host but I am unsure how else to do this here is an example of what I am talking about:

subroutine foo(example_device_array)
    
    real,Dimension(:,:,:),intent(inout),device::example_device_array


    !! Error happens in here upon trying to read in variable
    !! Happens with or withour allocatable keyword inside here
    
end subroutine foo

program main
    use example_kernels
    implicit none
    integer, parameter:: nx =20 , ny = 20, nz = 20


    real, Dimension(:,:,:), allocatable, device :: example_device_array
    allocate(example_device_array(64,64,64))

    call fill_out_device_array<<<256,256>>>(example_device_array)
    call cudaDeviceSynchronize()

    call foo(example_device_array)

end program main

I hope i havent oversimplified, I realize theres a chance it gets messed up in the previous kernel call but its fairly complex so I first wanted to check whether it was an issue with how im passing the variable.

Really appreciate any insights.

[MatColgrove]

Hi novus and welcome!

CUDA Fortran requires F90 interfaces in order to pass device arrays. Does it work if you put “foo” in a module so an implicit interface is created?

-Mat

[novus]

Hi Matt!

Sorry it took me a while to get to it, that worked! Thanks for the advice.