在CUDA中，如何高效地判断一个半精度浮点数是否有限？

后端开发 2026-07-09

CUDA提供一个（NVIDIA GPU）函数，用于检查半精度值是否为无穷大： __hisinf()。但——没有对应的 __hisfinite() 函数（就像用于双精度值的函数，确实存在：int isfinite(double x)）。

一个 bool isfinite(half x) 函数的最有效实现会是什么？具体来说，

我是不是就用 not __hisinf(x) and not __hisnan(x)，希望编译器能把它优化得足够好？
我是否需要进行位运算，确保指数部分不是全为1？
还有其他办法吗？

“Efficient” 并不是一个精确的术语，原理上可能存在延迟和吞吐量的权衡，但我猜在这种情况下，这并不是问题。

解决方案

（在 njuffa 的回答之后编辑。）

tl;dr: 最快的方法似乎是直接在半精度值的 unsigned short 表示中，将指数位与全1 的模式进行比较。但——这个结果可能取决于你如何衡量。

需要考虑的多种方法

到目前为止，在我们这些人之间，似乎有6 种建议的方法来检查半精度值的有限性/有限性：

（来自 @njuffa）按IEEE 754，当且仅当与正无穷相比为更小的数才是有限的。现在只需要得到正无穷的值——我们可以通过从已知的整数表示得到 __ushort2half()，或者通过对 INFINITY 宏进行简单的强制类型转换来获得：

bool isfinite(half x) { return __habs(x) < ((half)INFINITY); } 2. （来自 @njuffa）CUDA提供用于检查 float 值有限性的内建函数；因此我们直接把它强制转换为 float 并使用它。

bool isfinite(half x) { return isfinite(__half2float(x)); } 3. （来自 @wim）如果把表示的最低15位（即忽略符号）作为一个整数值来考虑，那么其中的前5 位需要全部为1；因此有限数的值要小于非有限数。

bool isfinite(half x) { unsigned short shifted_plus_infinity_rep = 0b0111110000000000u << 1; unsigned short shifted_input_rep = __half_as_ushort (x) << 1; return shifted_input_rep < shifted_plus_infinity_rep; } 4.直接检查指数位：非有限值的指数为全1；有限值则是其他值。

bool isfinite(half x){ auto exponent_bits_mask = 0b111110000000000u; return (__half_as_ushort(x) & exponent_bits_mask) != exponent_bits_mask; } 5.通过位域检查指数位（本质上与前一种方法相同）

bool isfinite(half x){ union destructured_half { half value; struct { unsigned short mantissa: 10; unsigned short exponent : 5; unsigned short sign : 1; } components; }; destructured_half d { .value = x }; return d.components.exponent != 0b11111u; // all ones } 6. （@ericpostpischil的想法）用自我相减。只有有限值在与自身相减时才会得到 0（对于无穷大和NaN，结果是NaN）。这种方法在代码上最简短，具有美学吸引力。

bool isfinite(half x){ return x - x == (half) 0; }

编译结果分析

改编自这里（GodBolt.org，CUDA 13.2，SM 89目标）。

前置、签名和最终存储指令省略。加载指令本应是通用的——但实际上并非如此！因此只去除了通用的 R1 设置。

方法1：

 LDC.U16 R0, c[0x0][0x168] 
 MOV R2, c[0x0][0x160] 
 IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] 
 ULDC.64 UR4, c[0x0][0x118] 
 HSETP2.GEU.AND P0, PT, |R0|.H0_H0, +INF , +INF , PT 
 SEL R5, RZ, 0x1, P0

方法2：

 HADD2.F32 R0, -RZ, c[0x0] [0x168].H0_H0 
 MOV R2, c[0x0][0x160] 
 IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] 
 ULDC.64 UR4, c[0x0][0x118] 
 FSETP.GEU.AND P0, PT, |R0|, +INF , PT 
 SEL R5, RZ, 0x1, P0

方法3：

 LDC.U16 R0, c[0x0][0x168] 
 MOV R2, c[0x0][0x160] 
 IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] 
 ULDC.64 UR4, c[0x0][0x118] 
 IMAD.SHL.U32 R0, R0, 0x2, RZ 
 LOP3.LUT R0, R0, 0xffff, RZ, 0xc0, !PT 
 ISETP.GE.U32.AND P0, PT, R0, 0xf800, PT 
 SEL R5, RZ, 0x1, P0

方法4：

 LDC.U16 R0, c[0x0][0x168] 
 IMAD.MOV.U32 R2, RZ, RZ, c[0x0][0x160] 
 ULDC.64 UR4, c[0x0][0x118] 
 IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] 
 LOP3.LUT R0, R0, 0x7c00, RZ, 0xc0, !PT 
 ISETP.NE.AND P0, PT, R0, 0x7c00, PT 
 SEL R5, RZ, 0x1, !P0

