局限性#
本节总结了 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 中,如果 cond 为 False,则对 print(out) 的调用将引发 UnboundLocalError,因为 out 仅在 if 块内定义。
在 Warp 中,行为有所不同。 即使 cond 为 False,对 print(out) 的调用也*不会*引发错误。 Warp 实际上使 out 可以在 if 块之外访问。 但是,如果 cond 为 False,则 out 将未初始化,从而导致未定义的行为。