局限性#

本节总结了 Warp 中的各种局限性和当前不支持的功能。问题、疑问和功能请求可以在GitHub Issues上提出。

不支持的功能#

为了在 GPU 上实现良好的性能,某些动态语言功能不受支持

  • Lambda 函数

  • 列表推导式

  • 异常

  • 递归

  • 表达式的运行时评估,例如:eval()

  • 动态结构,例如列表、集合、字典等。

内核和用户函数#

  • 字符串不能传递到内核中。

  • 不支持短路求值

  • wp.atomic_add() 不支持 wp.int64

  • wp.tid() 不能从用户函数中调用。

  • 如果在模块已经加载(例如通过 wp.launch()wp.load_module())的情况下,在运行时修改 wp.constant() 的值不会触发受影响内核的重新编译。

  • 如果 wp.constant()wp.float64 一起使用,可能会出现精度损失,因为它最初被分配给生成的代码中的 wp.float32 变量。

Warp 的一个限制是,用于启动内核的网格的每个维度都必须可以表示为 32 位有符号整数。 因此,网格的任何单个维度都不应超过 \(2^{31}-1\)

Warp 目前还使用每个块 256 个(CUDA)线程的固定块大小。 默认情况下,Warp 将尝试在一个 CUDA 线程中处理来自 Warp 网格的一个元素。 对于使用多维网格边界启动的内核,这并非总是可行,因为 CUDA 块维度存在硬件限制

当 CUDA 线程不可能仅处理 Warp 网格中的一个元素时,Warp 将自动回退到使用网格步幅循环。 发生这种情况时,某些 CUDA 线程可能会处理来自 Warp 网格的多个元素。 用户还可以设置 max_blocks 参数来微调内核的网格步幅行为,即使对于原本能够每个 CUDA 线程处理一个 Warp 网格元素的内核也是如此。

可微性#

有关自动微分的限制,请参阅“可微性”页面中的局限性和解决方法部分。

数组#

  • 数组最多可以有四个维度。

  • Warp 数组的每个维度不能大于 32 位有符号整数可以表示的最大值,\(2^{31}-1\)

  • 目前没有支持复数的数据类型。

结构体#

  • 结构体不能具有泛型成员,即类型为 typing.Any

体数据#

  • 在分配 Volume 的瓦片之后,无法更改稀疏体数据的 *拓扑*。

多个进程#

  • 在父进程中创建的 CUDA 上下文不能在 *forked* 子进程中使用。 请改用 spawn 启动方法,或避免在父进程中创建 CUDA 上下文。

  • 使用多个进程运行时,使用同一用户内核缓存目录可能会出现问题。一种解决方法是为每个进程使用单独的缓存目录。 有关如何更改缓存目录的信息,请参阅配置部分。

标量数学函数#

本节详细介绍了一些标量数学函数的局限性和与 CPython 语义的差异。

求模运算符#

当求模运算符 (%) 与负被除数或除数一起使用时,可能会出现与 Python 行为的偏差(另请参阅wp.mod())。 Warp 内核中求模运算符的行为遵循 C++11:结果的符号遵循 *被除数* 的符号。 在 Python 中,结果的符号遵循 *除数* 的符号

@wp.kernel
def modulus_test():
    # Kernel-scope behavior:
    a = -3 % 2 # a is -1
    b = 3 % -2 # b is 1
    c = 3 % 0  # Undefined behavior

# Python-scope behavior:
a = -3 % 2 # a is 1
b = 3 % -2 # b is -1
c = 3 % 0  # ZeroDivisionError

幂运算符#

Warp 内核中的幂运算符 (**) 仅适用于浮点数(另请参阅wp.pow())。 在 Python 中,幂运算符也可以用于整数。

反正弦和反余弦#

wp.asin()wp.acos() 自动将输入限制在 [-1, 1] 范围内。 在 Python 中,使用超出 [-1, 1] 范围的输入来使用math.asin()math.acos() 会引发 ValueError 异常。

舍入#

wp.round() 将中点情况舍入为远离零,但 Python 的 round() 将中点情况舍入为最接近的偶数选择(银行家舍入)。 如果需要银行家舍入,请使用 wp.rint()。 与 Python 不同,这两个舍入函数在 Warp 中的返回类型与输入类型相同

@wp.kernel
def halfway_rounding_test():
    # Kernel-scope behavior:
    a = wp.round(0.5) # a is 1.0
    b = wp.rint(0.5)  # b is 0.0
    c = wp.round(1.5) # c is 2.0
    d = wp.rint(1.5)  # d is 2.0

# Python-scope behavior:
a = round(0.5) # a is 0
c = round(1.5) # c is 2

变量作用域#

编写 Warp 内核时,变量作用域的行为可能与标准 Python 不同。 有时这会导致意外的结果。

在标准 Python 中,变量只能在定义它们的块内访问。 考虑以下示例

@wp.func
def foo(cond: bool):
    if cond:
        out = 123
    else:
        out = 234

    print(out)

此代码在标准 Python 中按预期工作。 无论 cond 的值如何,out 都会在打印之前定义。

但是,考虑一个稍作修改的示例

@wp.func
def foo(cond: bool):
    if cond:
        out = 123

    print(out) # No error even when `cond` is `False`.

在标准 Python 中,如果 condFalse,则对 print(out) 的调用将引发 UnboundLocalError,因为 out 仅在 if 块内定义。

在 Warp 中,行为有所不同。 即使 condFalse,对 print(out) 的调用也*不会*引发错误。 Warp 实际上使 out 可以在 if 块之外访问。 但是,如果 condFalse,则 out 将未初始化,从而导致未定义的行为。