我的计算密集型 Numba 函数在 GPU 上的运行速度比在 CPU 上慢 10 倍.我错过了什么吗?



我测试了Numba 的性能,并以两种方式尝试了一个虚拟但计算密集型的功能:在启用了并行性的 CPU 上(和范围(,以及在 GPU 上,我看到它在运行时占据了 100% 的 GPU。两者都运行良好,但 CPU 完成所需的时间减少了 10 倍。在这种情况下,我要求GPU更快,即使我的GPU不是很强大(Geforce 1050 ti(并且我的CPU很强大(Threadripper 3970x(。

这是我的 CPU 基准测试:

from numba import *
import numpy as np
from numba import cuda
import time

def benchmark():
input_list = np.random.randint(10, size=320000)
out = np.zeros(len(input_list))
cpu_run_test(input_list, out)
print('Result size: ' + str(len(out)) + ' ' + str(out))

@njit(parallel=True, fastmath=True, nogil=True)
def cpu_run_test(input_list, out):
for step in range(len(input_list)):
for j in range(10):
count = 0
for item2 in input_list:
if input_list[step] == item2:
count = count + 1
out[step] = count

if __name__ == '__main__':
import timeit
print(timeit.timeit("benchmark()", setup="from __main__ import benchmark", number=1))

这是我的 GPU 基准测试(相同的计算,只是分区工作方式不同,以适当地利用 GPU 块和线程:

from numba import *
import numpy as np
from numba import cuda
import time

def benchmark():
for xx in range(1):
new_array_duration = time.time()
input_list = np.random.randint(10, size=320000)
new_array_duration = time.time() - new_array_duration
print('New array duration: ' + str(new_array_duration))
to_device_duration = time.time()
d_array = cuda.to_device(input_list)
to_device_duration = time.time() - to_device_duration
print('To device duration: ' + str(to_device_duration))
kernel_duration = time.time()
run_test[16, 768](d_array)
kernel_duration = time.time() - kernel_duration
print('Kernel duration: ' + str(kernel_duration))
to_host_duration = time.time()
out = d_array.copy_to_host()
to_host_duration = time.time() - to_host_duration
print('To host duration: ' + str(to_host_duration))

@cuda.jit(fastmath=True)
def run_test(d_array):
array_slice_len = len(d_array) / cuda.blockDim.x
slice_start = (cuda.threadIdx.x * (cuda.blockIdx.x + 1)) * array_slice_len
for step in prange(slice_start, slice_start + array_slice_len):
if step > len(d_array) - 1:
return
for j in range(10):
count = 0
for item2 in d_array:
if d_array[step] == item2:
count = count + 1
d_array[step] = count

if __name__ == '__main__':
import timeit
# make_multithread(benchmark, 64)
print(timeit.timeit("benchmark()", setup="from __main__ import benchmark", number=1))

任何人都可以复制和粘贴并运行其中任何一个。我使用的是Linux Mint 20,Python 3.7,最新的Numba(0.51(和最新的cudatoolkit。

结果如下:

中央处理器:15.20秒

图形处理器:145.54 秒

这是正确的还是我缺少一些优化以使 GPU 代码运行得更快?

我错过了什么?

不一定是答案,但评论太长了。

我会尝试一些事情;改变块和线程,看看会发生什么。还利用 GPU 上的本地、共享和全局内存。

全局内存是唯一可以由主机设置的内存,并且所有线程都可以访问。它也是访问速度最慢的内存类型,因此您通常希望最大程度地减少传输量,并在可能的情况下使用内存合并。

共享内存是可用于一个块中所有线程的内存,比全局内存快。本地内存是最快的,仅对一个线程可见。

在 GPU 上启动 prange 实际上是使用 GPU 线程吗?因为正常(非 python(您不必在 GPU 上并行调用东西。也许该循环以串行方式运行?

您还应该尝试找到一种不涉及分支的算法(如果有的话(。如果我正确理解您的代码,您有一个包含 320.000 个项目的列表,其中包含 [0;10(. 对于每个项目,计算该项目出现的频率,并将该项目替换为该项目的计数。和 10 倍的良好衡量。这将导致结果集的项在值 [0;10( 和介于 [0;320.000(,但一般可能在 32.000 左右。

为了通过消除这种令人不快的振荡来更轻松地转换算法,将列表大小增加 10 倍并摆脱循环也应该这样做。

由于我不会说python,这里有一些非分支矢量化实现的伪代码。虽然我不知道这是否会更快,但我会尝试。

{l} local memory
{s} shared memory
{g} global memory
Host:
input_list{g} = (0..9)[3.200.000]
accumulator{g} = (0..0)[10]
run _1
run _2

Device Kernel _1:   
let blocksize = get_blocksize()
let threadId = get_threadId()
let globalId = get_globalId()
let localSum{s} = arr[10]
let localCopy{s} = arr[blocksize.x]
if(globalId.x < size(input_list)) {
localCopy[threadId] = input_list[globalId]

localSum[localCopy[threadId]]++
}
if(threadId.x < 10) {
accumulator[threadId] += localSum[threadId]
}

Device Kernel _2:
let blocksize = get_blocksize()
let threadId = get_threadId()
let globalId = get_globalId()
let localAcc{l} = arr[10]
for(i = 0 to 10) {
localAcc[i] = accumulator[i]
}
if(globalId < size(input_list)) {
let val = input_list[globalId]
input_list[globalId] = localAcc[val]
}

相关内容

最新更新