Cuda out of resources error when running python numbapro

Question

I am trying to run a cuda kernel in numbapro python, but I keep getting an out of resources error. I then tried to execute the kernel into a loop and send smaller arrays, but that still gave me the same error.

Here is my error message:

Traceback (most recent call last):
File "./predict.py", line 418, in <module>
predict[griddim, blockdim, stream](callResult_d, catCount, numWords, counts_d, indptr_d, indices_d, probtcArray_d, priorC_d)
File "/home/mhagen/Developer/anaconda/lib/python2.7/site-packages/numba/cuda/compiler.py", line 228, in __call__
sharedmem=self.sharedmem)
File "/home/mhagen/Developer/anaconda/lib/python2.7/site-packages/numba/cuda/compiler.py", line 268, in _kernel_call
cu_func(*args)
File "/home/mhagen/Developer/anaconda/lib/python2.7/site-packages/numba/cuda/cudadrv/driver.py", line 1044, in __call__
self.sharedmem, streamhandle, args)
File "/home/mhagen/Developer/anaconda/lib/python2.7/site-packages/numba/cuda/cudadrv/driver.py", line 1088, in launch_kernel
None)
File "/home/mhagen/Developer/anaconda/lib/python2.7/site-packages/numba/cuda/cudadrv/driver.py", line 215, in safe_cuda_api_call
self._check_error(fname, retcode)
File "/home/mhagen/Developer/anaconda/lib/python2.7/site-packages/numba/cuda/cudadrv/driver.py", line 245, in _check_error
raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: Call to cuLaunchKernel results in CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES

Here is my source code:

from numbapro.cudalib import cusparse
from numba import *
from numbapro import cuda

@cuda.jit(argtypes=(double[:], int64, int64, double[:], int64[:], int64[:], double[:,:], double[:] ))
def predict( callResult, catCount, wordCount, counts, indptr, indices, probtcArray, priorC ):

    i = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x

    correct = 0
    wrong = 0
    lastDocIndex = -1
    maxProb = -1e6
    picked = -1

    for cat in range(catCount):

        probSum = 0.0

        for j in range(indptr[i],indptr[i+1]):
            wordIndex = indices[j]
            probSum += (counts[j]*math.log(probtcArray[cat,wordIndex]))

        probSum += math.log(priorC[cat])
        if probSum > maxProb:
            maxProb = probSum
            picked = cat

    callResult[i] = picked

predictions = []
counter = 1000
for i in range(int(math.ceil(numDocs/(counter*1.0)))):
    docTestSliceList = docTestList[i*counter:(i+1)*counter]
    numDocsSlice = len(docTestSliceList)
    docTestArray = np.zeros((numDocsSlice,numWords))
    for j,doc in enumerate(docTestSliceList):
        for ind in doc:
            docTestArray[j,ind['term']] = ind['count']
    docTestArraySparse = cusparse.ss.csr_matrix(docTestArray)

    start = time.time()
    OPT_N = numDocsSlice
    blockdim = 1024, 1
    griddim = int(math.ceil(float(OPT_N)/blockdim[0])), 1 

    catCount = len(music_categories)
    callResult = np.zeros(numDocsSlice)
    stream = cuda.stream()
    with stream.auto_synchronize():
        probtcArray_d = cuda.to_device(numpy.asarray(probtcArray),stream)
        priorC_d = cuda.to_device(numpy.asarray(priorC),stream)
        callResult_d = cuda.to_device(callResult, stream)
        counts_d = cuda.to_device(docTestArraySparse.data, stream)
        indptr_d = cuda.to_device(docTestArraySparse.indptr, stream)
        indices_d = cuda.to_device(docTestArraySparse.indices, stream)
        predict[griddim, blockdim, stream](callResult_d, catCount, numWords, counts_d, indptr_d, indices_d, probtcArray_d, priorC_d)

        callResult_d.to_host(stream)
    #stream.synchronize()
    predictions += list(callResult)

    print "prediction %d: %f" % (i,time.time()-start)

Answer 1

I found out this was in the cuda procedure.

When you call predict the blockdim is set to 1024. predict[griddim, blockdim, stream](callResult_d, catCount, numWords, counts_d, indptr_d, indices_d, probtcArray_d, priorC_d)

But the procedure is called iteratively with slice sizes of 1000 elements not 1024. So, in the procedure it will attempt to write 24 elements that are out of bounds in the return array.

Sending a number of elements parameter (n_el) and placing an error checking call in the cuda procedure solves it.

@cuda.jit(argtypes=(double[:], int64, int64, int64, double[:], int64[:], int64[:], double[:,:], double[:] ))
def predict( callResult, n_el, catCount, wordCount, counts, indptr, indices, probtcArray, priorC ):

     i = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x

     if i < n_el:

         ....