Index: clang/include/clang/Driver/Options.td =================================================================== --- clang/include/clang/Driver/Options.td +++ clang/include/clang/Driver/Options.td @@ -241,9 +241,10 @@ // Args.hasArg(OPT_ffoo) is used to check that the flag is enabled. // This is useful if the option is usually disabled. multiclass OptInFFlag flags=[]> { + string help="", list flags=[], code keypath=""> { def f#NAME : Flag<["-"], "f"#name>, Flags, - Group, HelpText; + Group, HelpText, + MarshallingInfoFlag; def fno_#NAME : Flag<["-"], "fno-"#name>, Flags, Group, HelpText; } @@ -251,11 +252,12 @@ // A boolean option which is opt-out in CC1. The negative option exists in CC1 and // Args.hasArg(OPT_fno_foo) is used to check that the flag is disabled. multiclass OptOutFFlag flags=[]> { + string help="", list flags=[], code keypath=""> { def f#NAME : Flag<["-"], "f"#name>, Flags, Group, HelpText; def fno_#NAME : Flag<["-"], "fno-"#name>, Flags, - Group, HelpText; + Group, HelpText, + MarshallingInfoFlag; } ///////// @@ -559,27 +561,33 @@ def cl_strict_aliasing : Flag<["-"], "cl-strict-aliasing">, Group, Flags<[CC1Option]>, HelpText<"OpenCL only. This option is added for compatibility with OpenCL 1.0.">; def cl_single_precision_constant : Flag<["-"], "cl-single-precision-constant">, Group, Flags<[CC1Option]>, - HelpText<"OpenCL only. Treat double precision floating-point constant as single precision constant.">; -def cl_finite_math_only : Flag<["-"], "cl-finite-math-only">, Group, Flags<[CC1Option]>, + HelpText<"OpenCL only. Treat double precision floating-point constant as single precision constant.">, + MarshallingInfoFlag<"LangOpts->SinglePrecisionConstants", "false">; +def cl_finite_math_only : Flag<["-"], "cl-finite-math-only">, Group, HelpText<"OpenCL only. Allow floating-point optimizations that assume arguments and results are not NaNs or +-Inf.">; def cl_kernel_arg_info : Flag<["-"], "cl-kernel-arg-info">, Group, Flags<[CC1Option]>, - HelpText<"OpenCL only. Generate kernel argument metadata.">; -def cl_unsafe_math_optimizations : Flag<["-"], "cl-unsafe-math-optimizations">, Group, Flags<[CC1Option]>, + HelpText<"OpenCL only. Generate kernel argument metadata.">, + MarshallingInfoFlag<"CodeGenOpts.EmitOpenCLArgMetadata", "false">; +def cl_unsafe_math_optimizations : Flag<["-"], "cl-unsafe-math-optimizations">, Group, HelpText<"OpenCL only. Allow unsafe floating-point optimizations. Also implies -cl-no-signed-zeros and -cl-mad-enable.">; def cl_fast_relaxed_math : Flag<["-"], "cl-fast-relaxed-math">, Group, Flags<[CC1Option]>, - HelpText<"OpenCL only. Sets -cl-finite-math-only and -cl-unsafe-math-optimizations, and defines __FAST_RELAXED_MATH__.">; + HelpText<"OpenCL only. Sets -cl-finite-math-only and -cl-unsafe-math-optimizations, and defines __FAST_RELAXED_MATH__.">, + MarshallingInfoFlag<"LangOpts->FastRelaxedMath", "false">; def cl_mad_enable : Flag<["-"], "cl-mad-enable">, Group, Flags<[CC1Option]>, - HelpText<"OpenCL only. Allow use of less precise MAD computations in the generated binary.">; -def cl_no_signed_zeros : Flag<["-"], "cl-no-signed-zeros">, Group, Flags<[CC1Option]>, + HelpText<"OpenCL only. Allow use of less precise MAD computations in the generated binary.">, + MarshallingInfoFlag<"CodeGenOpts.LessPreciseFPMAD", "false">; +def cl_no_signed_zeros : Flag<["-"], "cl-no-signed-zeros">, Group, HelpText<"OpenCL only. Allow use of less precise no signed zeros computations in the generated binary.">; def cl_std_EQ : Joined<["-"], "cl-std=">, Group, Flags<[CC1Option]>, HelpText<"OpenCL language standard to compile for.">, Values<"cl,CL,cl1.1,CL1.1,cl1.2,CL1.2,cl2.0,CL2.0,cl3.0,CL3.0,clc++,CLC++">; def cl_denorms_are_zero : Flag<["-"], "cl-denorms-are-zero">, Group, HelpText<"OpenCL only. Allow denormals to be flushed to zero.">; def cl_fp32_correctly_rounded_divide_sqrt : Flag<["-"], "cl-fp32-correctly-rounded-divide-sqrt">, Group, Flags<[CC1Option]>, - HelpText<"OpenCL only. Specify that single precision floating-point divide and sqrt used in the program source are correctly rounded.">; + HelpText<"OpenCL only. Specify that single precision floating-point divide and sqrt used in the program source are correctly rounded.">, + MarshallingInfoFlag<"CodeGenOpts.CorrectlyRoundedDivSqrt", "false">; def cl_uniform_work_group_size : Flag<["-"], "cl-uniform-work-group-size">, Group, Flags<[CC1Option]>, - HelpText<"OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel">; + HelpText<"OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel">, + MarshallingInfoFlag<"CodeGenOpts.UniformWGSize", "false">; def client__name : JoinedOrSeparate<["-"], "client_name">; def combine : Flag<["-", "--"], "combine">, Flags<[DriverOption, Unsupported]>; def compatibility__version : JoinedOrSeparate<["-"], "compatibility_version">; @@ -998,7 +1006,7 @@ HelpText<"Controls the semantics of floating-point calculations.">; def ffp_exception_behavior_EQ : Joined<["-"], "ffp-exception-behavior=">, Group, Flags<[CC1Option]>, HelpText<"Specifies the exception behavior of floating-point operations.">; -defm fast_math : OptInFFlag<"fast-math", "Allow aggressive, lossy floating-point optimizations">; +defm fast_math : OptInFFlag<"fast-math", "Allow aggressive, lossy floating-point, optimizations", "", "", [], "LangOpts->FastMath">; defm math_errno : OptInFFlag<"math-errno", "Require math functions to indicate errors by setting errno">; def fbracket_depth_EQ : Joined<["-"], "fbracket-depth=">, Group, Flags<[CoreOption]>; def fsignaling_math : Flag<["-"], "fsignaling-math">, Group; @@ -1208,16 +1216,11 @@ Group; def fassociative_math : Flag<["-"], "fassociative-math">, Group; def fno_associative_math : Flag<["-"], "fno-associative-math">, Group; -def freciprocal_math : - Flag<["-"], "freciprocal-math">, Group, Flags<[CC1Option]>, - HelpText<"Allow division operations to be reassociated">; -def fno_reciprocal_math : Flag<["-"], "fno-reciprocal-math">, Group; -def ffinite_math_only : Flag<["-"], "ffinite-math-only">, Group, Flags<[CC1Option]>; -def fno_finite_math_only : Flag<["-"], "fno-finite-math-only">, Group; -def fsigned_zeros : Flag<["-"], "fsigned-zeros">, Group; -def fno_signed_zeros : - Flag<["-"], "fno-signed-zeros">, Group, Flags<[CC1Option]>, - HelpText<"Allow optimizations that ignore the sign of floating point zeros">; +defm reciprocal_math : OptInFFlag< "reciprocal-math", "Allow division operations to be reassociated", "", "", [], "LangOpts->AllowRecip">; +def fapprox_func : Flag<["-"], "fapprox-func">, Group, Flags<[CC1Option, NoDriverOption]>, + MarshallingInfoFlag<"LangOpts->ApproxFunc", "false">; +defm finite_math_only : OptInFFlag<"finite-math-only", "", "", "", [], "LangOpts->FiniteMathOnly">; +defm signed_zeros : OptOutFFlag<"signed-zeros", "Allow optimizations that ignore the sign of floating point zeros", "", "", [], "LangOpts->NoSignedZero">; def fhonor_nans : Flag<["-"], "fhonor-nans">, Group; def fno_honor_nans : Flag<["-"], "fno-honor-nans">, Group; def fhonor_infinities : Flag<["-"], "fhonor-infinities">, Group; @@ -3823,14 +3826,18 @@ def mdisable_tail_calls : Flag<["-"], "mdisable-tail-calls">, HelpText<"Disable tail call optimization, keeping the call stack accurate">; def menable_no_infinities : Flag<["-"], "menable-no-infs">, - HelpText<"Allow optimization to assume there are no infinities.">; + HelpText<"Allow optimization to assume there are no infinities.">, + MarshallingInfoFlag<"LangOpts->NoHonorInfs", "false">; def menable_no_nans : Flag<["-"], "menable-no-nans">, - HelpText<"Allow optimization to assume there are no NaNs.">; + HelpText<"Allow optimization to assume there are no NaNs.">, + MarshallingInfoFlag<"LangOpts->NoHonorNaNs", "false">; def menable_unsafe_fp_math : Flag<["-"], "menable-unsafe-fp-math">, HelpText<"Allow unsafe floating-point math optimizations which may decrease " - "precision">; + "precision">, + MarshallingInfoFlag<"LangOpts->UnsafeFPMath", "false">; def mreassociate : Flag<["-"], "mreassociate">, - HelpText<"Allow reassociation transformations for floating-point instructions">; + HelpText<"Allow reassociation transformations for floating-point instructions">, + MarshallingInfoFlag<"LangOpts->AllowFPReassoc", "false">; def mabi_EQ_ieeelongdouble : Flag<["-"], "mabi=ieeelongdouble">, HelpText<"Use IEEE 754 quadruple-precision for long double">; def mfloat_abi : Separate<["-"], "mfloat-abi">, Index: clang/lib/Driver/ToolChains/Clang.cpp =================================================================== --- clang/lib/Driver/ToolChains/Clang.cpp +++ clang/lib/Driver/ToolChains/Clang.cpp @@ -2544,7 +2544,7 @@ llvm::DenormalMode DenormalFP32Math = DefaultDenormalFP32Math; StringRef FPContract = ""; bool StrictFPModel = false; - + bool ApproxFunc = false; if (const Arg *A = Args.getLastArg(options::OPT_flimited_precision_EQ)) { CmdArgs.push_back("-mlimit-float-precision"); @@ -2770,6 +2770,7 @@ RoundingFPMath = false; // If fast-math is set then set the fp-contract mode to fast. FPContract = "fast"; + ApproxFunc = true; break; case options::OPT_fno_fast_math: HonorINFs = true; @@ -2824,8 +2825,13 @@ CmdArgs.push_back("-fmath-errno"); if (!MathErrno && AssociativeMath && ReciprocalMath && !SignedZeros && - !TrappingMath) + !TrappingMath) { CmdArgs.push_back("-menable-unsafe-fp-math"); + ApproxFunc = true; + } + + if (ApproxFunc) + CmdArgs.push_back("-fapprox-func"); if (!SignedZeros) CmdArgs.push_back("-fno-signed-zeros"); @@ -3153,30 +3159,74 @@ } static void RenderOpenCLOptions(const ArgList &Args, ArgStringList &CmdArgs) { + const auto &OptTable = getDriverOptTable(); + + auto GetArgString = [&Args, &OptTable](OptSpecifier Specifier) { + return Args.MakeArgString(OptTable.getOption(Specifier).getPrefixedName()); + }; + + auto RenderArg = [&](const Arg *Arg) { + CmdArgs.push_back(Args.MakeArgString(Arg->getOption().getPrefixedName())); + }; + + auto RenderFiniteMathOnlyArg = [&](const Arg *Arg) { + CmdArgs.push_back(GetArgString(options::OPT_ffinite_math_only)); + CmdArgs.push_back(GetArgString(options::OPT_menable_no_nans)); + CmdArgs.push_back(GetArgString(options::OPT_menable_no_infinities)); + }; + + auto RenderUnsafeMathOptsArg = [&](const Arg *Arg) { + CmdArgs.push_back(GetArgString(options::OPT_cl_mad_enable)); + CmdArgs.push_back(GetArgString(options::OPT_menable_unsafe_fp_math)); + CmdArgs.push_back(GetArgString(options::OPT_mreassociate)); + CmdArgs.push_back(GetArgString(options::OPT_fno_signed_zeros)); + CmdArgs.push_back(GetArgString(options::OPT_freciprocal_math)); + CmdArgs.push_back(GetArgString(options::OPT_fapprox_func)); + }; + + auto RenderFastRelaxedMathArg = [&](const Arg *Arg) { + RenderArg(Arg); + + CmdArgs.push_back(GetArgString(options::OPT_cl_mad_enable)); + CmdArgs.push_back(GetArgString(options::OPT_ffast_math)); + CmdArgs.push_back(GetArgString(options::OPT_ffinite_math_only)); + CmdArgs.push_back(GetArgString(options::OPT_menable_unsafe_fp_math)); + CmdArgs.push_back(GetArgString(options::OPT_mreassociate)); + CmdArgs.push_back(GetArgString(options::OPT_menable_no_nans)); + CmdArgs.push_back(GetArgString(options::OPT_menable_no_infinities)); + CmdArgs.push_back(GetArgString(options::OPT_fno_signed_zeros)); + CmdArgs.push_back(GetArgString(options::OPT_freciprocal_math)); + CmdArgs.push_back(GetArgString(options::OPT_fapprox_func)); + }; + + auto RenderNoSignedZerosArg = [&](const Arg *Arg) { + CmdArgs.push_back(GetArgString(options::OPT_fno_signed_zeros)); + }; + // cl-denorms-are-zero is not forwarded. It is translated into a generic flag // for denormal flushing handling based on the target. - const unsigned ForwardedArguments[] = { - options::OPT_cl_opt_disable, - options::OPT_cl_strict_aliasing, - options::OPT_cl_single_precision_constant, - options::OPT_cl_finite_math_only, - options::OPT_cl_kernel_arg_info, - options::OPT_cl_unsafe_math_optimizations, - options::OPT_cl_fast_relaxed_math, - options::OPT_cl_mad_enable, - options::OPT_cl_no_signed_zeros, - options::OPT_cl_fp32_correctly_rounded_divide_sqrt, - options::OPT_cl_uniform_work_group_size - }; + const std::pair> + CLArguments[] = { + {options::OPT_cl_opt_disable, RenderArg}, + {options::OPT_cl_strict_aliasing, RenderArg}, + {options::OPT_cl_single_precision_constant, RenderArg}, + {options::OPT_cl_finite_math_only, RenderFiniteMathOnlyArg}, + {options::OPT_cl_kernel_arg_info, RenderArg}, + {options::OPT_cl_unsafe_math_optimizations, RenderUnsafeMathOptsArg}, + {options::OPT_cl_fast_relaxed_math, RenderFastRelaxedMathArg}, + {options::OPT_cl_mad_enable, RenderArg}, + {options::OPT_cl_no_signed_zeros, RenderNoSignedZerosArg}, + {options::OPT_cl_fp32_correctly_rounded_divide_sqrt, RenderArg}, + {options::OPT_cl_uniform_work_group_size, RenderArg}}; if (Arg *A = Args.getLastArg(options::OPT_cl_std_EQ)) { std::string CLStdStr = std::string("-cl-std=") + A->getValue(); CmdArgs.push_back(Args.MakeArgString(CLStdStr)); } - for (const auto &Arg : ForwardedArguments) - if (const auto *A = Args.getLastArg(Arg)) - CmdArgs.push_back(Args.MakeArgString(A->getOption().getPrefixedName())); + for (const auto &Arg : CLArguments) + if (const auto *A = Args.getLastArg(Arg.first)) + Arg.second(A); } static void RenderARCMigrateToolOptions(const Driver &D, const ArgList &Args, Index: clang/lib/Frontend/CompilerInvocation.cpp =================================================================== --- clang/lib/Frontend/CompilerInvocation.cpp +++ clang/lib/Frontend/CompilerInvocation.cpp @@ -939,15 +939,8 @@ Opts.NoEscapingBlockTailCalls = Args.hasArg(OPT_fno_escaping_block_tail_calls); Opts.FloatABI = std::string(Args.getLastArgValue(OPT_mfloat_abi)); - Opts.LessPreciseFPMAD = Args.hasArg(OPT_cl_mad_enable) || - Args.hasArg(OPT_cl_unsafe_math_optimizations) || - Args.hasArg(OPT_cl_fast_relaxed_math); Opts.LimitFloatPrecision = std::string(Args.getLastArgValue(OPT_mlimit_float_precision)); - Opts.CorrectlyRoundedDivSqrt = - Args.hasArg(OPT_cl_fp32_correctly_rounded_divide_sqrt); - Opts.UniformWGSize = - Args.hasArg(OPT_cl_uniform_work_group_size); Opts.Reciprocals = Args.getAllArgValues(OPT_mrecip_EQ); Opts.StrictFloatCastOverflow = !Args.hasArg(OPT_fno_strict_float_cast_overflow); @@ -3021,8 +3014,6 @@ Opts.DumpVTableLayouts = Args.hasArg(OPT_fdump_vtable_layouts); Opts.SpellChecking = !Args.hasArg(OPT_fno_spell_checking); Opts.NoBitFieldTypeAlign = Args.hasArg(OPT_fno_bitfield_type_align); - Opts.SinglePrecisionConstants = Args.hasArg(OPT_cl_single_precision_constant); - Opts.FastRelaxedMath = Args.hasArg(OPT_cl_fast_relaxed_math); if (Opts.FastRelaxedMath) Opts.setDefaultFPContractMode(LangOptions::FPM_Fast); Opts.HexagonQdsp6Compat = Args.hasArg(OPT_mqdsp6_compat); @@ -3290,47 +3281,6 @@ if (InlineArg->getOption().matches(options::OPT_fno_inline)) Opts.NoInlineDefine = true; - Opts.FastMath = - Args.hasArg(OPT_ffast_math) || Args.hasArg(OPT_cl_fast_relaxed_math); - Opts.FiniteMathOnly = Args.hasArg(OPT_ffinite_math_only) || - Args.hasArg(OPT_ffast_math) || - Args.hasArg(OPT_cl_finite_math_only) || - Args.hasArg(OPT_cl_fast_relaxed_math); - Opts.UnsafeFPMath = Args.hasArg(OPT_menable_unsafe_fp_math) || - Args.hasArg(OPT_ffast_math) || - Args.hasArg(OPT_cl_unsafe_math_optimizations) || - Args.hasArg(OPT_cl_fast_relaxed_math); - Opts.AllowFPReassoc = Args.hasArg(OPT_mreassociate) || - Args.hasArg(OPT_menable_unsafe_fp_math) || - Args.hasArg(OPT_ffast_math) || - Args.hasArg(OPT_cl_unsafe_math_optimizations) || - Args.hasArg(OPT_cl_fast_relaxed_math); - Opts.NoHonorNaNs = - Args.hasArg(OPT_menable_no_nans) || Args.hasArg(OPT_ffinite_math_only) || - Args.hasArg(OPT_ffast_math) || Args.hasArg(OPT_cl_finite_math_only) || - Args.hasArg(OPT_cl_fast_relaxed_math); - Opts.NoHonorInfs = Args.hasArg(OPT_menable_no_infinities) || - Args.hasArg(OPT_ffinite_math_only) || - Args.hasArg(OPT_ffast_math) || - Args.hasArg(OPT_cl_finite_math_only) || - Args.hasArg(OPT_cl_fast_relaxed_math); - Opts.NoSignedZero = Args.hasArg(OPT_fno_signed_zeros) || - Args.hasArg(OPT_menable_unsafe_fp_math) || - Args.hasArg(OPT_ffast_math) || - Args.hasArg(OPT_cl_no_signed_zeros) || - Args.hasArg(OPT_cl_unsafe_math_optimizations) || - Args.hasArg(OPT_cl_fast_relaxed_math); - Opts.AllowRecip = Args.hasArg(OPT_freciprocal_math) || - Args.hasArg(OPT_menable_unsafe_fp_math) || - Args.hasArg(OPT_ffast_math) || - Args.hasArg(OPT_cl_unsafe_math_optimizations) || - Args.hasArg(OPT_cl_fast_relaxed_math); - // Currently there's no clang option to enable this individually - Opts.ApproxFunc = Args.hasArg(OPT_menable_unsafe_fp_math) || - Args.hasArg(OPT_ffast_math) || - Args.hasArg(OPT_cl_unsafe_math_optimizations) || - Args.hasArg(OPT_cl_fast_relaxed_math); - if (Arg *A = Args.getLastArg(OPT_ffp_contract)) { StringRef Val = A->getValue(); if (Val == "fast") Index: clang/test/CodeGen/complex-math.c =================================================================== --- clang/test/CodeGen/complex-math.c +++ clang/test/CodeGen/complex-math.c @@ -5,7 +5,9 @@ // RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple armv7-none-linux-gnueabi -o - | FileCheck %s --check-prefix=ARM // RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple armv7-none-linux-gnueabihf -o - | FileCheck %s --check-prefix=ARMHF // RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple thumbv7k-apple-watchos2.0 -o - -target-abi aapcs16 | FileCheck %s --check-prefix=ARM7K -// RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple aarch64-unknown-unknown -ffast-math -ffp-contract=fast -o - | FileCheck %s --check-prefix=AARCH64-FASTMATH +// RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple aarch64-unknown-unknown -ffast-math \ +// RUN: -menable-unsafe-fp-math -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math -fapprox-func \ +// RUN: -mreassociate -ffinite-math-only -ffp-contract=fast -o - | FileCheck %s --check-prefix=AARCH64-FASTMATH float _Complex add_float_rr(float a, float b) { // X86-LABEL: @add_float_rr( Index: clang/test/CodeGen/fast-math.c =================================================================== --- clang/test/CodeGen/fast-math.c +++ clang/test/CodeGen/fast-math.c @@ -1,4 +1,7 @@ -// RUN: %clang_cc1 -ffast-math -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -ffast-math -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math \ +// RUN: -menable-unsafe-fp-math -ffinite-math-only -fapprox-func -mreassociate -ffp-contract=fast \ +// RUN: -emit-llvm -o - %s | FileCheck %s + float f0, f1, f2; void foo(void) { Index: clang/test/CodeGen/finite-math.c =================================================================== --- clang/test/CodeGen/finite-math.c +++ clang/test/CodeGen/finite-math.c @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -ffinite-math-only -emit-llvm -o - %s | FileCheck %s -check-prefix=CHECK -check-prefix=FINITE +// RUN: %clang_cc1 -ffinite-math-only -menable-no-nans -menable-no-infs -emit-llvm -o - \ +// RUN: %s | FileCheck %s -check-prefix=CHECK -check-prefix=FINITE // RUN: %clang_cc1 -fno-signed-zeros -emit-llvm -o - %s | FileCheck %s -check-prefix=CHECK -check-prefix=NSZ // RUN: %clang_cc1 -freciprocal-math -emit-llvm -o - %s | FileCheck %s -check-prefix=CHECK -check-prefix=RECIP // RUN: %clang_cc1 -mreassociate -emit-llvm -o - %s | FileCheck %s -check-prefix=CHECK -check-prefix=REASSOC Index: clang/test/CodeGen/fp-floatcontrol-stack.cpp =================================================================== --- clang/test/CodeGen/fp-floatcontrol-stack.cpp +++ clang/test/CodeGen/fp-floatcontrol-stack.cpp @@ -1,6 +1,8 @@ // RUN: %clang_cc1 -triple x86_64-linux-gnu -ffp-contract=on -DDEFAULT=1 -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DDEFAULT %s // RUN: %clang_cc1 -triple x86_64-linux-gnu -ffp-contract=on -DEBSTRICT=1 -ffp-exception-behavior=strict -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DEBSTRICT %s -// RUN: %clang_cc1 -triple x86_64-linux-gnu -DFAST=1 -ffast-math -ffp-contract=fast -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-FAST %s +// RUN: %clang_cc1 -triple x86_64-linux-gnu -DFAST=1 -ffast-math -ffinite-math-only -menable-unsafe-fp-math \ +// RUN: -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math -fapprox-func -mreassociate \ +// RUN: -ffp-contract=fast -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-FAST %s // RUN: %clang_cc1 -triple x86_64-linux-gnu -ffp-contract=on -DNOHONOR=1 -menable-no-infs -menable-no-nans -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-NOHONOR %s #define FUN(n) \ Index: clang/test/CodeGen/fp-function-attrs.cpp =================================================================== --- clang/test/CodeGen/fp-function-attrs.cpp +++ clang/test/CodeGen/fp-function-attrs.cpp @@ -1,4 +1,6 @@ -// RUN: %clang_cc1 -triple x86_64-linux-gnu -ffast-math -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-linux-gnu -ffast-math -ffinite-math-only -menable-unsafe-fp-math \ +// RUN: -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math \ +// RUN: -fapprox-func -mreassociate -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s float test_default(float a, float b, float c) { float tmp = a; Index: clang/test/CodeGen/fp-options-to-fast-math-flags.c =================================================================== --- clang/test/CodeGen/fp-options-to-fast-math-flags.c +++ clang/test/CodeGen/fp-options-to-fast-math-flags.c @@ -1,12 +1,15 @@ // RUN: %clang_cc1 -triple x86_64-unknown-unknown -emit-llvm -o - %s | FileCheck -check-prefix CHECK-PRECISE %s // RUN: %clang_cc1 -triple x86_64-unknown-unknown -menable-no-nans -emit-llvm -o - %s | FileCheck -check-prefix CHECK-NO-NANS %s // RUN: %clang_cc1 -triple x86_64-unknown-unknown -menable-no-infs -emit-llvm -o - %s | FileCheck -check-prefix CHECK-NO-INFS %s -// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ffinite-math-only -emit-llvm -o - %s | FileCheck -check-prefix CHECK-FINITE %s +// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ffinite-math-only -menable-no-infs -menable-no-nans -emit-llvm -o - %s | FileCheck -check-prefix CHECK-FINITE %s // RUN: %clang_cc1 -triple x86_64-unknown-unknown -fno-signed-zeros -emit-llvm -o - %s | FileCheck -check-prefix CHECK-NO-SIGNED-ZEROS %s // RUN: %clang_cc1 -triple x86_64-unknown-unknown -mreassociate -emit-llvm -o - %s | FileCheck -check-prefix CHECK-REASSOC %s // RUN: %clang_cc1 -triple x86_64-unknown-unknown -freciprocal-math -emit-llvm -o - %s | FileCheck -check-prefix CHECK-RECIP %s -// RUN: %clang_cc1 -triple x86_64-unknown-unknown -menable-unsafe-fp-math -emit-llvm -o - %s | FileCheck -check-prefix CHECK-UNSAFE %s -// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ffast-math -emit-llvm -o - %s | FileCheck -check-prefix CHECK-FAST %s +// RUN: %clang_cc1 -triple x86_64-unknown-unknown -menable-unsafe-fp-math -freciprocal-math \ +// RUN: -mreassociate -fno-signed-zeros -fapprox-func -emit-llvm -o - %s | FileCheck -check-prefix CHECK-UNSAFE %s +// RUN: %clang_cc1 -triple x86_64-unknown-unknown -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans \ +// RUN: -fno-signed-zeros -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -emit-llvm \ +// RUN: -o - %s | FileCheck -check-prefix CHECK-FAST %s float fn(float); Index: clang/test/CodeGen/fpconstrained.c =================================================================== --- clang/test/CodeGen/fpconstrained.c +++ clang/test/CodeGen/fpconstrained.c @@ -1,11 +1,23 @@ // RUN: %clang_cc1 -ftrapping-math -frounding-math -ffp-exception-behavior=strict -fexperimental-strict-floating-point -emit-llvm -o - %s | FileCheck %s -check-prefix=FPMODELSTRICT // RUN: %clang_cc1 -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s -check-prefix=PRECISE -// RUN: %clang_cc1 -ffast-math -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST -// RUN: %clang_cc1 -ffast-math -emit-llvm -o - %s | FileCheck %s -check-prefix=FASTNOCONTRACT -// RUN: %clang_cc1 -ffast-math -ffp-contract=fast -ffp-exception-behavior=ignore -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST -// RUN: %clang_cc1 -ffast-math -ffp-contract=fast -ffp-exception-behavior=strict -fexperimental-strict-floating-point -emit-llvm -o - %s | FileCheck %s -check-prefix=EXCEPT -// RUN: %clang_cc1 -ffast-math -ffp-contract=fast -ffp-exception-behavior=maytrap -fexperimental-strict-floating-point -emit-llvm -o - %s | FileCheck %s -check-prefix=MAYTRAP - +// RUN: %clang_cc1 -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -ffp-contract=fast \ +// RUN: -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST +// RUN: %clang_cc1 -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate \ +// RUN: -emit-llvm -o - %s | FileCheck %s -check-prefix=FASTNOCONTRACT +// RUN: %clang_cc1 -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -ffp-contract=fast \ +// RUN: -ffp-exception-behavior=ignore -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST +// RUN: %clang_cc1 -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -ffp-contract=fast \ +// RUN: -ffp-exception-behavior=strict -fexperimental-strict-floating-point -emit-llvm -o - %s \ +// RUN: | FileCheck %s -check-prefix=EXCEPT +// RUN: %clang_cc1 -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -ffp-contract=fast \ +// RUN: -ffp-exception-behavior=maytrap -fexperimental-strict-floating-point -emit-llvm -o - %s \ +// RUN: | FileCheck %s -check-prefix=MAYTRAP + float f0, f1, f2; void foo() { Index: clang/test/CodeGen/fpconstrained.cpp =================================================================== --- clang/test/CodeGen/fpconstrained.cpp +++ clang/test/CodeGen/fpconstrained.cpp @@ -1,11 +1,23 @@ // RUN: %clang_cc1 -x c++ -ftrapping-math -fexceptions -fcxx-exceptions -frounding-math -ffp-exception-behavior=strict -fexperimental-strict-floating-point -emit-llvm -o - %s | FileCheck %s -check-prefix=FPMODELSTRICT // RUN: %clang_cc1 -x c++ -ffp-contract=fast -fexceptions -fcxx-exceptions -emit-llvm -o - %s | FileCheck %s -check-prefix=PRECISE -// RUN: %clang_cc1 -x c++ -ffast-math -fexceptions -fcxx-exceptions -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST -// RUN: %clang_cc1 -x c++ -ffast-math -fexceptions -fcxx-exceptions -emit-llvm -o - %s | FileCheck %s -check-prefix=FASTNOCONTRACT -// RUN: %clang_cc1 -x c++ -ffast-math -fexceptions -fcxx-exceptions -ffp-contract=fast -ffp-exception-behavior=ignore -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST -// RUN: %clang_cc1 -x c++ -ffast-math -fexceptions -fcxx-exceptions -ffp-contract=fast -ffp-exception-behavior=strict -fexperimental-strict-floating-point -emit-llvm -o - %s | FileCheck %s -check-prefix=EXCEPT -// RUN: %clang_cc1 -x c++ -ffast-math -fexceptions -fcxx-exceptions -ffp-contract=fast -ffp-exception-behavior=maytrap -fexperimental-strict-floating-point -emit-llvm -o - %s | FileCheck %s -check-prefix=MAYTRAP - +// RUN: %clang_cc1 -x c++ -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -fexceptions -fcxx-exceptions \ +// RUN: -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST +// RUN: %clang_cc1 -x c++ -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -fexceptions -fcxx-exceptions \ +// RUN: -emit-llvm -o - %s | FileCheck %s -check-prefix=FASTNOCONTRACT +// RUN: %clang_cc1 -x c++ -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -fexceptions -fcxx-exceptions \ +// RUN: -ffp-contract=fast -ffp-exception-behavior=ignore -emit-llvm -o - %s | FileCheck %s -check-prefix=FAST +// RUN: %clang_cc1 -x c++ -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -fexceptions -fcxx-exceptions \ +// RUN: -ffp-contract=fast -ffp-exception-behavior=strict -fexperimental-strict-floating-point -emit-llvm \ +// RUN: -o - %s | FileCheck %s -check-prefix=EXCEPT +// RUN: %clang_cc1 -x c++ -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros \ +// RUN: -freciprocal-math -menable-unsafe-fp-math -fapprox-func -mreassociate -fexceptions -fcxx-exceptions \ +// RUN: -ffp-contract=fast -ffp-exception-behavior=maytrap -fexperimental-strict-floating-point -emit-llvm \ +// RUN: -o - %s | FileCheck %s -check-prefix=MAYTRAP + float f0, f1, f2; template Index: clang/test/CodeGen/libcalls.c =================================================================== --- clang/test/CodeGen/libcalls.c +++ clang/test/CodeGen/libcalls.c @@ -1,6 +1,7 @@ // RUN: %clang_cc1 -fmath-errno -emit-llvm -o - %s -triple i386-unknown-unknown | FileCheck -check-prefix CHECK-YES %s // RUN: %clang_cc1 -emit-llvm -o - %s -triple i386-unknown-unknown | FileCheck -check-prefix CHECK-NO %s -// RUN: %clang_cc1 -menable-unsafe-fp-math -emit-llvm -o - %s -triple i386-unknown-unknown | FileCheck -check-prefix CHECK-FAST %s +// RUN: %clang_cc1 -menable-unsafe-fp-math -fapprox-func -fno-signed-zeros -mreassociate -freciprocal-math \ +// RUN: -emit-llvm -o - %s -triple i386-unknown-unknown | FileCheck -check-prefix CHECK-FAST %s // CHECK-YES-LABEL: define void @test_sqrt // CHECK-NO-LABEL: define void @test_sqrt Index: clang/test/CodeGenOpenCL/no-signed-zeros.cl =================================================================== --- clang/test/CodeGenOpenCL/no-signed-zeros.cl +++ clang/test/CodeGenOpenCL/no-signed-zeros.cl @@ -1,5 +1,5 @@ // RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s -check-prefix=NORMAL -// RUN: %clang_cc1 %s -emit-llvm -cl-no-signed-zeros -o - | FileCheck %s -check-prefix=NO-SIGNED-ZEROS +// RUN: %clang_cc1 %s -emit-llvm -fno-signed-zeros -o - | FileCheck %s -check-prefix=NO-SIGNED-ZEROS float signedzeros(float a) { return a; Index: clang/test/CodeGenOpenCL/relaxed-fpmath.cl =================================================================== --- clang/test/CodeGenOpenCL/relaxed-fpmath.cl +++ clang/test/CodeGenOpenCL/relaxed-fpmath.cl @@ -1,18 +1,25 @@ // RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s -check-prefix=NORMAL -// RUN: %clang_cc1 %s -emit-llvm -cl-fast-relaxed-math -o - | FileCheck %s -check-prefix=FAST -// RUN: %clang_cc1 %s -emit-llvm -cl-finite-math-only -o - | FileCheck %s -check-prefix=FINITE -// RUN: %clang_cc1 %s -emit-llvm -cl-unsafe-math-optimizations -o - | FileCheck %s -check-prefix=UNSAFE +// RUN: %clang_cc1 %s -emit-llvm -cl-fast-relaxed-math -cl-mad-enable -ffast-math -ffinite-math-only \ +// RUN: -menable-unsafe-fp-math -mreassociate -menable-no-nans -menable-no-infs -fno-signed-zeros \ +// RUN: -freciprocal-math -fapprox-func -o - | FileCheck %s -check-prefix=FAST +// RUN: %clang_cc1 %s -emit-llvm -ffinite-math-only -menable-no-nans -menable-no-infs -o - | FileCheck %s -check-prefix=FINITE +// RUN: %clang_cc1 %s -emit-llvm -cl-mad-enable -menable-unsafe-fp-math -mreassociate -fno-signed-zeros \ +// RUN: -freciprocal-math -fapprox-func -o - | FileCheck %s -check-prefix=UNSAFE // RUN: %clang_cc1 %s -emit-llvm -cl-mad-enable -o - | FileCheck %s -check-prefix=MAD -// RUN: %clang_cc1 %s -emit-llvm -cl-no-signed-zeros -o - | FileCheck %s -check-prefix=NOSIGNED +// RUN: %clang_cc1 %s -emit-llvm -fno-signed-zeros -o - | FileCheck %s -check-prefix=NOSIGNED // Check the fp options are correct with PCH. // RUN: %clang_cc1 %s -DGEN_PCH=1 -finclude-default-header -triple spir-unknown-unknown -emit-pch -o %t.pch // RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -o - | FileCheck %s -check-prefix=NORMAL -// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -cl-fast-relaxed-math -o - | FileCheck %s -check-prefix=FAST -// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -cl-finite-math-only -o - | FileCheck %s -check-prefix=FINITE -// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -cl-unsafe-math-optimizations -o - | FileCheck %s -check-prefix=UNSAFE +// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -cl-fast-relaxed-math -cl-mad-enable -ffast-math \ +// RUN: -ffinite-math-only -menable-unsafe-fp-math -mreassociate -menable-no-nans -menable-no-infs -fno-signed-zeros \ +// RUN: -freciprocal-math -fapprox-func -o - | FileCheck %s -check-prefix=FAST +// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -ffinite-math-only -menable-no-nans -menable-no-infs -o - \ +// RUN: | FileCheck %s -check-prefix=FINITE +// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -cl-mad-enable -menable-unsafe-fp-math -mreassociate \ +// RUN: -fno-signed-zeros -freciprocal-math -fapprox-func -o - | FileCheck %s -check-prefix=UNSAFE // RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -cl-mad-enable -o - | FileCheck %s -check-prefix=MAD -// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -cl-no-signed-zeros -o - | FileCheck %s -check-prefix=NOSIGNED +// RUN: %clang_cc1 %s -include-pch %t.pch -fno-validate-pch -emit-llvm -fno-signed-zeros -o - | FileCheck %s -check-prefix=NOSIGNED #if !GEN_PCH Index: clang/test/Driver/opencl.cl =================================================================== --- clang/test/Driver/opencl.cl +++ clang/test/Driver/opencl.cl @@ -28,12 +28,12 @@ // CHECK-OPT-DISABLE: "-cc1" {{.*}} "-cl-opt-disable" // CHECK-STRICT-ALIASING: "-cc1" {{.*}} "-cl-strict-aliasing" // CHECK-SINGLE-PRECISION-CONST: "-cc1" {{.*}} "-cl-single-precision-constant" -// CHECK-FINITE-MATH-ONLY: "-cc1" {{.*}} "-cl-finite-math-only" +// CHECK-FINITE-MATH-ONLY: "-cc1" {{.*}} "-ffinite-math-only" "-menable-no-nans" "-menable-no-infs" // CHECK-KERNEL-ARG-INFO: "-cc1" {{.*}} "-cl-kernel-arg-info" -// CHECK-UNSAFE-MATH-OPT: "-cc1" {{.*}} "-cl-unsafe-math-optimizations" -// CHECK-FAST-RELAXED-MATH: "-cc1" {{.*}} "-cl-fast-relaxed-math" +// CHECK-UNSAFE-MATH-OPT: "-cc1" {{.*}} "-cl-mad-enable" "-menable-unsafe-fp-math" "-mreassociate" "-fno-signed-zeros" "-freciprocal-math" "-fapprox-func" +// CHECK-FAST-RELAXED-MATH: "-cc1" {{.*}} "-cl-fast-relaxed-math" "-cl-mad-enable" "-ffast-math" "-ffinite-math-only" "-menable-unsafe-fp-math" "-mreassociate" "-menable-no-nans" "-menable-no-infs" "-fno-signed-zeros" "-freciprocal-math" "-fapprox-func" // CHECK-MAD-ENABLE: "-cc1" {{.*}} "-cl-mad-enable" -// CHECK-NO-SIGNED-ZEROS: "-cc1" {{.*}} "-cl-no-signed-zeros" +// CHECK-NO-SIGNED-ZEROS: "-cc1" {{.*}} "-fno-signed-zeros" // This is not forwarded // CHECK-DENORMS-ARE-ZERO-NOT: "-cl-denorms-are-zero" Index: clang/test/Headers/nvptx_device_math_sin.c =================================================================== --- clang/test/Headers/nvptx_device_math_sin.c +++ clang/test/Headers/nvptx_device_math_sin.c @@ -1,8 +1,15 @@ // REQUIRES: nvptx-registered-target // RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc // RUN: %clang_cc1 -x c -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=SLOW -// RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast -// RUN: %clang_cc1 -x c -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=FAST +// RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown \ +// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffinite-math-only \ +// RUN: -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math -menable-unsafe-fp-math \ +// RUN: -fapprox-func -mreassociate -ffp-contract=fast +// RUN: %clang_cc1 -x c -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers \ +// RUN: -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown \ +// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc \ +// RUN: -o - -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math \ +// RUN: -menable-unsafe-fp-math -fapprox-func -mreassociate -ffp-contract=fast | FileCheck %s --check-prefix=FAST // expected-no-diagnostics #include Index: clang/test/Headers/nvptx_device_math_sin.cpp =================================================================== --- clang/test/Headers/nvptx_device_math_sin.cpp +++ clang/test/Headers/nvptx_device_math_sin.cpp @@ -1,8 +1,15 @@ // REQUIRES: nvptx-registered-target // RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc // RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=SLOW -// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast -// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=FAST +// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown \ +// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffinite-math-only \ +// RUN: -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math -menable-unsafe-fp-math \ +// RUN: -fapprox-func -mreassociate -ffp-contract=fast +// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers \ +// RUN: -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown \ +// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc \ +// RUN: -o - -ffast-math -ffinite-math-only -menable-no-infs -menable-no-nans -fno-signed-zeros -freciprocal-math \ +// RUN: -menable-unsafe-fp-math -fapprox-func -mreassociate -ffp-contract=fast | FileCheck %s --check-prefix=FAST // expected-no-diagnostics #include Index: clang/test/SemaOpenCL/fp-options.cl =================================================================== --- clang/test/SemaOpenCL/fp-options.cl +++ clang/test/SemaOpenCL/fp-options.cl @@ -1,4 +1,4 @@ // RUN: %clang_cc1 %s -finclude-default-header -triple spir-unknown-unknown -emit-pch -o %t.pch -// RUN: %clang_cc1 %s -finclude-default-header -cl-no-signed-zeros -triple spir-unknown-unknown -include-pch %t.pch -fsyntax-only -verify +// RUN: %clang_cc1 %s -finclude-default-header -fno-signed-zeros -triple spir-unknown-unknown -include-pch %t.pch -fsyntax-only -verify // expected-no-diagnostics Index: llvm/utils/TableGen/OptParserEmitter.cpp =================================================================== --- llvm/utils/TableGen/OptParserEmitter.cpp +++ llvm/utils/TableGen/OptParserEmitter.cpp @@ -443,7 +443,8 @@ OS << "OPTION("; WriteOptRecordFields(OS, R); OS << ")\n"; - if (!isa(R.getValueInit("MarshallingKind"))) + if (!isa(R.getValueInit("MarshallingKind")) && + !R.getValueAsString("KeyPath").empty()) OptsWithMarshalling.push_back(MarshallingKindInfo::create(R)); } OS << "#endif // OPTION\n";