Cuda illegal memory access error when using array indexes stored in another array

Question

I'm using cuda fortran and I've been struggling with this problem in one simple kernel and I couldn't find the solution. Isn't it possible to use integer values stored in an array as the indexes for another array?

Here's a simple example (edited to include also the main program):

program test
  use cudafor

  integer:: ncell, i
  integer, allocatable:: values(:)
  integer, allocatable, device :: values_d(:)


  ncell = 10

  allocate(values(ncell), values_d(ncell))

  do i=1,ncell
        values(i) = i
  enddo


  values_d = values


  call multipleindices_kernel<<< ncell/1024+1,1024 >>> (values_d,
 + ncell)

  values = values_d

  write (*,*) values


  end program test

!////////////////////////////////////////////////////

attributes(global) subroutine multipleindices_kernel(valu, ncell)
use cudafor
  implicit none
  integer, value:: ncell   ! ncell = 10
  integer :: valu(ncell)
  integer :: tempind(10)
  integer:: i

  tempind(1)=10
  tempind(2)=3
  tempind(3)=5
  tempind(4)=7
  tempind(5)=9
  tempind(6)=2
  tempind(7)=4
  tempind(8)=6
  tempind(9)=8
  tempind(10)=1

  i = (blockidx%x - 1 ) * blockdim%x + threadidx%x

  if (i .LE. ncell) then
        valu(tempind(i))= 1
  endif


  end subroutine

I understand that if there were repeated values in the tempind array different threads could be accessing the same memory location for reading or writting, but that is not the case. Even though, this gives the error "0: copyout Memcpy (host=0x303610, dev=0x3e20000, size=40) FAILED: 77(an illegal memory access was encountered).

Does anyone know if it is possible to use this indexes coming from another array in cuda?

After some additional tests, I've noticed that the problem occurs not while running the kernel itself, but on the transfer of the data back to CPU (if I remove "values = values_d" then no error is displayed). Also, if I substitute in the kernel valu(tempind(i)) by valu(i) it works fine, but I want to have the indexes coming from an array since the purpose of this test is to make a parallelization of a CFD code where the indexes are stored like that.

Answer 1

The problem appears to be that the generated executable doesn't pass the variable ncell to the kernel correctly. Running the application through cuda-memcheck shows that threads outside of the 1-10 are passing through the branch statement, and adding a print statement to print ncell inside the kernel also gives strange answers.

It used to be a requirement that all attributes(global) subroutines had to reside within a module. This requirement seems to have been relaxed in more recent versions of CUDA Fortran (I cannot find references to it in the programming guide). I believe the code outside of the module is causing the error here. By placing multipleindices_kernel within a module and using that module in test I consistantly get correct answers with no errors. The code for this is below:

module testmod
contains
attributes(global) subroutine multipleindices_kernel(valu, ncell)
  use cudafor
  implicit none
  integer, value:: ncell   ! ncell = 10
  integer :: valu(ncell)
  integer :: tempind(10)
  integer:: i

  tempind(1)=10
  tempind(2)=3
  tempind(3)=5
  tempind(4)=7
  tempind(5)=9
  tempind(6)=2
  tempind(7)=4
  tempind(8)=6
  tempind(9)=8
  tempind(10)=1

  i = (blockidx%x - 1 ) * blockdim%x + threadidx%x

  if (i .LE. ncell) then
        valu(tempind(i))= 1
  endif


  end subroutine
end module testmod

  program test
  use cudafor
  use testmod

  integer:: ncell, i
  integer, allocatable:: values(:)
  integer, allocatable, device :: values_d(:)


  ncell = 10

  allocate(values(ncell), values_d(ncell))

  do i=1,ncell
        values(i) = i
  enddo


  values_d = values


  call multipleindices_kernel<<< ncell/1024+1,1024 >>> (values_d, ncell)

  values = values_d

  write (*,*) values


  end program test

!////////////////////////////////////////////////////