PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

Cannot allocate variables from used modules

 
Post new topic   Reply to topic    PGI User Forum Forum Index -> Accelerator Programming
View previous topic :: View next topic  
Author Message
cknaus



Joined: 18 Feb 2010
Posts: 4

PostPosted: Tue Nov 09, 2010 8:37 am    Post subject: Cannot allocate variables from used modules Reply with quote

Hi,

Here is a test case where I allocate a device variable which is declared in a separate module.

--8<-- compile with pgf90 -Mcuda mod1.cuf kernel.cuf main.f90

PROGRAM main
USE test_mod
CALL test
END PROGRAM

--

MODULE mod1
IMPLICIT NONE
INTEGER, DEVICE, ALLOCATABLE :: foo
CONTAINS
END MODULE

--

module test_mod
use cudafor
use mod1 ! this fails
! INTEGER, DEVICE, ALLOCATABLE :: foo ! this works
contains

attributes(global) subroutine test_kernel( )
foo = 42
end subroutine

subroutine test
REAL*8 temp
integer r
type(dim3) :: dimGrid, dimBlock

allocate(foo)

dimGrid = dim3(1, 1, 1)
dimBlock = dim3(1, 1, 1)
call test_kernel<<<dimGrid,dimBlock>>>()

r = cudathreadsynchronize()
write(0,*) "Value of cudathreadsynchronize = ", r

temp = foo
write(0,*) temp

end subroutine
end module

--8<--

When I run the program, it produces the following output:

--
Value of cudathreadsynchronize = 4
copyout Memcpy (host=0x7ffffa7af254, dev=0x110000, size=4) FAILED:4
--

However, if I declare the device variable inside the same module where it is allocated, then the test succeeds. My system configuration is: Ubuntu 10.04 (64-bit), GeForce GTX 260, original NVIDIA driver packaged by Ubuntu, PGI Accelerator Fortran Workstation 10.5.

Is this a compiler bug? If so, is there workaround while still allocating variables declared in separate modules?

Regards,

Claude Knaus
Back to top
View user's profile
mkcolg



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

PostPosted: Tue Nov 09, 2010 2:16 pm    Post subject: Reply with quote

Hi Claude,

In CUDA Fortran, device module data is only accessible by device routines within the same module or from host code that uses the module. Accessing device data declared from other modules is not allowed. (See the "Variable Qualifier" section of the CUDA Fortran Reference Guide)The problem being that there isn't a linker for device code, hence no way to associate external device symbols.

However, this has been our most requested feature and our engineers have been hard at work trying to find a way to support this. With Fermi and CUDA 3.0, we have found a way to perform this association at run time. The feature is still in development but will be available some time early next year. Full details can be found at http://www.pgroup.com/lit/articles/insider/v2n3a1.htm

- Mat
Back to top
View user's profile
cknaus



Joined: 18 Feb 2010
Posts: 4

PostPosted: Tue Nov 09, 2010 4:21 pm    Post subject: Reply with quote

Hi Mat,

Thanks for the prompt and clarifying answer!

Cheers,
-- Claude
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
Page 1 of 1

 
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