1

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

NVCC 似乎没有添加或更改预处理器定义--use_fast_math的时间。我通过比较以下输出来检查这一点：

nvcc -Xcompiler -dM -E -x cu -

与输出

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

它们完全一样；所以那条大道似乎被封锁了。现在，如果编译用户会调用 NVCC，--use_fast_math -DUSING_FAST_MATH那么我也可以检测到；但假设它是库代码，我们不能将这些限制强加给用户。

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

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

4

2 回答 2

1

这是一个丑陋的拼凑而不是答案：

在您的代码中放置如下内容：

#ifndef I_ENABLED_FAST_MATH_REALLY_I_DID_PLEASE_BELIEVE_ME
#warning "Foo library is possibly being compiled without `--use_fast_math`. Please enable that switch, and define I_ENABLED_FAST_MATH_REALLY_I_DID_PLEASE_BELIEVE_ME, to avoid this message.
#endif

虽然这不会强迫你的用户做你想做的事，但它会很烦人，至少会鼓励他们去做。（当然，他们可以简单地使用预处理器定义禁用警告。）

于 2020-02-11T12:46:06.443 回答
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 黑客来拦截或识别您正在寻找的字节码，但我非常怀疑这就是您的想法。

于 2020-02-11T08:24:36.023 回答