PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

CUDA-x86.

Pointer Arrays in CUDA Fortran

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



Joined: 27 Nov 2012
Posts: 12

PostPosted: Thu Dec 12, 2013 6:04 am    Post subject: Pointer Arrays in CUDA Fortran Reply with quote

Hi, all

I need to deal with lots of data which consists of column vectors with different length. As a result, it's not appropriate to use the two-dimension array. The pointer arrays seems can solve the problem perfectly. Although, it's not possible to declare an array of pointers in Fortran, we can create an pointer array by using derived data types as the sample routine does.

Code:
PROGRAM test_array_pointer

  IMPLICIT NONE

  TYPE :: ptr
    REAL, POINTER :: p(:)
  END TYPE

  TYPE(ptr), ALLOCATABLE :: x(:)

  REAL, TARGET :: a(1)
  REAL, TARGET :: b(2)
  REAL, TARGET :: c(3)


  a=(/1.0/)
  b=(/1.0,2.0/)
  c=(/1.0,2.0,3.0/)

  ALLOCATE(x(3))

  x(1)%p => a
  x(2)%p => b
  x(3)%p => c

  WRITE(*,*)  x(1)%p
  WRITE(*,*)  x(2)%p
  WRITE(*,*)  x(3)%p

  DEALLOCATE(x)

  STOP

END PROGRAM test_array_pointer


The program output:
Code:
    1.000000   
    1.000000        2.000000   
    1.000000        2.000000        3.000000   
FORTRAN STOP


It seems that the roundabout method works fine. Then I want to use it in CUDA Fortran, I write a test routine as follows.

Code:
PROGRAM test

  USE cudafor

  IMPLICIT NONE

  TYPE :: ptr
    REAL, POINTER :: p(:)
  END TYPE

  TYPE(ptr), DEVICE, ALLOCATABLE :: dev_x(:)

  REAL, DEVICE, TARGET :: a(1)
  REAL, DEVICE, TARGET :: b(2)
  REAL, DEVICE, TARGET :: c(3)

  REAL :: x(1)

  INTEGER :: istat

  a=(/1.0/)
  b=(/1.0,2.0/)
  c=(/1.0,2.0,3.0/)

  ALLOCATE(dev_x(3),STAT=istat)

  WRITE(*,*) istat

  dev_x(1)%p => a
  dev_x(2)%p => b
  dev_x(3)%p => c

  x = dev_x(1)%p

  WRITE(*,*) x

  DEALLOCATE(dev_x)

  STOP

END PROGRAM test


The program output:
Code:
           0
Segmentation fault (core dumped)


Errors occur on the memory operations, unfortunately.

BTW, the NVIDIA card is K20C and the compile enviroment shows below.
    PGI Fortran 13.10 X64, Cent OS 6.4 X64
    -Mcuda=cuda5.5,cc35

Can someone tell me where it goes wrong?
Or is there any other ways to deal with the column vectors with different length?
Back to top
View user's profile
mkcolg



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

PostPosted: Mon Dec 16, 2013 4:33 pm    Post subject: Reply with quote

Hi OceanCloud,

The basic problem here is with "dev_x(1)%p". You're trying to access a device data structure on the host. While the base device address of "dev_x" is stored on the host, any internal addresses are not. Hence, to make this work, you need to move these pointer assignments over to the device.

Something like the following:

Code:
% cat testp.cuf
module foo
  TYPE :: ptr
    REAL, POINTER :: p(:)
  END TYPE

  TYPE(ptr), DEVICE, ALLOCATABLE    :: dev_x(:)
  REAL, DEVICE                      :: xD(3)
  REAL, DEVICE, allocatable, TARGET :: a(:)
  REAL, DEVICE, allocatable, TARGET :: b(:)
  REAL, DEVICE, allocatable, TARGET :: c(:)

contains

  attributes(global) subroutine assignPtr()
     dev_x(1)%p => a
     dev_x(2)%p => b
     dev_x(3)%p => c
  end subroutine assign

  attributes(global) subroutine getA()
     xD(1:1)= dev_x(1)%p
  end subroutine getA

  attributes(global) subroutine getB()
     xD(1:2) = dev_x(2)%p
  end subroutine getB

  attributes(global) subroutine getC()
     xD = dev_x(3)%p
  end subroutine getC


end module foo

PROGRAM test

  USE cudafor
  use foo
  IMPLICIT NONE


  REAL :: x(3)
  INTEGER :: istat

  allocate(a(1), b(2), c(3))
  a=(/1.0/)
  b=(/1.0,2.0/)
  c=(/1.0,2.0,3.0/)
  x=0.0

  ALLOCATE(dev_x(3),STAT=istat)

  WRITE(*,*) istat
  call assignPtr<<<1,1>>>()
  xD=0.0
  call getA<<<1,1>>>()
  x=xD
  WRITE(*,*) x
  xD=0.0
  call getB<<<1,1>>>()
  x=xD
  WRITE(*,*) x
  xD=0.0
  call getC<<<1,1>>>()
  x=xD
  WRITE(*,*) x

  DEALLOCATE(dev_x,a,b,c)

  STOP

END PROGRAM test

% pgf90 -Mcuda testp.cuf -V13.9 ; a.out
            0
    1.000000        0.000000        0.000000
    1.000000        2.000000        0.000000
    1.000000        2.000000        3.000000
Warning: ieee_inexact is signaling
FORTRAN STOP


- Mat
Back to top
View user's profile
OceanCloud



Joined: 27 Nov 2012
Posts: 12

PostPosted: Mon Dec 16, 2013 7:11 pm    Post subject: Reply with quote

I know exactly where the problem is

You are a great help, Mat.
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