我正在尝试如何在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循环运行速度非常快。但是在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))
输出是:
Lick = b'GTX 1050'
maxThreadsPerBlock = 1024
maxBlockDimX = 1024
maxBlockDimY = 1024
maxBlockDimZ = 64
maxGridDimX = 2147483647
maxGridDimY = 65535
maxGridDimZ = 65535
maxSharedMemoryPerBlock = 49152
asyncEngineCount = 2
canMapHostMemory = 1
multiProcessorCount = 6
warpSize = 32
unifiedAddressing = 1
pciBusID = 3
pciDeviceID = 0
这是由与GPU内核启动相关联的异步启动队列引起的。
当你告诉numba提交GPU内核时:
matmul[(256,256),(16,16)](a1,b1,c1)
该请求进入队列,即使GPU内核尚未完成甚至尚未启动,发出该内核调用的CPU线程(即python)也可以继续。
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工作完成的实际速率继续进行,而不是“队列填充”速率。