nvvm parse error

This is an error I get which I don’t fully understand;

This are the code files (the ones where the error is emerging from)

https://dl.dropboxusercontent.com/u/59996494/host_subs_m.CUF
https://dl.dropboxusercontent.com/u/59996494/precision_m.F90

make gpu
pgf90 -g -Mcuda -Minfo -c precision_m.F90
pgf90 -g -Mcuda -Minfo -c cpurandom_m.F90
pgf90 -g -Mcuda -Minfo -c host_gen_m.F90
pgf90 -g -Mcuda -Minfo -c host_subs_m.CUF

nvvmCompileProgram error: 9.
Error: /tmp/pgcudafortgNe52bKRKvA.gpu (42, 0): parse error: global variable reference must have pointer type
PGF90-F-0155-Compiler failed to translate accelerator region (see -Minfo messages): Device compiler exited with error status code (host_subs_m.CUF: 1)
PGF90/x86-64 Linux 14.10-0: compilation aborted
make: *** [host_subs_m.o] Error 2

I have the same error with this small piece if test code:

make test
pgf90 -g -Mcuda -Minfo	 -c precision_m.F90
pgf90 -g -Mcuda -Minfo	 -c test_host_gen.CUF
pgf90 -g -Mcuda -Minfo	 -c test_global.CUF
nvvmCompileProgram error: 9.
Error: /tmp/pgcudaforf6KgpUvKTkXq.gpu (19, 0): parse error: global variable reference must have pointer type
PGF90-F-0155-Compiler failed to translate accelerator region (see -Minfo messages): Device compiler exited with error status code (test_global.CUF: 1)
PGF90/x86-64 Linux 14.10-0: compilation aborted
make: *** [test_global.o] Error 2



module precision_m
      integer , parameter :: singlePrecision = kind (0.0)
      integer , parameter :: doublePrecision = kind (0.0d0)
      ! Comment out one of the lines below
      	integer , parameter :: fp_kind = singlePrecision
      !integer , parameter :: fp_kind = doublePrecision
      integer , parameter :: hb_kind = doublePrecision
      end module precision_m



module test_host_gen

	use precision_m

	implicit none

	real(fp_kind) , device :: D_M_he
	real(fp_kind) :: M_he


	contains

	subroutine read_control_file(ctrl_file)

			character(len=100) , intent(in) :: ctrl_file
			character(len=100) :: buffer, label
			integer :: pos , swch
  			integer, parameter :: fh = 15
  			integer :: ios = 0
  			integer :: line = 0

  			open(fh, file=ctrl_file)

  			do while (ios == 0)
     			read(fh, '(A)', iostat=ios) buffer
		     if (ios == 0) then
		        line = line + 1

		        ! Find the first instance of whitespace.  Split label and data.
		        pos = scan(buffer, '    ')
		        label = buffer(1:pos)
		        buffer = buffer(pos+1:)

		        select case (label)
		        case ('M_he')
		           read(buffer, *, iostat=ios) M_he
		           print *, 'Read Mass of Helium M_he : ', M_he			       			      				       			       			      				       		
		        case default
		           print *, 'Skipping invalid label at line', line
		        end select
		     end if
		  end do
		  
	end subroutine read_control_file

end module test_host_gen



module test_global

	use test_host_gen

	contains

	attributes(global) subroutine rnd12(rnd0)

		integer , device ::	is,js,ks,ns
		real(fp_kind) , intent(out):: rnd0
		save

		mzran = is - ks
		if (mzran .lt. 0) then
		    mzran = mzran +  2147483579
		endif

		is = js
		js = ks
		ks = mzran
		ns = 69069*ns + 1013904243
		mzran = mzran + ns
	    
	    rnd0 = 0.5d0 + mzran/2.d0**32 

	end subroutine rnd12

end module test_global



program test

use test_global

implicit none

	real(fp_kind) , device :: k_M_he , d_random
	character(len=100) :: ctrl_file
	real(fp_kind) :: tempo , random
	integer :: tmp_int

	ctrl_file = "input.txt"
	rnd_file = "random.txt"

	open(unit=1, file=name, iostat=ios, status="new", action="write")
	if ( ios /= 0 ) stop "Error opening file name"


	! test Random numbers

	do i = 1 : 1000000

		call rnd12<<<1,1>>>(d_random)

		random = d_random

		write(1,*) random
	enddo

end program test

Compiler:

pgf90 14.10-0 64-bit target on x86-64 Linux -tp istanbul
The Portland Group - PGI Compilers and Tools
Copyright (c) 2014, NVIDIA CORPORATION. All rights reserved.

Hi egodfred,

The problem here is that the compiler isn’t catching that you’re trying to use a “SAVE” statement in the device routines. I’ve added TPR#21520 requesting we emit a more informative error message when encountering a SAVE statement in device code.

A SAVE statement causes local variables to get global storage so that the values persistent from call to call. This doesn’t make sense in a parallel context since all threads would then be accessing the same data. Hence, it’s not supported in CUDA Fortran device code.

My suggestion here would be create a module device array with one element per thread to store these variables.

Thanks!
Mat

TPR 21520 - CUDA Fortran: Need to issue error when “save” is used in a device routine.

Has been fixed in the current 15.5 release.

thanks,
dave