PGI User Forum

[msgross42]

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?

[mkcolg]

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

- Mat

[msgross42]

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?

[mkcolg]

[msgross42]

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?