一个 CUDA 错误 当一个大数组被用作输入数据时
a CUDA error When a large array is used as input data
我有一个代码可以通过 python3.5 使用 numba 和 CUDA8.0 在 GPU 中进行一些计算。当输入大小为（50,27）的数组时，它 运行 成功并得到正确的结果。我把输入数据改成size(200,340)，报错
我在我的代码中使用了共享内存。是不是共享内存不够？还是grid size和block size不好？我不知道如何识别它并为网格和块选择合适的大小。
我设置了小格子大小和块大小，错误是一样的。
我应该怎么做才能解决这个问题？感谢您的一些建议。
我简化了我的代码，它有同样的错误。这里方便设置输入数据的大小：df = np.random.random_sample((300, 200)) + 10.
代码：
import os,sys,time,math
import pandas as pd
import numpy as np
from numba import cuda, float32
os.environ['NUMBAPRO_NVVM']=r'D:\NVIDIA GPU Computing Toolkit\CUDA\v8.0\nvvm\bin\nvvm64_31_0.dll'
os.environ['NUMBAPRO_LIBDEVICE']=r'D:\NVIDIA GPU Computing Toolkit\CUDA\v8.0\nvvm\libdevice'
bpg = 8
tpb = (4,32)
tsize = (3,4)
hsize = (1,4)
@cuda.jit
def calcu_T(D, T):
gw = cuda.gridDim.x
bx = cuda.blockIdx.x
tx = cuda.threadIdx.x
bw = cuda.blockDim.x
ty = cuda.threadIdx.y
bh = cuda.blockDim.y
c_num = D.shape[1]
c_index = bx
while c_index<c_num*c_num:
c_x = int(c_index/c_num)
c_y = c_index%c_num
if c_x==c_y:
T[c_x,c_y] = 0.0
else:
X = D[:,c_x]
Y = D[:,c_y]
hbuf = cuda.shared.array(hsize, float32)
h = tx
Xi = X[h:]
Xi1 = X[:-h]
Yih = Y[:-h]
sbuf = cuda.shared.array(tsize, float32)
L = len(Xi)
#mean
if ty==0:
Xi_m = 0.0
Xi1_m = 0.0
Yih_m = 0.0
for i in range(L):
Xi_m += Xi[i]
Xi1_m += Xi1[i]
Yih_m += Yih[i]
Xi_m = Xi_m/L
Xi1_m = Xi1_m/L
Yih_m = Yih_m/L
sbuf[0,tx] = Xi_m
sbuf[1,tx] = Xi1_m
sbuf[2,tx] = Yih_m
cuda.syncthreads()
sl = cuda.shared.array(tpb, float32)
r_index = ty
s_l = 0.0
while r_index<L:
s1 = 0.0
for i in range(L):
s1 += (Xi[r_index]+Xi1[i])/sbuf[0,tx]
s_l += s1
r_index +=bh
sl[tx,ty] = s_l
cuda.syncthreads()
#
if ty==0:
ht = 0.0
for i in range(bh):
ht += sl[tx,i]
hbuf[0,tx] = ht/L
cuda.syncthreads()
#max
if tx==0 and ty==0:
m_t = 0.0
for index,ele in enumerate(hbuf[0]):
if index==0:
m_t = ele
elif ele>m_t:
m_t = ele
T[c_x,c_y] = m_t
c_index +=gw
df = np.random.random_sample((300, 200)) + 10
D = np.array(df, dtype=np.float32)
r,c = D.shape
T = np.empty([c,c])
dD = cuda.to_device(D)
dT = cuda.device_array_like(T)
calcu_T[bpg, tpb](dD,dT)
dT.copy_to_host(T)
错误：
Traceback (most recent call last):
File "G:\myworkspace\python3.5\forte\forte170327\test10fortest8.py", line 118, in <module>
dT.copy_to_host(T)
File "D:\python3.5.3\lib\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 "D:\python3.5.3\lib\site-packages\numba\cuda\cudadrv\driver.py", line 1481, in device_to_host
fn(host_pointer(dst), device_pointer(src), size, *varargs)
File "D:\python3.5.3\lib\site-packages\numba\cuda\cudadrv\driver.py", line 259, in safe_cuda_api_call
self._check_error(fname, retcode)
File "D:\python3.5.3\lib\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
我的设备信息：
Device 0:
CUDA Driver Version / Runtime Version 8.0 / 8.0
CUDA Capability Major/Minor version number: 5.0
Total amount of global memory: 2048 MBytes (2147483648 bytes)
( 5) Multiprocessors, (128) CUDA Cores/MP: 640 CUDA Cores
Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
Maximum Layered 1D Texture Size, (num) layers 1D=(16384), 2048 layers
Maximum Layered 2D Texture Size, (num) layers 2D=(16384, 16384), 2048 layers
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 49152 bytes
Total number of registers available per block: 65536
Warp size: 32
Maximum number of threads per multiprocessor: 2048
Maximum number of threads per block: 1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)
Maximum memory pitch: 2147483647 bytes
Texture alignment: 512 bytes
您的代码没有任何问题。如果我 运行 你的代码在我的 GTX970 上，我得到这个：
In [11]: main??
Signature: main()
Source:
def main():
df = np.random.random_sample((300, 200)) + 10
D = np.array(df, dtype=np.float32)
r,c = D.shape
T = np.empty([c,c])
dD = cuda.to_device(D)
dT = cuda.device_array_like(T)
calcu_T[bpg, tpb](dD,dT)
dT.copy_to_host(T)
File: ~/SO/crash.py
Type: function
In [12]: %timeit -n 3 -r 3 main()
3 loops, best of 3: 6.61 s per loop
即没有 运行 时间错误，但是包含内核的 python 代码需要 6.6 秒才能 运行。如果我使用 CUDA 分析器分析代码：
$ nvprof python crash.py
==13828== NVPROF is profiling process 13828, command: python crash.py
All finished
==13828== Profiling application: python crash.py
==13828== Profiling result:
Time(%) Time Calls Avg Min Max Name
100.00% 6.59109s 1 6.59109s 6.59109s 6.59109s cudapy::__main__::calcu_T1(Array<float, int=2, A, mutable, aligned>, Array<double, int=2, A, mutable, aligned>)
0.00% 26.271us 1 26.271us 26.271us 26.271us [CUDA memcpy DtoH]
0.00% 21.279us 1 21.279us 21.279us 21.279us [CUDA memcpy HtoD]
==13828== API calls:
Time(%) Time Calls Avg Min Max Name
98.51% 6.59118s 1 6.59118s 6.59118s 6.59118s cuMemcpyDtoH
1.42% 94.890ms 1 94.890ms 94.890ms 94.890ms cuDevicePrimaryCtxRetain
0.05% 3.4116ms 1 3.4116ms 3.4116ms 3.4116ms cuModuleLoadDataEx
0.01% 417.96us 1 417.96us 417.96us 417.96us cuLinkCreate
0.00% 227.57us 1 227.57us 227.57us 227.57us cuLinkAddData
0.00% 195.72us 2 97.859us 95.710us 100.01us cuMemAlloc
0.00% 190.10us 1 190.10us 190.10us 190.10us cuLinkComplete
0.00% 139.04us 1 139.04us 139.04us 139.04us cuMemGetInfo
0.00% 53.193us 1 53.193us 53.193us 53.193us cuMemcpyHtoD
0.00% 29.538us 1 29.538us 29.538us 29.538us cuDeviceGetName
0.00% 17.895us 1 17.895us 17.895us 17.895us cuLaunchKernel
0.00% 2.0250us 1 2.0250us 2.0250us 2.0250us cuCtxPushCurrent
0.00% 2.0150us 5 403ns 255ns 752ns cuFuncGetAttribute
0.00% 1.6260us 2 813ns 547ns 1.0790us cuDeviceGetCount
0.00% 1.1430us 1 1.1430us 1.1430us 1.1430us cuModuleGetFunction
0.00% 951ns 2 475ns 372ns 579ns cuDeviceGet
0.00% 796ns 1 796ns 796ns 796ns cuLinkDestroy
0.00% 787ns 1 787ns 787ns 787ns cuDeviceComputeCapability
可以看到你贴出的内核需要6.5秒才能运行。
您没有提供任何详细信息，但我猜您 运行 正在 Windows，您的 GPU 是显示 GPU，您的代码 运行s 足够慢以至于它正在达到 WDDM 显示管理器看门狗超时限制。这是非常有据可查的，之前已经被问过数百次——例如 here.
您选择的搜索引擎和 CUDA Windows 入门指南将为您提供有关从操作系统和硬件角度改善情况的备选方案的信息。然而，最明显的是简单地改进您的代码以使其 运行 更快。
我有一个代码可以通过 python3.5 使用 numba 和 CUDA8.0 在 GPU 中进行一些计算。当输入大小为（50,27）的数组时，它 运行 成功并得到正确的结果。我把输入数据改成size(200,340)，报错
我在我的代码中使用了共享内存。是不是共享内存不够？还是grid size和block size不好？我不知道如何识别它并为网格和块选择合适的大小。
我设置了小格子大小和块大小，错误是一样的。
我应该怎么做才能解决这个问题？感谢您的一些建议。
我简化了我的代码，它有同样的错误。这里方便设置输入数据的大小：df = np.random.random_sample((300, 200)) + 10.
代码：
import os,sys,time,math
import pandas as pd
import numpy as np
from numba import cuda, float32
os.environ['NUMBAPRO_NVVM']=r'D:\NVIDIA GPU Computing Toolkit\CUDA\v8.0\nvvm\bin\nvvm64_31_0.dll'
os.environ['NUMBAPRO_LIBDEVICE']=r'D:\NVIDIA GPU Computing Toolkit\CUDA\v8.0\nvvm\libdevice'
bpg = 8
tpb = (4,32)
tsize = (3,4)
hsize = (1,4)
@cuda.jit
def calcu_T(D, T):
gw = cuda.gridDim.x
bx = cuda.blockIdx.x
tx = cuda.threadIdx.x
bw = cuda.blockDim.x
ty = cuda.threadIdx.y
bh = cuda.blockDim.y
c_num = D.shape[1]
c_index = bx
while c_index<c_num*c_num:
c_x = int(c_index/c_num)
c_y = c_index%c_num
if c_x==c_y:
T[c_x,c_y] = 0.0
else:
X = D[:,c_x]
Y = D[:,c_y]
hbuf = cuda.shared.array(hsize, float32)
h = tx
Xi = X[h:]
Xi1 = X[:-h]
Yih = Y[:-h]
sbuf = cuda.shared.array(tsize, float32)
L = len(Xi)
#mean
if ty==0:
Xi_m = 0.0
Xi1_m = 0.0
Yih_m = 0.0
for i in range(L):
Xi_m += Xi[i]
Xi1_m += Xi1[i]
Yih_m += Yih[i]
Xi_m = Xi_m/L
Xi1_m = Xi1_m/L
Yih_m = Yih_m/L
sbuf[0,tx] = Xi_m
sbuf[1,tx] = Xi1_m
sbuf[2,tx] = Yih_m
cuda.syncthreads()
sl = cuda.shared.array(tpb, float32)
r_index = ty
s_l = 0.0
while r_index<L:
s1 = 0.0
for i in range(L):
s1 += (Xi[r_index]+Xi1[i])/sbuf[0,tx]
s_l += s1
r_index +=bh
sl[tx,ty] = s_l
cuda.syncthreads()
#
if ty==0:
ht = 0.0
for i in range(bh):
ht += sl[tx,i]
hbuf[0,tx] = ht/L
cuda.syncthreads()
#max
if tx==0 and ty==0:
m_t = 0.0
for index,ele in enumerate(hbuf[0]):
if index==0:
m_t = ele
elif ele>m_t:
m_t = ele
T[c_x,c_y] = m_t
c_index +=gw
df = np.random.random_sample((300, 200)) + 10
D = np.array(df, dtype=np.float32)
r,c = D.shape
T = np.empty([c,c])
dD = cuda.to_device(D)
dT = cuda.device_array_like(T)
calcu_T[bpg, tpb](dD,dT)
dT.copy_to_host(T)
错误：
Traceback (most recent call last):
File "G:\myworkspace\python3.5\forte\forte170327\test10fortest8.py", line 118, in <module>
dT.copy_to_host(T)
File "D:\python3.5.3\lib\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 "D:\python3.5.3\lib\site-packages\numba\cuda\cudadrv\driver.py", line 1481, in device_to_host
fn(host_pointer(dst), device_pointer(src), size, *varargs)
File "D:\python3.5.3\lib\site-packages\numba\cuda\cudadrv\driver.py", line 259, in safe_cuda_api_call
self._check_error(fname, retcode)
File "D:\python3.5.3\lib\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
我的设备信息：
Device 0:
CUDA Driver Version / Runtime Version 8.0 / 8.0
CUDA Capability Major/Minor version number: 5.0
Total amount of global memory: 2048 MBytes (2147483648 bytes)
( 5) Multiprocessors, (128) CUDA Cores/MP: 640 CUDA Cores
Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
Maximum Layered 1D Texture Size, (num) layers 1D=(16384), 2048 layers
Maximum Layered 2D Texture Size, (num) layers 2D=(16384, 16384), 2048 layers
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 49152 bytes
Total number of registers available per block: 65536
Warp size: 32
Maximum number of threads per multiprocessor: 2048
Maximum number of threads per block: 1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)
Maximum memory pitch: 2147483647 bytes
Texture alignment: 512 bytes
您的代码没有任何问题。如果我 运行 你的代码在我的 GTX970 上，我得到这个：
In [11]: main??
Signature: main()
Source:
def main():
df = np.random.random_sample((300, 200)) + 10
D = np.array(df, dtype=np.float32)
r,c = D.shape
T = np.empty([c,c])
dD = cuda.to_device(D)
dT = cuda.device_array_like(T)
calcu_T[bpg, tpb](dD,dT)
dT.copy_to_host(T)
File: ~/SO/crash.py
Type: function
In [12]: %timeit -n 3 -r 3 main()
3 loops, best of 3: 6.61 s per loop
即没有 运行 时间错误，但是包含内核的 python 代码需要 6.6 秒才能 运行。如果我使用 CUDA 分析器分析代码：
$ nvprof python crash.py
==13828== NVPROF is profiling process 13828, command: python crash.py
All finished
==13828== Profiling application: python crash.py
==13828== Profiling result:
Time(%) Time Calls Avg Min Max Name
100.00% 6.59109s 1 6.59109s 6.59109s 6.59109s cudapy::__main__::calcu_T1(Array<float, int=2, A, mutable, aligned>, Array<double, int=2, A, mutable, aligned>)
0.00% 26.271us 1 26.271us 26.271us 26.271us [CUDA memcpy DtoH]
0.00% 21.279us 1 21.279us 21.279us 21.279us [CUDA memcpy HtoD]
==13828== API calls:
Time(%) Time Calls Avg Min Max Name
98.51% 6.59118s 1 6.59118s 6.59118s 6.59118s cuMemcpyDtoH
1.42% 94.890ms 1 94.890ms 94.890ms 94.890ms cuDevicePrimaryCtxRetain
0.05% 3.4116ms 1 3.4116ms 3.4116ms 3.4116ms cuModuleLoadDataEx
0.01% 417.96us 1 417.96us 417.96us 417.96us cuLinkCreate
0.00% 227.57us 1 227.57us 227.57us 227.57us cuLinkAddData
0.00% 195.72us 2 97.859us 95.710us 100.01us cuMemAlloc
0.00% 190.10us 1 190.10us 190.10us 190.10us cuLinkComplete
0.00% 139.04us 1 139.04us 139.04us 139.04us cuMemGetInfo
0.00% 53.193us 1 53.193us 53.193us 53.193us cuMemcpyHtoD
0.00% 29.538us 1 29.538us 29.538us 29.538us cuDeviceGetName
0.00% 17.895us 1 17.895us 17.895us 17.895us cuLaunchKernel
0.00% 2.0250us 1 2.0250us 2.0250us 2.0250us cuCtxPushCurrent
0.00% 2.0150us 5 403ns 255ns 752ns cuFuncGetAttribute
0.00% 1.6260us 2 813ns 547ns 1.0790us cuDeviceGetCount
0.00% 1.1430us 1 1.1430us 1.1430us 1.1430us cuModuleGetFunction
0.00% 951ns 2 475ns 372ns 579ns cuDeviceGet
0.00% 796ns 1 796ns 796ns 796ns cuLinkDestroy
0.00% 787ns 1 787ns 787ns 787ns cuDeviceComputeCapability
可以看到你贴出的内核需要6.5秒才能运行。
您没有提供任何详细信息，但我猜您 运行 正在 Windows，您的 GPU 是显示 GPU，您的代码 运行s 足够慢以至于它正在达到 WDDM 显示管理器看门狗超时限制。这是非常有据可查的，之前已经被问过数百次——例如 here.
您选择的搜索引擎和 CUDA Windows 入门指南将为您提供有关从操作系统和硬件角度改善情况的备选方案的信息。然而，最明显的是简单地改进您的代码以使其 运行 更快。