如何在 CUDA 中高效判断 half-precision 值是否有限？

结论

优先用位运算检查指数位：通过 __half_as_ushort() 取得 half 的原始编码，判断指数是否全为 1。这种写法直接，不涉及浮点格式转换，适合作为性能优化的基准版本。

!__hisinf(x) && !__hisnan(x) 也正确，但编译器未必会合并这两个分类操作。哪种写法最快，要结合目标 GPU、CUDA 编译器和实际调用位置检查生成代码，单看源码无法确定。

推荐实现

#include <cuda_fp16.h>

__device__ __forceinline__
bool isfinite_half(__half x)
{
    const unsigned short bits = __half_as_ushort(x);
    return (bits & 0x7C00u) != 0x7C00u;
}

__half_as_ushort() 按位重解释数据，保留原始的 16 位编码。它与 __half2ushort_rz() 的含义不同，后者会把浮点数转换成整数。相关声明位于 cuda_fp16.h。CUDA 数据转换与移动函数文档

函数命名为 isfinite_half，便于与已有的 isfinite 重载区分。

为什么只检查指数就够了？

CUDA 的 __half 使用 binary16 格式，包括 1 位符号、5 位指数和 10 位尾数字段。CUDA __half 类型文档

位编号：15   14       10   9                 0
        符号    指数             尾数
掩码：        0x7C00

binary16 编码与数值类别的对应关系如下：

判断有限性时，无需区分无穷与 NaN，也无需检查符号位，因此可以写成：

(bits & 0x7C00u) != 0x7C00u

这个表达式能正确处理正负零、次正规数、所有有限正规数，以及不同符号和尾数编码的 NaN。

也可以先清除符号位，再进行整数比较，结果等价：

return (__half_as_ushort(x) & 0x7FFFu) < 0x7C00u;

清除符号位后，有限值的编码小于 0x7C00，无穷等于它，NaN 大于它。两种表达式都很简短，但简短并不代表其中一种在所有 GPU 上都更快。

组合 __hisinf() 和 __hisnan() 可以吗？

可以，下面的写法能正确判断有限性：

__device__ __forceinline__
bool isfinite_half_checks(__half x)
{
    return !__hisinf(x) && !__hisnan(x);
}

__hisinf() 对负无穷返回 -1，对正无穷返回 1，其他情况返回 0；__hisnan() 返回是否为 NaN。因此，逻辑非 ! 能正确处理正负无穷。CUDA half 比较函数文档

这段源码包含两个分类操作和一次逻辑组合。编译器可能将它们合并、简化，也可能保留更多指令。&& 有短路语义，但最终机器码未必会出现分支。

如果这段代码位于热点路径，直接检查指数位更便于检查优化结果。仅把 && 改成 & 也不能保证更快。

还有其他可选方案吗？

还可以用浮点比较判断 |x| < +∞，对应的 half 写法如下：

__device__ __forceinline__
bool isfinite_half_compare(__half x)
{
    const __half positive_inf = __ushort_as_half(0x7C00u);
    return __hlt(__habs(x), positive_inf);
}

有限值满足这个条件；正负无穷取绝对值后，都不满足严格小于。__hlt() 遇到 NaN 时返回 false，也符合有限性判断的要求。这里不能换成 __hltu()，因为后者遇到 NaN 会返回 true。CUDA half 比较函数文档

这种实现可以纳入性能比较。它是否比整数位运算更快，取决于绝对值、比较和常量在目标架构上最终编译成什么机器指令。

先通过 __half2float(x) 转成 float，再调用 isfinite()，同样能正确判断有限性。如果后续本来就需要这个 float 值，可以复用转换结果；如果只是为了分类，则没有明显理由优先选择这种写法。

怎样确认实际效率？

先把指数位版本放到实际调用位置，使用实际的优化选项和目标架构编译，再检查最终 SASS。cuobjdump 和 nvdisasm 都能查看机器指令。PTX 仍是中间表示，不能直接用来判断最终指令成本。CUDA Binary Utilities 文档

比较时，主要检查是否有额外的格式转换、位提取、比较和分支，以及判定结果能否直接用于后续条件。源码中的一次按位与加一次比较，未必恰好对应两条机器指令。

实测时应保留计算结果，防止编译器删除整个判断，并尽量在实际 kernel 中比较。单独测量延迟和吞吐量时，需要分别设计依赖链测试和包含多个独立操作的测试。读取数组后逐项判断的 kernel，也可能主要受内存访问限制。

正确性检查可以遍历全部 65,536 种 16 位编码，重点覆盖正负零、最小次正规数、最大有限值、正负无穷和不同 NaN 编码。在没有目标硬件的生成代码和测量结果时，可以优先选择指数位检查，但没有充分依据认定它在所有架构上都最快。

备注：内容仅供参考。