代码之家 › 专栏 › 技术社区 › ZHANG Juenjie

cuda的正确用法是什么。地方的numba中的数组?

numba cuda python

ZHANG Juenjie · 技术社区 · 7 年前

我使用Nuba用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的用法感到困惑。地方的大堆

它有两个论点。一个是shape,另一个是dtype。

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

我不理解这种争论的形式。

有人知道吗?

1 回复 | 直到 7 年前

talonmies 7 年前

直接从 documentation :

本地内存是每个线程专用的内存区域。使用本地当标量局部变量这还不够。在内核,不同于传统的动态内存管理。
numba.cuda.local.array(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

i、 e.使用不正确的本地数组大小运行会产生预期的内存访问错误。然而,代码实际上仍在运行。另一方面,如果我修改您的代码以使用 shape=3 :

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

内存访问错误消失。因此,您不应该混淆正确工作和未定义的行为(这可能包括意外工作,但抛出错误,如本例所示)。发生这种情况的确切原因将隐藏在Nuba运行时及其编译器生成的代码中。我没有兴趣对此进行详细解释。