PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

CUDA Fortran : device variable in module
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
Tuan



Joined: 11 Jun 2009
Posts: 233

PostPosted: Thu Feb 18, 2010 12:53 pm    Post subject: Reply with quote

goblinsqueen wrote:
Hi,
I receive the same warning than RTLEE when I compile the following code

Code:

module cudamod

  use cudafor

  real(kind=8), dimension(:), allocatable, device :: vec_dev

contains

  attributes(device) subroutine sub1(i)
   
    use cudafor
   
    implicit none
   
    integer, value :: i
    real(kind=8), dimension(32) :: vec0
    integer :: tid
   
    tid = threadidx%x
    vec0(tid) = vec_dev(i)
   
  end subroutine

  attributes(global) subroutine kernel_test(N,nblocks,nthreads)
   
    use cudafor
   
    implicit none
   
    integer, value :: N, nthreads, nblocks
    integer :: i, idx

    idx = (blockidx%x-1)*blockdim%x + threadidx%x
    do i=idx,N,nthreads*nblocks
        call sub1(i)
    end do

  end subroutine

end module

program test

  use cudafor
  use cudamod

  integer, parameter :: N = 1000
  real(kind=8), dimension(N) :: vector 
  integer :: nblocks, nthreads

  vector = 1.D0
  allocate(vec_dev(N))
  vec_dev = vector(1:N)
  nblocks = 30
  nthreads = 256
  call kernel_test<<<nblocks,nthreads>>>(N,nblocks,nthreads)
  deallocate(vec_dev)

end program


From the PGI CUDA Fortran documentation and from Mat's answer, I was assuming that it is not allowed for the kernel to access device data declared in a module, but that it is possible for any device subprogram in that module and by any host program that uses the module.

Am I wrong?

Thank you in advance for any explanation!


I experience the same issue. Don't access the device data from module.

Tuan.
Back to top
View user's profile
mkcolg



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

PostPosted: Thu Feb 18, 2010 9:46 pm    Post subject: Reply with quote

Hi Goblinqueen,

Currently only fixed size device arrays are allowed in a Module. However, we are actively working on adding this (it's our most requested feature). It should be available in 10.6 or possibly sooner.

- Mat
Back to top
View user's profile
goblinsqueen



Joined: 04 Feb 2010
Posts: 14

PostPosted: Fri Feb 19, 2010 2:03 am    Post subject: Reply with quote

Thank you, it would be very important!
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 © phpBB Group