在numba中cuda.local.array的正确用法是什么？

Question

我使用numba在python中编写了一个测试代码。

from numba import cuda
import numpy as np
import numba
@cuda.jit
def function(output, size, random_array):
    i_p, i_k1, i_k2 = cuda.grid(3)
    a=cuda.local.array(shape=1,dtype=numba.float64)
    if i_p<size and i_k1<size and i_k2<size:
        a1=i_p
        a2=i_k1+1
        a3=i_k2+2

        a[0]=a1
        a[1]=a2
        a[2]=a3
        for i in range(len(random_array)):
            output[i_p,i_k1,i_k2,i] = a[int(random_array[i])]
output=cuda.device_array((2,2,2,5))

random_array=np.array([np.random.random()*3 for i in range(5)])
print(random_array)
random_array0=cuda.to_device(random_array)
size=2
threadsperblock = (8, 8, 8)
blockspergridx=(size + (threadsperblock[0] - 1)) // threadsperblock[0]
blockspergrid = ((blockspergridx, blockspergridx, blockspergridx))

# Start the kernel 
function[blockspergrid, threadsperblock](output, size, random_array0)
print(output.copy_to_host())

# test if it is consistent with non gpu case
output=np.zeros([2,2,2,5])
for i in range(size):
    for j in range(size):
        for k in range(size):
            a=[i,j+1,k+2]
            for ii in range(len(random_array)):
                output[i,j,k,ii] = a[int(random_array[ii])]
print(output)

我对cuda.local.array的用法感到困惑。

它有两个论点。一个是形状，另一个是dtype。

但是，结果不会随着形状的不同而改变。例如，shape = 0或shape = 1或shape = 100。

我不明白这个论点的形状。

有谁知道这个？

Answer 1

直接从documentation引用：

本地内存是每个线程专用的内存区域。当标量局部变量不足时，使用本地内存有助于分配一些暂存区域。与传统的动态内存管理不同，内存在内核持续时间内分配一次。

numba.cuda.local.array(shape, type) 

在设备上分配给定形状和类型的本地数组。 shape是整数或表示数组维度的整数元组，必须是简单的常量表达式。 type是需要存储在数组中的Numba类型的元素。该数组对当前线程是私有的。返回类似于数组的对象，可以像任何标准数组一样读取和写入（例如通过索引）。

因此，在这种情况下，如果您希望本地内存至少包含三个元素，则必须使用shape >= 3才能使代码正常工作。

您编码似乎与shape=1一起使用的事实应该被视为未定义的行为。如果我使用cuda-memcheck运行你的代码，我得到这个：

$ cuda-memcheck python indexing.py 
========= CUDA-MEMCHECK
[ 1.99261914  1.91166157  2.85454532  1.64078385  1.9576766 ]
========= Invalid __local__ write of size 8
=========     at 0x000001b0 in cudapy::__main__::function$241(Array<double, int=4, A, mutable, aligned>, __int64, Array<double, int=1, A, mutable, aligned>)
=========     by thread (1,1,1) in block (0,0,0)
=========     Address 0x00fffc80 is out of bounds

[SNIPPED for brevity]


=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so (cuLaunchKernel + 0x2cd) [0x23c06d]
Traceback (most recent call last):
  File "indexing.py", line 42, in <module>
    outputd = output.copy_to_host()
  File "/opt/miniconda3/lib/python3.6/site-packages/numba/cuda/cudadrv/devicearray.py", line 198, in copy_to_host
    _driver.device_to_host(hostary, self, self.alloc_size, stream=stream)
  File "/opt/miniconda3/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py", line 1481, in device_to_host
    fn(host_pointer(dst), device_pointer(src), size, *varargs)
  File "/opt/miniconda3/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py", line 259, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "/opt/miniconda3/lib/python3.6/site-packages/numba/cuda/cudadrv/driver.py", line 296, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [719] Call to cuMemcpyDtoH results in UNKNOWN_CUDA_ERROR
========= ERROR SUMMARY: 9 errors

即，使用不正确的本地数组大小运行会产生内存访问错误，如您所料。但是，代码仍然实际运行。另一方面，如果我修改你的代码使用shape=3：

$ cuda-memcheck python indexing.py 
========= CUDA-MEMCHECK
[ 1.98532356  1.53822652  0.69376061  2.22448278  0.76800584]
True
========= ERROR SUMMARY: 0 errors

内存访问错误消失。所以你不应该混淆正确的工作和未定义的行为（可能包括意外工作，但抛出错误，就像在这种情况下）。发生这种情况的确切原因将隐藏在numba运行时及其编译器生成的代码中。我没有兴趣仔细研究它以进一步解释。