PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

Device Derived Types

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



Joined: 19 Feb 2013
Posts: 6

PostPosted: Mon Feb 25, 2013 11:20 am    Post subject: Device Derived Types Reply with quote

Is there a way to define a derived type as device variables so the entire type can be passed to the GPU kernel in the GPU subroutine call in the host program?
Back to top
View user's profile
mkcolg



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

PostPosted: Mon Feb 25, 2013 12:36 pm    Post subject: Reply with quote

Provided that the UDT is fixed size (i.e. no allocatable) then, yes, you can uses these with both CUDA Fortran and OpenACC.

- Mat
Back to top
View user's profile
msgross42



Joined: 19 Feb 2013
Posts: 6

PostPosted: Mon Feb 25, 2013 12:56 pm    Post subject: Reply with quote

Here is the code I am trying to test this on:
MODULE GPUModule
use cudafor
implicit none
type globalvar
real*8:: a
real*8:: x(10)
real*8:: y(10)
real*8:: z(10)
end type globalvar

CONTAINS
attributes(global) subroutine sumv(gvar)
type(globalvar)::gvar
integer i
i = (blockidx%x-1) * blockdim%x + threadidx%x
gvar%z(i)=gvar%a*gvar%x(i)+gvar%y(i)
end subroutine sumv
end MODULE GPUMODULE

PROGRAM testgpu
use GPUModule
use cudafor
implicit none
type(globalvar) gvar
integer,parameter :: n=10
integer i,istat
real*8 zn
gvar%a=5
do i=1,n
gvar%x(i)=i
gvar%y(i)=i
enddo
call sumv<<<1,n>>>(gvar)
zn=gvar%z(2)
write(*,*) zn
end PROGRAM

As written, the compiler gives me an error of device attribute mismatch when calling the sumv routine. What do I need to do to define my type such that it is passed to the kernel as a device variable?
Back to top
View user's profile
mkcolg



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

PostPosted: Mon Feb 25, 2013 2:26 pm    Post subject: Reply with quote

Quote:
As written, the compiler gives me an error of device attribute mismatch when calling the sumv routine. What do I need to do to define my type such that it is passed to the kernel as a device variable?


gvar needs to be declared with the device attribute otherwise it's a host variable.

Code:
PROGRAM testgpu
use GPUModule
use cudafor
implicit none
type(globalvar), device :: gvar
integer,parameter :: n=10


Hope this helps,
Mat
Back to top
View user's profile
msgross42



Joined: 19 Feb 2013
Posts: 6

PostPosted: Wed Feb 27, 2013 10:52 am    Post subject: Reply with quote

Is there a way to define a type or a type within another type as constant or shared or can they only be defined as device variables?
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