Я потратил 2 дня, пытаясь выяснить, почему я получаю ошибку «Недостаточно ресурсов» в программе Numba Cuda.
Моя проблема: Ошибка возникает при изменениикод из потока обрабатывает один элемент матрицы, который работает, чтобы поток обрабатывал весь столбец (как, например, в матрице умножается).Поскольку все матрицы помещаются в память независимо от того, каким образом я это делаю, и поскольку я не использую разделяемую память, я не могу понять, где я использую больше ресурсов.Если что-то зациклено на столбце, кажется, что это может быть меньше, а не больше ресурсов.
вот две версии ядра
# h,n, u, v f and out are NxM scale matricies where N~1000 M~10000.
def cuda_onethread_per_element (h, n, u, v, f, dx, dy, out):
gj,gi = cuda.grid(2)
gj,gi = sj,si
if (gj < out_u.shape[1]-1): # discard threads beyond last index.
out[gi,gj] =calc(gi,gj,h, n, u, v, f, dx, dy)
def cuda_1000_elements_per_thread(h, n, u, v, f, dx, dy, out):
gj,gi = cuda.grid(2)
for gi in range(out.shape[0]-1): # this loop is the change
if (gj < out_u.shape[1]-1):
out[gi,gj] =calc(gi,gj,h, n, u, v, f, dx, dy)
что происходит втораяпроисходит сбой, когда размер матрицы обычно составляет от 1000x10000 до 1000x20000, который начинает заполнять более половины графического процессора.
Traceback (most recent call last):
File "testerr.py", line 105, in <module>
testit((1,1),(1,1001),cu_u_driver_global_funcs_col)
File "testerr.py", line 102, in testit
cu_u_driver[blockspergrid,threadsperblock](h1, n1, u1, v1, f1, dx, dy, out_u1)
File "/home/rrs/anaconda3/envs/shallowwater/lib/python3.7/site-packages/numba/cuda/compiler.py", line 497, in __call__
sharedmem=self.sharedmem)
File "/home/rrs/anaconda3/envs/shallowwater/lib/python3.7/site-packages/numba/cuda/compiler.py", line 571, in _kernel_call
cu_func(*kernelargs)
File "/home/rrs/anaconda3/envs/shallowwater/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py", line 1514, in __call__
self.sharedmem, streamhandle, args)
File "/home/rrs/anaconda3/envs/shallowwater/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py", line 1558, in launch_kernel
None)
File "/home/rrs/anaconda3/envs/shallowwater/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py", line 290, in safe_cuda_api_call
self._check_error(fname, retcode)
File "/home/rrs/anaconda3/envs/shallowwater/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py", line 325, in _check_error
raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [701] Call to cuLaunchKernel results in CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES
Я не получаю сообщений об ошибках, если запускаю его в симуляторе Cuda GPU, только на графическом процессореRTX2070
Как должен работать код Первый из них заполняет только один элемент матрицы одним потоком, поэтому требуется сетка для покрытия всей матрицы блоками.второй требует только одного блока, чтобы покрыть всю матрицу, так как он зацикливается 10000 раз в каждом потоке, и каждый поток обрабатывает один из 1000 столбцов.
Регрессия Коды кажутся логически непротиворечивыми: когда код работает, оба метода дают идентичные числовые результаты.Там нет Нанс или Инфс.До тех пор, пока я не достигну предела ресурсов, изменение размера массива или размера блока не приведет к ошибке или численно отличному результату.
оба кода отлично работают на маленьких матрицах
Подсказка? Как ни странно, оба кода запускаются, когда я уменьшаю количество числовых вычислений в вычислениях, которые я делаю даже при использовании больших матриц.
Если я изменяю сложность вычисления, вызываемого небольшими суммами, я могу найтиместо, где это не подведет.в этой граничной точке в месте сбоя воспроизводимость сбоя исчезает - я предполагаю, что это может иметь отношение к другим незначительным использованиям графического процессора, который использует ОС.Но так как это продолжалось в течение многих дней, нет никакого систематического или сильного использования графического процессора, и только 10% его памяти связано с другими процессами для запуска дисплея.Когда код запускается, NVIDI-SMI сообщает, что используется 2/3 из 8 ГБ памяти на карте.
Мысли о том, что может быть причиной этого: Я полагаю, что второй метод, имеющийОдин блок размером от 1000 до 1024 может работать только на одном СМ, а первый может распространить нашу сетку, если захочет.Это может иметь значение, если используется много общей памяти, но мой код не использует ее.Просто глобальная память и локальные переменные.
У меня есть предположение, в котором я тоже сомневаюсь, что, возможно, Numba одновременно развертывает циклы и выстраивает вычисления.Если бы кто-то абсурдно развернул все 10000 циклов и включил все вычисления, то, может быть, просто, может быть, в итоге получилось бы 30000 локальных переменных.Но это было бы чокнутым.И я думаю, что коды умножения матриц потерпят неудачу, если бы это было правдой.
Более вероятной ошибкой было бы то, что я допустил какую-то глупую ошибку.но я не могу определить это.и я переписал код множеством способов, чтобы попытаться немного поколебать дерево и выявить ошибку.
Вот полная версия исполняемого кода
import numpy as np
import time
import numba as nb
import math
sqrt = math.sqrt
from numba import cuda
device_ii_compiler = cuda.jit('float32(int32,int32,float32[:,:], float32[:,:], float32[:,:],float32[:,:],float32[:,:],float32,float32)',device=True,inline=False)
compiler = cuda.jit('void(float32[:,:], float32[:,:], float32[:,:],float32[:,:],float32[:,:],float32,float32,float32[:,:])' ,inline=False)
def compileit (fn,blockdim=None,border=(0,0)):
name = fn.__name__
fc = compiler(fn)
fc.__name__ = name+'cuda'
fc.BLOCKDIM = blockdim
fc.BORDER = border
return fc
def dudt2_dummy_ret_py(i,j,h, n, f, u, v, dx, dy) :
p5 = np.float32(0.5)
last_v_n_i = (v[2,j]+v[2,5])*p5
v_n = (v[i,j]+v[i,5])*p5 # stradles n[i,j]
coriolis_u = (f[99,j]*last_v_n_i+f[i,j]*v_n)*p5 # coriolis force F is n-centered
return coriolis_u
cuda_dudt2_dummy_ret=device_ii_compiler(dudt2_dummy_ret_py)
def cu_u_driver_global_funcs_col_py(h, n, u, v, f, dx, dy, out_u):
sj,si = cuda.grid(2) # cuda.x cuda.y skow,fast
gj = sj
for gi in range(1, out_u.shape[0]-2):
if (gj < out_u.shape[1]-2) and gj>0 :
tmp = cuda_dudt2_dummy_ret(gi,gj,h, n, f, u, v, dx, dy) #
# now I just add on some more dummy calucaltions to titrate the
# place where the error gets tripped. removing the
# following lines removes the error. And the number
# of these additional calcs I need to trigger the error
# depends on the array size
tmp += sqrt(abs(np.float32(tmp)*dy+np.float32(gj)*dx)) # comment or un comment this line
tmp += sqrt(abs(np.float32(gi)*dx+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(tmp)*dy+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(gi)*dx+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(tmp)*dy+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(gi)*dx+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(gi)*dy+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(gi)*dx+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(gi)*dy+np.float32(gj)*dx)) # comment or un comment this line
# tmp += sqrt(abs(np.float32(gi)*dx+np.float32(gj)*dx)) # comment or un comment this line
# # tmp += sqrt(abs(gi*dy+gj*dx)) # comment or un comment this line
out_u[gi,gj] = tmp #funcs(si,sj,h, n, u, v, f, dx, dy)
cu_u_driver_global_funcs_col = compileit(cu_u_driver_global_funcs_col_py,blockdim=(1024,1))
M=20000
N=1000
h = np.ones((N,M),dtype=np.float32) #np.asarray( np.float32(2)+np.random.random((N,M)) ,dtype=np.float32)
n = np.ones((N,M),dtype=np.float32) #np.asarray(np.float32(2)+np.random.random((N,M)) ,dtype=np.float32)
u =np.ones((N+1,M),dtype=np.float32) # np.asarray(np.random.random((N+1,M)) ,dtype=np.float32)
v = np.ones((N,M+1),dtype=np.float32) #np.asarray(np.random.random((N,M+1)) ,dtype=np.float32)
f = np.ones((N,M),dtype=np.float32) #np.asarray(np.random.random((N,M)) ,dtype=np.float32)
dx = np.float32(0.1)
dy = np.float32(0.2)
out_u = np.zeros_like(u)
def testit(blockspergrid,threadsperblock,cu_u_driver,res1=None ):
print (cu_u_driver.__name__)
h1 = cuda.to_device(h)
n1 = cuda.to_device(n)
u1 = cuda.to_device(u)
v1 = cuda.to_device(v)
f1 = cuda.to_device(f)
out_u1 =cuda.to_device(out_u)
print( "blocks per grid", blockspergrid)
print("threads per block",threadsperblock)
cu_u_driver[blockspergrid,threadsperblock](h1, n1, u1, v1, f1, dx, dy, out_u1)
testit((1,1),(1,1001),cu_u_driver_global_funcs_col)