diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst --- a/clang/docs/LanguageExtensions.rst +++ b/clang/docs/LanguageExtensions.rst @@ -3045,6 +3045,38 @@ section of the code. This can be useful when fast contraction is otherwise enabled for the translation unit with the ``-ffp-contract=fast`` flag. +The ``#pragma float_control`` pragma allows precise floating-point +semantics and floating-point exception behavior to be specified +for a section of the source code. This pragma can only appear at file scope or +at the start of a compound statement (excluding comments). When using within a +compound statement, the pragma is active within the scope of the compound +statement. This pragma is modeled after a Microsoft pragma with the +same spelling and syntax. For pragmas specified at file scope, a stack +is supported so that the pragma float_control settings can be pushed or popped. + +When ``float_control(precise, on)`` is enabled, the section of code governed +by the pragma behaves as though the command-line option ``ffp-model=precise`` +is enabled. That is, fast-math is disabled and fp-contract=on (fused +multiply add) is enabled. + +When ``float_control(except, on)`` is enabled, the section of code governed +by the pragma behaves as though the command-line + ``ffp-exception-behavior=strict`` is enabled, ``float-control(precise, off)`` +selects ``ffp-exception-behavior=ignore``. + +The full syntax this pragma supports is +``float_control(except|precise, on|off [, push])`` and +``float_control(push|pop)``. +The ``push`` and ``pop`` forms can only occur at file scope. + +.. code-block:: c++ + + for(...) { + // This block will be compiled with fno-fast-math and ffp-contract=on + #pragma float_control(precise, on) + a = b[i] * c[i] + e; + } + Specifying an attribute for multiple declarations (#pragma clang attribute) =========================================================================== diff --git a/clang/docs/UsersManual.rst b/clang/docs/UsersManual.rst --- a/clang/docs/UsersManual.rst +++ b/clang/docs/UsersManual.rst @@ -1190,8 +1190,50 @@ Controlling Floating Point Behavior ----------------------------------- -Clang provides a number of ways to control floating point behavior. The options -are listed below. +Clang provides a number of ways to control floating point behavior, including +with command line options and source pragmas. This section +describes the various floating point semantic modes and the corresponding options. + +.. csv-table:: Floating Point Semantic Modes + :header: "Mode", "Values" + :widths: 15, 30, 30 + + "except_behavior", "{ignore, strict, may_trap}", "ffp-exception-behavior" + "fenv_access", "{off, on}", "(none)" + "rounding_mode", "{dynamic, tonearest, downward, upward, towardzero}", "frounding-math" + "contract", "{on, off, fast}", "ffp-contract" + "denormal_fp_math", "{IEEE, PreserveSign, PositiveZero}", "fdenormal-fp-math" + "denormal_fp32_math", "{IEEE, PreserveSign, PositiveZero}", "fdenormal-fp-math-fp32" + "support_math_errno", "{on, off}", "fmath-errno" + "no_honor_nans", "{on, off}", "fhonor-nans" + "no_honor_infinities", "{on, off}", "fhonor-infinities" + "no_signed_zeros", "{on, off}", "fsigned-zeros" + "allow_reciprocal", "{on, off}", "freciprocal-math" + "allow_approximate_fns", "{on, off}", "(none)" + "allow_reassociation", "{on, off}", "fassociative-math" + + +This table describes the option settings that correspond to the three +floating point semantic models: precise (the default), strict, and fast. + + +.. csv-table:: Floating Point Models + :header: "Mode", "Precise", "Strict", "Fast" + :widths: 25, 15, 15, 15 + + "except_behavior", "ignore", "strict", "ignore" + "fenv_access", "off", "on", "off" + "rounding_mode", "tonearest", "dynamic", "tonearest" + "contract", "on", "off", "fast" + "denormal_fp_math", "IEEE", "IEEE", "PreserveSign" + "denormal_fp32_math", "IEEE","IEEE", "PreserveSign" + "support_math_errno", "on", "on", "off" + "no_honor_nans", "off", "off", "on" + "no_honor_infinities", "off", "off", "on" + "no_signed_zeros", "off", "off", "on" + "allow_reciprocal", "off", "off", "on" + "allow_approximate_fns", "off", "off", "on" + "allow_reassociation", "off", "off", "on" .. option:: -ffast-math @@ -1385,7 +1427,7 @@ and ``fast``. Details: - * ``precise`` Disables optimizations that are not value-safe on floating-point data, although FP contraction (FMA) is enabled (``-ffp-contract=fast``). This is the default behavior. + * ``precise`` Disables optimizations that are not value-safe on floating-point data, although FP contraction (FMA) is enabled (``-ffp-contract=on``). This is the default behavior. * ``strict`` Enables ``-frounding-math`` and ``-ffp-exception-behavior=strict``, and disables contractions (FMA). All of the ``-ffast-math`` enablements are disabled. * ``fast`` Behaves identically to specifying both ``-ffast-math`` and ``ffp-contract=fast`` diff --git a/clang/include/clang/AST/Stmt.h b/clang/include/clang/AST/Stmt.h --- a/clang/include/clang/AST/Stmt.h +++ b/clang/include/clang/AST/Stmt.h @@ -533,7 +533,7 @@ /// This is only meaningful for operations on floating point /// types and 0 otherwise. - unsigned FPFeatures : 8; + unsigned FPFeatures : 14; SourceLocation OpLoc; }; @@ -604,7 +604,7 @@ unsigned OperatorKind : 6; // Only meaningful for floating point types. - unsigned FPFeatures : 8; + unsigned FPFeatures : 14; }; class CXXRewrittenBinaryOperatorBitfields { @@ -1100,7 +1100,7 @@ Stmt &operator=(Stmt &&) = delete; Stmt(StmtClass SC) { - static_assert(sizeof(*this) <= 8, + static_assert(sizeof(*this) <= 16, "changing bitfields changed sizeof(Stmt)"); static_assert(sizeof(*this) % alignof(void *) == 0, "Insufficient alignment!"); diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td --- a/clang/include/clang/Basic/DiagnosticParseKinds.td +++ b/clang/include/clang/Basic/DiagnosticParseKinds.td @@ -1109,9 +1109,9 @@ "'#pragma init_seg' is only supported when targeting a " "Microsoft environment">, InGroup; -// - #pragma fp_contract -def err_pragma_fp_contract_scope : Error< - "'#pragma fp_contract' can only appear at file scope or at the start of a " +// - #pragma restricted to file scope or start of compound statement +def err_pragma_file_or_compound_scope : Error< + "'#pragma %0' can only appear at file scope or at the start of a " "compound statement">; // - #pragma stdc unknown def ext_stdc_pragma_ignored : ExtWarn<"unknown pragma in STDC namespace">, @@ -1130,6 +1130,12 @@ def err_pragma_detect_mismatch_malformed : Error< "pragma detect_mismatch is malformed; it requires two comma-separated " "string literals">; +// - #pragma float_control +def err_pragma_float_control_malformed : Error< + "pragma float_control is malformed; use 'float_control({push|pop})' or " + "'float_control({precise|except}, {on|off} [,push])'">; +def err_pragma_float_control_unknown_kind : Error< + "unknown kind of pragma float_control">; // - #pragma pointers_to_members def err_pragma_pointers_to_members_unknown_kind : Error< "unexpected %0, expected to see one of %select{|'best_case', 'full_generality', }1" @@ -1297,9 +1303,8 @@ def err_pragma_fp_invalid_argument : Error< "unexpected argument '%0' to '#pragma clang fp %1'; " "expected 'on', 'fast' or 'off'">; -def err_pragma_fp_scope : Error< - "'#pragma clang fp' can only appear at file scope or at the start of a " - "compound statement">; +def err_pragma_fc_pp_scope : Error< + "'#pragma float_control push/pop' can only appear at file scope">; def err_pragma_invalid_keyword : Error< "invalid argument; expected 'enable'%select{|, 'full'}0%select{|, 'assume_safety'}1 or 'disable'">; diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h --- a/clang/include/clang/Basic/LangOptions.h +++ b/clang/include/clang/Basic/LangOptions.h @@ -359,25 +359,48 @@ FPOptions() : fp_contract(LangOptions::FPC_Off), fenv_access(LangOptions::FEA_Off), rounding(LangOptions::FPR_ToNearest), - exceptions(LangOptions::FPE_Ignore) - {} + exceptions(LangOptions::FPE_Ignore), + allow_reassoc(0), + no_nans(0), + no_infs(0), + no_signed_zeros(0), + allow_reciprocal(0), + approx_func(0) + {} // Used for serializing. explicit FPOptions(unsigned I) : fp_contract(static_cast(I & 3)), fenv_access(static_cast((I >> 2) & 1)), rounding(static_cast((I >> 3) & 7)), - exceptions(static_cast((I >> 6) & 3)) + exceptions(static_cast((I >> 6) & 3)), + allow_reassoc((I>>8) & 1), + no_nans((I>>9) & 1), + no_infs((I>>10) & 1), + no_signed_zeros((I>>11) & 1), + allow_reciprocal((I>>12) & 1), + approx_func((I>>13) & 1) {} explicit FPOptions(const LangOptions &LangOpts) : fp_contract(LangOpts.getDefaultFPContractMode()), fenv_access(LangOptions::FEA_Off), - rounding(LangOptions::FPR_ToNearest), - exceptions(LangOptions::FPE_Ignore) + rounding(LangOpts.getFPRoundingMode()), + exceptions(LangOpts.getFPExceptionMode()), + allow_reassoc(LangOpts.FastMath || LangOpts.AllowFPReassoc), + no_nans(LangOpts.FastMath || LangOpts.NoHonorNaNs), + no_infs(LangOpts.FastMath || LangOpts.NoHonorInfs), + no_signed_zeros(LangOpts.FastMath || LangOpts.NoSignedZero), + allow_reciprocal(LangOpts.FastMath || LangOpts.AllowRecip), + approx_func(LangOpts.FastMath || LangOpts.ApproxFunc) {} // FIXME: Use getDefaultFEnvAccessMode() when available. + void setFastMath(bool B = true) { + allow_reassoc = no_nans = no_infs = no_signed_zeros = approx_func = + allow_reciprocal = B; + } + bool allowFPContractWithinStatement() const { return fp_contract == LangOptions::FPC_On; } @@ -404,6 +427,18 @@ fenv_access = LangOptions::FEA_On; } + void setFPPreciseEnabled(bool Value) { + if (Value) { + /* Precise mode implies fp_contract=on and disables ffast-math */ + setFastMath(false); + setAllowFPContractWithinStatement(); + } else { + /* Precise mode implies fp_contract=fast and enables ffast-math */ + setFastMath(true); + setAllowFPContractAcrossStatement(); + } + } + void setDisallowFEnvAccess() { fenv_access = LangOptions::FEA_Off; } LangOptions::FPRoundingModeKind getRoundingMode() const { @@ -422,6 +457,34 @@ exceptions = EM; } + /// Flag queries + bool allowReassoc() const { return allow_reassoc; } + bool noNaNs() const { return no_nans; } + bool noInfs() const { return no_infs; } + bool noSignedZeros() const { return no_signed_zeros; } + bool allowReciprocal() const { return allow_reciprocal; } + bool approxFunc() const { return approx_func; } + + /// Flag setters + void setAllowReassoc(bool B = true) { + allow_reassoc = B; + } + void setNoNaNs(bool B = true) { + no_nans = B; + } + void setNoInfs(bool B = true) { + no_infs = B; + } + void setNoSignedZeros(bool B = true) { + no_signed_zeros = B; + } + void setAllowReciprocal(bool B = true) { + allow_reciprocal = B; + } + void setApproxFunc(bool B = true) { + approx_func = B; + } + bool isFPConstrained() const { return getRoundingMode() != LangOptions::FPR_ToNearest || getExceptionMode() != LangOptions::FPE_Ignore || @@ -431,7 +494,24 @@ /// Used to serialize this. unsigned getInt() const { return fp_contract | (fenv_access << 2) | (rounding << 3) - | (exceptions << 6); + | (exceptions << 6) + | (allow_reassoc << 8) | (no_nans << 9) + | (no_infs << 10) | (no_signed_zeros << 11) + | (allow_reciprocal << 12) | (approx_func << 13); + } + + /// Used with getInt() to manage the float_control pragma stack. + void Restore(unsigned I) { + fp_contract = (static_cast(I & 3)); + fenv_access = (static_cast((I >> 2) & 1)); + rounding = (static_cast((I >> 3) & 7)); + exceptions = (static_cast((I >> 6) & 3)); + allow_reassoc = ((I>>8) & 1); + no_nans = ((I>>9) & 1); + no_infs = ((I>>10) & 1); + no_signed_zeros = ((I>>11) & 1); + allow_reciprocal = ((I>>12) & 1); + approx_func = ((I>>13) & 1); } private: @@ -442,6 +522,12 @@ unsigned fenv_access : 1; unsigned rounding : 3; unsigned exceptions : 2; + unsigned allow_reassoc : 1; + unsigned no_nans : 1; + unsigned no_infs : 1; + unsigned no_signed_zeros : 1; + unsigned allow_reciprocal : 1; + unsigned approx_func : 1; }; /// Describes the kind of translation unit being processed. diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -187,6 +187,12 @@ COMPATIBLE_LANGOPT(FastMath , 1, 0, "fast FP math optimizations, and __FAST_MATH__ predefined macro") COMPATIBLE_LANGOPT(FiniteMathOnly , 1, 0, "__FINITE_MATH_ONLY__ predefined macro") COMPATIBLE_LANGOPT(UnsafeFPMath , 1, 0, "Unsafe Floating Point Math") +COMPATIBLE_LANGOPT(AllowFPReassoc , 1, 0, "Permit Floating Point reassociation") +COMPATIBLE_LANGOPT(NoHonorNaNs , 1, 0, "Permit Floating Point optimization without regard to NaN") +COMPATIBLE_LANGOPT(NoHonorInfs , 1, 0, "Permit Floating Point optimization without regard to infinities") +COMPATIBLE_LANGOPT(NoSignedZero , 1, 0, "Permit Floating Point optimization without regard to signed zeros") +COMPATIBLE_LANGOPT(AllowRecip , 1, 0, "Permit Floating Point reciprocal") +COMPATIBLE_LANGOPT(ApproxFunc , 1, 0, "Permit Floating Point approximation") BENIGN_LANGOPT(ObjCGCBitmapPrint , 1, 0, "printing of GC's bitmap layout for __weak/__strong ivars") diff --git a/clang/include/clang/Basic/PragmaKinds.h b/clang/include/clang/Basic/PragmaKinds.h --- a/clang/include/clang/Basic/PragmaKinds.h +++ b/clang/include/clang/Basic/PragmaKinds.h @@ -25,6 +25,16 @@ PMSST_ON // #pragms ms_struct on }; +enum PragmaFloatControlKind { + PFC_Unknown, + PFC_Precise, // #pragma float_control(precise, [,on]) + PFC_NoPrecise, // #pragma float_control(precise, off) + PFC_Except, // #pragma float_control(except [,on]) + PFC_NoExcept, // #pragma float_control(except, off) + PFC_Push, // #pragma float_control(push) + PFC_Pop // #pragma float_control(pop) +}; + } #endif diff --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def --- a/clang/include/clang/Basic/TokenKinds.def +++ b/clang/include/clang/Basic/TokenKinds.def @@ -806,6 +806,11 @@ // handles them. PRAGMA_ANNOTATION(pragma_fenv_access) +// Annotation for #pragma float_control +// The lexer produces these so that they only take effect when the parser +// handles them. +PRAGMA_ANNOTATION(pragma_float_control) + // Annotation for #pragma pointers_to_members... // The lexer produces these so that they only take effect when the parser // handles them. diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -178,6 +178,7 @@ std::unique_ptr PCSectionHandler; std::unique_ptr MSCommentHandler; std::unique_ptr MSDetectMismatchHandler; + std::unique_ptr FloatControlHandler; std::unique_ptr MSPointersToMembers; std::unique_ptr MSVtorDisp; std::unique_ptr MSInitSeg; @@ -721,6 +722,10 @@ /// #pragma STDC FENV_ACCESS... void HandlePragmaFEnvAccess(); + /// Handle the annotation token produced for + /// #pragma float_control + void HandlePragmaFloatControl(); + /// \brief Handle the annotation token produced for /// #pragma clang fp ... void HandlePragmaFP(); diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -552,6 +552,9 @@ PragmaStack ConstSegStack; PragmaStack CodeSegStack; + // This stacks the current state of Sema.FPFeatures + PragmaStack FpPragmaStack; + // RAII object to push / pop sentinel slots for all MS #pragma stacks. // Actions should be performed only if we enter / exit a C++ method body. class PragmaStackSentinelRAII { @@ -9395,6 +9398,11 @@ void ActOnPragmaDetectMismatch(SourceLocation Loc, StringRef Name, StringRef Value); + /// ActOnPragmaFloatControl - Call on well-formed \#pragma float_control + void ActOnPragmaFloatControl(SourceLocation Loc, + PragmaMsStackAction Action, + PragmaFloatControlKind Value); + /// ActOnPragmaUnused - Called on well-formed '\#pragma unused'. void ActOnPragmaUnused(const Token &Identifier, Scope *curScope, diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp --- a/clang/lib/CodeGen/CGExprScalar.cpp +++ b/clang/lib/CodeGen/CGExprScalar.cpp @@ -216,7 +216,14 @@ /// Update the FastMathFlags of LLVM IR from the FPOptions in LangOptions. static void updateFastMathFlags(llvm::FastMathFlags &FMF, FPOptions FPFeatures) { - FMF.setAllowContract(FPFeatures.allowFPContractAcrossStatement()); + FMF.setAllowReassoc(FPFeatures.allowReassoc()); + FMF.setNoNaNs(FPFeatures.noNaNs()); + FMF.setNoInfs(FPFeatures.noInfs()); + FMF.setNoSignedZeros(FPFeatures.noSignedZeros()); + FMF.setAllowReciprocal(FPFeatures.allowReciprocal()); + FMF.setApproxFunc(FPFeatures.approxFunc()); + FMF.setAllowContract(FPFeatures.allowFPContractAcrossStatement() || + FPFeatures.allowFPContractWithinStatement()); } /// Propagate fast-math flags from \p Op to the instruction in \p V. @@ -423,6 +430,36 @@ Value *Visit(Expr *E) { ApplyDebugLocation DL(CGF, E); + if (BinaryOperator * BinOp = dyn_cast(E)) { + // Preserve the old values + llvm::IRBuilder<>::FastMathFlagGuard FMFG(Builder); + auto FPFeatures = BinOp->getFPFeatures(); + auto NewRoundingBehavior = ToConstrainedRoundingMD( + FPFeatures.getRoundingMode()); + Builder.setDefaultConstrainedRounding(NewRoundingBehavior); + auto NewExceptionBehavior = ToConstrainedExceptMD( + FPFeatures.getExceptionMode()); + Builder.setDefaultConstrainedExcept(NewExceptionBehavior); + auto FMF = Builder.getFastMathFlags(); + FMF.setAllowReassoc(FPFeatures.allowReassoc()); + FMF.setNoNaNs(FPFeatures.noNaNs()); + FMF.setNoInfs(FPFeatures.noInfs()); + FMF.setNoSignedZeros(FPFeatures.noSignedZeros()); + FMF.setAllowReciprocal(FPFeatures.allowReciprocal()); + FMF.setApproxFunc(FPFeatures.approxFunc()); + FMF.setAllowContract(FPFeatures.allowFPContractAcrossStatement() || + FPFeatures.allowFPContractWithinStatement()); + Builder.setFastMathFlags(FMF); + assert((CGF.CurFuncDecl==nullptr || + Builder.getIsFPConstrained() || + isa(CGF.CurFuncDecl) || + isa(CGF.CurFuncDecl) || + (NewExceptionBehavior == llvm::fp::ebIgnore && + NewRoundingBehavior == llvm::fp::rmToNearest)) && + "FPConstrained should be enabled on entire function"); + + return StmtVisitor::Visit(E); + } return StmtVisitor::Visit(E); } diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -4419,6 +4419,16 @@ } } // end namespace CodeGen + +// Map the LangOption for floating point rounding mode into +// the corresponding enum in the IR. +llvm::fp::RoundingMode ToConstrainedRoundingMD( + LangOptions::FPRoundingModeKind Kind); + +// Map the LangOption for floating point exception behavior into +// the corresponding enum in the IR. +llvm::fp::ExceptionBehavior ToConstrainedExceptMD( + LangOptions::FPExceptionModeKind Kind); } // end namespace clang #endif diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -108,7 +108,7 @@ // Map the LangOption for rounding mode into // the corresponding enum in the IR. -static llvm::fp::RoundingMode ToConstrainedRoundingMD( +llvm::fp::RoundingMode clang::ToConstrainedRoundingMD( LangOptions::FPRoundingModeKind Kind) { switch (Kind) { @@ -123,7 +123,7 @@ // Map the LangOption for exception behavior into // the corresponding enum in the IR. -static llvm::fp::ExceptionBehavior ToConstrainedExceptMD( +llvm::fp::ExceptionBehavior clang::ToConstrainedExceptMD( LangOptions::FPExceptionModeKind Kind) { switch (Kind) { @@ -140,14 +140,14 @@ auto fpExceptionBehavior = ToConstrainedExceptMD( getLangOpts().getFPExceptionMode()); + Builder.setDefaultConstrainedRounding(fpRoundingMode); + Builder.setDefaultConstrainedExcept(fpExceptionBehavior); if (fpExceptionBehavior == llvm::fp::ebIgnore && fpRoundingMode == llvm::fp::rmToNearest) // Constrained intrinsics are not used. - ; + Builder.setIsFPConstrained(false); else { Builder.setIsFPConstrained(true); - Builder.setDefaultConstrainedRounding(fpRoundingMode); - Builder.setDefaultConstrainedExcept(fpExceptionBehavior); } } @@ -912,9 +912,11 @@ if (FD->isMain()) Fn->addFnAttr(llvm::Attribute::NoRecurse); - if (const FunctionDecl *FD = dyn_cast_or_null(D)) + if (const FunctionDecl *FD = dyn_cast_or_null(D)) { + Builder.setIsFPConstrained(FD->usesFPIntrin()); if (FD->usesFPIntrin()) Fn->addFnAttr(llvm::Attribute::StrictFP); + } // If a custom alignment is used, force realigning to this alignment on // any main function which certainly will need it. diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -2520,10 +2520,9 @@ llvm::DenormalMode DenormalFPMath = DefaultDenormalFPMath; llvm::DenormalMode DenormalFP32Math = DefaultDenormalFP32Math; - StringRef FPContract = ""; + StringRef FPContract = "on"; bool StrictFPModel = false; - if (const Arg *A = Args.getLastArg(options::OPT_flimited_precision_EQ)) { CmdArgs.push_back("-mlimit-float-precision"); CmdArgs.push_back(A->getValue()); @@ -2560,12 +2559,10 @@ // ffp-model= is a Driver option, it is entirely rewritten into more // granular options before being passed into cc1. // Use the gcc option in the switch below. - if (!FPModel.empty() && !FPModel.equals(Val)) { + if (!FPModel.empty() && !FPModel.equals(Val)) D.Diag(clang::diag::warn_drv_overriding_flag_option) << Args.MakeArgString("-ffp-model=" + FPModel) << Args.MakeArgString("-ffp-model=" + Val); - FPContract = ""; - } if (Val.equals("fast")) { optID = options::OPT_ffast_math; FPModel = Val; @@ -2573,13 +2570,14 @@ } else if (Val.equals("precise")) { optID = options::OPT_ffp_contract; FPModel = Val; - FPContract = "fast"; + FPContract = "on"; PreciseFPModel = true; } else if (Val.equals("strict")) { StrictFPModel = true; optID = options::OPT_frounding_math; FPExceptionBehavior = "strict"; FPModel = Val; + FPContract = "off"; TrappingMath = true; } else D.Diag(diag::err_drv_unsupported_option_argument) @@ -2658,7 +2656,7 @@ case options::OPT_ffp_contract: { StringRef Val = A->getValue(); if (PreciseFPModel) { - // -ffp-model=precise enables ffp-contract=fast as a side effect + // -ffp-model=precise enables ffp-contract=on as a side effect // the FPContract value has already been set to a string literal // and the Val string isn't a pertinent value. ; @@ -2757,7 +2755,7 @@ // -fno_fast_math restores default denormal and fpcontract handling DenormalFPMath = DefaultDenormalFPMath; DenormalFP32Math = DefaultDenormalFP32Math; - FPContract = ""; + FPContract = "on"; break; } if (StrictFPModel) { @@ -2768,7 +2766,7 @@ !AssociativeMath && !ReciprocalMath && SignedZeros && TrappingMath && RoundingFPMath && DenormalFPMath != llvm::DenormalMode::getIEEE() && - FPContract.empty()) + (FPContract.equals("off") || FPContract.empty())) // OK: Current Arg doesn't conflict with -ffp-model=strict ; else { diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2476,6 +2476,7 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, const TargetOptions &TargetOpts, PreprocessorOptions &PPOpts, + CodeGenOptions &CGOpts, DiagnosticsEngine &Diags) { // FIXME: Cleanup per-file based stuff. LangStandard::Kind LangStd = LangStandard::lang_unspecified; @@ -3187,6 +3188,15 @@ Opts.UnsafeFPMath = Args.hasArg(OPT_menable_unsafe_fp_math) || Args.hasArg(OPT_cl_unsafe_math_optimizations) || Args.hasArg(OPT_cl_fast_relaxed_math); + Opts.AllowFPReassoc = Opts.FastMath || CGOpts.Reassociate; + Opts.NoHonorNaNs = Opts.FastMath || CGOpts.NoNaNsFPMath || + Opts.FiniteMathOnly; + Opts.NoHonorInfs = Opts.FastMath || CGOpts.NoInfsFPMath || + Opts.FiniteMathOnly; + Opts.NoSignedZero = Opts.FastMath || CGOpts.NoSignedZeros; + Opts.AllowRecip = Opts.FastMath || CGOpts.ReciprocalMath; + // Currently there's no clang option to enable this individually + Opts.ApproxFunc = Opts.FastMath; if (Arg *A = Args.getLastArg(OPT_ffp_contract)) { StringRef Val = A->getValue(); @@ -3607,7 +3617,7 @@ // Other LangOpts are only initialized when the input is not AST or LLVM IR. // FIXME: Should we really be calling this for an Language::Asm input? ParseLangArgs(LangOpts, Args, DashX, Res.getTargetOpts(), - Res.getPreprocessorOpts(), Diags); + Res.getPreprocessorOpts(), Res.getCodeGenOpts(), Diags); if (Res.getFrontendOpts().ProgramAction == frontend::RewriteObjC) LangOpts.ObjCExceptions = 1; if (T.isOSDarwin() && DashX.isPreprocessed()) { diff --git a/clang/lib/Parse/ParseDeclCXX.cpp b/clang/lib/Parse/ParseDeclCXX.cpp --- a/clang/lib/Parse/ParseDeclCXX.cpp +++ b/clang/lib/Parse/ParseDeclCXX.cpp @@ -3361,6 +3361,14 @@ // are complete and we can parse the delayed portions of method // declarations and the lexed inline method definitions, along with any // delayed attributes. + + // Save the state of Sema.FPFeatures, and change the setting + // to the levels specified on the command line. Previous level + // will be restored when the RAII object is destroyed. + Sema::FPFeaturesStateRAII SaveFPFeaturesState(Actions); + FPOptions fpOptions(getLangOpts()); + Actions.FPFeatures.Restore(fpOptions.getInt()); + SourceLocation SavedPrevTokLocation = PrevTokLocation; ParseLexedPragmas(getCurrentClass()); ParseLexedAttributes(getCurrentClass()); diff --git a/clang/lib/Parse/ParsePragma.cpp b/clang/lib/Parse/ParsePragma.cpp --- a/clang/lib/Parse/ParsePragma.cpp +++ b/clang/lib/Parse/ParsePragma.cpp @@ -184,6 +184,16 @@ Sema &Actions; }; +struct PragmaFloatControlHandler : public PragmaHandler { + PragmaFloatControlHandler(Sema &Actions) + : PragmaHandler("float_control"), Actions(Actions) {} + void HandlePragma(Preprocessor &PP, PragmaIntroducer Introducer, + Token &FirstToken) override; + +private: + Sema &Actions; +}; + struct PragmaMSPointersToMembers : public PragmaHandler { explicit PragmaMSPointersToMembers() : PragmaHandler("pointers_to_members") {} void HandlePragma(Preprocessor &PP, PragmaIntroducer Introducer, @@ -334,6 +344,9 @@ PP.AddPragmaHandler(MSCommentHandler.get()); } + FloatControlHandler = + std::make_unique(Actions); + PP.AddPragmaHandler(FloatControlHandler.get()); if (getLangOpts().MicrosoftExt) { MSDetectMismatchHandler = std::make_unique(Actions); @@ -438,6 +451,8 @@ PP.RemovePragmaHandler("clang", PCSectionHandler.get()); PCSectionHandler.reset(); + PP.RemovePragmaHandler(FloatControlHandler.get()); + FloatControlHandler.reset(); if (getLangOpts().MicrosoftExt) { PP.RemovePragmaHandler(MSDetectMismatchHandler.get()); MSDetectMismatchHandler.reset(); @@ -646,6 +661,18 @@ ConsumeAnnotationToken(); } +void Parser::HandlePragmaFloatControl() { + assert(Tok.is(tok::annot_pragma_float_control)); + + uintptr_t Value = reinterpret_cast(Tok.getAnnotationValue()); + Sema::PragmaMsStackAction Action = + static_cast((Value >> 16) & 0xFFFF); + PragmaFloatControlKind Kind = + PragmaFloatControlKind(Value & 0xFFFF); + SourceLocation PragmaLoc = ConsumeAnnotationToken(); + Actions.ActOnPragmaFloatControl(PragmaLoc, Action, Kind); +} + void Parser::HandlePragmaFEnvAccess() { assert(Tok.is(tok::annot_pragma_fenv_access)); tok::OnOffSwitch OOS = @@ -2489,6 +2516,139 @@ PP.EnterToken(AnnotTok, /*IsReinject*/ false); } +/// Handle the \#pragma float_control extension. +/// +/// The syntax is: +/// \code +/// #pragma float_control(keyword[, setting] [,push]) +/// \endcode +/// Where 'keyword' and 'setting' are identifiers. +// 'keyword' can be: precise, except, push, pop +// 'setting' can be: on, off +/// The optional arguments 'setting' and 'push' are supported only +/// when the keyword is 'precise' or 'except'. +void PragmaFloatControlHandler::HandlePragma(Preprocessor &PP, + PragmaIntroducer Introducer, + Token &Tok) { + Sema::PragmaMsStackAction Action = Sema::PSK_Set; + SourceLocation FloatControlLoc = Tok.getLocation(); + PP.Lex(Tok); + if (Tok.isNot(tok::l_paren)) { + PP.Diag(FloatControlLoc, diag::err_expected) << tok::l_paren; + return; + } + + // Read the identifier. + PP.Lex(Tok); + if (Tok.isNot(tok::identifier)) { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } + + // Verify that this is one of the float control options. + IdentifierInfo *II = Tok.getIdentifierInfo(); + PragmaFloatControlKind Kind = + llvm::StringSwitch( + II->getName()) + .Case("precise", PFC_Precise) + .Case("except", PFC_Except) + .Case("push", PFC_Push) + .Case("pop", PFC_Pop) + .Default(PFC_Unknown); + PP.Lex(Tok); // the identifier + if (Kind == PFC_Unknown) { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_unknown_kind); + return; + } else if (Kind == PFC_Push || + Kind == PFC_Pop) { + if (Actions.getCurScope()->getParent() != nullptr) { + // FIXME this doesn't detect file-scope: + // the token immediately following a function definition + // returns false, but the token immediately following a forward + // class declaration returns true + //FIXME PP.Diag(Tok.getLocation(), diag::err_pragma_fc_pp_scope); + return; + } + if (Tok.isNot(tok::r_paren)) { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } + PP.Lex(Tok); // Eat the r_paren + Action = (Kind == PFC_Pop) ? Sema::PSK_Pop : Sema::PSK_Push; + } else { + if (Tok.is(tok::r_paren)) + // Selecting Precise or Except + PP.Lex(Tok); // the r_paren + else if (Tok.isNot(tok::comma)) { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } else { + PP.Lex(Tok); // , + if (!Tok.isAnyIdentifier()) { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } + StringRef PushOnOff = Tok.getIdentifierInfo()->getName(); + if (PushOnOff == "on") + // Kind is set correctly + ; + else if (PushOnOff == "off") { + if (Kind == PFC_Precise ) + Kind = PFC_NoPrecise ; + if (Kind == PFC_Except ) + Kind = PFC_NoExcept ; + } else if (PushOnOff == "push") { + if (!Actions.CurContext->isTranslationUnit()) { + //FIXME PP.Diag(Tok.getLocation(), diag::err_pragma_fc_pp_scope); + return; + } + Action = Sema::PSK_Push_Set; + } else { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } + PP.Lex(Tok); // the identifier + if (Tok.is(tok::comma)) { + PP.Lex(Tok); // , + if (!Tok.isAnyIdentifier()) { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } + StringRef ExpectedPush = Tok.getIdentifierInfo()->getName(); + if (ExpectedPush == "push") { + if (!Actions.CurContext->isTranslationUnit()) { + //FIXME PP.Diag(Tok.getLocation(), diag::err_pragma_fc_pp_scope); + return; + } + Action = Sema::PSK_Push_Set; + } else { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } + PP.Lex(Tok); // the push identifier + } + if (Tok.isNot(tok::r_paren)) { + PP.Diag(Tok.getLocation(), diag::err_pragma_float_control_malformed); + return; + } + PP.Lex(Tok); // the r_paren + } + } + SourceLocation EndLoc = Tok.getLocation(); + + // Note: there is no accomodation for PP callback for this pragma. + + // Enter the annotation. + Token AnnotTok; + AnnotTok.startToken(); + AnnotTok.setKind(tok::annot_pragma_float_control); + AnnotTok.setLocation(FloatControlLoc); + AnnotTok.setAnnotationEndLoc(EndLoc); + AnnotTok.setAnnotationValue(reinterpret_cast( + static_cast((Action << 16) | (Kind & 0xFFFF)))); + PP.EnterToken(AnnotTok, /*IsReinject=*/false); +} + /// Handle the Microsoft \#pragma detect_mismatch extension. /// /// The syntax is: diff --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp --- a/clang/lib/Parse/ParseStmt.cpp +++ b/clang/lib/Parse/ParseStmt.cpp @@ -353,13 +353,13 @@ case tok::annot_pragma_fp_contract: ProhibitAttributes(Attrs); - Diag(Tok, diag::err_pragma_fp_contract_scope); + Diag(Tok, diag::err_pragma_file_or_compound_scope) << "fp_contract"; ConsumeAnnotationToken(); return StmtError(); case tok::annot_pragma_fp: ProhibitAttributes(Attrs); - Diag(Tok, diag::err_pragma_fp_scope); + Diag(Tok, diag::err_pragma_file_or_compound_scope) << "clang fp"; ConsumeAnnotationToken(); return StmtError(); @@ -368,6 +368,12 @@ HandlePragmaFEnvAccess(); return StmtEmpty(); + case tok::annot_pragma_float_control: + ProhibitAttributes(Attrs); + Diag(Tok, diag::err_pragma_file_or_compound_scope) << "float_control"; + ConsumeAnnotationToken(); + return StmtError(); + case tok::annot_pragma_opencl_extension: ProhibitAttributes(Attrs); HandlePragmaOpenCLExtension(); @@ -936,6 +942,9 @@ case tok::annot_pragma_fenv_access: HandlePragmaFEnvAccess(); break; + case tok::annot_pragma_float_control: + HandlePragmaFloatControl(); + break; case tok::annot_pragma_ms_pointers_to_members: HandlePragmaMSPointersToMembers(); break; diff --git a/clang/lib/Parse/Parser.cpp b/clang/lib/Parse/Parser.cpp --- a/clang/lib/Parse/Parser.cpp +++ b/clang/lib/Parse/Parser.cpp @@ -763,6 +763,9 @@ case tok::annot_pragma_fenv_access: HandlePragmaFEnvAccess(); return nullptr; + case tok::annot_pragma_float_control: + HandlePragmaFloatControl(); + return nullptr; case tok::annot_pragma_fp: HandlePragmaFP(); break; diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -157,8 +157,9 @@ LangOpts.getMSPointerToMemberRepresentationMethod()), VtorDispStack(LangOpts.getVtorDispMode()), PackStack(0), DataSegStack(nullptr), BSSSegStack(nullptr), ConstSegStack(nullptr), - CodeSegStack(nullptr), CurInitSeg(nullptr), VisContext(nullptr), - PragmaAttributeCurrentTargetDecl(nullptr), + CodeSegStack(nullptr), FpPragmaStack(FPFeatures.getInt()), + CurInitSeg(nullptr), + VisContext(nullptr), PragmaAttributeCurrentTargetDecl(nullptr), IsBuildingRecoveryCallExpr(false), Cleanup{}, LateTemplateParser(nullptr), LateTemplateParserCleanup(nullptr), OpaqueParser(nullptr), IdResolver(pp), StdExperimentalNamespaceCache(nullptr), StdInitializerList(nullptr), diff --git a/clang/lib/Sema/SemaAttr.cpp b/clang/lib/Sema/SemaAttr.cpp --- a/clang/lib/Sema/SemaAttr.cpp +++ b/clang/lib/Sema/SemaAttr.cpp @@ -407,6 +407,49 @@ Consumer.HandleTopLevelDecl(DeclGroupRef(PDMD)); } +void Sema::ActOnPragmaFloatControl(SourceLocation Loc, + PragmaMsStackAction Action, + PragmaFloatControlKind Value) { + auto NewValue = FpPragmaStack.CurrentValue; + switch(Value) { + default: + llvm_unreachable("invalid pragma float_control kind"); + case PFC_Precise: + case PFC_NoPrecise: + case PFC_Except: + case PFC_NoExcept: + switch(Value) { + default: + llvm_unreachable("invalid pragma float_control kind"); + case PFC_Precise: + FPFeatures.setFPPreciseEnabled(true); + break; + case PFC_NoPrecise: + FPFeatures.setFPPreciseEnabled(false); + break; + case PFC_Except: + FPFeatures.setExceptionMode(LangOptions::FPE_Strict); + break; + case PFC_NoExcept: + FPFeatures.setExceptionMode(LangOptions::FPE_Ignore); + break; + } + NewValue = FPFeatures.getInt(); + FpPragmaStack.Act(Loc, Action, StringRef(), NewValue); + break; + case PFC_Push: + // Just duplicate the top entry on the stack + FpPragmaStack.Act(Loc, Action, StringRef(), NewValue); + break; + case PFC_Pop: + // Pop the value from the stack and restore the settings + FpPragmaStack.Act(Loc, Action, StringRef(), NewValue); + NewValue = FpPragmaStack.CurrentValue; + FPFeatures.Restore(NewValue); + break; + } +} + void Sema::ActOnPragmaMSPointersToMembers( LangOptions::PragmaMSPointersToMembersKind RepresentationMethod, SourceLocation PragmaLoc) { diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -13130,14 +13130,6 @@ if (ResultTy.isNull() || LHS.isInvalid() || RHS.isInvalid()) return ExprError(); - if (ResultTy->isRealFloatingType() && - (getLangOpts().getFPRoundingMode() != LangOptions::FPR_ToNearest || - getLangOpts().getFPExceptionMode() != LangOptions::FPE_Ignore)) - // Mark the current function as usng floating point constrained intrinsics - if (FunctionDecl *F = dyn_cast(CurContext)) { - F->setUsesFPIntrin(true); - } - // Some of the binary operations require promoting operands of half vector to // float vectors and truncating the result back to half vector. For now, we do // this only when HalfArgsAndReturn is set (that is, when the target is arm or diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp --- a/clang/lib/Sema/SemaStmt.cpp +++ b/clang/lib/Sema/SemaStmt.cpp @@ -374,6 +374,13 @@ } void Sema::ActOnStartOfCompoundStmt(bool IsStmtExpr) { + if (getFPOptions().isFPConstrained()) { + // Mark the current function as usng floating point constrained intrinsics + if (FunctionDecl *F = dyn_cast(CurContext)) { + F->setUsesFPIntrin(true); + } + } + PushCompoundScope(IsStmtExpr); } diff --git a/clang/test/CodeGen/constrained-math-builtins.c b/clang/test/CodeGen/constrained-math-builtins.c --- a/clang/test/CodeGen/constrained-math-builtins.c +++ b/clang/test/CodeGen/constrained-math-builtins.c @@ -154,9 +154,9 @@ (double)f * f - f; (long double)-f * f + f; -// CHECK: call float @llvm.experimental.constrained.fmuladd.f32 +// CHECK: call contract float @llvm.experimental.constrained.fmuladd.f32 // CHECK: fneg -// CHECK: call double @llvm.experimental.constrained.fmuladd.f64 +// CHECK: call contract double @llvm.experimental.constrained.fmuladd.f64 // CHECK: fneg -// CHECK: call x86_fp80 @llvm.experimental.constrained.fmuladd.f80 +// CHECK: call contract x86_fp80 @llvm.experimental.constrained.fmuladd.f80 }; diff --git a/clang/test/CodeGen/fast-math.c b/clang/test/CodeGen/fast-math.c --- a/clang/test/CodeGen/fast-math.c +++ b/clang/test/CodeGen/fast-math.c @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -ffast-math -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -ffast-math -ffp-contract=fast -emit-llvm -o - %s | FileCheck %s float f0, f1, f2; void foo(void) { diff --git a/clang/test/CodeGen/fp-contract-on-pragma.cpp b/clang/test/CodeGen/fp-contract-on-pragma.cpp --- a/clang/test/CodeGen/fp-contract-on-pragma.cpp +++ b/clang/test/CodeGen/fp-contract-on-pragma.cpp @@ -3,7 +3,7 @@ // Is FP_CONTRACT honored in a simple case? float fp_contract_1(float a, float b, float c) { // CHECK: _Z13fp_contract_1fff -// CHECK: tail call float @llvm.fmuladd +// CHECK: tail call contract float @llvm.fmuladd #pragma clang fp contract(on) return a * b + c; } @@ -31,7 +31,7 @@ float fp_contract_3(float a, float b, float c) { // CHECK: _Z13fp_contract_3fff - // CHECK: tail call float @llvm.fmuladd + // CHECK: tail call contract float @llvm.fmuladd return template_muladd(a, b, c); } @@ -45,13 +45,13 @@ template class fp_contract_4; // CHECK: _ZN13fp_contract_4IiE6methodEfff -// CHECK: tail call float @llvm.fmuladd +// CHECK: tail call contract float @llvm.fmuladd // Check file-scoped FP_CONTRACT #pragma clang fp contract(on) float fp_contract_5(float a, float b, float c) { // CHECK: _Z13fp_contract_5fff - // CHECK: tail call float @llvm.fmuladd + // CHECK: tail call contract float @llvm.fmuladd return a * b + c; } @@ -69,8 +69,8 @@ float fp_contract_7(float a, float b, float c) { // CHECK: _Z13fp_contract_7fff -// CHECK: %[[M:.+]] = fmul float %b, 2.000000e+00 -// CHECK-NEXT: fsub float %[[M]], %c +// CHECK: %[[M:.+]] = fmul contract float %b, 2.000000e+00 +// CHECK-NEXT: fsub contract float %[[M]], %c #pragma clang fp contract(on) return (a = 2 * b) - c; } diff --git a/clang/test/CodeGen/fp-contract-pragma.cpp b/clang/test/CodeGen/fp-contract-pragma.cpp --- a/clang/test/CodeGen/fp-contract-pragma.cpp +++ b/clang/test/CodeGen/fp-contract-pragma.cpp @@ -3,7 +3,7 @@ // Is FP_CONTRACT honored in a simple case? float fp_contract_1(float a, float b, float c) { // CHECK: _Z13fp_contract_1fff -// CHECK: tail call float @llvm.fmuladd +// CHECK: tail call contract float @llvm.fmuladd #pragma STDC FP_CONTRACT ON return a * b + c; } @@ -31,7 +31,7 @@ float fp_contract_3(float a, float b, float c) { // CHECK: _Z13fp_contract_3fff -// CHECK: tail call float @llvm.fmuladd +// CHECK: tail call contract float @llvm.fmuladd return template_muladd(a, b, c); } @@ -44,13 +44,13 @@ template class fp_contract_4; // CHECK: _ZN13fp_contract_4IiE6methodEfff -// CHECK: tail call float @llvm.fmuladd +// CHECK: tail call contract float @llvm.fmuladd // Check file-scoped FP_CONTRACT #pragma STDC FP_CONTRACT ON float fp_contract_5(float a, float b, float c) { // CHECK: _Z13fp_contract_5fff -// CHECK: tail call float @llvm.fmuladd +// CHECK: tail call contract float @llvm.fmuladd return a * b + c; } @@ -68,24 +68,24 @@ float fp_contract_7(float a, float b, float c) { // CHECK: _Z13fp_contract_7fff -// CHECK: %[[M:.+]] = fmul float %b, 2.000000e+00 -// CHECK-NEXT: fsub float %[[M]], %c +// CHECK: %[[M:.+]] = fmul contract float %b, 2.000000e+00 +// CHECK-NEXT: fsub contract float %[[M]], %c #pragma STDC FP_CONTRACT ON return (a = 2 * b) - c; } float fp_contract_8(float a, float b, float c) { // CHECK: _Z13fp_contract_8fff -// CHECK: fneg float %c -// CHECK: tail call float @llvm.fmuladd +// CHECK: fneg contract float %c +// CHECK: tail call contract float @llvm.fmuladd #pragma STDC FP_CONTRACT ON return a * b - c; } float fp_contract_9(float a, float b, float c) { // CHECK: _Z13fp_contract_9fff -// CHECK: fneg float %a -// CHECK: tail call float @llvm.fmuladd +// CHECK: fneg contract float %a +// CHECK: tail call contract float @llvm.fmuladd #pragma STDC FP_CONTRACT ON return c - a * b; } diff --git a/clang/test/CodeGen/fp-floatcontrol-class.cpp b/clang/test/CodeGen/fp-floatcontrol-class.cpp new file mode 100644 --- /dev/null +++ b/clang/test/CodeGen/fp-floatcontrol-class.cpp @@ -0,0 +1,19 @@ +// RUN: %clang -c -Xclang -emit-llvm -o - %s | FileCheck %s +// Verify that float_control does not pertain to initializer expressions + +float y(); +float z(); +#pragma float_control(except, on) +class ON { +float w = 2 + y() * z(); +// CHECK-LABEL: define {{.*}} void @_ZN2ONC2Ev{{.*}} +//CHECK: call contract float {{.*}}llvm.fmuladd +}; +ON on; +#pragma float_control( except, off) +class OFF { +float w = 2 + y() * z(); +// CHECK-LABEL: define {{.*}} void @_ZN3OFFC2Ev{{.*}} +//CHECK: call contract float {{.*}}llvm.fmuladd +}; +OFF off; diff --git a/clang/test/CodeGen/fp-floatcontrol-pragma.cpp b/clang/test/CodeGen/fp-floatcontrol-pragma.cpp new file mode 100644 --- /dev/null +++ b/clang/test/CodeGen/fp-floatcontrol-pragma.cpp @@ -0,0 +1,47 @@ +// RUN: %clang_cc1 -emit-llvm -o - %s | FileCheck %s + +float fff(float x, float y) { +// CHECK-LABEL: define float @_Z3fffff{{.*}} +// CHECK: entry +#pragma float_control(except, on) + float z; + z = z*z; +//CHECK: llvm.experimental.constrained.fmul{{.*}} + { + z = x*y; +//CHECK: llvm.experimental.constrained.fmul{{.*}} + } + { +// This pragma has no effect since if there are any fp intrin in the +// function then all the operations need to be fp intrin +#pragma float_control(except, off) + z = z + x*y; +//CHECK: llvm.experimental.constrained.fmul{{.*}} + } + z = z*z; +//CHECK: llvm.experimental.constrained.fmul{{.*}} + return z; +} +float check_precise(float x, float y) { +// CHECK-LABEL: define float @_Z13check_preciseff{{.*}} + float z; + { +#pragma float_control(precise, on) + z = x*y + z; +//CHECK: llvm.fmuladd{{.*}} + } + { +#pragma float_control(precise, off) + z = x*y + z; +//CHECK: fmul fast float +//CHECK: fadd fast float + } + return z; +} +float fma_test1(float a, float b, float c) { +// CHECK-LABEL define float @_Z9fma_test1fff{{.*}} +#pragma float_control(precise, on) + float x = a * b + c; +//CHECK: fmuladd + return x; +} diff --git a/clang/test/CodeGen/fp-floatcontrol-stack.cpp b/clang/test/CodeGen/fp-floatcontrol-stack.cpp new file mode 100644 --- /dev/null +++ b/clang/test/CodeGen/fp-floatcontrol-stack.cpp @@ -0,0 +1,256 @@ +// RUN: %clang -c -DDEFAULT=1 -Xclang -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DDEFAULT %s +// RUN: %clang -c -DEBSTRICT=1 -ffp-exception-behavior=strict -Xclang -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DEBSTRICT %s +// RUN: %clang -c -DFAST=1 -ffast-math -Xclang -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-FAST %s +// RUN: %clang -c -DNOHONOR=1 -fno-honor-nans -fno-honor-infinities -Xclang -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-NOHONOR %s + +#define FUN(n) (float z) { return n * z + n; } + +float fun_default FUN(1) +//CHECK-LABEL: define {{.*}} @_Z11fun_defaultf{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: call contract float @llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +// Note that backend wants constrained intrinsics used +// throughout the function if they are needed anywhere in the function. +// In that case, operations are built with constrained intrinsics operator +// but using default settings for exception behavior and rounding mode. +//CHECK-DEBSTRICT: llvm.experimental.constrained.fmul{{.*}}tonearest{{.*}}strict +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif +class ResetScope; + +#pragma float_control(except, on, push) +float exc_on FUN(2) +//CHECK-LABEL: define {{.*}} @_Z6exc_onf{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: llvm.experimental.constrained.fmul{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: llvm.experimental.constrained.fmuladd{{.*}}tonearest{{.*}}strict +#endif +#if NOHONOR +//CHECK-NOHONOR: nnan ninf contract float {{.*}}llvm.experimental.constrained.fmuladd{{.*}}tonearest{{.*}}strict +#endif +#if FAST +//CHECK-FAST: fast float {{.*}}llvm.experimental.constrained.fmul{{.*}}tonearest{{.*}}strict +//CHECK-FAST: fast float {{.*}}llvm.experimental.constrained.fadd{{.*}}tonearest{{.*}}strict +#endif + +class ResetScope; +#pragma float_control(pop) +float exc_pop FUN(5) +//CHECK-LABEL: define {{.*}} @_Z7exc_popf{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: call contract float @llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: llvm.experimental.constrained.fmuladd{{.*}}tonearest{{.*}}strict +#endif +#if NOHONOR +//CHECK-NOHONOR: call nnan ninf contract float @llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif + +class ResetScope; +#pragma float_control(except, off) +float exc_off FUN(5) +//CHECK-LABEL: define {{.*}} @_Z7exc_offf{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: call contract float @llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: call contract float @llvm.fmuladd{{.*}} +#endif +#if NOHONOR +//CHECK-NOHONOR: call nnan ninf contract float @llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif + +class ResetScope; +#pragma float_control(precise, on, push) +float precise_on FUN(3) +//CHECK-LABEL: define {{.*}} @_Z10precise_onf{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if NOHONOR +// If precise is pushed then all fast-math should be off! +//CHECK-NOHONOR: call contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: contract float {{.*}}llvm.fmuladd{{.*}} +#endif + +class ResetScope; +#pragma float_control(pop) +float precise_pop FUN(3) +//CHECK-LABEL: define {{.*}} @_Z11precise_popf{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if NOHONOR +//CHECK-NOHONOR: call nnan ninf contract float @llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif +class ResetScope; +#pragma float_control(precise, off) +float precise_off FUN(4) +//CHECK-LABEL: define {{.*}} @_Z11precise_offf{{.*}} +#if DEFAULT +// Note: precise_off enables fp_contract=fast and the instructions +// generated do not include the contract flag, although it was enabled +// in IRBuilder. +//CHECK-DDEFAULT: fmul fast float +//CHECK-DDEFAULT: fadd fast float +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: fmul fast float +//CHECK-DEBSTRICT: fadd fast float +#endif +#if NOHONOR +// fast math should be enabled, and contract should be fast +//CHECK-NOHONOR: fmul fast float +//CHECK-NOHONOR: fadd fast float +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif + +class ResetScope; +#pragma float_control(precise, on) +float precise_on2 FUN(3) +//CHECK-LABEL: define {{.*}} @_Z11precise_on2f{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if NOHONOR +// fast math should be off, and contract should be on +//CHECK-NOHONOR: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: contract float {{.*}}llvm.fmuladd{{.*}} +#endif + +class ResetScope; +#pragma float_control(push) +float precise_push FUN(3) +//CHECK-LABEL: define {{.*}} @_Z12precise_pushf{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if NOHONOR +//CHECK-NOHONOR: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: contract float {{.*}}llvm.fmuladd{{.*}} +#endif + +class ResetScope; +#pragma float_control(precise, off) +float precise_off2 FUN(4) +//CHECK-LABEL: define {{.*}} @_Z12precise_off2f{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: fmul fast float +//CHECK-DDEFAULT: fadd fast float +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: fmul fast float +//CHECK-DEBSTRICT: fadd fast float +#endif +#if NOHONOR +// fast math settings since precise is off +//CHECK-NOHONOR: fmul fast float +//CHECK-NOHONOR: fadd fast float +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif + +class ResetScope; +#pragma float_control(pop) +float precise_pop2 FUN(3) +//CHECK-LABEL: define {{.*}} @_Z12precise_pop2f{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: llvm.fmuladd{{.*}} +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if NOHONOR +//CHECK-NOHONOR: contract float {{.*}}llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: contract float {{.*}}llvm.fmuladd{{.*}} +#endif + +class ResetScope; +// --------- end of push pop test +#pragma float_control(except, on) +float y(); +class ON { +// Settings for top level class initializer revert to command line +// source pragma's do not pertain. +float z = 2 + y() * 7; +//CHECK-LABEL: define {{.*}} void @_ZN2ONC2Ev{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: call contract float {{.*}}llvm.fmuladd +#endif +#if EBSTRICT +//Currently, same as default [command line options not considered] +//CHECK-DEBSTRICT: call contract float {{.*}}llvm.fmuladd +#endif +#if NOHONOR +//CHECK-NOHONOR: call nnan ninf contract float @llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif +}; +ON on; +#pragma float_control( except, off) +class OFF { +float w = 2 + y() * 7; +//CHECK-LABEL: define {{.*}} void @_ZN3OFFC2Ev{{.*}} +#if DEFAULT +//CHECK-DDEFAULT: call contract float {{.*}}llvm.fmuladd +#endif +#if EBSTRICT +//CHECK-DEBSTRICT: call contract float {{.*}}llvm.fmuladd +#endif +#if NOHONOR +//CHECK-NOHONOR: call nnan ninf contract float @llvm.fmuladd{{.*}} +#endif +#if FAST +//CHECK-FAST: fmul fast float +//CHECK-FAST: fadd fast float +#endif +}; +OFF off; diff --git a/clang/test/CodeGen/fpconstrained.c b/clang/test/CodeGen/fpconstrained.c --- a/clang/test/CodeGen/fpconstrained.c +++ b/clang/test/CodeGen/fpconstrained.c @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -ftrapping-math -frounding-math -ffp-exception-behavior=strict -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=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 -emit-llvm -o - %s | FileCheck %s -check-prefix=EXCEPT // RUN: %clang_cc1 -ffast-math -ffp-contract=fast -ffp-exception-behavior=maytrap -emit-llvm -o - %s | FileCheck %s -check-prefix=MAYTRAP @@ -17,6 +17,7 @@ // STRICTNOEXCEPT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.dynamic", metadata !"fpexcept.ignore") // PRECISE: fadd contract float %{{.*}}, %{{.*}} // FAST: fadd fast + // FASTNOCONTRACT: fadd reassoc nnan ninf nsz arcp afn float f0 = f1 + f2; // CHECK: ret diff --git a/clang/test/CodeGen/fpconstrained.cpp b/clang/test/CodeGen/fpconstrained.cpp --- a/clang/test/CodeGen/fpconstrained.cpp +++ b/clang/test/CodeGen/fpconstrained.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -x c++ -ftrapping-math -fexceptions -fcxx-exceptions -frounding-math -ffp-exception-behavior=strict -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=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 -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 -emit-llvm -o - %s | FileCheck %s -check-prefix=MAYTRAP @@ -27,6 +27,7 @@ // STRICTNOEXCEPT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.dynamic", metadata !"fpexcept.ignore") // PRECISE: fadd contract float %{{.*}}, %{{.*}} // FAST: fadd fast + // FASTNOCONTRACT: fadd reassoc nnan ninf nsz arcp afn float f0 = f1 + f2; // CHECK: ret void diff --git a/clang/test/CodeGen/ppc-emmintrin.c b/clang/test/CodeGen/ppc-emmintrin.c --- a/clang/test/CodeGen/ppc-emmintrin.c +++ b/clang/test/CodeGen/ppc-emmintrin.c @@ -2,9 +2,9 @@ // REQUIRES: powerpc-registered-target // RUN: %clang -S -emit-llvm -target powerpc64-unknown-linux-gnu -mcpu=pwr8 -ffreestanding -DNO_WARN_X86_INTRINSICS %s \ -// RUN: -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-BE +// RUN: -ffp-contract=off -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-BE // RUN: %clang -S -emit-llvm -target powerpc64le-unknown-linux-gnu -mcpu=pwr8 -ffreestanding -DNO_WARN_X86_INTRINSICS %s \ -// RUN: -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-LE +// RUN: -ffp-contract=off -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-LE // CHECK-BE-DAG: @_mm_movemask_pd.perm_mask = internal constant <4 x i32> , align 16 // CHECK-BE-DAG: @_mm_shuffle_epi32.permute_selectors = internal constant [4 x i32] [i32 66051, i32 67438087, i32 134810123, i32 202182159], align 4 diff --git a/clang/test/CodeGen/ppc-xmmintrin.c b/clang/test/CodeGen/ppc-xmmintrin.c --- a/clang/test/CodeGen/ppc-xmmintrin.c +++ b/clang/test/CodeGen/ppc-xmmintrin.c @@ -2,9 +2,9 @@ // REQUIRES: powerpc-registered-target // RUN: %clang -S -emit-llvm -target powerpc64-unknown-linux-gnu -mcpu=pwr8 -ffreestanding -DNO_WARN_X86_INTRINSICS %s \ -// RUN: -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-BE +// RUN: -ffp-contract=off -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-BE // RUN: %clang -S -emit-llvm -target powerpc64le-unknown-linux-gnu -mcpu=pwr8 -ffreestanding -DNO_WARN_X86_INTRINSICS %s \ -// RUN: -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-LE +// RUN: -ffp-contract=off -fno-discard-value-names -mllvm -disable-llvm-optzns -o - | llvm-cxxfilt -n | FileCheck %s --check-prefixes=CHECK,CHECK-LE #include diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl @@ -10,8 +10,8 @@ typedef unsigned short __attribute__((ext_vector_type(2))) ushort2; // CHECK-LABEL: @builtins_amdgcn_dl_insts -// CHECK: call float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 false) -// CHECK: call float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 true) +// CHECK: call contract float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 false) +// CHECK: call contract float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 true) // CHECK: call i32 @llvm.amdgcn.sdot2(<2 x i16> %v2ssA, <2 x i16> %v2ssB, i32 %siC, i1 false) // CHECK: call i32 @llvm.amdgcn.sdot2(<2 x i16> %v2ssA, <2 x i16> %v2ssB, i32 %siC, i1 true) diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl @@ -5,7 +5,7 @@ #pragma OPENCL EXTENSION cl_khr_fp16 : enable // CHECK-LABEL: @test_fmed3_f16 -// CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c) +// CHECK: call contract half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c) void test_fmed3_f16(global half* out, half a, half b, half c) { *out = __builtin_amdgcn_fmed3h(a, b, c); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl @@ -19,7 +19,7 @@ // CHECK-LABEL: test_interp_f32 // CHECK: call float @llvm.amdgcn.interp.p1 -// CHECK: call float @llvm.amdgcn.interp.p2 +// CHECK: call contract float @llvm.amdgcn.interp.p2 void test_interp_f32(global float* out, float i, float j, int m0) { float p1 = __builtin_amdgcn_interp_p1(i, 1, 4, m0); @@ -27,7 +27,7 @@ } // CHECK-LABEL: test_interp_mov -// CHECK: call float @llvm.amdgcn.interp.mov +// CHECK: call contract float @llvm.amdgcn.interp.mov void test_interp_mov(global float* out, float i, float j, int m0) { *out = __builtin_amdgcn_interp_mov(2, 3, 4, m0); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl @@ -20,70 +20,70 @@ // CHECK-LABEL: @test_mfma_f32_32x32x1f32 -// CHECK: call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float %a, float %b, <32 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float %a, float %b, <32 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_32x32x1f32(global v32f* out, float a, float b, v32f c) { *out = __builtin_amdgcn_mfma_f32_32x32x1f32(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_16x16x1f32 -// CHECK: call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %a, float %b, <16 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %a, float %b, <16 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_16x16x1f32(global v16f* out, float a, float b, v16f c) { *out = __builtin_amdgcn_mfma_f32_16x16x1f32(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_4x4x1f32 -// CHECK: call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float %a, float %b, <4 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float %a, float %b, <4 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_4x4x1f32(global v4f* out, float a, float b, v4f c) { *out = __builtin_amdgcn_mfma_f32_4x4x1f32(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_32x32x2f32 -// CHECK: call <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float %a, float %b, <16 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float %a, float %b, <16 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_32x32x2f32(global v16f* out, float a, float b, v16f c) { *out = __builtin_amdgcn_mfma_f32_32x32x2f32(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_16x16x4f32 -// CHECK: call <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float %a, float %b, <4 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float %a, float %b, <4 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_16x16x4f32(global v4f* out, float a, float b, v4f c) { *out = __builtin_amdgcn_mfma_f32_16x16x4f32(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_32x32x4f16 -// CHECK: call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half> %a, <4 x half> %b, <32 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half> %a, <4 x half> %b, <32 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_32x32x4f16(global v32f* out, v4h a, v4h b, v32f c) { *out = __builtin_amdgcn_mfma_f32_32x32x4f16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_16x16x4f16 -// CHECK: call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half> %a, <4 x half> %b, <16 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half> %a, <4 x half> %b, <16 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_16x16x4f16(global v16f* out, v4h a, v4h b, v16f c) { *out = __builtin_amdgcn_mfma_f32_16x16x4f16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_4x4x4f16 -// CHECK: call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half> %a, <4 x half> %b, <4 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half> %a, <4 x half> %b, <4 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_4x4x4f16(global v4f* out, v4h a, v4h b, v4f c) { *out = __builtin_amdgcn_mfma_f32_4x4x4f16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_32x32x8f16 -// CHECK: call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> %a, <4 x half> %b, <16 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> %a, <4 x half> %b, <16 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_32x32x8f16(global v16f* out, v4h a, v4h b, v16f c) { *out = __builtin_amdgcn_mfma_f32_32x32x8f16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_16x16x16f16 -// CHECK: call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> %a, <4 x half> %b, <4 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> %a, <4 x half> %b, <4 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_16x16x16f16(global v4f* out, v4h a, v4h b, v4f c) { *out = __builtin_amdgcn_mfma_f32_16x16x16f16(a, b, c, 0, 0, 0); @@ -125,35 +125,35 @@ } // CHECK-LABEL: @test_mfma_f32_32x32x2bf16 -// CHECK: call <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16> %a, <2 x i16> %b, <32 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16> %a, <2 x i16> %b, <32 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_32x32x2bf16(global v32f* out, v2s a, v2s b, v32f c) { *out = __builtin_amdgcn_mfma_f32_32x32x2bf16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_16x16x2bf16 -// CHECK: call <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_16x16x2bf16(global v16f* out, v2s a, v2s b, v16f c) { *out = __builtin_amdgcn_mfma_f32_16x16x2bf16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_4x4x2bf16 -// CHECK: call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_4x4x2bf16(global v4f* out, v2s a, v2s b, v4f c) { *out = __builtin_amdgcn_mfma_f32_4x4x2bf16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_32x32x4bf16 -// CHECK: call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_32x32x4bf16(global v16f* out, v2s a, v2s b, v16f c) { *out = __builtin_amdgcn_mfma_f32_32x32x4bf16(a, b, c, 0, 0, 0); } // CHECK-LABEL: @test_mfma_f32_16x16x8bf16 -// CHECK: call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %c, i32 0, i32 0, i32 0) +// CHECK: call contract <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %c, i32 0, i32 0, i32 0) void test_mfma_f32_16x16x8bf16(global v4f* out, v2s a, v2s b, v4f c) { *out = __builtin_amdgcn_mfma_f32_16x16x8bf16(a, b, c, 0, 0, 0); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -9,49 +9,49 @@ typedef unsigned long ulong; // CHECK-LABEL: @test_div_fixup_f16 -// CHECK: call half @llvm.amdgcn.div.fixup.f16 +// CHECK: call contract half @llvm.amdgcn.div.fixup.f16 void test_div_fixup_f16(global half* out, half a, half b, half c) { *out = __builtin_amdgcn_div_fixuph(a, b, c); } // CHECK-LABEL: @test_rcp_f16 -// CHECK: call half @llvm.amdgcn.rcp.f16 +// CHECK: call contract half @llvm.amdgcn.rcp.f16 void test_rcp_f16(global half* out, half a) { *out = __builtin_amdgcn_rcph(a); } // CHECK-LABEL: @test_rsq_f16 -// CHECK: call half @llvm.amdgcn.rsq.f16 +// CHECK: call contract half @llvm.amdgcn.rsq.f16 void test_rsq_f16(global half* out, half a) { *out = __builtin_amdgcn_rsqh(a); } // CHECK-LABEL: @test_sin_f16 -// CHECK: call half @llvm.amdgcn.sin.f16 +// CHECK: call contract half @llvm.amdgcn.sin.f16 void test_sin_f16(global half* out, half a) { *out = __builtin_amdgcn_sinh(a); } // CHECK-LABEL: @test_cos_f16 -// CHECK: call half @llvm.amdgcn.cos.f16 +// CHECK: call contract half @llvm.amdgcn.cos.f16 void test_cos_f16(global half* out, half a) { *out = __builtin_amdgcn_cosh(a); } // CHECK-LABEL: @test_ldexp_f16 -// CHECK: call half @llvm.amdgcn.ldexp.f16 +// CHECK: call contract half @llvm.amdgcn.ldexp.f16 void test_ldexp_f16(global half* out, half a, int b) { *out = __builtin_amdgcn_ldexph(a, b); } // CHECK-LABEL: @test_frexp_mant_f16 -// CHECK: call half @llvm.amdgcn.frexp.mant.f16 +// CHECK: call contract half @llvm.amdgcn.frexp.mant.f16 void test_frexp_mant_f16(global half* out, half a) { *out = __builtin_amdgcn_frexp_manth(a); @@ -65,7 +65,7 @@ } // CHECK-LABEL: @test_fract_f16 -// CHECK: call half @llvm.amdgcn.fract.f16 +// CHECK: call contract half @llvm.amdgcn.fract.f16 void test_fract_f16(global half* out, half a) { *out = __builtin_amdgcn_fracth(a); @@ -107,19 +107,19 @@ } // CHECK-LABEL: @test_ds_fadd -// CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) +// CHECK: call contract float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) void test_ds_faddf(local float *out, float src) { *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false); } // CHECK-LABEL: @test_ds_fmin -// CHECK: call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) +// CHECK: call contract float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) void test_ds_fminf(local float *out, float src) { *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false); } // CHECK-LABEL: @test_ds_fmax -// CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) +// CHECK: call contract float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) void test_ds_fmaxf(local float *out, float src) { *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false); } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -61,133 +61,133 @@ } // CHECK-LABEL: @test_div_fmas_f32 -// CHECK: call float @llvm.amdgcn.div.fmas.f32 +// CHECK: call contract float @llvm.amdgcn.div.fmas.f32 void test_div_fmas_f32(global float* out, float a, float b, float c, int d) { *out = __builtin_amdgcn_div_fmasf(a, b, c, d); } // CHECK-LABEL: @test_div_fmas_f64 -// CHECK: call double @llvm.amdgcn.div.fmas.f64 +// CHECK: call contract double @llvm.amdgcn.div.fmas.f64 void test_div_fmas_f64(global double* out, double a, double b, double c, int d) { *out = __builtin_amdgcn_div_fmas(a, b, c, d); } // CHECK-LABEL: @test_div_fixup_f32 -// CHECK: call float @llvm.amdgcn.div.fixup.f32 +// CHECK: call contract float @llvm.amdgcn.div.fixup.f32 void test_div_fixup_f32(global float* out, float a, float b, float c) { *out = __builtin_amdgcn_div_fixupf(a, b, c); } // CHECK-LABEL: @test_div_fixup_f64 -// CHECK: call double @llvm.amdgcn.div.fixup.f64 +// CHECK: call contract double @llvm.amdgcn.div.fixup.f64 void test_div_fixup_f64(global double* out, double a, double b, double c) { *out = __builtin_amdgcn_div_fixup(a, b, c); } // CHECK-LABEL: @test_trig_preop_f32 -// CHECK: call float @llvm.amdgcn.trig.preop.f32 +// CHECK: call contract float @llvm.amdgcn.trig.preop.f32 void test_trig_preop_f32(global float* out, float a, int b) { *out = __builtin_amdgcn_trig_preopf(a, b); } // CHECK-LABEL: @test_trig_preop_f64 -// CHECK: call double @llvm.amdgcn.trig.preop.f64 +// CHECK: call contract double @llvm.amdgcn.trig.preop.f64 void test_trig_preop_f64(global double* out, double a, int b) { *out = __builtin_amdgcn_trig_preop(a, b); } // CHECK-LABEL: @test_rcp_f32 -// CHECK: call float @llvm.amdgcn.rcp.f32 +// CHECK: call contract float @llvm.amdgcn.rcp.f32 void test_rcp_f32(global float* out, float a) { *out = __builtin_amdgcn_rcpf(a); } // CHECK-LABEL: @test_rcp_f64 -// CHECK: call double @llvm.amdgcn.rcp.f64 +// CHECK: call contract double @llvm.amdgcn.rcp.f64 void test_rcp_f64(global double* out, double a) { *out = __builtin_amdgcn_rcp(a); } // CHECK-LABEL: @test_rsq_f32 -// CHECK: call float @llvm.amdgcn.rsq.f32 +// CHECK: call contract float @llvm.amdgcn.rsq.f32 void test_rsq_f32(global float* out, float a) { *out = __builtin_amdgcn_rsqf(a); } // CHECK-LABEL: @test_rsq_f64 -// CHECK: call double @llvm.amdgcn.rsq.f64 +// CHECK: call contract double @llvm.amdgcn.rsq.f64 void test_rsq_f64(global double* out, double a) { *out = __builtin_amdgcn_rsq(a); } // CHECK-LABEL: @test_rsq_clamp_f32 -// CHECK: call float @llvm.amdgcn.rsq.clamp.f32 +// CHECK: call contract float @llvm.amdgcn.rsq.clamp.f32 void test_rsq_clamp_f32(global float* out, float a) { *out = __builtin_amdgcn_rsq_clampf(a); } // CHECK-LABEL: @test_rsq_clamp_f64 -// CHECK: call double @llvm.amdgcn.rsq.clamp.f64 +// CHECK: call contract double @llvm.amdgcn.rsq.clamp.f64 void test_rsq_clamp_f64(global double* out, double a) { *out = __builtin_amdgcn_rsq_clamp(a); } // CHECK-LABEL: @test_sin_f32 -// CHECK: call float @llvm.amdgcn.sin.f32 +// CHECK: call contract float @llvm.amdgcn.sin.f32 void test_sin_f32(global float* out, float a) { *out = __builtin_amdgcn_sinf(a); } // CHECK-LABEL: @test_cos_f32 -// CHECK: call float @llvm.amdgcn.cos.f32 +// CHECK: call contract float @llvm.amdgcn.cos.f32 void test_cos_f32(global float* out, float a) { *out = __builtin_amdgcn_cosf(a); } // CHECK-LABEL: @test_log_clamp_f32 -// CHECK: call float @llvm.amdgcn.log.clamp.f32 +// CHECK: call contract float @llvm.amdgcn.log.clamp.f32 void test_log_clamp_f32(global float* out, float a) { *out = __builtin_amdgcn_log_clampf(a); } // CHECK-LABEL: @test_ldexp_f32 -// CHECK: call float @llvm.amdgcn.ldexp.f32 +// CHECK: call contract float @llvm.amdgcn.ldexp.f32 void test_ldexp_f32(global float* out, float a, int b) { *out = __builtin_amdgcn_ldexpf(a, b); } // CHECK-LABEL: @test_ldexp_f64 -// CHECK: call double @llvm.amdgcn.ldexp.f64 +// CHECK: call contract double @llvm.amdgcn.ldexp.f64 void test_ldexp_f64(global double* out, double a, int b) { *out = __builtin_amdgcn_ldexp(a, b); } // CHECK-LABEL: @test_frexp_mant_f32 -// CHECK: call float @llvm.amdgcn.frexp.mant.f32 +// CHECK: call contract float @llvm.amdgcn.frexp.mant.f32 void test_frexp_mant_f32(global float* out, float a) { *out = __builtin_amdgcn_frexp_mantf(a); } // CHECK-LABEL: @test_frexp_mant_f64 -// CHECK: call double @llvm.amdgcn.frexp.mant.f64 +// CHECK: call contract double @llvm.amdgcn.frexp.mant.f64 void test_frexp_mant_f64(global double* out, double a) { *out = __builtin_amdgcn_frexp_mant(a); @@ -208,14 +208,14 @@ } // CHECK-LABEL: @test_fract_f32 -// CHECK: call float @llvm.amdgcn.fract.f32 +// CHECK: call contract float @llvm.amdgcn.fract.f32 void test_fract_f32(global int* out, float a) { *out = __builtin_amdgcn_fractf(a); } // CHECK-LABEL: @test_fract_f64 -// CHECK: call double @llvm.amdgcn.fract.f64 +// CHECK: call contract double @llvm.amdgcn.fract.f64 void test_fract_f64(global int* out, double a) { *out = __builtin_amdgcn_fract(a); @@ -417,25 +417,25 @@ } // CHECK-LABEL: @test_cubeid( -// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c) +// CHECK: call contract float @llvm.amdgcn.cubeid(float %a, float %b, float %c) void test_cubeid(global float* out, float a, float b, float c) { *out = __builtin_amdgcn_cubeid(a, b, c); } // CHECK-LABEL: @test_cubesc( -// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c) +// CHECK: call contract float @llvm.amdgcn.cubesc(float %a, float %b, float %c) void test_cubesc(global float* out, float a, float b, float c) { *out = __builtin_amdgcn_cubesc(a, b, c); } // CHECK-LABEL: @test_cubetc( -// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c) +// CHECK: call contract float @llvm.amdgcn.cubetc(float %a, float %b, float %c) void test_cubetc(global float* out, float a, float b, float c) { *out = __builtin_amdgcn_cubetc(a, b, c); } // CHECK-LABEL: @test_cubema( -// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c) +// CHECK: call contract float @llvm.amdgcn.cubema(float %a, float %b, float %c) void test_cubema(global float* out, float a, float b, float c) { *out = __builtin_amdgcn_cubema(a, b, c); } @@ -528,7 +528,7 @@ } // CHECK-LABEL: @test_fmed3_f32 -// CHECK: call float @llvm.amdgcn.fmed3.f32( +// CHECK: call contract float @llvm.amdgcn.fmed3.f32( void test_fmed3_f32(global float* out, float a, float b, float c) { *out = __builtin_amdgcn_fmed3f(a, b, c); @@ -620,7 +620,7 @@ } // CHECK-LABEL: @test_cvt_pkrtz( -// CHECK: tail call <2 x half> @llvm.amdgcn.cvt.pkrtz(float %src0, float %src1) +// CHECK: tail call contract <2 x half> @llvm.amdgcn.cvt.pkrtz(float %src0, float %src1) kernel void test_cvt_pkrtz(global half2* out, float src0, float src1) { *out = __builtin_amdgcn_cvt_pkrtz(src0, src1); } diff --git a/clang/test/CodeGenOpenCL/builtins-f16.cl b/clang/test/CodeGenOpenCL/builtins-f16.cl --- a/clang/test/CodeGenOpenCL/builtins-f16.cl +++ b/clang/test/CodeGenOpenCL/builtins-f16.cl @@ -6,66 +6,66 @@ void test_half_builtins(half h0, half h1, half h2) { volatile half res; - // CHECK: call half @llvm.copysign.f16(half %h0, half %h1) + // CHECK: call contract half @llvm.copysign.f16(half %h0, half %h1) res = __builtin_copysignf16(h0, h1); - // CHECK: call half @llvm.fabs.f16(half %h0) + // CHECK: call contract half @llvm.fabs.f16(half %h0) res = __builtin_fabsf16(h0); - // CHECK: call half @llvm.ceil.f16(half %h0) + // CHECK: call contract half @llvm.ceil.f16(half %h0) res = __builtin_ceilf16(h0); - // CHECK: call half @llvm.cos.f16(half %h0) + // CHECK: call contract half @llvm.cos.f16(half %h0) res = __builtin_cosf16(h0); - // CHECK: call half @llvm.exp.f16(half %h0) + // CHECK: call contract half @llvm.exp.f16(half %h0) res = __builtin_expf16(h0); - // CHECK: call half @llvm.exp2.f16(half %h0) + // CHECK: call contract half @llvm.exp2.f16(half %h0) res = __builtin_exp2f16(h0); - // CHECK: call half @llvm.floor.f16(half %h0) + // CHECK: call contract half @llvm.floor.f16(half %h0) res = __builtin_floorf16(h0); - // CHECK: call half @llvm.fma.f16(half %h0, half %h1, half %h2) + // CHECK: call contract half @llvm.fma.f16(half %h0, half %h1, half %h2) res = __builtin_fmaf16(h0, h1 ,h2); - // CHECK: call half @llvm.maxnum.f16(half %h0, half %h1) + // CHECK: call contract half @llvm.maxnum.f16(half %h0, half %h1) res = __builtin_fmaxf16(h0, h1); - // CHECK: call half @llvm.minnum.f16(half %h0, half %h1) + // CHECK: call contract half @llvm.minnum.f16(half %h0, half %h1) res = __builtin_fminf16(h0, h1); - // CHECK: frem half %h0, %h1 + // CHECK: frem contract half %h0, %h1 res = __builtin_fmodf16(h0, h1); - // CHECK: call half @llvm.pow.f16(half %h0, half %h1) + // CHECK: call contract half @llvm.pow.f16(half %h0, half %h1) res = __builtin_powf16(h0, h1); - // CHECK: call half @llvm.log10.f16(half %h0) + // CHECK: call contract half @llvm.log10.f16(half %h0) res = __builtin_log10f16(h0); - // CHECK: call half @llvm.log2.f16(half %h0) + // CHECK: call contract half @llvm.log2.f16(half %h0) res = __builtin_log2f16(h0); - // CHECK: call half @llvm.log.f16(half %h0) + // CHECK: call contract half @llvm.log.f16(half %h0) res = __builtin_logf16(h0); - // CHECK: call half @llvm.rint.f16(half %h0) + // CHECK: call contract half @llvm.rint.f16(half %h0) res = __builtin_rintf16(h0); - // CHECK: call half @llvm.round.f16(half %h0) + // CHECK: call contract half @llvm.round.f16(half %h0) res = __builtin_roundf16(h0); - // CHECK: call half @llvm.sin.f16(half %h0) + // CHECK: call contract half @llvm.sin.f16(half %h0) res = __builtin_sinf16(h0); - // CHECK: call half @llvm.sqrt.f16(half %h0) + // CHECK: call contract half @llvm.sqrt.f16(half %h0) res = __builtin_sqrtf16(h0); - // CHECK: call half @llvm.trunc.f16(half %h0) + // CHECK: call contract half @llvm.trunc.f16(half %h0) res = __builtin_truncf16(h0); - // CHECK: call half @llvm.canonicalize.f16(half %h0) + // CHECK: call contract half @llvm.canonicalize.f16(half %h0) res = __builtin_canonicalizef16(h0); } diff --git a/clang/test/CodeGenOpenCL/builtins-r600.cl b/clang/test/CodeGenOpenCL/builtins-r600.cl --- a/clang/test/CodeGenOpenCL/builtins-r600.cl +++ b/clang/test/CodeGenOpenCL/builtins-r600.cl @@ -2,7 +2,7 @@ // RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s // CHECK-LABEL: @test_recipsqrt_ieee_f32 -// CHECK: call float @llvm.r600.recipsqrt.ieee.f32 +// CHECK: call contract float @llvm.r600.recipsqrt.ieee.f32 void test_recipsqrt_ieee_f32(global float* out, float a) { *out = __builtin_r600_recipsqrt_ieeef(a); @@ -10,7 +10,7 @@ #if cl_khr_fp64 // XCHECK-LABEL: @test_recipsqrt_ieee_f64 -// XCHECK: call double @llvm.r600.recipsqrt.ieee.f64 +// XCHECK: call contract double @llvm.r600.recipsqrt.ieee.f64 void test_recipsqrt_ieee_f64(global double* out, double a) { *out = __builtin_r600_recipsqrt_ieee(a); diff --git a/clang/test/CodeGenOpenCL/relaxed-fpmath.cl b/clang/test/CodeGenOpenCL/relaxed-fpmath.cl --- a/clang/test/CodeGenOpenCL/relaxed-fpmath.cl +++ b/clang/test/CodeGenOpenCL/relaxed-fpmath.cl @@ -8,12 +8,12 @@ float spscalardiv(float a, float b) { // CHECK: @spscalardiv( - // NORMAL: fdiv float + // NORMAL: fdiv contract float // FAST: fdiv fast float - // FINITE: fdiv nnan ninf float - // UNSAFE: fdiv nnan nsz float - // MAD: fdiv float - // NOSIGNED: fdiv nsz float + // FINITE: fdiv nnan ninf contract float + // UNSAFE: fdiv nnan nsz contract float + // MAD: fdiv contract float + // NOSIGNED: fdiv nsz contract float return a / b; } // CHECK: attributes diff --git a/clang/test/CodeGenOpenCL/single-precision-constant.cl b/clang/test/CodeGenOpenCL/single-precision-constant.cl --- a/clang/test/CodeGenOpenCL/single-precision-constant.cl +++ b/clang/test/CodeGenOpenCL/single-precision-constant.cl @@ -1,6 +1,6 @@ // RUN: %clang_cc1 %s -cl-single-precision-constant -emit-llvm -o - | FileCheck %s float fn(float f) { - // CHECK: tail call float @llvm.fmuladd.f32(float %f, float 2.000000e+00, float 1.000000e+00) + // CHECK: tail call contract float @llvm.fmuladd.f32(float %f, float 2.000000e+00, float 1.000000e+00) return f*2. + 1.; } diff --git a/clang/test/Driver/fp-model.c b/clang/test/Driver/fp-model.c --- a/clang/test/Driver/fp-model.c +++ b/clang/test/Driver/fp-model.c @@ -27,9 +27,9 @@ // RUN: | FileCheck --check-prefix=WARN5 %s // WARN5: warning: overriding '-ffp-model=strict' option with '-ffp-contract=fast' [-Woverriding-t-option] -// RUN: %clang -### -ffp-model=strict -ffp-contract=off -c %s 2>&1 \ +// RUN: %clang -### -ffp-model=strict -ffp-contract=fast -c %s 2>&1 \ // RUN: | FileCheck --check-prefix=WARN6 %s -// WARN6: warning: overriding '-ffp-model=strict' option with '-ffp-contract=off' [-Woverriding-t-option] +// WARN6: warning: overriding '-ffp-model=strict' option with '-ffp-contract=fast' [-Woverriding-t-option] // RUN: %clang -### -ffp-model=strict -ffp-contract=on -c %s 2>&1 \ // RUN: | FileCheck --check-prefix=WARN7 %s @@ -100,13 +100,14 @@ // RUN: %clang -### -nostdinc -ffp-model=precise -c %s 2>&1 \ // RUN: | FileCheck --check-prefix=CHECK-FPM-PRECISE %s // CHECK-FPM-PRECISE: "-cc1" -// CHECK-FPM-PRECISE: "-ffp-contract=fast" +// CHECK-FPM-PRECISE: "-ffp-contract=on" // CHECK-FPM-PRECISE: "-fno-rounding-math" // RUN: %clang -### -nostdinc -ffp-model=strict -c %s 2>&1 \ // RUN: | FileCheck --check-prefix=CHECK-FPM-STRICT %s // CHECK-FPM-STRICT: "-cc1" // CHECK-FPM-STRICT: "-ftrapping-math" +// CHECK-FPM-STRICT: "-ffp-contract=off" // CHECK-FPM-STRICT: "-frounding-math" // CHECK-FPM-STRICT: "-ffp-exception-behavior=strict" diff --git a/clang/test/Parser/fp-floatcontrol-syntax.cpp b/clang/test/Parser/fp-floatcontrol-syntax.cpp new file mode 100644 --- /dev/null +++ b/clang/test/Parser/fp-floatcontrol-syntax.cpp @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 -fsyntax-only -verify -DCHECK_ERROR %s +// XFAIL: * + +float function_scope(float a) { + return a; +} + +// There seems to be a bug in Actions.CurContext->isTranslationUnit() +// unless this dummy class is used. FIXME Fix the issue then remove this +// workaround. +#define TU_WORKAROUND +#ifdef TU_WORKAROUND +class ResetTUScope; +#endif +#ifdef CHECK_ERROR +# pragma float_control(push) +# pragma float_control(pop) +# pragma float_control(precise,on,push) +void check_stack() { +#pragma float_control(push) // expected-error {{can only appear at file scope}} +#pragma float_control(pop) // expected-error {{can only appear at file scope}} +#pragma float_control(precise,on,push) // expected-error {{can only appear at file scope}} + return; +} +#endif diff --git a/llvm/include/llvm/IR/IRBuilder.h b/llvm/include/llvm/IR/IRBuilder.h --- a/llvm/include/llvm/IR/IRBuilder.h +++ b/llvm/include/llvm/IR/IRBuilder.h @@ -215,6 +215,8 @@ /// Get the flags to be applied to created floating point ops FastMathFlags getFastMathFlags() const { return FMF; } + FastMathFlags& getFastMathFlags() { return FMF; } + /// Clear the fast-math flags. void clearFastMathFlags() { FMF.clear(); } @@ -299,10 +301,16 @@ IRBuilderBase &Builder; FastMathFlags FMF; MDNode *FPMathTag; + bool IsFPConstrained; + fp::ExceptionBehavior DefaultConstrainedExcept; + fp::RoundingMode DefaultConstrainedRounding; public: FastMathFlagGuard(IRBuilderBase &B) - : Builder(B), FMF(B.FMF), FPMathTag(B.DefaultFPMathTag) {} + : Builder(B), FMF(B.FMF), FPMathTag(B.DefaultFPMathTag), + IsFPConstrained(B.IsFPConstrained), + DefaultConstrainedExcept(B.DefaultConstrainedExcept), + DefaultConstrainedRounding(B.DefaultConstrainedRounding) {} FastMathFlagGuard(const FastMathFlagGuard &) = delete; FastMathFlagGuard &operator=(const FastMathFlagGuard &) = delete; @@ -310,6 +318,9 @@ ~FastMathFlagGuard() { Builder.FMF = FMF; Builder.DefaultFPMathTag = FPMathTag; + Builder.IsFPConstrained = IsFPConstrained; + Builder.DefaultConstrainedExcept = DefaultConstrainedExcept; + Builder.DefaultConstrainedRounding = DefaultConstrainedRounding; } };