cuda - 我可以在编译时确定是否设置了 --use_fast_math 吗？

Question

我正在编写一些 CUDA 代码，我希望它根据 --use_fast_math 是否设置而表现不同。而且 - 我想在编译时做出决定，而不是在运行时。

当设置 --use_fast_math 时，NVCC 似乎不添加或更改预处理器定义。我通过比较以下输出来检查这一点:

nvcc -Xcompiler -dM -E -x cu -

输出为

nvcc -Xcompiler -dM -E --use_fast_math -x cu -

而且它们完全一样；所以那条路似乎被堵住了。现在，如果编译用户使用 --use_fast_math -DUSING_FAST_MATH 调用 NVCC，那么我也可以检测到这一点；但假设它是库代码，我们不能对用户施加这些限制。

是否有其他方法可以让正在编译的代码注意到 --use_fast_math 已打开？

注意:“注意到”可能意味着使用预处理器 #if 或 #ifdef 指令、使用 SFINAE、使用编译器内置值或 constexpr 函数 - 编译时可用的任何内容时间。

Answer 1

答案几乎肯定是否定的。快速数学函数是硬件指令，它们被 CUDA 设备代码编译器内的代码生成所取代。一个例子:

$ cat nonsense.cu

__global__ void kernel(float* in, float* out)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    out[idx] = sqrtf(cosf(in[idx]));
}

$ nvcc -v -arch=sm_60 --keep -c --use_fast_math nonsense.cu 
#$ _SPACE_= 
#$ _CUDART_=cudart
#$ _HERE_=/opt/cuda-10.1/bin
#$ _THERE_=/opt/cuda-10.1/bin
#$ _TARGET_SIZE_=
#$ _TARGET_DIR_=
#$ _TARGET_DIR_=targets/x86_64-linux
#$ TOP=/opt/cuda-10.1/bin/..
#$ NVVMIR_LIBRARY_DIR=/opt/cuda-10.1/bin/../nvvm/libdevice
#$ LD_LIBRARY_PATH=/opt/cuda-10.1/bin/../lib:/opt/cuda-10.1/lib64
#$ PATH=/opt/cuda-10.1/bin/../nvvm/bin:/opt/cuda-10.1/bin:/opt/miniconda3/bin:/opt/miniconda3/condabin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/snap/bin:/opt/cuda-10.1/bin
#$ INCLUDES="-I/opt/cuda-10.1/bin/../targets/x86_64-linux/include"  
#$ LIBRARIES=  "-L/opt/cuda-10.1/bin/../targets/x86_64-linux/lib/stubs" "-L/opt/cuda-10.1/bin/../targets/x86_64-linux/lib"
#$ CUDAFE_FLAGS=
#$ PTXAS_FLAGS=
#$ gcc -std=c++14 -D__CUDA_ARCH__=600 -E -x c++  -DCUDA_DOUBLE_MATH_FUNCTIONS -D__CUDACC__ -D__NVCC__  "-I/opt/cuda-10.1/bin/../targets/x86_64-linux/include"    -D__CUDACC_VER_MAJOR__=10 -D__CUDACC_VER_MINOR__=1 -D__CUDACC_VER_BUILD__=105 -include "cuda_runtime.h" -m64 "nonsense.cu" > "nonsense.cpp1.ii" 
#$ cicc --c++14 --gnu_version=70400 --allow_managed   -arch compute_60 -m64 -ftz=1 -prec_div=0 -prec_sqrt=0 -fmad=1 -fast-math --gen_div_approx_ftz --include_file_name "nonsense.fatbin.c" -tused -nvvmir-library "/opt/cuda-10.1/bin/../nvvm/libdevice/libdevice.10.bc" --gen_module_id_file --module_id_file_name "nonsense.module_id" --orig_src_file_name "nonsense.cu" --gen_c_file_name "nonsense.cudafe1.c" --stub_file_name "nonsense.cudafe1.stub.c" --gen_device_file_name "nonsense.cudafe1.gpu"  "nonsense.cpp1.ii" -o "nonsense.ptx"
#$ ptxas -arch=sm_60 -m64  "nonsense.ptx"  -o "nonsense.sm_60.cubin" 
#$ fatbinary --create="nonsense.fatbin" -64 "--image=profile=sm_60,file=nonsense.sm_60.cubin" "--image=profile=compute_60,file=nonsense.ptx" --embedded-fatbin="nonsense.fatbin.c" 
#$ gcc -std=c++14 -E -x c++ -D__CUDACC__ -D__NVCC__  "-I/opt/cuda-10.1/bin/../targets/x86_64-linux/include"    -D__CUDACC_VER_MAJOR__=10 -D__CUDACC_VER_MINOR__=1 -D__CUDACC_VER_BUILD__=105 -include "cuda_runtime.h" -m64 "nonsense.cu" > "nonsense.cpp4.ii" 
#$ cudafe++ --c++14 --gnu_version=70400 --allow_managed  --m64 --parse_templates --gen_c_file_name "nonsense.cudafe1.cpp" --stub_file_name "nonsense.cudafe1.stub.c" --module_id_file_name "nonsense.module_id" "nonsense.cpp4.ii" 
#$ gcc -std=c++14 -D__CUDA_ARCH__=600 -c -x c++  -DCUDA_DOUBLE_MATH_FUNCTIONS "-I/opt/cuda-10.1/bin/../targets/x86_64-linux/include"   -m64 -o "nonsense.o" "nonsense.cudafe1.cpp" 

$ cat nonsense.ptx
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-25769353
// Cuda compilation tools, release 10.1, V10.1.105
// Based on LLVM 3.4svn
//

.version 6.4
.target sm_60
.address_size 64

    // .globl   _Z6kernelPfS_

.visible .entry _Z6kernelPfS_(
    .param .u64 _Z6kernelPfS__param_0,
    .param .u64 _Z6kernelPfS__param_1
)
{
    .reg .f32   %f<4>;
    .reg .b32   %r<5>;
    .reg .b64   %rd<8>;


    ld.param.u64    %rd1, [_Z6kernelPfS__param_0];
    ld.param.u64    %rd2, [_Z6kernelPfS__param_1];
    cvta.to.global.u64  %rd3, %rd2;
    cvta.to.global.u64  %rd4, %rd1;
    mov.u32     %r1, %tid.x;
    mov.u32     %r2, %ctaid.x;
    mov.u32     %r3, %ntid.x;
    mad.lo.s32  %r4, %r3, %r2, %r1;
    mul.wide.s32    %rd5, %r4, 4;
    add.s64     %rd6, %rd4, %rd5;
    ld.global.f32   %f1, [%rd6];
    cos.approx.ftz.f32  %f2, %f1;
    sqrt.approx.ftz.f32     %f3, %f2;
    add.s64     %rd7, %rd3, %rd5;
    st.global.f32   [%rd7], %f3;
    ret;
}

您可以看到，没有 nvcc 引导的预处理器魔法，只是传递给设备编译器的参数，这导致了带有必要指令的 PTX 代码。这意味着，理论上，您可能能够使用 LLVM hack 来拦截或识别您正在寻找的字节码，但我非常怀疑这就是您的想法。