當使用存儲在另一個陣列中的數組索引時，Cuda非法內存訪問錯誤

Question

我正在使用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 

我明白，如果有所述tempind陣列中重複的值不同的線程可以訪問相同的存儲器位置讀取或書面方式， 但事實並非如此。 儘管如此，這給出錯誤「0：拷貝Memcpy（主機= 0x303610，dev = 0x3e20000，大小= 40）失敗：77（遇到非法內存訪問）

有誰知道是否有可能使用來自cuda中的另一個數組的這些索引？

經過一些額外的測試後，我發現問題不在運行內核本身時發生，而是在將數據傳輸回CPU（如果刪除「值= values_d「，那麼不會顯示錯誤）另外，如果我通過valu（i）替換內核值（tempind（i）），它可以正常工作，但我希望索引來自數組，因爲這樣做的目的測試的目的是對索引存儲的CFD代碼進行並行化處理。

Answer 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 

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