Могу ли я определить во время компиляции, был ли установлен параметр --use_fast_math?

Question

Я пишу некоторый код CUDA и хочу, чтобы он вел себя по-разному в зависимости от того, был установлен --use_fast_math или нет. И - я хочу принять это решение во время компиляции, а не во время выполнения.

Кажется, что NV CC не не добавляет или изменяет определение препроцессора, когда установлено --use_fast_math , Я проверил это, сравнив вывод:

nvcc -Xcompiler -dM -E -x cu -

с выводом

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

, и они абсолютно одинаковы; так что проспект кажется заблокированным. Теперь, если пользователь компиляции вызовет NV CC с --use_fast_math -DUSING_FAST_MATH, то я также могу это обнаружить; но предположим, что это код библиотеки, и мы не можем наложить эти ограничения на пользователя.

Есть ли другой способ для кода, проходящего компиляцию, заметить, что --use_fast_math включен?

Примечание: " Заметить "может означать использование директив препроцессора #if или #ifdef, использование SFINAE, использование встроенных в компилятор значений или функций constexpr - все, что доступно во время компиляции.

Answer 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

Хотя это не заставит ваших пользователей делать то, что вы хотите, это будет достаточно раздражающим, чтобы хотя бы побудить их сделать это. (Естественно, они могут просто отключить предупреждение с помощью определения препроцессора.)

Answer 2

Ответ почти наверняка нет. Быстрые математические функции являются аппаратными инструкциями и заменяются генерацией кода в компиляторе кода устройства 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;
}

Вы можете видеть, что нет управляемых препроцессоров magi c nv cc, только аргументы, передаваемые компилятору устройства, который привел к коду PTX с необходимыми инструкциями в место. Это означает, что теоретически вы могли бы возиться с хакерами LLVM, чтобы перехватить или идентифицировать искомый байт-код, но я очень сомневаюсь, что это то, что вы имели в виду.