为什么 numba cuda 召回几次后运行缓慢？

Question

我正在尝试如何在 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))

输出是：

name = b'GeForce GTX 1050 Ti'

maxThreadsPerBlock = 1024

maxBlockDimX = 1024

maxBlockDimY = 1024

maxBlockDimZ = 64

maxGridDimX = 2147483647

maxGridDimY = 65535

maxGridDimZ = 65535

maxSharedMemoryPerBlock = 49152

asyncEngineCount = 2

canMapHostMemory = 1

multiProcessorCount = 6

warpSize = 32

统一地址 = 1

pciBusID = 3

pciDeviceID = 0

Answer 1

这是由与 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() 所以如果你修改你的for循环如下：

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

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