Bry*_*oze 3 python performance cuda numba nvprof
值得注意的细节
首先是尝试将 guvectorize 与以下函数一起使用。我正在传递一堆 numpy 数组并尝试使用它们在两个数组之间进行乘法运算。如果使用 cuda 以外的目标运行，则此方法有效。但是，当切换到 cuda 时，会导致未知错误：
文件“C:\ProgramData\Anaconda3\lib\site-packages\numba\cuda\decorators.py”，>第 82 行，在 jitwrapper 中 debug=debug）
TypeError: init () 得到了一个意外的关键字参数 'debug'
在按照我可以从这个错误中找到的所有内容之后，我除了死胡同什么也没找到。我猜这是一个非常简单的修复，我完全没有，但哦，好吧。还应该说，此错误仅在运行一次并由于内存过载而崩溃后才会发生。
os.environ["NUMBA_ENABLE_CUDASIM"] = "1"
os.environ["CUDA_VISIBLE_DEVICES"] = "10DE 1B06 63933842"
...
所有的数组都是 numpy
@guvectorize(['void(int64, float64[:,:], float64[:,:], float64[:,:,:],
int64, int64, float64[:,:,:])'], '(),(m,o),(m,o),(n,m,o),(),() -> (n,m,o)',
target='cuda', nopython=True)
def cVestDiscount (ed, orCV, vals, discount, n, rowCount, cv):
for as_of_date in range(0,ed):
for ID in range(0,rowCount):
for num in range(0,n):
cv[as_of_date][ID][num] = orCV[ID][num] * discount[as_of_date][ID][num]
尝试在命令行中使用 nvprofiler 运行代码会导致以下错误：
警告：当前配置不支持统一内存分析，因为在此 ?multi-GPU 设置上检测到一对没有对等支持的设备。当对等映射不可用时，系统回退到使用零拷贝内存。它会导致访问统一内存的内核运行速度变慢。可以在以下位置找到更多详细信息：http : //docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-managed-memory
我意识到我使用的是支持 SLI 的显卡（两张卡都是相同的，evga gtx 1080ti，并且具有相同的设备 ID），所以我禁用了 SLI 并添加了“CUDA_VISIBLE_DEVICES”行以尝试限制其他一张卡，但我留下了相同的结果。
我仍然可以使用 nvprof 运行代码，但与 njit(parallel=True) 和 prange 相比，cuda 函数很慢。通过使用较小的数据大小，我们可以运行代码，但它比 target='parallel' 和 target='cpu' 慢。
为什么 cuda 这么慢，这些错误是什么意思？
谢谢您的帮助！
编辑：这是代码的一个工作示例：
import numpy as np
from numba import guvectorize
import time
from timeit import default_timer as timer
@guvectorize(['void(int64, float64[:,:], float64[:,:,:], int64, int64, float64[:,:,:])'], '(),(m,o),(n,m,o),(),() -> (n,m,o)', target='cuda', nopython=True)
def cVestDiscount (countRow, multBy, discount, n, countCol, cv):
for as_of_date in range(0,countRow):
for ID in range(0,countCol):
for num in range(0,n):
cv[as_of_date][ID][num] = multBy[ID][num] * discount[as_of_date][ID][num]
countRow = np.int64(100)
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(countRow, multBy, discount, n, countCol, cv)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
我可以使用 gtx 1080ti 在 cuda 中运行代码，但是，它比并行或 CPU 运行要慢得多。我查看了与 guvectorize 相关的其他帖子，但它们都没有帮助我了解在 guvectorize 中运行什么是最佳的，什么不是最佳的。有什么方法可以使此代码“对 cuda 友好”，或者只是在数组之间进行乘法太简单而无法看到任何好处？
首先，您展示的基本操作是取两个矩阵，将它们传输到 GPU，进行一些元素乘法以生成第三个数组，然后将该第三个数组传回主机。
有可能制作一个 numba/cuda guvectorize（或 cuda.jit 内核）实现，它可能比简单的串行 python 实现运行得更快，但我怀疑它是否有可能超过编写良好的主机代码的性能（例如使用一些并行化方法，例如 guvectorize）来做同样的事情。这是因为主机和设备之间传输的每个字节的算术强度太低了。这个操作太简单了。
其次，我相信，从了解 numbavectorize和guvectorize打算做什么开始是很重要的。基本原则是从“worker会做什么”的角度来写ufunc定义。然后允许 numba 从中启动多个工人。您指示 numba 启动多个 worker 的方法是传递一个大于您提供的签名的数据集。应该注意的是，numba 不知道如何在 ufunc 定义中并行化 for 循环。它通过获取您的 ufunc 定义并在并行工作人员之间运行它来获得并行“强度”，其中每个工作人员处理数据的“切片”，但运行您的整个 ufunc 定义在那片上。作为一些额外的阅读，我也在这里介绍了一些这方面的内容。
因此，我们在您的实现中遇到的一个问题是您编写了一个签名（和 ufunc），它将整个输入数据集映射到单个工作人员。正如@talonmies 所示，您的底层内核总共有 64 个线程/工作线程（即使除了上述关于算术强度的陈述之外，这对于 GPU 来说也太小了），但我怀疑实际上64 实际上只是 numba 的最小线程块大小，实际上该线程块中只有 1 个线程在执行任何有用的计算工作。该线程以串行方式执行您的整个 ufunc，包括所有 for 循环。
这显然不是任何人想要合理使用vectorizeor 的目的guvectorize。
因此，让我们重新审视您正在尝试做的事情。最终，您的 ufunc 想要将一个数组的输入值乘以另一个数组的输入值，并将结果存储在第三个数组中。我们想多次重复这个过程。如果所有 3 个数组大小都相同，我们实际上可以实现这一点，vectorize甚至不必求助于更复杂的guvectorize. 让我们将该方法与您的原始方法进行比较，重点关注 CUDA 内核执行。这是一个有效的示例，其中 t14.py 是您的原始代码，使用分析器运行，而 t15.py 是vectorize它的一个版本，确认我们已更改您的multBy数组大小以匹配cv和discount：
$ nvprof --print-gpu-trace python t14.py
==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 --print-gpu-trace python t15.py
==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 秒，而 vectorize 内核在大约 842 微秒（即小于 1 毫秒）内执行。我们还注意到计算内核时间现在比将 3 个数组传输到/从 GPU 所需的时间小得多（总共需要大约 20 毫秒），我们注意到内核尺寸现在是 15625 个块，每块 128每个线程的总线程/工作线程数为 2000000，完全匹配要完成的乘法运算的总数，
鉴于上述vectorize方法的简单性，如果您真正想要做的是这种逐元素乘法，那么您可以考虑只进行复制，multBy以便在维度上匹配其他两个数组，然后就可以完成了。
但问题仍然存在：如何处理不同的输入数组大小，就像在原始问题中一样？为此，我认为我们需要去guvectorize（或者，正如@talonmies 指出的那样，编写您自己的@cuda.jit内核，这可能是最好的建议，尽管这些方法都无法克服向/从设备传输数据的开销，如前所述）。
为了解决这个问题guvectorize，我们需要更仔细地考虑已经提到的“切片”概念。让我们重新编写您的guvectorize内核，使其仅对整体数据的“切片”进行操作，然后允许guvectorize启动函数启动多个工作程序来处理它，每个切片一个工作程序。
在 CUDA 中，我们喜欢有很多工人；你真的不能有太多。所以这将影响我们如何“切片”我们的数组，以便为多个工人提供行动的可能性。如果我们沿着第三个（最后一个n）维度切片，我们将只有 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 循环以及对函数签名进行适当更改以反映这一点有关。我们还将签名中的3 维函数修改为二维函数。毕竟，我们正在对 3 维数据进行二维“切片”，并让每个工人在切片上工作。
分析器报告的内核尺寸现在是 2 个块而不是 1 个。这是有道理的，因为在最初的实现中，实际上只提供了 1 个“切片”，因此需要 1 个工人，因此需要 1 个线程（但 numba 旋转1 个线程块（共 64 个线程））。在这个实现中，有 100 个切片，numba 选择启动 64 个工人/线程的 2 个线程块，以提供所需的 100 个工人/线程。
分析器报告的内核性能为 47.4 毫秒，现在介于原始版本（~1.224 秒）和大规模并行vectorize版本（~0.001 秒）之间。因此，从 1 名工作人员增加到 100 名工作人员大大加快了速度，但可能会有更多的性能提升。如果您弄清楚如何在countCol维度上切片，您可能会更接近vectorize版本，性能方面（见下文）。请注意，我们现在所处的位置（~47ms）和矢量化版本（~1ms）之间的差异足以弥补将稍大的multBy矩阵传输到设备的额外传输成本（~5ms 或更少），以方便vectorize简单。
关于 python 计时的一些附加评论：我相信 python 如何为原始、vectorize 和 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)
#warm-up run
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 个工作人员时的约 47 毫秒下降到分割成 4000 个工作人员时的约 9 毫秒。类似地，我们观察到 numba 选择启动 63 个块，每个块 64 个线程，总共 4032 个线程，以处理此“切片”所需的 4000 个工人。
仍然不如 ~1msvectorize内核（它为工作人员提供更多可用的并行“切片”）快，但比原始问题中提出的 ~1.2s 内核快得多。即使包含所有 Python 开销，python 代码的整体挂载时间也快了大约 2 倍。
最后，让我们回顾一下我之前发表的声明（类似于评论和其他答案中的声明）：
“我怀疑是否有可能超过编写良好的主机代码的性能（例如使用一些并行化方法，例如 guvectorize）来做同样的事情。”
我们现在在 t16.py 或 t17.py 中有方便的测试用例，我们可以用来测试它。为简单起见，我将选择 t16.py。我们可以通过从guvectorizeufunc 中删除目标名称来“将其转换回 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 复制数据。整体walltime 也是我们的最佳测量值，约为 0.5 秒，而我们最好的 GPU 情况下约为 1.0 秒。所以这个特殊的问题，由于每字节数据传输的算术强度低，可能不太适合 GPU 计算。
| 归档时间： |
|
| 查看次数： |
1989 次 |
| 最近记录： |