方法5：

 LDC.U16 R0, c[0x0][0x168] 
 IMAD.MOV.U32 R2, RZ, RZ, c[0x0][0x160] 
 ULDC.64 UR4, c[0x0][0x118] 
 IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] 
 LOP3.LUT R0, R0, 0x7c00, RZ, 0xc0, !PT 
 ISETP.NE.AND P0, PT, R0, 0x7c00, PT 
 SEL R5, RZ, 0x1, !P0

方法6：

 LDC.U16 R0, c[0x0][0x168] 
 MOV R2, c[0x0][0x160] 
 IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] 
 ULDC.64 UR4, c[0x0][0x118] 
 HADD2 R0, R0.H0_H0, -c[0x0] [0x168].H0_H0 
 HSETP2.NEU.AND P0, PT, R0.H0_H0, RZ.H0_H0, PT 
 SEL R5, RZ, 0x1, P0

方法2 中的 HADD2.F32 非常奇怪！它把半精度值加到负零上。为什么？

性能比较

基准测试方法

我采用了 @njuffa的微基准内核写法：遍历所有 unsigned short 值，将每个值视为半精度值，并检查其是否有限。不同的线程从不同的 unsigned short 值开始。内核被调度多次。仅运行一个相同的块网格尺寸、但不包含有限性检查逻辑的内核开销可以忽略不计——只有几微秒；因此忽略它的影响。

不过我用的计时方式与 @njuffa不同：

使用CUDA事件而不是系统时间。我也使用了cuda-api-wrappers。
取内核调用的中位数而不是最小值，重复10次。
运行100次迭代

（如果你认为这些选择有问题，请在评论中说明。）

结果

在NVIDIA RTX 4600上：

序号	方法名称	时间（毫秒）
1	小于无穷大字面量	6.7031
2	使用float类型的CUDA内置函数	7.25789
3	小于无符号短整型阈值值	8.26982
4	直接比较指数位	6.56589
5	通过结构化解构比较指数位	6.56835
6	自我相减判定法	7.25725

看起来获胜的方法是将指数位与无穷大与NaN的模式进行对比。没有明显差异

我很快会在其他GPU上也试试，敬请关注...

基准测试程序：

#include <cuda/api.hpp>
#include <cuda_fp16.h>
#include <iostream>
#include <iomanip>

enum FinityCheckMethod {
    smaller_than_infinity_literal,                        // Method 1
    using_cuda_intrinsic_for_float_type,                  // Method 2
    less_than_ushort_threshold_value_for_non_finite,      // Method 3
    bit_pattern_comparison,                               // Method 4
    bit_pattern_comparison_via_destructuring,             // Method 5
    self_subtraction,                                     // Method 6
};

#define __fd__ __forceinline__ __device__

template <FinityCheckMethod FCM>
__fd__ bool isfinite (half x);

template <> __fd__ bool isfinite<smaller_than_infinity_literal>(half x)
{
    return __habs(x) < ((half)INFINITY);
}

template <> __fd__ bool isfinite<using_cuda_intrinsic_for_float_type>(half x)
{
    return isfinite(__half2float(x));
}

template <> __fd__ bool isfinite<less_than_ushort_threshold_value_for_non_finite>(half x)
{
    static constexpr unsigned short shifted_plus_infinity_rep = 0b0111110000000000u << 1;
    return (unsigned short)(__half_as_ushort (x) << 1) < shifted_plus_infinity_rep;
}

template <> __fd__ bool isfinite<bit_pattern_comparison>(half x)
{
    auto exponent_bits_mask = 0b0111110000000000u; // 10 least-significant bits are the mantissa
    return (__half_as_ushort(x) & exponent_bits_mask) != exponent_bits_mask;
}

template <> __fd__ bool isfinite<bit_pattern_comparison_via_destructuring>(half x)
{
    union destructured_half {
        half value;
        struct {
            unsigned short mantissa: 10;
            unsigned short exponent : 5;
            unsigned short sign : 1;
        } components;
    };
    destructured_half d { .value = x };
    return d.components.exponent != 0b11111u; // all ones
}

template <> __fd__ bool isfinite<self_subtraction>(half x)
{
    return x - x == (half) 0;
}

using count_t = unsigned int;

template <FinityCheckMethod FCM>
__global__ void kernel(count_t *finity_counts)
{
    auto tid = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned short arg = tid;
    count_t res = 0;
    // one iteration for every one of the 2^16 possible ushort values
    do {
        res = res + isfinite<FCM>(__ushort_as_half(arg));
        arg++;
    } while (arg != tid);
    finity_counts[tid] = res;
}

