为什么 NVCC 不优化 ceilf() 的文字？

Question

_{（关于编译时上限函数的后续问题，对于文字，在 C 中？）}

考虑以下 CUDA 函数：

__device__ int foo_f() { return ceilf(1007.1111); }

应该很容易对其进行优化以生成仅返回 1008 的设备函数：

mov.u32         %r1, 1008;
st.param.b32    [func_retval0+0], %r1;
ret;

但相反，它编译（使用 NVCC 11.5）成本更高：

mov.f32          %f1, 0f447C0000;
cvt.rzi.s32.f32  %r1, %f1;
st.param.b32     [func_retval0+0], %r1;
ret;

如果代码是：

static __device__ int poor_mans_ceilf(float x)
{
    return (int) x + ( ((float)(int) x < x) ? 1 : 0);
}

__device__ int foo_pf() { return poor_mans_ceilf(1007.1111); }

编译器应该更容易“注意到”这是一个优化机会。

那么，为什么 NVCC 未能在这里进行优化（而典型的 C/C++ 编译器确实会进行优化）？是否存在一些微妙的问题阻止（编辑）PTX 代码的优化？我意识到ptxas最终有机会优化它，但这不是特定于微架构的优化。

在 上查看这一切GodBolt。

_{PS：我知道这可以通过使用来规避constexpr。}

Answer 1

您无疑完全清楚，PTX 是一种虚拟汇编语言，并不由 GPU 运行。如果我们将您的代码编译为机器代码，我们会看到：

$ cat bogogogo.cu
__device__ int foo_f() { return ceilf(1007.1111); }

$ nvcc -dc -Xptxas='-v' bogogogo.cu
ptxas info    : 0 bytes gmem
ptxas info    : Function properties for _Z5foo_fv
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads

$ cuobjdump -sass bogogogo.o

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
compressed

        code for sm_52
                Function : _Z5foo_fv
        .headerflags    @"EF_CUDA_SM52 EF_CUDA_PTX_SM(EF_CUDA_SM52)"
                                                       /* 0x001ffc00ffe007f0 */
        /*0008*/         {         MOV32I R4, 0x3f0 ;  /* 0x010000003f07f004 */
        /*0010*/                   RET         }
                                                       /* 0xe32000000007000f */
        /*0018*/                   BRA 0x18 ;          /* 0xe2400fffff87000f */
                                                       /* 0x001f8000fc0007e0 */
        /*0028*/                   NOP;                /* 0x50b0000000070f00 */
        /*0030*/                   NOP;                /* 0x50b0000000070f00 */
        /*0038*/                   NOP;                /* 0x50b0000000070f00 */

您可以清楚地看到，在最终的汇编器输出中，转换已被优化为立即常数 (0x3f0 = 1008 = ceilf(1007.1111))。因此，您要求的优化是由 PTX 汇编器执行的，而不是由前端 C++ 编译器执行的。

如果您想知道为什么 NVIDIA 工具链设计者决定以这种方式在编译器和汇编器之间执行优化工作，您必须直接询问他们。