title: Why numba cuda is running slow after recalling it several times?
body: <p>I am experimenting how to use cuda inside numba. However I have encountered something different from my expectation. Here is my code</p>
  
  <pre><code>from numba import cuda
  @cuda.jit
  def matmul(A, B, C):
  """Perform square matrix multiplication of C = A * B
  """
  d=cuda.local.array((3,3),dtype=numba.float64)
  i, j = cuda.grid(2)
  if i &lt; C.shape[0] and j &lt; C.shape[1]:
      tmp = 0.
      for k in range(A.shape[1]):
          tmp += A[i, k] * B[k, j]
      C[i, j] = tmp
  </code></pre>
  
  <p>This is the matrix function I self defined for testing using numba.cuda. Before running the tests, I also loaded the arrays in the following code:</p>
  
  <pre><code>import numpy as np
  a=np.random.rand(2000,2000)
  b=np.random.rand(2000,2000)
  c=np.empty((2000,2000))
  a1=cuda.to_device(a)
  b1=cuda.to_device(b)
  c1=cuda.to_device(c)
  </code></pre>
  
  <p>Then I used the following code for experiment:</p>
  
  <pre><code>from time import time
  count =0
  start=time()
  for i in range(2000):
    matmul[(256,256),(16,16)](a1,b1,c1)
    count +=1
    print(count)
  </code></pre>
  
  <p>The for loops ran very fast for the first 1028 runs. However it ran very slow after the 1028th.What exactly caused this and how do I fix it. I am running on win10 by the way.</p>
  
  <p>Here is my cuda information called from numba.cuda</p>
  
  <pre><code>from numba import cuda
  gpu = cuda.get_current_device()
  print("name = %s" % gpu.name)
  print("maxThreadsPerBlock = %s" % str(gpu.MAX_THREADS_PER_BLOCK))
  print("maxBlockDimX = %s" % str(gpu.MAX_BLOCK_DIM_X))
  print("maxBlockDimY = %s" % str(gpu.MAX_BLOCK_DIM_Y))
  print("maxBlockDimZ = %s" % str(gpu.MAX_BLOCK_DIM_Z))
  print("maxGridDimX = %s" % str(gpu.MAX_GRID_DIM_X))
  print("maxGridDimY = %s" % str(gpu.MAX_GRID_DIM_Y))
  print("maxGridDimZ = %s" % str(gpu.MAX_GRID_DIM_Z))
  print("maxSharedMemoryPerBlock = %s" % 
  str(gpu.MAX_SHARED_MEMORY_PER_BLOCK))
  print("asyncEngineCount = %s" % str(gpu.ASYNC_ENGINE_COUNT))
  print("canMapHostMemory = %s" % str(gpu.CAN_MAP_HOST_MEMORY))
  print("multiProcessorCount = %s" % str(gpu.MULTIPROCESSOR_COUNT))
  print("warpSize = %s" % str(gpu.WARP_SIZE))
  print("unifiedAddressing = %s" % str(gpu.UNIFIED_ADDRESSING))
  print("pciBusID = %s" % str(gpu.PCI_BUS_ID))
  print("pciDeviceID = %s" % str(gpu.PCI_DEVICE_ID))
  </code></pre>
  
  <p>and the output is:</p>
  
  <p>name = b'GeForce GTX 1050 Ti'</p>
  
  <p>maxThreadsPerBlock = 1024</p>
  
  <p>maxBlockDimX = 1024</p>
  
  <p>maxBlockDimY = 1024</p>
  
  <p>maxBlockDimZ = 64</p>
  
  <p>maxGridDimX = 2147483647</p>
  
  <p>maxGridDimY = 65535</p>
  
  <p>maxGridDimZ = 65535</p>
  
  <p>maxSharedMemoryPerBlock = 49152</p>
  
  <p>asyncEngineCount = 2</p>
  
  <p>canMapHostMemory = 1</p>
  
  <p>multiProcessorCount = 6</p>
  
  <p>warpSize = 32</p>
  
  <p>unifiedAddressing = 1</p>
  
  <p>pciBusID = 3</p>
  
  <p>pciDeviceID = 0</p>