PGF90-S-0528: Device attribute mismatch

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.

  • Mat