<p>gufunc Numba发射和运行如此缓慢的原因在剖析时立即变得显而易见(Numba 0.38.1与cuda8.0)</p>
<pre><code>==24691== Profiling application: python slowvec.py
==24691== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput Device Context Stream Name
271.33ms 1.2800us - - - - - 8B 5.9605MB/s GeForce GTX 970 1 7 [CUDA memcpy HtoD]
271.65ms 14.591us - - - - - 156.25KB 10.213GB/s GeForce GTX 970 1 7 [CUDA memcpy HtoD]
272.09ms 2.5868ms - - - - - 15.259MB 5.7605GB/s GeForce GTX 970 1 7 [CUDA memcpy HtoD]
274.98ms 992ns - - - - - 8B 7.6909MB/s GeForce GTX 970 1 7 [CUDA memcpy HtoD]
275.17ms 640ns - - - - - 8B 11.921MB/s GeForce GTX 970 1 7 [CUDA memcpy HtoD]
276.33ms 657.28ms (1 1 1) (64 1 1) 40 0B 0B - - GeForce GTX 970 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>) [38]
933.62ms 3.5128ms - - - - - 15.259MB 4.2419GB/s GeForce GTX 970 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.
</code></pre>
<p>运行代码的最终内核启动使用64个线程的单个块。在一个GPU上,理论上每mp2048个线程,23mp,这意味着你的GPU理论处理能力的99.9%没有被使用。这看起来像是numba开发人员的一个荒谬的设计选择,如果你被它阻碍了,我会把它作为一个bug来报告(看起来你是这样)。在</p>
<p>显而易见的解决方案是将函数重写为cudapython内核方言中的<code>@cuda.jit</code>函数,并显式地控制执行参数。这样,您至少可以确保代码运行时有足够的线程来潜在地使用您的硬件的所有容量。它仍然是一个内存受限的操作,因此您可以实现的加速可能会被限制在远低于您的GPU的内存带宽与CPU的比率。而且,这可能不足以分摊主机到设备内存传输的成本,因此在最好的情况下,性能可能不会提高,尽管这还远远不够。在</p>
<p>简而言之,要小心automagic编译器生成的并行性的危险。。。。在</p>
<p>Postscript补充说,我设法弄清楚了如何获得numba发出的代码的PTX,并且足以说明这绝对是一个废话(而且我不能真正发布所有这些东西):</p>
^{pr2}$
<p>所有这些整型运算都只执行一个双精度乘法!在</p>