Device Derived Types

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?

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

  • Mat

Here is the code I am trying to test this on:
MODULE GPUModule
use cudafor
implicit none
type globalvar
real8:: a
real
8:: x(10)
real8:: 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
real8 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?

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.

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

Hope this helps,
Mat

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?