CUDA Python 中支持的 Python 特性

本页列出了 CUDA Python 中支持的 Python 功能。这包括使用 @cuda.jit 编译的所有内核和设备函数，以及其他面向 CUDA GPU 的高级 Numba 装饰器。

语言

执行模型

CUDA Python 直接映射到 CUDA 的 单指令多线程 执行 (SIMT) 模型。每条指令都隐式地由多个线程并行执行。在这种执行模型下，数组表达式的作用较小，因为我们不希望多个线程执行相同的任务。相反，我们希望线程以协作的方式执行任务。

详情请参阅 CUDA 编程指南。

浮点错误模型

默认情况下，CUDA Python 内核使用 NumPy 错误模型执行。在此模型中，除以零不会引发异常，而是产生 inf、-inf 或 nan 的结果。这与正常的 Python 错误模型不同，在正常模型中，除以零会引发 ZeroDivisionError。

当启用调试时（通过向 @cuda.jit 装饰器传递 debug=True），将使用 Python 错误模型。这允许在核函数执行期间识别除以零的错误。

构造

以下 Python 结构不受支持：

• 异常处理 (try .. except, try .. finally)

• 上下文管理（with 语句）

• 推导式（列表、字典、集合或生成器推导式）

• 生成器 (任何 yield 语句)

raise 和 assert 语句是被支持的，但有以下限制：

• 它们只能在内核中使用，不能在设备函数中使用。

• 只有在向 @cuda.jit 装饰器传递 debug=True 时，它们才会生效。这与 CUDA C/C++ 中的 assert 关键字的行为类似，除非在启用设备调试的情况下编译，否则该关键字将被忽略。

支持字符串、整数和浮点数的打印，但打印是一个异步操作——为了确保在内核启动后所有输出都被打印，需要调用 numba.cuda.synchronize()。省略对 synchronize 的调用是可以接受的，但内核的输出可能会在其他后续驱动操作期间出现（例如后续内核启动、内存传输等），或者在程序执行完成之前未能出现。最多可以向 print 函数传递 32 个参数——如果传递的参数更多，则会发出格式字符串并生成警告。这是由于 CUDA 打印中的一般限制，如 CUDA C++ 编程指南中 关于打印限制的章节 所述。

递归

支持自递归设备函数，但递归调用的参数类型必须与函数初始调用的参数类型相同。例如，以下形式的递归是支持的：

@cuda.jit("int64(int64)", device=True)
def fib(n):
    if n < 2:
        return n
    return fib(n - 1) + fib(n - 2)

(fib 函数总是有一个 int64 参数), 而以下是不支持的:

# Called with x := int64, y := float64
@cuda.jit
def type_change_self(x, y):
    if x > 1 and y > 0:
        return x + type_change_self(x - y, y)
    else:
        return y

外部调用 type_change_self 提供了 (int64, float64) 参数，但内部调用使用了 (float64, float64) 参数（因为 x - y / int64 - float64 结果是一个 float64 类型）。因此，此函数不受支持。

函数之间的相互递归（例如，函数 func1() 调用 func2()，而 func2() 又调用 func1()）是不支持的。

备注

CUDA 中的调用堆栈大小通常非常有限，因此在 CUDA 设备上通过递归调用更容易使其溢出，相比之下在 CPU 上则不容易。

栈溢出将导致内核执行期间发生未指定的启动失败（ULF）。为了确定ULF是否由栈溢出引起，可以在 Compute Sanitizer 下运行程序，它会明确指出何时发生了栈溢出。

内置类型

以下内置类型的支持继承自CPU nopython模式。

• 整数

• 浮动

• 复杂

• 布尔

• 无

• 元组

• Enum, IntEnum

请参阅 nopython 内置类型。

对于NumPy数组中使用的字符序列（字节和Unicode字符串），也有一些非常有限的支持。请注意，这种支持只能用于CUDA 11.2及更高版本。

内置函数

支持以下内置函数：

• abs()

• bool

• complex

• enumerate()

• float

• int：仅支持单参数形式

• len()

• min()：仅支持多参数形式

• max(): 仅支持多参数形式

• pow()

• range

• round()

• zip()

标准库模块

cmath

以下是 cmath 模块中支持的函数：

• cmath.acos()

• cmath.acosh()

