首先,您展示的基本操作是将两个矩阵传输到GPU,执行一些逐元素乘法以产生第三个数组,并将该第三个数组传回主机。可能可以实现numba/cuda guvectorize(或cuda.jit kernel)实现,它可能比朴素的串行python实现运行得更快,但我怀疑无法超过编写良好的主机代码(例如使用某些并行化方法,如guvectorize)来完成相同的操作的性能。这是因为每个字节在主机和设备之间传输的算术强度太低。这个操作太简单了。
其次,我认为重要的是要先了解numba中
vectorize
和
guvectorize
的用途。基本原则是从“工作人员将做什么?”的角度编写ufunc定义,然后允许numba从中启动多个工作人员。指导numba启动多个工作人员的方法是
传递一个大于您给出的签名的数据集。需要注意的是,
numba不知道如何并行化ufunc定义内部的for循环。它通过将您的ufunc定义在并行工作人员之间运行,并在每个工作人员处理“切片”数据时运行
整个ufunc定义来获得并行“强度”。如果需要更多阅读,我也在
这里涵盖了一些相关内容。
在您的实现中存在一个问题,即您编写了一个将整个输入数据集映射到单个工作程序(和ufunc)的签名。正如@talonmies所示,您的基础内核正在使用总共64个线程/工作程序(这对于GPU来说太小而不足为奇,即使与上述算术强度无关)。但是我怀疑事实上64实际上只是numba最小线程块大小,而实际上只有1个线程在该线程块中执行任何有用的计算工作。这一个线程以串行方式执行您的整个ufunc,包括所有for循环。
显然,这不是任何人对vectorize
或guvectorize
合理使用的意图。
让我们回顾一下你要做什么。最终,您的ufunc想要将一个数组中的输入值乘以另一个数组中的输入值,并将结果存储在第三个数组中。我们希望重复执行这个过程很多次。如果所有3个数组大小相同,我们实际上可以使用
vectorize
来实现这一点,甚至不必诉诸更复杂的
guvectorize
。让我们比较这种方法和您的原始方法,着重于CUDA内核执行。这是一个完整的示例,其中t14.py是您的原始代码,使用分析器运行,而t15.py是它的
vectorize
版本,承认我们已经改变了
multBy
数组的大小来匹配
cv
和
discount
:
$ nvprof
==4145== NVPROF is profiling process 4145, command: python t14.py
Function: discount factor cumVest duration (seconds):1.24354910851
==4145== Profiling application: python t14.py
==4145== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
312.36ms 1.2160us - - - - - 8B 6.2742MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
312.81ms 27.392us - - - - - 156.25KB 5.4400GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
313.52ms 5.8696ms - - - - - 15.259MB 2.5387GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
319.74ms 1.0880us - - - - - 8B 7.0123MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
319.93ms 896ns - - - - - 8B 8.5149MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
321.40ms 1.22538s (1 1 1) (64 1 1) 63 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__gufunc_cVestDiscount$242(Array<__int64, int=1, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>, Array<double, int=4, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<double, int=4, A, mutable, aligned>) [37]
1.54678s 7.1816ms - - - - - 15.259MB 2.0749GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$ cat t15.py
import numpy as np
from numba import guvectorize,vectorize
import time
from timeit import default_timer as timer
@vectorize(['float64(float64, float64)'], target='cuda')
def cVestDiscount (a, b):
return a * b
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
multBy = np.full_like(discount, 1)
cv = np.empty_like(discount)
func_start = timer()
cv = cVestDiscount(multBy, discount)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
$ nvprof
==4167== NVPROF is profiling process 4167, command: python t15.py
Function: discount factor cumVest duration (seconds):0.37507891655
==4167== Profiling application: python t15.py
==4167== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
193.92ms 6.2729ms - - - - - 15.259MB 2.3755GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
201.09ms 5.7101ms - - - - - 15.259MB 2.6096GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
364.92ms 842.49us (15625 1 1) (128 1 1) 13 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__vectorized_cVestDiscount$242(Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>) [31]
365.77ms 7.1528ms - - - - - 15.259MB 2.0833GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$
我们看到你的应用程序报告了大约1.244秒的运行时间,而向量化版本报告了大约0.375秒的运行时间。但这两个数字都包含Python开销。如果我们查看分析器中生成的CUDA内核持续时间,差异甚至更加明显。我们看到原始内核花费了大约1.225秒,而向量化内核在大约842微秒(即不到1毫秒)内执行。我们还注意到,计算内核时间现在比将3个数组传输到/从GPU所需的时间要小得多(总共需要约20毫秒),并且我们注意到内核尺寸现在为15625个128线程块,每个块有一个总线程/工作数为2000000,正好匹配要执行的总乘法操作数量,并且比原始代码中可怜的64个线程(可能只有1个线程)要多得多。
考虑到上述“向量化”方法的简单性,如果您真正想做的是逐元素乘法,则可以考虑复制“multBy”,使其在尺寸上与其他两个数组匹配,并完成它。
但问题仍然存在:如何处理不同的输入数组大小,就像原始问题一样?为此,我认为我们需要使用
guvectorize
(或者,正如@talonmies所指出的那样,编写自己的
@cuda.jit
内核,这可能是最好的建议,尽管可能没有这些方法中的任何一种能够克服数据传输到/从设备的开销,正如已经提到的那样)。
为了使用guvectorize
解决这个问题,我们需要更加仔细地考虑已经提到的“切片”概念。让我们重新编写您的guvectorize
内核,以便它仅对整体数据的“切片”进行操作,然后允许guvectorize
启动函数为其启动多个工作器来处理它,每个切片一个工作器。
在CUDA中,我们喜欢有很多的工作者;实际上你不能有太多。因此,这将影响我们如何“切割”数组,以便让多个工作者可以行动。如果我们沿着第三(最后一个)维度切片,我们只有5个切片可用,所以最多只能有5个工作者。同样,如果我们沿着第一或
countRow
维度切片,我们将有100个切片,所以最多只能有100个工作者。理想情况下,我们应该沿着第二或
countCol
维度进行切片。然而为了简单起见,我将沿着第一或
countRow
维度进行切片。这显然是不优化的,但请参见下面的示例,了解如何处理按第二维度切片问题。按照第一维进行切片意味着我们将从guvectorize内核中去掉第一个for循环,并允许ufunc系统沿着那个维度并行化(基于我们传递的数组的大小)。代码可能看起来像这样:
$ cat t16.py
import numpy as np
from numba import guvectorize
import time
from timeit import default_timer as timer
@guvectorize(['void(float64[:,:], float64[:,:], int64, int64, float64[:,:])'], '(m,o),(m,o),(),() -> (m,o)', target='cuda', nopython=True)
def cVestDiscount (multBy, discount, n, countCol, cv):
for ID in range(0,countCol):
for num in range(0,n):
cv[ID][num] = multBy[ID][num] * discount[ID][num]
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
n = np.int64(5)
countCol = np.int64(4000)
cv = np.zeros(shape=(100,4000,5), dtype=np.float64)
func_start = timer()
cv = cVestDiscount(multBy, discount, n, countCol, cv)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
$ nvprof --print-gpu-trace python t16.py
==4275== NVPROF is profiling process 4275, command: python t16.py
Function: discount factor cumVest duration (seconds):0.0670170783997
==4275== Profiling application: python t16.py
==4275== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
307.05ms 27.392us - - - - - 156.25KB 5.4400GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
307.79ms 5.9293ms - - - - - 15.259MB 2.5131GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
314.34ms 1.3440us - - - - - 8B 5.6766MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
314.54ms 896ns - - - - - 8B 8.5149MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
317.27ms 47.398ms (2 1 1) (64 1 1) 63 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__gufunc_cVestDiscount$242(Array<double, int=3, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>) [35]
364.67ms 7.3799ms - - - - - 15.259MB 2.0192GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$
观察:
这些代码更改与删除 countCol
参数、从 guvectorize 内核中删除第一个 for 循环以及相应地更改函数签名有关。我们还在签名中将我们的三维函数修改为二维。毕竟,我们正在获取三维数据的二维“切片”,让每个工作进程处理一个切片。
分析器报告的内核维度现在是 2 个块而不是 1 个。这是有道理的,因为在原始实现中,实际上只有 1 个“切片”被呈现,因此只需要 1 个工作进程,因此只有 1 个线程(但 numba 启动了 64 个线程块)。在这个实现中,有 100 个切片,numba 选择启动 2 个包含 64 个工作进程/线程的线程块,以提供所需的 100 个工作进程/线程。
分析器报告的内核性能为 47.4ms,介于原始版本(约为 1.224s)和大规模并行的 vectorize
版本(约为 0.001s)之间。因此,从 1 个到 100 个工作进程的转变使事情加快了很多,但还有更多的性能提升可能。如果您找出如何在 countCol
维度上切片,您可能可以在性能方面更接近 vectorize
版本(请参见下文)。请注意,我们现在的状态(约 47ms)和向量化版本(约 1ms)之间的差异足以弥补传输稍大的 multBy
矩阵的附加传输成本(~5ms 或更少),以便实现 vectorize
的简单性。
关于Python计时的一些补充说明:我认为Python编译原始版本、向量化版本和guvectorize改进版本所需内核的确切行为是不同的。如果我们修改t15.py代码以运行“预热”运行,那么至少Python计时在趋势上与总体墙时和仅内核计时保持一致:
$ cat t15.py
import numpy as np
from numba import guvectorize,vectorize
import time
from timeit import default_timer as timer
@vectorize(['float64(float64, float64)'], target='cuda')
def cVestDiscount (a, b):
return a * b
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
multBy = np.full_like(discount, 1)
cv = np.empty_like(discount)
cv = cVestDiscount(multBy, discount)
func_start = timer()
cv = cVestDiscount(multBy, discount)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
[bob@cluster2 python]$ time python t14.py
Function: discount factor cumVest duration (seconds):1.24376320839
real 0m2.522s
user 0m1.572s
sys 0m0.809s
$ time python t15.py
Function: discount factor cumVest duration (seconds):0.0228319168091
real 0m1.050s
user 0m0.473s
sys 0m0.445s
$ time python t16.py
Function: discount factor cumVest duration (seconds):0.0665760040283
real 0m1.252s
user 0m0.680s
sys 0m0.441s
$
现在,回答评论中的一个问题,有效地说:“我应该如何重构问题以沿着4000(
countCol
或“中间”)维度切片?”我们可以参考切片第一维的方法。一种可能的方法是重新排列数组的形状,使4000维成为第一维,然后移除它,类似于我们在之前处理
guvectorize
时所做的。下面是一个实例:
$ cat t17.py
import numpy as np
from numba import guvectorize
import time
from timeit import default_timer as timer
@guvectorize(['void(int64, float64[:], float64[:,:], int64, float64[:,:])'], '(),(o),(m,o),() -> (m,o)', target='cuda', nopython=True)
def cVestDiscount (countCol, multBy, discount, n, cv):
for ID in range(0,countCol):
for num in range(0,n):
cv[ID][num] = multBy[num] * discount[ID][num]
countRow = np.int64(100)
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(4000,100,5))
n = np.int64(5)
countCol = np.int64(4000)
cv = np.zeros(shape=(4000,100,5), dtype=np.float64)
func_start = timer()
cv = cVestDiscount(countRow, multBy, discount, n, cv)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
[bob@cluster2 python]$ python t17.py
Function: discount factor cumVest duration (seconds):0.0266749858856
$ nvprof --print-gpu-trace python t17.py
==8544== NVPROF is profiling process 8544, command: python t17.py
Function: discount factor cumVest duration (seconds):0.0268459320068
==8544== Profiling application: python t17.py
==8544== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
304.92ms 1.1840us - - - - - 8B 6.4437MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
305.36ms 27.392us - - - - - 156.25KB 5.4400GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
306.08ms 6.0208ms - - - - - 15.259MB 2.4749GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
312.44ms 1.0880us - - - - - 8B 7.0123MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
313.59ms 8.9961ms (63 1 1) (64 1 1) 63 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__gufunc_cVestDiscount$242(Array<__int64, int=1, A, mutable, aligned>, Array<double, int=2, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>) [35]
322.59ms 7.2772ms - - - - - 15.259MB 2.0476GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$
相对可预测的是,我们观察到当我们将数据分成100个worker时,执行时间从约47毫秒降至当我们将数据分成4000个worker时的约9毫秒。同样地,我们观察到numba选择启动63个块,每个块有64个线程,总共4032个线程,以处理这个“slicing”所需的4000个worker。
仍然不如速度约为1毫秒的
vectorize
内核(其具有更多可用的并行“slices”),但比原始问题中提出的约1.2秒的内核快得多。即使考虑到所有Python开销,Python代码的总体墙时也快了约2倍。
最后观察一下我之前所做的陈述(与评论和其他答案中所做的陈述类似):
现在我们有方便的测试用例可以在t16.py或t17.py中使用来进行测试。为了简单起见,我将选择t16.py。我们可以通过从
guvectorize
ufunc中删除目标指定来“将其转换回CPU代码”:
$ cat t16a.py
import numpy as np
from numba import guvectorize
import time
from timeit import default_timer as timer
@guvectorize(['void(float64[:,:], float64[:,:], int64, int64, float64[:,:])'], '(m,o),(m,o),(),() -> (m,o)')
def cVestDiscount (multBy, discount, n, countCol, cv):
for ID in range(0,countCol):
for num in range(0,n):
cv[ID][num] = multBy[ID][num] * discount[ID][num]
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
n = np.int64(5)
countCol = np.int64(4000)
cv = np.zeros(shape=(100,4000,5), dtype=np.float64)
func_start = timer()
cv = cVestDiscount(multBy, discount, n, countCol, cv)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
$ time python t16a.py
Function: discount factor cumVest duration (seconds):0.00657796859741
real 0m0.528s
user 0m0.474s
sys 0m0.047s
$
因此,我们可以看到这个仅使用CPU的版本大约在6毫秒内运行函数,并且它没有GPU“开销”,例如CUDA初始化和数据传输到/从GPU。总体的墙时也是我们最好的测量,大约为0.5秒,而我们最好的GPU情况下则为1.0秒左右。因此,由于每个数据传输字节的低算术强度,这个特定问题可能不适合GPU计算。