PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

PGF90-S-0528: Device attribute mismatch

 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling
View previous topic :: View next topic  
Author Message
dcwarren



Joined: 18 Jun 2012
Posts: 29

PostPosted: Mon Oct 08, 2012 1:31 pm    Post subject: PGF90-S-0528: Device attribute mismatch Reply with quote

What follows is a minimal working example:

Code:
!*********************************************************
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?
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6129
Location: The Portland Group Inc.

PostPosted: Mon Oct 08, 2012 1:46 pm    Post subject: Reply with quote

Hi dcwarren,

Quote:
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.

Code:

!*********************************************************
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
Back to top
View user's profile
dcwarren



Joined: 18 Jun 2012
Posts: 29

PostPosted: Mon Oct 08, 2012 9:55 pm    Post subject: Reply with quote

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.
Back to top
View user's profile
mkcolg



Joined: 30 Jun 2004
Posts: 6129
Location: The Portland Group Inc.

PostPosted: Tue Oct 09, 2012 6:58 am    Post subject: Reply with quote

Quote:
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.

Code:
   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


Quote:
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
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Programming and Compiling All times are GMT - 7 Hours
Page 1 of 1

 
Jump to:  
You cannot post new topics in this forum
You cannot reply to topics in this forum
You cannot edit your posts in this forum
You cannot delete your posts in this forum
You cannot vote in polls in this forum


Powered by phpBB © phpBB Group