3.13。 CUDA Ufuncs 和广义 Ufuncs

原文： http://numba.pydata.org/numba-doc/latest/cuda/ufunc.html

此页面描述了类似 CUDA ufunc 的对象。

为了支持 CUDA 程序的编程模式，CUDA Vectorize 和 GUVectorize 不能生成传统的 ufunc。相反，返回类似 ufunc 的对象。此对象是一个非常模拟但与常规 NumPy ufunc 不完全兼容的对象。 CUDA ufunc 增加了对传递设备内阵列（已在 GPU 设备上）的支持，以减少 PCI-express 总线上的流量。它还接受<cite>流</cite>关键字以在异步模式下启动。

3.13.1。示例：基本示例

import math
from numba import vectorize, cuda
import numpy as np
@vectorize(['float32(float32, float32, float32)',
            'float64(float64, float64, float64)'],
           target='cuda')
def cu_discriminant(a, b, c):
    return math.sqrt(b ** 2 - 4 * a * c)
N = 10000
dtype = np.float32
# prepare the input
A = np.array(np.random.sample(N), dtype=dtype)
B = np.array(np.random.sample(N) + 10, dtype=dtype)
C = np.array(np.random.sample(N), dtype=dtype)
D = cu_discriminant(A, B, C)
print(D)  # print result

3.13.2。示例：调用设备功能

所有 CUDA ufunc 内核都能够调用其他 CUDA 设备功能：

from numba import vectorize, cuda
# define a device function
@cuda.jit('float32(float32, float32, float32)', device=True, inline=True)
def cu_device_fn(x, y, z):
    return x ** y / z
# define a ufunc that calls our device function
@vectorize(['float32(float32, float32, float32)'], target='cuda')
def cu_ufunc(x, y, z):
    return cu_device_fn(x, y, z)

3.13.3。广义 CUDA ufuncs

可以使用 CUDA 在 GPU 上执行通用 ufunc，类似于 CUDA ufunc 功能。这可以通过以下方式完成：

from numba import guvectorize
@guvectorize(['void(float32[:,:], float32[:,:], float32[:,:])'],
             '(m,n),(n,p)->(m,p)', target='cuda')
def matmulcore(A, B, C):
    ...

有时候 gufunc 内核会使用太多的 GPU 资源，这会导致内核启动失败。用户可以通过在已编译的 gufunc 对象上设置 <cite>max_blocksize</cite> 属性来显式控制线程块的最大大小。

from numba import guvectorize
@guvectorize(..., target='cuda')
def very_complex_kernel(A, B, C):
    ...
very_complex_kernel.max_blocksize = 32  # limits to 32 threads per block