• cmath.asin()

• cmath.asinh()

• cmath.atan()

• cmath.atanh()

• cmath.cos()

• cmath.cosh()

• cmath.exp()

• cmath.isfinite()

• cmath.isinf()

• cmath.isnan()

• cmath.log()

• cmath.log10()

• cmath.phase()

• cmath.polar()

• cmath.rect()

• cmath.sin()

• cmath.sinh()

• cmath.sqrt()

• cmath.tan()

• cmath.tanh()

math

以下是 math 模块中支持的函数：

• math.acos()

• math.asin()

• math.atan()

• math.acosh()

• math.asinh()

• math.atanh()

• math.cos()

• math.sin()

• math.tan()

• math.hypot()

• math.cosh()

• math.sinh()

• math.tanh()

• math.atan2()

• math.erf()

• math.erfc()

• math.exp()

• math.expm1()

• math.fabs()

• math.frexp()

• math.ldexp()

• math.gamma()

• math.lgamma()

• math.log()

• math.log2()

• math.log10()

• math.log1p()

• math.sqrt()

• math.remainder()

• math.pow()

• math.ceil()

• math.floor()

• math.copysign()

• math.fmod()

• math.modf()

• math.isnan()

• math.isinf()

• math.isfinite()

operator

以下 operator 模块中的函数是被支持的：

• operator.add()

• operator.and_()

• operator.eq()

• operator.floordiv()

• operator.ge()

• operator.gt()

• operator.iadd()

• operator.iand()

• operator.ifloordiv()

• operator.ilshift()

• operator.imod()

• operator.imul()

• operator.invert()

• operator.ior()

• operator.ipow()

• operator.irshift()

• operator.isub()

• operator.itruediv()

• operator.ixor()

• operator.le()

• operator.lshift()

• operator.lt()

• operator.mod()

• operator.mul()

• operator.ne()

• operator.neg()

• operator.not_()

• operator.or_()

• operator.pos()

• operator.pow()

• operator.rshift()

• operator.sub()

• operator.truediv()

• operator.xor()

NumPy 支持

由于CUDA编程模型，内核内部的动态内存分配是低效的，通常并不需要。Numba禁止任何内存分配功能。这禁用了大量的NumPy API。为了获得最佳性能，用户应编写代码，使得每个线程一次只处理一个元素。

支持的 NumPy 功能：

• 访问 ndarray 属性 .shape、.strides、.ndim、.size 等。

• 索引和切片功能正常。

• 支持ufuncs的一个子集，但输出数组必须作为位置参数传入（参见 调用 NumPy UFunc）。请注意，ufuncs在每个线程中顺序执行——输入数组的元素之间没有自动的ufuncs并行化。

  以下ufuncs受支持：

  • numpy.sin()

  • numpy.cos()

  • numpy.tan()

  • numpy.arcsin()

  • numpy.arccos()

  • numpy.arctan()

  • numpy.arctan2()

  • numpy.hypot()

  • numpy.sinh()

  • numpy.cosh()

  • numpy.tanh()

  • numpy.arcsinh()

  • numpy.arccosh()

  • numpy.arctanh()

  • numpy.deg2rad()

  • numpy.radians()

  • numpy.rad2deg()

  • numpy.degrees()

  • numpy.greater()

  • numpy.greater_equal()

  • numpy.less()

  • numpy.less_equal()

  • numpy.not_equal()

  • numpy.equal()

  • numpy.log()

  • numpy.log2()

  • numpy.log10()

  • numpy.logical_and()

  • numpy.logical_or()

  • numpy.logical_xor()

  • numpy.logical_not()

  • numpy.maximum()

  • numpy.minimum()

  • numpy.fmax()

  • numpy.fmin()

  • numpy.bitwise_and()

  • numpy.bitwise_or()

  • numpy.bitwise_xor()

  • numpy.invert()

  • numpy.bitwise_not()

  • numpy.left_shift()

  • numpy.right_shift()

不支持的 NumPy 功能：

• 数组创建API。

• 数组方法。

• 返回新数组的函数。

CFFI 支持

cffi.FFI 对象的 from_buffer() 方法是被支持的。这对于获取可以传递给外部 C / C++ / PTX 函数的指针非常有用（参见 CUDA FFI 文档）。