为什么numba cuda调用几次后运行速度变慢？

我正在尝试如何在 numba 中使用 cuda。然而我却遇到了与我预想不同的事情。这是我的代码

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 < C.shape[0] and j < C.shape[1]:
    tmp = 0.
    for k in range(A.shape[1]):
        tmp += A[i, k] * B[k, j]
    C[i, j] = tmp

这是我自己定义的矩阵函数，用于使用numba.cuda进行测试。在运行测试之前，我还在以下代码中加载了数组：

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)

然后我用下面的代码进行实验：

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)

for 循环在前 1028 次运行中运行得非常快。然而，在 1028 号之后，它运行得非常慢。到底是什么导致了这个问题，我该如何解决它。顺便说一句，我是在win10上运行的。

这是我从 numba.cuda 调用的 cuda 信息

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))

输出是：

名称 = b'GeForce GTX 1050 Ti'

每个块的最大线程数 = 1024

最大块尺寸X = 1024

最大块尺寸 = 1024

最大块尺寸Z = 64

maxGridDimX = 2147483647

最大网格尺寸 = 65535

最大网格尺寸Z = 65535

每个块的最大共享内存 = 49152

异步引擎计数 = 2

canMapHostMemory = 1

多处理器计数 = 6

扭曲大小 = 32

统一寻址 = 1

pci总线ID = 3

pci设备ID = 0

这是由与 GPU 内核启动相关的异步启动队列引起的。

当您告诉 numba 提交 GPU 内核时：

matmul[(256,256),(16,16)](a1,b1,c1)

该请求进入队列，发出该内核调用的 CPU 线程（即 python）可以继续，即使 GPU 内核尚未完成或什至尚未启动。

CUDA 运行时将这些请求排队并在 GPU 准备好执行更多工作时发出它们。

在 for 循环快速递增期间，您最初看到的是队列填满了工作请求。这并不代表 GPU 执行工作所需的实际时间。

最终队列填满，CUDA 运行时在内核启动时暂停 CPU 线程（即 python），直到队列槽打开。此时，for 循环可以继续进行一次迭代。正是在这一点上（可能在 1028 次迭代左右），您开始看到“速度变慢”。此后，for 循环继续进行大约GPU 内核执行和从处理队列中删除的速率。

这里没有什么需要修复的；这是预期的行为。

如果您希望 for 循环仅按照 GPU 内核实际执行的速率进行，那么您应该在 for 循环中插入一个同步函数。

例如，numba 提供numba.cuda.synchronize() https://numba.pydata.org/numba-doc/dev/cuda-reference/host.html因此，如果您按如下方式修改 for 循环：

for i in range(2000):
  matmul[(256,256),(16,16)](a1,b1,c1)
  cuda.synchronize()
  count +=1
  print(count)

您将看到 for 循环以 GPU 工作完成的实际速率进行，而不是“队列填充”速率。