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

Posted

tags:

篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了在numba中cuda.local.array的正确用法是什么?相关的知识,希望对你有一定的参考价值。

我使用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。

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

有谁知道这个?

答案

直接从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运行时及其编译器生成的代码中。我没有兴趣仔细研究它以进一步解释。

以上是关于在numba中cuda.local.array的正确用法是什么?的主要内容,如果未能解决你的问题,请参考以下文章

可以在用户创建的 numba CUDA 设备函数中调用 numba.cuda.random 设备函数吗?

numba.errors.TypingError:在 nopython 模式管道中失败(步骤:nopython 前端)

在 numba 中使用 numpy.vstack

如何使用 numba 在 GPU 上泛化快速矩阵乘法

如何在 Numba 中使用指针包装 CFFI 函数

如何使用 python 和 numba 在 RTX GPU 中对 NVIDIA 的张量核心进行编程?