const char* name_of(FinityCheckMethod FCM)
{
    static constexpr const char* names[] = {
        "smaller than infinity literal",
        "using cuda intrinsic for float type",
        "less than ushort threshold value for non-finite",
        "direct exponent bits comparison",
        "exponent bits comparison via destructuring",
        "self-substraction criterion",
    };
    static constexpr auto num_names = sizeof(names)/sizeof(char const*);
    return (FCM >= num_names) ? nullptr : names[FCM];
}

template <FinityCheckMethod FCM>
void sanity_check(cuda::span<count_t> d_finity_counts, cuda::span<count_t> finity_counts)
{
    for (size_t tid = 0; tid < d_finity_counts.size(); tid++) {
        auto num_encodings_of_finite_half_values { 0xf800u };
        if (finity_counts[tid] != num_encodings_of_finite_half_values) {
            std::cout
                << "sanity check for check variant " << name_of(FCM) << " failed: "
                << "Thread " << tid << "'s counted " << finity_counts[tid]
                << " finite representations rather than " <<  num_encodings_of_finite_half_values;
            exit(EXIT_FAILURE);
        }
    }
}

template <FinityCheckMethod FCM>
void time_method(cuda::device_t const& device)
{
    auto threads_per_block { 128 };
    auto num_ushort_values { 1lu << sizeof(unsigned short) * CHAR_BIT };
    assert(num_ushort_values % threads_per_block == 0);
    auto num_iterations { 100 };

    auto launch_config = cuda::launch_config_builder()
        .overall_size(num_ushort_values)
        .block_size(threads_per_block)
        .build();

    auto total_threads = num_ushort_values;

    auto d_finity_counts = cuda::memory::make_unique_span<count_t>(device, total_threads);
    auto finity_counts = cuda::make_unique_span<count_t>(total_threads);

    cuda::memory::zero(d_finity_counts);
    auto events = std::make_pair(cuda::event::create(device), cuda::event::create(device));
    auto timings = cuda::generate_unique_span<cuda::event::duration_t>(num_iterations,
        [&](size_t) {
            events.first.record();
            cuda::launch(kernel<FCM>, device, launch_config, d_finity_counts.data());
            events.second.record();
            device.synchronize();
            return cuda::event::time_elapsed_between(events.first, events.second);
        });
    std::nth_element(timings.begin(), timings.begin() + num_iterations / 2, timings.end());
    std::cout << "isfinite check method  " << std::left << std::setw(50) << name_of(FCM)
        << " : " << std::setw(12) << timings[num_iterations / 2].count() << " msec\n";
    cuda::memory::copy(finity_counts, d_finity_counts);

    sanity_check<FCM>(d_finity_counts, finity_counts);
}

int main()
{
    auto device = cuda::device::current::get();
    time_method<smaller_than_infinity_literal>(device);
    time_method<using_cuda_intrinsic_for_float_type>(device);
    time_method<less_than_ushort_threshold_value_for_non_finite>(device);
    time_method<bit_pattern_comparison>(device);
    time_method<bit_pattern_comparison_via_destructuring>(device);
    time_method<self_subtraction>(device);
}

GPU

站内所有文章版权归属LeftHeroAI导航站，无授权禁止任何主体转载、抄袭、复制内容，亦不得私自架设镜像站点。一经侵权，本站将通过法律途径追责。

如何让我的轮播切换更平滑？

2026-07-12

在同一运行环境中，为什么一个算法能比另一个算法更快地把相同的分数加起来？

2026-07-12

WebView中的内容闪烁

2026-07-12

如何为Hibernate查询设置超时？

2026-07-12

实时服务器问题

2026-07-12

在Firebird SQL 4中，包含row_number() 窗口函数的公用表表达式的性能

2026-07-12

在CUDA中，如何高效地判断一个半精度浮点数是否有限？

解决方案

需要考虑的多种方法

编译结果分析

性能比较

基准测试方法

结果

基准测试程序：

在本地化的验证错误信息中嵌入字符串

使用pandoc API（writeEPUB3）生成EPUB时，注释会重复，且缺少索引编号

相关文章

最新文章

热门工具

超人AI导航

在CUDA中，如何高效地判断一个半精度浮点数是否有限？

解决方案

需要考虑的多种方法

编译结果分析

性能比较

基准测试方法

结果

基准测试程序：

在本地化的验证错误信息中嵌入字符串

使用pandoc API（writeEPUB3）生成EPUB时，注释会重复，且缺少索引编号

相关文章

最新文章

热门工具