gpu - 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: 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，因此这些功能是不可能的。但这是一件好事，因为向量化分配一个数组来放入值，这是一个不必要的中间数组，并且在 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)