为何NVCC未将带constexpr2次幂指数的ldexpf优化为简单fmul？

考虑以下CUDA代码：

enum { p = 5 };
__device__ float adjust_mul(float x) { return x * (1 << p); }
__device__ float adjust_ldexpf(float x) { return ldexpf(x, p); }

我原本期望NVCC能将第一个函数和第二个函数优化为完全相同的代码，因为ldexpf(x, p)的定义是“将x乘以2的p次幂”。但两个函数的PTX代码差异显著：

第一个函数编译后的PTX代码：

.visible .func  (.param .b32 func_retval0) adjust_mul(float)(
    .param .b32 adjust_mul(float)_param_0
)
{
    ld.param.f32    %f1, [adjust_mul(float)_param_0];
    mul.f32     %f2, %f1, 0f42000000;
    st.param.f32    [func_retval0+0], %f2;
    ret;
}

第二个函数编译后的PTX代码：

.visible .func  (.param .b32 func_retval0) adjust_ldexpf(float)(
    .param .b32 adjust_ldexpf(float)_param_0
)
{
    ld.param.f32    %f5, [adjust_ldexpf(float)_param_0];
    abs.f32     %f1, %f5;
    setp.eq.f32     %p1, %f1, 0f00000000;
    setp.eq.f32     %p2, %f1, 0f7F800000;
    or.pred     %p3, %p1, %p2;
    @%p3 bra    $L__BB1_2;
    bra.uni     $L__BB1_1;
$L__BB1_2:
    setp.gt.f32     %p4, %f1, 0f00000000;
    add.f32     %f8, %f5, %f5;
    selp.f32    %f9, %f5, %f8, %p4;
    bra.uni     $L__BB1_3;
$L__BB1_1:
    mov.f32     %f6, 0f40A00000;
    ex2.approx.ftz.f32  %f7, %f6;
    mul.f32     %f9, %f7, %f5;
$L__BB1_3:
    st.param.f32    [func_retval0+0], %f9;
    ret;
}

为何NVCC没有将第二个函数优化为与第一个函数相同的PTX代码？

这是因为ldexpf的行为和直接乘以2^p并不完全等价，NVCC需要保留ldexpf对特殊浮点值的标准处理逻辑：

• 特殊值处理：ldexpf必须严格遵循C标准对0、无穷大（Inf）、NaN的行为定义：
  • 输入为0时，ldexpf(0, p)必须返回0；输入为无穷大或NaN时，需原样返回这些值。虽然直接乘法在常规场景下结果一致，但ldexpf的实现会显式分支处理这类情况，确保完全符合标准语义。
• 实现逻辑差异：ldexpf本质是通过调整浮点数的指数位实现缩放，而非直接执行乘法。对于非特殊值，它会采用指数操作（比如PTX中的ex2.approx配合乘法）完成计算，这种路径和直接乘常数的逻辑不同，编译器不会随意替换，除非能100%证明两者语义完全一致且不违反标准。
• 标准合规性优先级：编译器会优先保证标准库函数的行为严格符合C/C++规范，即便代码看起来可以等价优化。ldexpf作为标准库函数，必须严格遵守规范，不能因“常规场景等价”就简化为乘法，避免边缘场景出现不符合预期的结果。

简言之，虽然常规输入下两者效果一致，但ldexpf的标准语义要求处理更多边缘情况，这导致编译器无法将其优化为简单的乘法操作。

内容的提问来源于stack exchange，提问作者einpoklum