Numba 中的组合向量化函数

Question

我正在使用 Numba（版本 0.37.0）优化 GPU 代码。我想使用组合矢量化函数（使用 Numba 的 @vectorize 装饰器）。

导入和数据：

import numpy as np
from math import sqrt
from numba import vectorize, guvectorize

angles = np.random.uniform(-np.pi, np.pi, 10)
coords = np.stack([np.cos(angles), np.sin(angles)], axis=1)

这按预期工作：

@guvectorize(['(float32[:], float32[:])'], '(i)->()', target='cuda')
def l2_norm(vec, out):
    acc = 0.0
    for value in vec:
        acc += value**2
    out[0] = sqrt(acc)

l2_norm(coords)

输出：

array([1., 1., 1., 1., 1., 1., 1., 1., 1., 1.], dtype=float32)

但我想通过调用另一个向量化函数来避免在 "l2_norm" 中使用那个 "for"。

我试过这个：

@vectorize(["float32(float32)"], target="cuda")
def power(value):
    return value**2

@guvectorize(['(float32[:], float32[:])'], '(i)->()', target='cuda')
def l2_norm_power(vec, out):
    acc = 0.0
    acc = power(vec)
    acc = acc.sum()
    out[0] = sqrt(acc)

l2_norm_power(coords)

但会引发 TypingError：

TypingError: Failed at nopython (nopython frontend)
Untyped global name 'power': cannot determine Numba type of <class  
'numba.cuda.dispatcher.CUDAUFuncDispatcher'>

知道如何执行此组合吗？

关于使用 Numba 优化 l2_norm 的其他方法有什么建议吗？

Answer 1

我认为你只能从其他 cuda 函数中调用 device=True 函数：

3.13.2. Example: Calling Device Functions

All CUDA ufunc kernels have the ability to call other CUDA device functions:

 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)

请注意，您可以使用 device:

调用 cuda.jit 函数

@cuda.jit(device=True)
def sum_of_squares(arr):
    acc = 0
    for item in arr:
        acc += item ** 2
    return acc

@nb.guvectorize(['(float32[:], float32[:])'], '(i)->()', target='cuda')
def l2_norm_power(vec, out):
    acc = sum_of_squares(vec)
    out[0] = sqrt(acc)

l2_norm_power(coords)

但这可能违背了拆分它的目的。

由于 numba.vectorize 不支持 device，因此这些功能是不可能的。但这是一件好事，因为 vectorize 分配一个数组来放入值，这是一个不必要的中间数组，并且在 GPU 上分配数组也非常低效（在 numba 中被禁止）：

3.5.5. Numpy support

Due to the CUDA programming model, dynamic memory allocation inside a kernel is inefficient and is often not needed. Numba disallows any memory allocating features. This disables a large number of NumPy APIs. For best performance, users should write code such that each thread is dealing with a single element at a time.

考虑到所有这些，我会简单地使用您原来的方法：

@guvectorize(['(float32[:], float32[:])'], '(i)->()', target='cuda')
def l2_norm(vec, out):
    acc = 0.0
    for value in vec:
        acc += value**2
    out[0] = sqrt(acc)

l2_norm(coords)

Numba 中的组合向量化函数

Combined vectorized functions in Numba

gpu

vectorization

dispatch

numba

numpy-ufunc