PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

error with derived types in PGI CUDA 10.4
Goto page Previous  1, 2
 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
mkcolg



Joined: 30 Jun 2004
Posts: 4996
Location: The Portland Group Inc.
PostPosted: Tue May 04, 2010 1:02 pm    Post subject: Reply with quote
Hi goblinqueen,

I just verified that TPR#16836 will be fixed in 10.5.

- Mat
Back to top
View user's profile
TroelsH



Joined: 24 Mar 2010
Posts: 9
PostPosted: Mon May 10, 2010 9:33 am    Post subject: Reply with quote
mkcolg wrote:
Hi goblinqueen,

I just verified that TPR#16836 will be fixed in 10.5.

- Mat


It is half-fixed. I can use derived type constant variables, which is very nice because I have packed all my small constants in one derived type, to save on memory transfer, but derived type device arrays do not work yet. E.g.

Code:
MODULE cuda_arrays
  USE cudafor
  implicit none
  type :: float2
    real x, y
  end type
  type(float2), dimension(:),  allocatable, device :: gf2
CONTAINS
!==============================
attributes(global) SUBROUTINE kernel(gf2)
  implicit none
  type(float2), dimension(100) :: gf2
  if (threadidx%x==1) gf2(blockidx%x)%x = 1.
END SUBROUTINE kernel
END MODULE
!==============================
PROGRAM test
  USE cuda_arrays
  implicit none
  type(dim3)      :: nt, nb
  allocate(gf2(100))
  nt = dim3(32, 1, 1)
  nb = dim3(100, 1, 1)
  call kernel<<<nb,nt>>>(gf2)
END


gives internal compiler error with v10.5 and compiles with v10.3

Troels
Back to top
View user's profile
TroelsH



Joined: 24 Mar 2010
Posts: 9
PostPosted: Tue May 11, 2010 1:42 am    Post subject: Reply with quote
I have just found a work-around. It works if the device array is not declared in the module, but passed as argument. The following compiles and runs with both v10.3, v10.4, and v10.5 :

Code:
MODULE cuda_arrays
  USE cudafor
  implicit none
CONTAINS
!==============================
attributes(global) SUBROUTINE kernel(gf2)
  implicit none
  type :: float2
    real x, y
  end type
  type(float2), dimension(100) :: gf2
  if (threadidx%x==1) gf2(blockidx%x)%x = 1.
END SUBROUTINE kernel
END MODULE
!==============================
PROGRAM test
  USE cuda_arrays
  implicit none
  type :: float2
    real x, y
  end type
  type(float2), dimension(:),  allocatable, device :: gf2
  type(dim3)      :: nt, nb
  allocate(gf2(100))
  nt = dim3(32, 1, 1)
  nb = dim3(100, 1, 1)
  call kernel<<<nb,nt>>>(gf2)
END


Troels

TroelsH wrote:


It is half-fixed. I can use derived type constant variables, which is very nice because I have packed all my small constants in one derived type, to save on memory transfer, but derived type device arrays do not work yet. E.g.

Code:

...


gives internal compiler error with v10.5 and compiles with v10.3

Troels
Back to top
View user's profile
goblinsqueen



Joined: 04 Feb 2010
Posts: 14
PostPosted: Tue May 11, 2010 6:49 am    Post subject: Reply with quote
Thank you Troels,
your comments are very useful!

I'll try to fix my code with your solution, but it will be not straightforward (many of my data are declared as derived type).

I hope it will be completely fixed in 10.6.
Back to top
View user's profile
Display posts from previous:   
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming All times are GMT - 7 Hours
Goto page Previous  1, 2
Page 2 of 2

 
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 © 2001, 2002 phpBB Group