为什么 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)

前 1028 次运行的 for 循环运行非常快。但是它运行非常慢 1028th.What 正是造成这个问题的原因，我该如何解决。顺便说一句，我在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'

maxThreadsPerBlock = 1024

maxBlockDimX = 1024

maxBlockDimY = 1024

maxBlockDimZ = 64

maxGridDimX = 2147483647

maxGridDimY = 65535

maxGridDimZ = 65535

每块最大共享内存 = 49152

asyncEngineCount = 2

canMapHostMemory = 1

多处理器计数 = 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 次迭代），您开始看到 "slow down"。此后，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 工作完成的实际速率进行，而不是 "queue-filling" 速率。

为什么 numba cuda 召回几次后就运行变慢了？

Why numba cuda is running slow after recalling it several times?

python-3.x

cuda

numba

pycuda

numba-pro

为什么 numba cuda 召回几次后就运行 变慢了？

Why numba cuda is running slow after recalling it several times?

python-3.x

cuda

numba

pycuda

numba-pro

为什么 numba cuda 召回几次后就运行变慢了？