使用存储在另一个数组中的数组索引时，Cuda 非法内存访问错误

Posted 2023-02-17

【中文标题】使用存储在另一个数组中的数组索引时，Cuda 非法内存访问错误

【英文标题】：Cuda illegal memory access error when using array indexes stored in another array

【发布时间】：2014-10-01 18:02:05

【问题描述】：

我正在使用 cuda fortran，我一直在一个简单的内核中努力解决这个问题，但我找不到解决方案。 难道不能用一个数组中存储的整数值作为另一个数组的索引吗？

这是一个简单的例子（编辑后还包括主程序）：

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

我知道如果临时数组中有重复的值，不同的线程可能会访问相同的内存位置以进行读取或写入，但事实并非如此。 尽管如此，这会给出错误“0: copyout Memcpy (host=0x303610, dev=0x3e20000, size=40) FAILED: 77（遇到非法内存访问）。

有谁知道是否可以使用来自 cuda 中另一个数组的索引？

经过一些额外的测试，我注意到问题不是在运行内核本身时出现，而是在将数据传输回 CPU 时出现（如果我删除“values = values_d”，则不会显示错误）。此外，如果我用 valu(i) 替换内核 value(tempind(i)) 它可以正常工作，但我希望索引来自数组，因为此测试的目的是使 CFD 代码并行化索引是这样存储的。

【问题讨论】：

有可能，所以报错很可能和value有关。自从我做任何 CUDA Fortran 以来已经有一段时间了，但它不应该有设备属性吗？无论如何，我建议您发布一个完整的复制示例，否则任何人都很难帮助您。

内核子例程定义中的变量声明（即使是参数）不需要device 属性。但是，在主程序代码中这是必需的（对于内核子程序使用的变量）。这看起来很奇怪：integer :: valu(ncell) 我本来期望的：integer :: valu(:) 但它可能没问题。我同意需要一个完整的复制示例。

您好，谢谢您的回答。我已经编辑发布了调用内核的简单程序测试，并添加了额外的 cmets 来帮助我们找到这个问题的原因。

【参考方案1】：

问题似乎是生成的可执行文件没有将变量ncell 正确传递给内核。通过cuda-memcheck运行应用程序显示1-10之外的线程正在通过分支语句，在内核内部添加打印语句打印ncell也会给出奇怪的答案。

过去要求所有attributes(global) 子例程都必须驻留在一个模块中。在更新的 CUDA Fortran 版本中，这个要求似乎已经放宽了（我在编程指南中找不到对它的引用）。我相信模块外部的代码导致了这里的错误。通过将multipleindices_kernel 放在一个模块中并在test 中使用该模块，我始终可以得到正确的答案而没有错误。代码如下：

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

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

【讨论】：

您好，杰兹，谢谢您的回答。确实解决了问题。我不知道它是否可以在您的编译器中这样工作，但就我而言，使用 PGI 编译器，我必须将“包含”语句放在模块的开头。不知道为什么这在没有模块的情况下不起作用，因为到目前为止我一直在使用其他类似的内核没有问题（不使用存储在其他数组中的索引）。

你是对的。 contains 粘贴时不知何故迷路了！固定。