PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

CUDA Fortran : device variable in module
Goto page 1, 2  Next
 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
RTLEE



Joined: 01 Mar 2006
Posts: 19

PostPosted: Wed Nov 18, 2009 4:56 pm    Post subject: CUDA Fortran : device variable in module Reply with quote

In CUDA Fortran, can device subroutines directly access device data in the same module? I thought the answer was yes, but I get a strange compiler warning (and the code crashes) when I try to do it. For example, this code compiles with the warning
Quote:

warning: cast to pointer from integer of different size
Warning: Cannot tell what pointer points to, assuming global memory space


Code:

module cudamod
   implicit none

   integer, device, allocatable, dimension(:)  :: int_d

   contains
   attributes(global) subroutine foo
      int_d(threadidx%x) = threadidx%x
   end subroutine foo
end module cudamod

program fcuda
   use cudafor
   use cudamod
   implicit none

   integer :: int_h(16)

   int_h = 0

   allocate(int_d(16))
   call foo<<<1,16>>>
   
   int_h = int_d
   
   print *,'int_h = ',int_h
   
   deallocate(int_d)
end program fcuda


When I run it, I get
Quote:

copyout Memcpy FAILED:4


Any help is appreciated!
Back to top
View user's profile
mkcolg



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

PostPosted: Fri Nov 20, 2009 11:41 am    Post subject: Reply with quote

Hi RTLEE,

Unfortunately, support for device data in a module didn't make it into 10.0. It's coming, but most likely not until early 2010. In the mean time, you'll need to pass the variable into the kernel.

Thanks,
Mat
Back to top
View user's profile
tty103



Joined: 19 Oct 2009
Posts: 8

PostPosted: Tue Jan 19, 2010 1:23 pm    Post subject: Reply with quote

Is this included in Jan 7 release?
Back to top
View user's profile
mkcolg



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

PostPosted: Tue Jan 19, 2010 3:05 pm    Post subject: Reply with quote

Sorry, no. I'll update this post once it's available.

- Mat
Back to top
View user's profile
goblinsqueen



Joined: 04 Feb 2010
Posts: 14

PostPosted: Thu Feb 18, 2010 8:05 am    Post subject: Reply with quote

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!
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 1, 2  Next
Page 1 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