Global device variables CUDA Fortran

Accelerated Computing HPC Compilers Legacy PGI Compilers
1

I will like to know if its possible to define some device global variables in a host module, which can later be used by device codes i.e global and device subroutines. I will like to know if this is possible normally , and what this compiler errors really mean, its a test code consisting of two modules and the main program as follows

The precision module:

    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
      ! This is to enable hybrid algorithms
      integer , parameter :: hb_kind = doublePrecision
    end module precision_m

The read control file

    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

The Kernels Module:

    module test_global

	   use test_host_gen

	   contains

	   attributes(global) subroutine test_kernel(k_M_he)

	      real(fp_kind) , device :: k_M_he
	
	      k_M_he = D_M_he

	   end subroutine test_kernel

    end module test_global

The main program:

    program test

       use test_global

       implicit none

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

	   ctrl_file = "input.txt"

	   call read_control_file(ctrl_file)

	   D_M_he = M_he

	   call test_kernel<<<1,1>>>(k_M_he)

	   tempo = M_he
	   ! Copy GPU Value
	   M_he = k_M_he
	   if (tempo.eq.M_he) print *, 'It works !!!!'

    end program test

compiler errors:

pgf90 -g -o testgpu precision_m.o test_host_gen.o test_global.o test.o
 
test_host_gen.o: In function `..cuda_fortran_constructor_1':
./test_host_gen.CUF:42: undefined reference to `__pgi_cuda_register_fat_binary'
./test_host_gen.CUF:42: undefined reference to `__cudaRegisterVar'
test_global.o: In function `test_global_test_kernel_':
./test_global.CUF:11: undefined reference to `cudaSetupArgument'
./test_global.CUF:11: undefined reference to `cudaLaunch'
test_global.o: In function `..cuda_fortran_constructor_1':
./test_global.CUF:13: undefined reference to `__pgi_cuda_register_fat_binary'
./test_global.CUF:13: undefined reference to `__cudaRegisterFunction'
test.o: In function `test':
./test.CUF:12: undefined reference to `pgf90_dev_alloc04'
./test.CUF:12: undefined reference to `pgf90_dev_symhandle_copyin'
./test.CUF:12: undefined reference to `pgf90_dev_configure_call'
./test.CUF:12: undefined reference to `pgf90_dev_copyout'
./test.CUF:23: undefined reference to `pgf90_dev_dealloc03'
make: *** [test] Error 2

Thanks.

2

Hi egodfred,

The problem here is that you need to add the flag “-Mcuda” to your link options so the compiler driver knows to add the CUDA libraries to your link.

I will like to know if its possible to define some device global variables in a host module, which can later be used by device codes i.e global and device subroutines.

Yes, this possible.

Hope this helps,
Mat

3

Thanks