a CUDA error When a large array is used as input data

Question

I have a code to do some calculation in GPU by python3.5 with numba and CUDA8.0. When an array with size(50,27) was input, it run successfully and get right result. I change the input data to size(200,340), it has an error.

I use shared memory in my code. Is there not enough shared memory? Or the grid size and block size are not good? I don't know how to identify it and choose appropriate size for grid and block.

I set small grid size and block size, the error is the same.

What should I do to solve this problem? Thanks for some advice.

I simplified my code and it has the same error. It is convenient to set the size of the input data here:df = np.random.random_sample((300, 200)) + 10.

The code:

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)

The error:

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

My device information:

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

You have already asked and deleted this question once. The problem with the question is the same -- there is not enough detail given to provide an answer. Please read [MCVE]. — talonmies, Apr 12 '17 at 14:21
Sorry, I thought my previous question was premature and I should solve it by myself, so I deleted it(I don't know deleting a question here is bad. Now I have known). During these days, I'm working for this problem. Until now, the problem still hasn't been resolved so that I ask this question again. — S.Jan, Apr 13 '17 at 02:08

score 4 · Accepted Answer · edited May 23 '17 at 12:32

There is nothing wrong with your code. If I run your code on my GTX970, I get this:

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

i.e. no runtime errors, but the python code including the kernel takes 6.6 seconds to run. If I profile the code with the CUDA profiler:

$ 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_T$241(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

you can see that the kernel you have posted takes 6.5 seconds to run.

You have provided no details, but I will guess that you are running on Windows, your GPU is a display GPU and your code runs slow enough that it is hitting the WDDM display manager watchdog timeout limit. This is extremely well documented and has been asked about literally hundreds of times before -- for example here.

Your search engine of choice and the CUDA Windows getting started guide will provide you with information on what your alternatives are to improve the situation from a operating system and hardware point of view. The most obvious, however, is simply to improve your code to make it run faster.

My GPU is GTX850M. As you said, It's a display GPU. I run this code on my notebook with Windows. Thank you for your helps recently. I'll learn more and try to make my code run faster. — S.Jan, Apr 13 '17 at 08:57

a CUDA error When a large array is used as input data

1 Answers1