passing device allocatable array to kernel subprogram

Accelerated Computing HPC Compilers Legacy PGI Compilers
1

when i use static array like

integer, dimension(10,10), device :: arr2

instead of

integer, dimension(:,:), allocatable, device :: arr2

the code below works fine.

module test_mod                                                                        
 use cudafor
 real, dimension(10,10), device :: arr1
  integer, dimension(:,:),allocatable, device :: arr2
! integer, dimension(10,10), device :: arr2
contains
  attributes(global) subroutine doOnGPU()
     integer, value :: i, j

     i = threadIdx%x
      do j = 1, 10
        arr1(i,j) = i +j/0.5d0
        arr1(i,j) =  i + j + arr1(arr2(i,i), j)
      enddo

   end subroutine doOnGPU

   subroutine doOnCPU()
       allocate(arr2(10,10))
       arr2 = 1;

       call doOnGPU<<<1, 10>>>()

   end subroutine doOnCPU

end module test_mod

program  test1
   use test_mod

   call doOnCPU()
end program test1

However, when I use allocatable, I get this warning

/tmp/pgcudaforVr5cnWS3-zwQ.gpu: In function `doongpu’:
/tmp/pgcudaforVr5cnWS3-zwQ.gpu:17: warning: cast to pointer from integer of different size
/tmp/pgcudaforVr5cnWS3-zwQ.gpu(17): Warning: Cannot tell what pointer points to, assuming global memory space

Could you please tell me why?

Thanks,
Tuan

2

Hi Tuan,

Currently only fixed size device arrays are allowed in a Module. However, we are actively working on adding this (it’s our most requested feature). It should be available in 10.6 or possibly sooner.

  • Mat