What follows is a minimal working example:
!*********************************************************
module gpu_kernel
use cudafor
contains
! *****************************************************
attributes(global) subroutine Cuda_Kernel(iseed)
implicit none
! Input/output variables
integer, intent(inout) :: iseed
iseed = iseed - (blockIdx%x-1)*blockDim%x - threadIdx%x
return
end subroutine
end module gpu_kernel
!*********************************************************
program mc_delB
use gpu_kernel
implicit none
integer :: iseed, BlocksPerGrid, ThreadsPerBlock
BlocksPerGrid = 16; ThreadsPerBlock = 1; iseed = -2255
call Cuda_Kernel<<<BlocksPerGrid,ThreadsPerBlock>>>(iseed)
stop
end program
When I attempt to compile (pgfortran -Mcuda -o mc_delB.exe mc_delB.f90) I get the error message in the title of this thread, pointing me to the first argument of Cuda_Kernel: iseed. This happens also if iseed in the subroutine is explicitly given the “device” attribute.
What am I doing incorrectly here?
Hi dcwarren,
What am I doing incorrectly here?
Iseed needs to have the device attribute applied in the host code. It’s implied within a device kernel.
!*********************************************************
module gpu_kernel
use cudafor
contains
! *****************************************************
attributes(global) subroutine Cuda_Kernel(iseed)
implicit none
! Input/output variables
integer, intent(inout) :: iseed
iseed = iseed - (blockIdx%x-1)*blockDim%x - threadIdx%x
return
end subroutine
end module gpu_kernel
!*********************************************************
program mc_delB
use gpu_kernel
implicit none
integer, device :: iseed
integer :: BlocksPerGrid, ThreadsPerBlock
BlocksPerGrid = 16; ThreadsPerBlock = 1; iseed = -2255
call Cuda_Kernel<<<BlocksPerGrid,ThreadsPerBlock>>>(iseed)
stop
end program
Hope this helps,
Mat
Okay. Thanks for the explanation. For some reason I thought we were doing the C equivalent of passing by value when moving from the host to the device, but Fortran is always pass by reference so I don’t know why it would be different here.
For some reason I thought we were doing the C equivalent of passing by value when moving from the host to the device,
Fortran is pass by reference by default, however you can add the “value” attribute to have the variable passed by value. Though, you would also need to remove the “intent(inout)” since the variable can’t be passed back.
attributes(global) subroutine Cuda_Kernel(iseed)
implicit none
! Input/output variables
integer, value :: iseed
iseed = iseed - (blockIdx%x-1)*blockDim%x - threadIdx%x
return
end subroutine
but Fortran is always pass by reference so I don’t know why it would be different here.
Iseed is being passed by reference but a reference to it’s location in host memory. Hence, you need to declare it as a device variable so you’re passing in an address in device memory.