在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 设备函数吗?