PGI User Forum
 SearchSearch   MemberlistMemberlist     RegisterRegister   ProfileProfile    Log inLog in 

Free OpenACC Webinar

Problems with the allocate statements in device subroutines

 
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: Wed Sep 25, 2013 7:14 am    Post subject: Problems with the allocate statements in device subroutines Reply with quote

Hi,

I have been troubling by the allocate statements in the device subroutines for quite a few days, and I can't figure it out, however.

I give the compile environment firstly, I guess the compile option is the key to this problem.

Compile Environment:
Quote:
PGI Visual Fortran 13.8 x64, Windows 7 x64, Visual Studio 2010

Fortran | Command Line
-gopt -Bstatic -Mbackslash -Mcuda=nofma,cuda5.0,cc35,cc3x,rdc -I"c:\program files\pgi\win64\13.8\include" -I"C:\Program Files\PGI\Microsoft Open Tools 11\include" -I"C:\Program Files (x86)\Windows Kits\8.0\Include\shared" -I"C:\Program Files (x86)\Windows Kits\8.0\Include\um" -Minform=inform

Linker | Command Line
-gopt -Bstatic -Mcuda=nofma,cuda5.0,cc35,cc3x,rdc -o "D:\Research\Programming\Routine\CUDA Fortran\test\x64\Debug\test.exe"


I wrote a test routine as follows:

Code:
module a_test

contains
 
  attributes(global) subroutine kernel(a,n)

    implicit none
    integer,value   :: n
    integer,device  :: a(n)
    integer :: i
    integer,device,allocatable :: b(:)
       
    if (threadidx%x==1) then
       allocate(b(n))
    end if
   
   call  syncthreads()
         
    i=threadIdx%x
    b(i)=i
    a(i)=b(i)
   return
  end subroutine kernel

end module

program prog

   use a_test
   use cudafor
   implicit none

   integer,parameter  :: n=128
   integer,device     :: a_d(n)
   integer                   :: a(n)

   call kernel<<<1,128>>>(a_d,n)
   a=a_d
 
   write(*,*)a(10:20)

end program


When I ran it, errors occured at link procedure.

error message:

Code:
Compiling Project  ...
Linking...
pgnvd-Error-Required tool nvlink was not found
pgnvd... looked for nvlink at c:\program files\pgi\win64/2013/cuda/5.0/bin\nvlink
child process exit with signal 1: c:\program files\pgi\win64\13.8\bin\pgnvd.exe
test build failed.


I searched the nvlink file on the whole disk and copyed it to the file folder, I ran it another time, and errors occured again.

error message:
Code:
Linking...
C:\Users\Administrator\AppData\Local\Temp\pgnvd3bty4b5LcTr5vB.s: Assembler messages:
C:\Users\Administrator\AppData\Local\Temp\pgnvd3bty4b5LcTr5vB.s:4: Warning: value 0xffed truncated to 0xed
C:\Users\Administrator\AppData\Local\Temp\pgnvd3bty4b5LcTr5vB.s:4: Warning: value 0xffba truncated to 0xba
C:\Users\Administrator\AppData\Local\Temp\pgnvd3bty4b5LcTr5vB.s:5: Warning: value 0xff80 truncated to 0x80
C:\Users\Administrator\AppData\Local\Temp\pgnvd3bty4b5LcTr5vB.s:6: Warning: value 0xffb3 truncated to 0xb3
C:\Users\Administrator\AppData\Local\Temp\pgnvd3bty4b5LcTr5vB.s:13: Warning: value 0xff80 truncated to 0x80
C:\Users\Administrator\AppData\Local\Temp\pgnvd3bty4b5LcTr5vB.s:15: Warning: value 0xffb7 truncated to 0xb7
test build failed.


Have I set the wrong compiler options and how can I do with the nvlink file?
Back to top
View user's profile
mkcolg



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

PostPosted: Fri Sep 27, 2013 9:50 am    Post subject: Reply with quote

Hi OceanCloud,

Apologies for the late response. I was getting input from engineering an wanting to test the released 13.9 (using CUDA 5.5) before responding. Unfortunately, I'm still seeing various errors when building with either CUDA 5.0 or 5.5. Hence, I added a problem report (TPR#19604) and sent it on to engineering. There is no work around at this point.

On a side note, the code isn't correct in that you have a local copy of "b" but only allocate it for one thread. I'm assuming you're either just testing allocation, or that you really want "b" to be a shared array with a length of "n". If it's the latter, then the better way to do this is to use a shared automatic array and then pass in the size of the array as the third argument in the kernel launch.

For example:
Code:
module a_test

 contains

   attributes(global) subroutine kernel(a,n)

     implicit none
     integer,value   :: n
     integer,device  :: a(n)
     integer :: i
     integer,shared :: b(n)


    call  syncthreads()

     i=threadIdx%x
     b(i)=i
     a(i)=b(i)
    return
   end subroutine kernel

 end module

 program prog

    use a_test
    use cudafor
    implicit none

    integer,parameter  :: n=128
    integer,device     :: a_d(n)
    integer                   :: a(n)

    call kernel<<<1,128,n*4>>>(a_d,n)
    a=a_d

    write(*,*)a(10:20)

 end program


- Mat
Back to top
View user's profile
OceanCloud



Joined: 27 Nov 2012
Posts: 12

PostPosted: Sat Sep 28, 2013 7:27 pm    Post subject: Reply with quote

Many thanks, Mat

I just want to test the allocate statement.

The reason for the allocation problem is the compile itself, and I'd better do it under the Linux system, isn't it?
Back to top
View user's profile
mkcolg



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

PostPosted: Mon Sep 30, 2013 10:02 am    Post subject: Reply with quote

We have better support for this on Linux (the nvlink error goes away), but I'm still seeing the same errors that I see on Windows. Why this particular small example fails while others do not, I'm not sure. I'll need to wait for engineering to investigate TPR#19604.

- Mat
Back to top
View user's profile
jtull



Joined: 30 Jun 2004
Posts: 445

PostPosted: Thu Jul 24, 2014 5:16 pm    Post subject: TPR 19604 - UF: CUDA Fortran device allocation - fixed Reply with quote

TPR 19604 - UF: CUDA Fortran device allocation fails to compile on Windows

should now work in the 14.6 and current 14.7 release.

regards,
dave
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