Add support for pragma float_control, to control precision and exception behavior at the source level

[mibintc]

Summary

Intel would like to support #pragma float_control which allows control over precision and exception behavior at the source level. This pragma is supported by both the Microsoft compiler and the Intel compiler, and our customers have found it useful. This message is to describe the pragma, provide the patch for review, and request feedback from the community.

As the pragma was originally defined in the Microsoft compiler, the pragma supports a stack of settings, so that you can use push and pop to save and restore the state. That functionality is already in clang and this patch piggy-backs on that support (reference PragmaStack and PragmaMsStackAction). The Microsoft compiler supports it only at file scope, but the Intel compiler, and this patch, supports the pragma at either file scope or at the beginning of a compound statement. Clang currently provides support for pragma fp_contract which enables floating point contraction--also at file scope and compound statement, and uses FPFeatures on certain expression nodes to hold the pragma settings. This patch piggy-backs FPFeatures, adding fields to show the setting of "float_control(precise)" and "float_control(except)"

This patch includes an update to the pragma documentation, to summarize, using pragma float_control(precise) is like enabling ffp-model=precise for a region of the program. pragma float_control(except) is like enabling ffp-exception-behavior=strict/ignore for a region of the program.

[andrew.w.kaylor]

clang/include/clang/Basic/DiagnosticParseKinds.td

1135

This isn't quite accurate. The pop case has no comma-separated arguments. It might be better to print the full syntax here if that's feasible.

clang/include/clang/Basic/LangOptions.h

362

It seems like fp_precise describes too many things to be a single option. Even within this set of options it overlaps with fp_contract.

clang/test/CodeGen/fp-floatcontrol-stack.cpp

3

Can you add run lines for -ffast-math and (separately) "-fno-honor-nans -fno-honor-infinities"?

18

There should be a constrained fadd here also, right?

53

Why is the constrained intrinsic generated in this case? If we've got both constraints set to the defaults at the file scope I would have expected that to turn off the constrained mode.

125

Are there also fast-math flags set here? If not, why not?

[mibintc]

In D72841#1842772, @andrew.w.kaylor wrote:

It's not clear to me from reading this how the "precise" control is going to work with relation to the fast math flags. I don't think MSVC allows the granularity of control over these flags that we have in clang, so maybe they don't have this problem.

You're right, MSVC only allows the pragma at file scope.

Consider this code: https://godbolt.org/z/mHiLCm

With the options "-ffp-model=precise -fno-honor-infinities -fno-honor-nans" the math operations here are generated with the "nnan ninf contract" flags. That's correct. What will happen when I use the pragma to turn precise off? Does it enable all fast math flags? Will the subsequent "pop" leave the "ninf nnan" fast math flags enabled?

This patch doesn't have support for turning precise off, that's a bug, I will revisit. This is my plan for how to handle enabling/disabling fast math: IRBuilder.h has settings FMF, and also supplies clearFastMathFlags, setFastMathFlags(flags) and FastMathFlagGuard. When the expression is walked that alters the FMF, use FastMathFlagGuard to save the current state of FMF, modify the settings using the clear or set functions, walk the expression. After the expression is walked, the FMF settings will be restored to previous state by the FastMathFlagGuard destructor.

As I said, I don't think you can get into this situation with MSVC. I believe that icc will go into full fast math mode with the "precise, off, push" pragma but will go back to "nnan ninf contract" mode with the pop. At least, that's what the design docs say. I haven't looked at the code to verify this. It seems like the correct behavior in any case. I think the clang FPOptions needs individual entries for all of the fast math flags to handle this case.

Thanks I'll investigate this and add test cases. I think possibly since IRBuilder has the FMF like I described above it might work with current support. Is there currently a way to modify nan and inf at the source level or only by compiler option? BTW I've been asked to implement a pragma to control fp "reassoc" at the source level. I'm planning to do that after this patch is complete.

clang/docs/LanguageExtensions.rst

3059

Oh, I looked back at the patch for -ffp-model and precise is documented to set ffp-contract=fast. Not sure why I thought that was right. I'll have to redo it.

3067

No, the stack that tracks the float control pragma settings is a pair, roughly (IsPreciseEnabled, IsExceptEnabled)

clang/include/clang/Basic/LangOptions.h

362

I see your point. I wanted it to reflect the current pragma setting that's why I kept it intact. I'll rethink this.

clang/test/CodeGen/fp-floatcontrol-stack.cpp

3

OK. i'll add pragma's to set precise off too.

18

yes there's a constrained add following. i can add that pattern into the file check.

53

The "run" line in this case uses ffp-exception-behavior=struct; i was trying to address https://bugs.llvm.org/show_bug.cgi?id=44571 by checking the command line options to see if strict was enabled. That's why constrained intrinsics are enabled. Evidently that's incorrect.

125

that's a bug. thanks

[sepavloff]

You could check the scope where the pragma appears is the same way as pragma clang fp does. The code

case tok::annot_pragma_fp:
  ProhibitAttributes(Attrs);
  Diag(Tok, diag::err_pragma_fp_scope);
  ConsumeAnnotationToken();
  return StmtError();

is put into Parser::ParseStatementOrDeclarationAfterAttributes and it produces errors on misplaced pragma. Probably the same way may work for you.

[mibintc]

This patch is a work in progress.

The problem that I want to work on next is that the scope is wrong when the pragma token is seen following the right brace of a function body. In this patch I worked around that problem by modifying the test case to insert bogus "class ResetScope;" to force the scope back to file-level before each pragma that occurs at file scope. You can see that workaround in the test case clang/test/CodeGen/fp-floatcontrol-stack.cpp

These are the main differences with previous version of this patch:

1. per Andy Kaylor's comments in the llvm-dev discussion "Floating Point semantic modes", I changed the -ffp-model=precise to be the default. I changed ffp-model=precise to set -ffp-contract=on [previously it was -ffp-contract=fast]. This caused the need to change some tests because the LLVM IR now has the "contract" bit enabled.
2. I put FMF into the FPOptions struct. As I mentioned in a previous comment, I have a request to allow changing the REASSOC bit via source pragma. I plan to submit that in a separate patch. Putting FMF into FPOptions allows the fast bits to be set or disabled via pragma float_control precise, Andy had suggested this change in his code review.
3. I put Andy's distillation of floating point options and floating point modes into UsersManual.rst
4. I changed how the strictfp attribute was set on functions, now it's set if the strict mode is enabled in the function body e.g. either by the option level or a float_control pragma

Still needs to be done:

1. The problem with the file scope I mentioned above
2. I assume that the 2 denormal settings need to be added to FPOptions.
3. Required diagnostics a la MSVC that Andy mentioned in llvm-dev discussion "Floating Point semantic modes"
4. What else?

[lebedev.ri]

I would think contract change can be separated from the rest of the changes, and therefore should be a separate review (to reduce noise)?

[mibintc]

In D72841#1869931, @lebedev.ri wrote:

I would think contract change can be separated from the rest of the changes, and therefore should be a separate review (to reduce noise)?

I split off that change to https://reviews.llvm.org/D74436 and added you as reviewer, thanks.

[mibintc]

I found the problem in the #pragma float_control (push/pop) stack, it was just a dumb bug.

I also added -include-pch test cases, and added code to ASTWriter ASTReader to preserve the floatcontrol pragma stack

For 2 of the new floatcontrol test cases, I temporarily added XFAIL because the patch to default -ffp-precise=fast and ffp-contract=on has been withdrawn (will re-commit soon).

[mibintc]

This patch is code complete and ready for your review and consideration.

[rjmccall]

clang/docs/LanguageExtensions.rst

3055

pragma float_control should be in backticks.

Throughout this documentation, when referring to command-line options, please spell them the way they're actually spelled on the command line, i.e. with a dash.

clang/include/clang/AST/Stmt.h

1104

What's happening here is exactly what this assertion is supposed to prevent. If you need more bits in one of these classes (I assume it's CXXOperatorCallExpr), you need to either make a field in the actual class or investigate more arcane mechanisms like trailing storage to reduce the normal impact. The latter is probably unnecessary for CXXOperatorCallExpr.

clang/include/clang/Basic/DiagnosticParseKinds.td

1138

Maybe "operation" would be a better user-facing name than "kind"? Also, this diagnostic is more specific but less helpful than the diagnostic just above.

clang/include/clang/Basic/LangOptions.def

196

Please align the commas.

Would it make more sense to just store an FPOptions in LangOptions instead of breaking all of the bits down separately?

We may need to reconsider at some point whether any of these are really "compatible" language options. Headers can contain inline code, and we shouldn't compile that incorrectly just because we reused a module we built under different language settings. Although... maybe we can figure out a way to store just the ways that an expression's context overrides the default semantics and then merge those semantics into the default set for the translation unit; that would make them actually compatible. Of course, it would also require more bits in expressions where it matters, and you might need to investigate trailing storage at that point.

clang/include/clang/Basic/LangOptions.h

466

Somewhere in this type, it should be obvious where I can go in order to understand what any of these flags means precisely. Ideally that would be reinforced by the method names, instead of using non-term-of-art abbreviations like "reassoc".

515

The more conventional method names here would an instance method called something like getAsOpaqueInt and then a static method called something like getFromOpaqueInt.

clang/lib/CodeGen/CGExprScalar.cpp

462

You can override VisitBinOp and just do this in that case. But why does it need to be done at this level at all, setting global state on the builder for all emissions, instead of in the leaves where we know we're emitting floating-point operations? This is adding a lot of overhead in some of the most commonly-exercised code paths in IRGen, but building FP expressions is relatively uncommon. I would definitely prefer a little bit of repetitive code over burdening the common case this much. It might also be nice to figure out when this is unnecessary.

Also, please extract a function to make FastMathFlags from FPOptions; you'll need it elsewhere, e.g. in CGExprComplex.

[mibintc]

Responding to @rjmccall 's review. Previously there was a FIXME about adding FPFeatures to UnaryOperator. This seems like the right time to do that, so that change is included there. I'll add inline replies too.

[mibintc]

some inline replies and comments

clang/include/clang/AST/Stmt.h

1104

@rjmccall The reason i changed the assertion is because FPOptions is now wider, so I had to change the assertion. See line 609 above. Is there something I need to do differently?

clang/include/clang/Basic/DiagnosticParseKinds.td

1138

I got rid of the diagnostic with the unhelpful string and just using the single diagnostic which has full information about how to form the pragma

clang/include/clang/Basic/LangOptions.def

196

I aligned the commas. I didn't put FPOptions into LangOptions, would you like me to make that change too? I don't know about trailing storage. I see that term in the code but I didn't see details about what that is/how that works.

clang/include/clang/Basic/LangOptions.h

466

I put the comments on the field declarations in the private part. I changed the names of the accessor methods to be more descriptive. (Previously I was using the same names as LLVM uses for those fields).

515

I changed the names like you suggested but not using the static method, is this OK?

clang/lib/CodeGen/CGExprScalar.cpp

462

I removed it from here and pushed this work towards the leaves. I decided that I should put FPFeatures onto UnaryOperator nodes which was left as a FIXME by an earlier author in this area. I added the FastMathFlags function like you suggested but i suppose it needs to be moved out of this file.

4081

In the previous rendition of this patch, when the Builder.FMF settings were modified at Visit(BinaryExpression), the assign is seen as a binary expression and so the FPFeatures was passed into IRBuilder. I'm not confident this patch is in the right place, I'd really like to put FPFeatures onto the CallExpr node, because if you call a builtin intrinsic function, and the mode is set to float_control(except, on), the call node for the intrinsic doesn't have the FPFeature bits, so it isn't marked as expected. Before I make that change I want @rjmccall to take another look; If FPFeatures was on CallExpr then I'd remove it here and modify IRBuilder.FMF when visiting CallExpr

clang/test/CodeGenOpenCL/builtins-r600.cl

5

OpenCL CompilerInvocation always sets fp_contract to "on"; inside clang I check if either fp_contract==on or fp_contract==fast, that expression is used to set the IRBuilder.FMF.contract bit.

CUDA CompilerInvocation always sets fp_contract to "fast"

[rjmccall]

clang/include/clang/AST/Stmt.h

1104

Because Stmt is a common base class for so much of the AST but only needs to store a small amount of state itself, we have a complicated system for optimizing space usage in subclasses by allocating bit-fields into Stmt. Letting an individual subclass's bit-field usage run over the expected size and therefore inflate Stmt for all subclasses would be counter-productive, hence the static_assert and why it shouldn't be changed. You need to move the storage of FPOptions into the appropriate subclass wherever it would cause the static_assert to fail.

[mibintc]

@rjmccall suggested that I needed to remove FPOptions from the Stmt class since the sizeof assertion failed. I moved that information into the relevant expression nodes and fixed a few small things that I didn't like in the previous version.

[mibintc]

clang/lib/CodeGen/CGExprScalar.cpp

620

The call expr might be a call to an intrinsic, the floating point intrinsic calls need to be marked properly with information from FPFeatures

4081

I got rid of the bogus code here and moved it into VisitCallExpr where it belongs.

[mibintc]

I got an email like this "Harbormaster failed remote builds in B48237: Diff 248536!" but there is no further information. It builds OK from my workstation. I did have to paste the review because an upload to Phabricator exceeded the size limit, could that be why? Is there a way to get more information about the build failure?

[rjmccall]

In D72841#1908084, @mibintc wrote:

@rjmccall suggested that I needed to remove FPOptions from the Stmt class since the sizeof assertion failed. I moved that information into the relevant expression nodes and fixed a few small things that I didn't like in the previous version.

You only need to do this for the expression nodes where it causes an overflow in the size of the bit-field; I don't think you're overflowing the capacity of UnaryOperatorBitfields, for example.

[mibintc]

In D72841#1911330, @rjmccall wrote:

In D72841#1908084, @mibintc wrote:

@rjmccall suggested that I needed to remove FPOptions from the Stmt class since the sizeof assertion failed. I moved that information into the relevant expression nodes and fixed a few small things that I didn't like in the previous version.

You only need to do this for the expression nodes where it causes an overflow in the size of the bit-field; I don't think you're overflowing the capacity of UnaryOperatorBitfields, for example.

I added unsigned FPFeatures : 14; to class UnaryOperatorBitfields but I encountered the same assertion failure in the Stmt constructor, In constructor ‘clang::Stmt::Stmt(clang::Stmt::StmtClass)’:
/iusers/sandbox/pragma-ws/llvm-project/clang/include/clang/AST/Stmt.h:1095:5: error: static assertion failed: changing bitfields changed sizeof(Stmt)

static_assert(sizeof(*this) <= 8

[rjmccall]

In D72841#1912416, @mibintc wrote:

In D72841#1911330, @rjmccall wrote:

In D72841#1908084, @mibintc wrote:

@rjmccall suggested that I needed to remove FPOptions from the Stmt class since the sizeof assertion failed. I moved that information into the relevant expression nodes and fixed a few small things that I didn't like in the previous version.

You only need to do this for the expression nodes where it causes an overflow in the size of the bit-field; I don't think you're overflowing the capacity of UnaryOperatorBitfields, for example.

I added unsigned FPFeatures : 14; to class UnaryOperatorBitfields but I encountered the same assertion failure in the Stmt constructor, In constructor ‘clang::Stmt::Stmt(clang::Stmt::StmtClass)’:
/iusers/sandbox/pragma-ws/llvm-project/clang/include/clang/AST/Stmt.h:1095:5: error: static assertion failed: changing bitfields changed sizeof(Stmt)

static_assert(sizeof(*this) <= 8

Oh, okay. It's actually pretty unfortunate that we need to increase the memory use of every UnaryOperator, BinaryOperator, and CallExpr in the AST just in case these pragmas are in use. Most instances of these expressions provably have nothing to do with floating-point, and in practice the pragmas are very rarely used.

I mentioned up-thread that it would be nice to make serialized AST actually agnostic about the basic language settings. That would require us to track not just the active option set, but what overrides were actually in place when parsing any given expression; we could then merge the two pretty easily in IRGen. (We could represent this with masks: we could then merge the information by taking the default-settings mask, bitwise-and'ing with the override mask to clear any values that will be overridden, and bitwise-or'ing the override values on top.) But if we did that, then it would be easy for us to identify when creating an AST node that no overrides are in effect; that would let us just store a single flag in the class indicating whether there were overrides, then use trailing storage for the actual overrides if present. This would relieve almost all of the pressure to keep the storage size down.

We could take it a step further by trying to recognize when building certain AST nodes that they're not floating-point operations and just building them with no overrides.

[mibintc]

@rjmccall Since CompoundAssignmentOperator derives from BinaryOperator, it's not simple to add Trailing storage here. I think I will have to fold CompoundAssignmentOperator into BinaryOperator and then add the 2 extra fields needed by CompoundAssignmentOperator into Trailing storage. Can you think of a better way? I worked on Trailing storage for UnaryOperator first and that wasn't too bad, but Binary is a different story.

[rjmccall]

In D72841#1917340, @mibintc wrote:

@rjmccall Since CompoundAssignmentOperator derives from BinaryOperator, it's not simple to add Trailing storage here. I think I will have to fold CompoundAssignmentOperator into BinaryOperator and then add the 2 extra fields needed by CompoundAssignmentOperator into Trailing storage. Can you think of a better way? I worked on Trailing storage for UnaryOperator first and that wasn't too bad, but Binary is a different story.

It's something we deal with occasionally, but it's definitely annoying. You basically have to test for which concrete class you have and then ask that class for its trailing storage.

Collapsing the types might be okay but could get involved.

[rjmccall]

In D72841#1917459, @rjmccall wrote:

In D72841#1917340, @mibintc wrote:

@rjmccall Since CompoundAssignmentOperator derives from BinaryOperator, it's not simple to add Trailing storage here. I think I will have to fold CompoundAssignmentOperator into BinaryOperator and then add the 2 extra fields needed by CompoundAssignmentOperator into Trailing storage. Can you think of a better way? I worked on Trailing storage for UnaryOperator first and that wasn't too bad, but Binary is a different story.

It's something we deal with occasionally, but it's definitely annoying. You basically have to test for which concrete class you have and then ask that class for its trailing storage.

Collapsing the types might be okay but could get involved.

To be clear, we do this in the Clang AST by doing a lot of hand-rolled trailing storage rather than by using llvm::TrailingStorage. If you're willing to do the refactors necessary to use the latter, though, that's great; but as mentioned, you'll need to either collapse CompoundAssignmentOperator into BinaryOperator or introduce a concrete subclass for the non-compound operators and make BinaryOperator an abstract node.

[mibintc]

The previous version of this patch was changed to move BinaryOperator.FPFeatures into TrailingStorage. This patch is rebased on that change. Looking for code review, (check-clang passes) thank you!

[mibintc]

added a couple inline explanatory comments

clang/include/clang/Basic/LangOptions.h

303

I added this boolean as part of validating the correctness of the pragma's that access the FP environment, according to the Microsoft checks.. Copying from the Microsoft doc: "There are restrictions on the ways you can use the fenv_access pragma in combination with other floating-point settings:

You can't enable fenv_access unless precise semantics are enabled. Precise semantics can be enabled either by the float_control pragma, or by using the /fp:precise or /fp:strict compiler options. The compiler defaults to /fp:precise if no other floating-point command-line option is specified.

You can't use float_control to disable precise semantics when fenv_access(on) is set."
This is copied from https://docs.microsoft.com/en-us/cpp/preprocessor/fenv-access?view=vs-2019

clang/test/CodeGen/constrained-math-builtins.c

161

Since this patch constructs the FPFeatures using the floating point settings from the command line versus the default FPOptions() constructor, several tests need to be changed. Some of the changes I made showing the flags on the IR, other tests I changed by adding ffp-contract to the RUN line to match the expected IR.

[mibintc]

noted bug in an inline comment. i will upload a fix

[mibintc]

corrected the bitfield width in Stmt.h for FPFeatures on cxx operator call

[mibintc]

clang/docs/LanguageExtensions.rst

3064

there's an extra whitespace here, i'll get rid of it

[erichkeane]

@rjmccall : can you comment on the CallExpr's trailing storage issue? And give advice as to what you meant in the last review?

clang/include/clang/AST/Expr.h

2122 ↗

(On Diff #260320)

needs a newline between functions.

2251 ↗

(On Diff #260320)

Is this really useful/usable at all? It seems to me that since this would require re-allocating this object that noone should be able to set this after construction.

2255 ↗

(On Diff #260320)

Why is this a separate function from getTrailingFPFeatures?

2256 ↗

(On Diff #260320)

If they have to be separate, the assert here doesn't really make sense, since getTrailingFPFeatures has the same assert.

2261 ↗

(On Diff #260320)

For the asserts, we should probably prefer using the hasStoredFPFeatures function.

2268 ↗

(On Diff #260320)

Is there use in having both this AND the get-stored, as opposed to just making everyone access via the same function? At least having 2 public versions aren't very clear what the difference is to me.

2701 ↗

(On Diff #260320)

This type already has trailing-storage type stuff. I think in the past review @rjmccall encouraged you to put this in the fake-trailing-storage like above.

clang/include/clang/Basic/LangOptions.h

193–209

Is this an unrelated change? What is the purpose for this?

372–382

Is this the same logic as getFromOpaqueInt? If so, we should probably just call that.

clang/include/clang/Sema/Sema.h

555

This comment is really oddly phrased and uses the 'stack'-noun as a verb? Something like: (please feel free to wordsmith): "This stack tracks the current state of Sema.CurFPFeatures."?

clang/lib/AST/Expr.cpp

1309 ↗

(On Diff #260320)

Is there a reason this isn't a member initializer?

[erichkeane]

clang/lib/Parse/ParsePragma.cpp

668

Oh boy these are some magic lookin numbers... Can you document these two lines?

2536

Replace this with BalancedDelimiterTracker instead, it gives consistent errors and are a bit easier to use. Additionally, I think it does some fixups that allow us to recover better.

I'd also suggest some refactoring with the PragmaFloatControlKind if/elses below. Perhaps handle the closing paren at the end, and do a switch-case for that handling.

clang/lib/Sema/SemaAttr.cpp

428

I guess I don't get why you're switching on both here? Can the two just be combined? I don't know if the 'NewValue = CurFPFeatures.getAsOpaqueInt();

FpPragmaStack.Act(Loc, Action, StringRef(), NewValue);'

part is sufficiently motivated to do 2 separate switches.

1004

Should we still be setting this even if there was an error?

clang/lib/Sema/SemaStmt.cpp

377

unrelated change?

400

Don't use curleys for single liners, both of these probably shouldn't need curleys at all. Comment could be at the top for clarity.

clang/lib/Serialization/ASTReaderStmt.cpp

684 ↗

(On Diff #260320)

Rather than this variable, why not just ask 'E' below?

[mibintc]

A couple replies to @erichkeane

clang/include/clang/AST/Expr.h

2251 ↗

(On Diff #260320)

It's only used during serialization (ASTReader); I guess the node has already been allocated by then so it's superfluous, because the allocation point could set this field.

2268 ↗

(On Diff #260320)

John suggested the name getStored hasStored as "less tempting" names. The getStored and hasStored are only used during Serialization. John suggested the getFPFeatures function as the public interface and it uses the LangOptions parameter. The features are saved in the node if they can't be recreated from the command line floating point options (due to presence of floating point pragma)

2701 ↗

(On Diff #260320)

whoops i meant that to go to the CallExprBits

clang/include/clang/Basic/LangOptions.h

193–209

it's a NFC the llvm:: prefix wasn't needed. maybe the clang formatter did that?

clang/lib/Serialization/ASTReaderStmt.cpp

684 ↗

(On Diff #260320)

yes i could do that. it would be a function call

[mibintc]

I dropped FPOptions from CallExpr -- if it's needed I will add it using TrailingStorage in a separate patch; I responded to @erichkeane 's code review

[mibintc]

clang/include/clang/AST/Expr.h

2701 ↗

(On Diff #260320)

Adding FPOptions to CallExprBits makes the field too large, I'm going to drop adding FPOptions to CallExpr, i'll add it via trailingstorage in a separate patch if it's needed.

clang/lib/Parse/ParsePragma.cpp

2536

BalancedDelimiterTracker doesn't work here because there's no access to the Parser object. Rewriting it would be an extensive change and I'm doubtful about making this change. PragmaHandler is defined in Lex. I think there are 60 pragma's that use the PragmaHandler.

clang/lib/Sema/SemaAttr.cpp

1004

Should we still be setting this even if there was an error?

It's not harmful to set it, if there's an error diagnostic then there is no codegen right?

clang/lib/Serialization/ASTReaderStmt.cpp

684 ↗

(On Diff #260320)

i made some changes in this area, eliminating setHasStoredFPFeatures

[erichkeane]

2 small things, @rjmccall and @sepavloff , anything else?

clang/include/clang/Basic/DiagnosticSemaKinds.td

866

The last 4 can be done via selects as well! Save a couple more spaces before we have to up the diagnostic id size :)

clang/include/clang/Sema/Sema.h

555

Just needs a period at the end.

clang/lib/Parse/ParsePragma.cpp

2536

Thats unfortunate :/ That type does some nice fixup work.

[mibintc]

I changed a comment to add a period, rebased, and used clang-format

[mibintc]

clang/include/clang/Basic/DiagnosticSemaKinds.td

866

The last 4 can be done via selects as well

Combining these 4 into 1 diagnostic is doable but it's ugly.

[erichkeane]

clang/include/clang/Basic/DiagnosticSemaKinds.td

866

Concur, I spent some time on it and don't really like it.

[erichkeane]

Looks OK to me. Please give @sepavloff and @rjmccall a day or so to make final comments before committing.

[uabelho]

Hi @mibintc ,

I think I'm seeing some oddities with this patch (current trunk, 0054c46095).
With

clang -S -Xclang -emit-llvm bbi-42553.c -o -

on the input

float rintf(float x);
#pragma STDC FENV_ACCESS ON

void t()
{
    (void)rintf(0.0f);
}

I get

bbi-42553.c:2:14: warning: pragma STDC FENV_ACCESS ON is not supported, ignoring pragma [-Wunknown-pragmas]
#pragma STDC FENV_ACCESS ON
             ^
; ModuleID = 'bbi-42553.c'
source_filename = "bbi-42553.c"
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

; Function Attrs: noinline nounwind optnone strictfp uwtable
define dso_local void @t() #0 {
entry:
  %0 = call float @llvm.experimental.constrained.rint.f32(float 0.000000e+00, metadata !"round.tonearest", metadata !"fpexcept.ignore") #2
  ret void
}

and if I remove the pragma I instead get

define dso_local void @t() #0 {
entry:
  %0 = call float @llvm.rint.f32(float 0.000000e+00)
  ret void
}

Before this patch I got the call to llvm.rint.f32 in both cases.

[uabelho]

It seems to be this change in SemaStmt.cpp that makes the FENV-pragma have some effect regardless of the warning saying that the pragma doesn't have any effect:

index aa0d89ac09c3..f76994a6dab3 100644
@@ -394,6 +394,11 @@ StmtResult Sema::ActOnCompoundStmt(SourceLocation L, SourceLocation R,
                                    ArrayRef<Stmt *> Elts, bool isStmtExpr) {
   const unsigned NumElts = Elts.size();
 
+  // Mark the current function as usng floating point constrained intrinsics
+  if (getCurFPFeatures().isFPConstrained())
+    if (FunctionDecl *F = dyn_cast<FunctionDecl>(CurContext))
+      F->setUsesFPIntrin(true);
+

[mibintc]

In D72841#2022354, @uabelho wrote:

It seems to be this change in SemaStmt.cpp that makes the FENV-pragma have some effect regardless of the warning saying that the pragma doesn't have any effect:

index aa0d89ac09c3..f76994a6dab3 100644
@@ -394,6 +394,11 @@ StmtResult Sema::ActOnCompoundStmt(SourceLocation L, SourceLocation R,
                                    ArrayRef<Stmt *> Elts, bool isStmtExpr) {
   const unsigned NumElts = Elts.size();
 
+  // Mark the current function as usng floating point constrained intrinsics
+  if (getCurFPFeatures().isFPConstrained())
+    if (FunctionDecl *F = dyn_cast<FunctionDecl>(CurContext))
+      F->setUsesFPIntrin(true);
+

Thank you I will follow up!

[rjmccall]

We're also seeing test failures in Apple's Clang fork, e.g.

test/CodeGen/finite-math.c:12:10: error: NSZ: expected string not found in input
 // NSZ: fadd nsz
         ^
<stdin>:11:20: note: scanning from here
define void @foo() #0 {
                   ^
<stdin>:15:9: note: possible intended match here
 %add = fadd float %0, %1
        ^

I haven't checked yet whether we can reproduce upstream.

[mibintc]

I posted a patch to fix the bug reported by @uabelho here https://reviews.llvm.org/D79510

@rjmccall I used check-clang and check-all on D71841 from my linux x86-64 server before submitting, and the testing was clear. Maybe your branch has some calls to create BinaryOperator that isn't passing FPFeatures with the correct value--related to D76384. My sandbox is showing this

%add = fadd nsz float %0, %1

[mibintc]

I got a report that this patch was causing a problem with Windows <numeric> header because #pragma float_control should be supported in namespace context. I've posted a patch for review here https://reviews.llvm.org/D79631

[plotfi]

Hi @rjmccall, I am also seeing similar failures. It is failing on apple's master (and the swift branches as well) because ParseLangArgs and ParseCodeGenArgs are getting called in the opposite order in apple/master from the order they are called in llvm/master. I posted a PR to fix those failures here: https://github.com/apple/llvm-project/pull/1202

but I don't know if this is the most correct approach.

In D72841#2023058, @rjmccall wrote:

We're also seeing test failures in Apple's Clang fork, e.g.

test/CodeGen/finite-math.c:12:10: error: NSZ: expected string not found in input
 // NSZ: fadd nsz
         ^
<stdin>:11:20: note: scanning from here
define void @foo() #0 {
                   ^
<stdin>:15:9: note: possible intended match here
 %add = fadd float %0, %1
        ^

I haven't checked yet whether we can reproduce upstream.

[rjmccall]

In D72841#2027740, @plotfi wrote:

Hi @rjmccall, I am also seeing similar failures. It is failing on apple's master (and the swift branches as well) because ParseLangArgs and ParseCodeGenArgs are getting called in the opposite order in apple/master from the order they are called in llvm/master. I posted a PR to fix those failures here: https://github.com/apple/llvm-project/pull/1202

but I don't know if this is the most correct approach.

Oh, thank you for figuring that out. Yeah, it's reasonable for code-gen option parsing to depend on language-option parsing, which means that this patch is wrong. The right fix is that we need to stop parsing these as code-gen options, which is reasonable since they have direct language-semantics impact. If we still need the code-gen option flags, we should be able to recreate them from the language options, I think.

[plotfi]

@rjmccall @mibintc Can we revert this patch for now then, and re-land when this patch is reworked? It would be good to get those bots passing. @rjmccall are the bots that you see failing on your end public?

In D72841#2028834, @rjmccall wrote:

In D72841#2027740, @plotfi wrote:

Hi @rjmccall, I am also seeing similar failures. It is failing on apple's master (and the swift branches as well) because ParseLangArgs and ParseCodeGenArgs are getting called in the opposite order in apple/master from the order they are called in llvm/master. I posted a PR to fix those failures here: https://github.com/apple/llvm-project/pull/1202

but I don't know if this is the most correct approach.

Oh, thank you for figuring that out. Yeah, it's reasonable for code-gen option parsing to depend on language-option parsing, which means that this patch is wrong. The right fix is that we need to stop parsing these as code-gen options, which is reasonable since they have direct language-semantics impact. If we still need the code-gen option flags, we should be able to recreate them from the language options, I think.

[mibintc]

This comment was removed by mibintc.

[mibintc]

Some inline replies/comments to @rjmccall and @plotfi

clang/lib/Frontend/CompilerInvocation.cpp

3186

@rjmccall @plotfi These earlier patches are also deriving the value of LangOpts from the settings of CG opts

3188

@rjmccall @plotfi here the codegen args are evaluated first. Perhaps we could have a "reconcile codegen and lang args" function which would resolve the floating point settings into a final setting? so that codegen or lang could be parsed in either order?

[mibintc]

@rjmccall Uncertain how to proceed, can you recommend? If I recall correctly, I added the lines in CompilerOptions because there were many failing lit tests, i could have fixed the lit fails by adding the lang options to the lit tests. (of course that change could have other repercussions)

[mibintc]

clang/lib/Frontend/CompilerInvocation.cpp

3193

@rjmccall I could set these by using Args.hasArg instead of CGOpts, would that be acceptabel?

[rjmccall]

clang/lib/Frontend/CompilerInvocation.cpp

3186

I don't know what you mean here; the code you're quoting just seems to be looking at Args. It's fine to re-parse arguments in both places if that makes something easier. The problem is that you're looking at the CodeGenOptions structure itself.

[rjmccall]

clang/lib/Frontend/CompilerInvocation.cpp

3193

I think so, yes. Ideally the CG options would then be set based on the earlier values, or replaced with uses of the language options structure, but not having a direct dependency at all may be simpler.

[plotfi]

@ab @rjmccall @mibintc Posted D79730 for consideration.
@mibintc can you produce a version of _this_ diff that works with D79730 applied. Currently the following fail, as they do on Apple Master:

Failing Tests (4):
  Clang :: CodeGen/finite-math.c
  Clang :: CodeGen/fp-floatcontrol-stack.cpp
  Clang :: CodeGenOpenCL/relaxed-fpmath.cl
  Clang :: Frontend/diagnostics-order.c

In D72841#2029309, @mibintc wrote:

@rjmccall Uncertain how to proceed, can you recommend? If I recall correctly, I added the lines in CompilerOptions because there were many failing lit tests, i could have fixed the lit fails by adding the lang options to the lit tests. (of course that change could have other repercussions)

[mibintc]

In D72841#2029821, @plotfi wrote:

@ab @rjmccall @mibintc Posted D79730 for consideration.
@mibintc can you produce a version of _this_ diff that works with D79730 applied. Currently the following fail, as they do on Apple Master:

@rjmccall accepted the proposed patch https://reviews.llvm.org/D79735, so I pushed that. I also tried your patch and the 3 CodeGen tests pass but the diagnostics-order.c test fails

Failing Tests (4):
  Clang :: CodeGen/finite-math.c
  Clang :: CodeGen/fp-floatcontrol-stack.cpp
  Clang :: CodeGenOpenCL/relaxed-fpmath.cl
  Clang :: Frontend/diagnostics-order.c

[plotfi]

That was a good fix. I am pretty sure this does mean the diagnostics-order.c will fail on apple's bots. The same diagnostics lines print, but in the wrong order. I haven't root caused that yet.

In D72841#2030099, @mibintc wrote:

In D72841#2029821, @plotfi wrote:

@ab @rjmccall @mibintc Posted D79730 for consideration.
@mibintc can you produce a version of _this_ diff that works with D79730 applied. Currently the following fail, as they do on Apple Master:

@rjmccall accepted the proposed patch https://reviews.llvm.org/D79735, so I pushed that. I also tried your patch and the 3 CodeGen tests pass but the diagnostics-order.c test fails

Failing Tests (4):
  Clang :: CodeGen/finite-math.c
  Clang :: CodeGen/fp-floatcontrol-stack.cpp
  Clang :: CodeGenOpenCL/relaxed-fpmath.cl
  Clang :: Frontend/diagnostics-order.c

[plotfi]

@rjmccall @mibintc I think the diagnostics-order.c test is still behaving correctly technically. The note lines are still printing with the associated error lines, it just happens that one of the warning lines prints at the end instead of in the middle. ie:

error: invalid value '-foo' in '-verify='
note: -verify prefixes must start with a letter and contain only alphanumeric characters, hyphens, and underscores
error: invalid value 'bogus' in '-std=bogus'
note: use 'c89', 'c90', or 'iso9899:1990' for 'ISO C 1990' standard
...
warning: optimization level '-O999' is not supported; using '-O3' instead

instead of

error: invalid value '-foo' in '-verify='
note: -verify prefixes must start with a letter and contain only alphanumeric characters, hyphens, and underscores
warning: optimization level '-O999' is not supported; using '-O3' instead
error: invalid value 'bogus' in '-std=bogus'
note: use 'c89', 'c90', or 'iso9899:1990' for 'ISO C 1990' standard
...

In D72841#2030468, @plotfi wrote:

That was a good fix. I am pretty sure this does mean the diagnostics-order.c will fail on apple's bots. The same diagnostics lines print, but in the wrong order. I haven't root caused that yet.

In D72841#2030099, @mibintc wrote:

In D72841#2029821, @plotfi wrote:

@ab @rjmccall @mibintc Posted D79730 for consideration.
@mibintc can you produce a version of _this_ diff that works with D79730 applied. Currently the following fail, as they do on Apple Master:

@rjmccall accepted the proposed patch https://reviews.llvm.org/D79735, so I pushed that. I also tried your patch and the 3 CodeGen tests pass but the diagnostics-order.c test fails

Failing Tests (4):
  Clang :: CodeGen/finite-math.c
  Clang :: CodeGen/fp-floatcontrol-stack.cpp
  Clang :: CodeGenOpenCL/relaxed-fpmath.cl
  Clang :: Frontend/diagnostics-order.c

[rjmccall]

I think we might have had to change that test on our fork when we changed the parsing order.

[michele.scandale]

clang/lib/CodeGen/CGExprScalar.cpp

226

I'm not convinced it correct to set contract when allowFPContractWithinStatement return true. Can someone clarify this?

If I compile the following example with -ffp-contract=on:

float test1(float a, float b, float c) {
  float x = a * b;
  return x + c;
}

float test2(float a, float b, float c) {
  return a * b + c;
}

Before this change the generated code was:

define float @test1(float %a, float %b, float %c) {
  %0 = fmul float %a, %b
  %1 = fadd float %0, %c
  ret float %1
}

define float @test2(float %a, float %b, float %c) {
  %0 = call float @llvm.fmuladd.f32(float %a, float%b, float %c)
  ret float %0
}

And my understanding is that the in-statement contraction is implemented by emitting the llvm.fmuladd call that a backend might decide to implement as fmul + fadd or as fma.

With this change the generated code is:

define float @test1(float %a, float %b, float %c) {
  %0 = fmul contract float %a, %b
  %1 = fadd contract float %0, %c
  ret float %1
}

define float @test2(float %a, float %b, float %c) {
  %0 = call contract float @llvm.fmuladd.f32(float %a, float%b, float %c)
  ret float %0
}

and it seems to me that in test1 (where multiple statements where explicitly used) the optimizer is now allowed to perform the contraction, violating the original program semantic where only "in-statement" contraction was allowed.

[rjmccall]

IIUC, the way within-statement contraction is supposed to work is that there are supposed to be blocking intrinsics inserted at various places. I don't remember the details, though, or if anyone's thought about how it interacts with cross-statement contraction, which this pragma permits within the same function (and could occur with inlining regardless).

[michele.scandale]

clang/include/clang/Basic/LangOptions.h

398–403

Same comment on LangOpts.FastMath || as the one for CompilerInvocation.cpp.

clang/lib/Frontend/CompilerInvocation.cpp

3191–3197

Why do we need Opts.FastMath || here? The code in the compiler driver clang/lib/Driver/ToolChains/Clang.cpp (https://github.com/llvm/llvm-project/blob/master/clang/lib/Driver/ToolChains/Clang.cpp#L2510) already takes care of generating the right flags for the CC1 to configure the floating point rules.

Moreover, if we ignore what the compiler driver does, the fact that Args.hasArg(OPT_ffast_math) is not considered in the definition of the codegen options such as NoInfsFPMath, NoNaNsFPMath, NoSignedZeros, Reassociate, so the you have already two distinct options for the same abstract property that might not match.

I think that at the CC1 level the reasoning should be done in terms of the fine grain options, and let the compiler driver makes life easy for the users -- i.e. LangOpts.FastMath should just control whether the macro __FAST_MATH__ is defined or not.

[michele.scandale]

In D72841#2030707, @rjmccall wrote:

IIUC, the way within-statement contraction is supposed to work is that there are supposed to be blocking intrinsics inserted at various places. I don't remember the details, though, or if anyone's thought about how it interacts with cross-statement contraction, which this pragma permits within the same function (and could occur with inlining regardless).

Prior to this change contract was never generated in the case of in-statement contraction only, instead clang was emitting llvm.fmuladd to inform the backend that only those were eligible for contraction. From a correctness perspective I think this was perfectly fine.

Currently I don't see any logic to generate "blocking intrinsics" (I guess to define a region around the instructions emitted for the given statement). Until such mechanism is in place, I think that generating the contract fast-math flag also for in-statement contraction is wrong because it breaks the original program semantic.

Am I missing something?

[rjmccall]

No, if that's how we handle that, then you're absolutely right that we shouldn't set the contract flag.

[scanon]

Prior to this change contract was never generated in the case of in-statement contraction only, instead clang was emitting llvm.fmuladd to inform the backend that only those were eligible for contraction. From a correctness perspective I think this was perfectly fine.

Currently I don't see any logic to generate "blocking intrinsics" (I guess to define a region around the instructions emitted for the given statement). Until such mechanism is in place, I think that generating the contract fast-math flag also for in-statement contraction is wrong because it breaks the original program semantic.

This is exactly right. If we are going to have new in-statement optimizations, then we probably do need to add some blocking intrinsic (which would be elidable given suitable fast-math flags); the system of generating fmuladd works well for FMA contraction, but doesn't really generalize to other optimizations of this sort.

[mibintc]

clang/lib/CodeGen/CGExprScalar.cpp

226

Thanks @michele.scandale i will work on a patch for this

[yaxunl]

clang/lib/Serialization/ASTReader.cpp

7871

This changes the behavior regarding AST reader and seems to be too hash restriction. Essentially this requires a pch can only be used with the same fp options with which the pch is generated. Since there are lots of fp options, it is impractical to generate pch for all the combinations.

We have seen regressions due to this assertion.

Can this assertion be dropped or done under some options?

Thanks.

[mibintc]

clang/lib/Serialization/ASTReader.cpp

7871

@yaxunl Can you please send me a reproducer, I'd like to see what's going on, not sure if just getting rid of the assertion will give the desired outcome.

[yaxunl]

clang/lib/Serialization/ASTReader.cpp

7871

diff.fp-options.1.1.txt444 BDownload

Pls apply the patch.

Thanks.

[mibintc]

clang/lib/CodeGen/CGExprScalar.cpp

226

@michele.scandale I posted a patch for 'contract' here, https://reviews.llvm.org/D79903

clang/lib/Serialization/ASTReader.cpp

7871

@rjmccall In the example supplied by @yaxunl, the floating point options in the pch file when created are default, and the floating point options in the use have no-signed-zeros flag. The discrepancy causes an error diagnostic when the pch is used. I added the FMF flags into FPFeatures in this patch, I made them COMPATIBLE_LANGOPT which is the encoding also being used for FastMath, FiniteMathOnly, and UnsafeFPMath. Do you have some advice about this issue?

[rjmccall]

clang/lib/Serialization/ASTReader.cpp

7871

A couple things are going on here.

First: a PCH can only end at the top level, not in the middle of a declaration, but otherwise Sema can be in an arbitrary semantic configuration. That definitely includes arbitrary pragmas being in effect, so in general the end state might not match the default FP state, so this assertion is bogus. When loading a PCH, you need to restore the pragma stack and current FP state to the configuration it was in at the end of the PCH.

Second: if you restore the pragma stack and FP state naively given the current representation of FP state, you will completely overwrite the FP settings of the current translation unit with the FP settings that were in effect when the PCH was built, which is obviously not okay. This is one way (among several) that the current representation is not really living up to the statement that these language options are "compatible". The better way to do this would be for the pragma stack and Expr nodes to record the current set of overrides in effect rather than the absolute current state; this could then be easily applied to an arbitrary global FP state.

[michele.scandale]

clang/lib/CodeGen/CGExprScalar.cpp

226

Thanks!

[mibintc]

I documented the issue reported by @yaxunl here, https://bugs.llvm.org/show_bug.cgi?id=46166, and take ownership of the bug. Thanks for the report.

[cfang]

-ffast-math flag got lost in the Builder after this change.

FMF.isFast() is true before updateFastMathFlags(FMF, FPFeatures), but turns false after.
It seems the Builder.FMF has been correctly set before, but I am not clear what FPFeatures should be at this point:

+static void setBuilderFlagsFromFPFeatures(CGBuilderTy &Builder,
+ CodeGenFunction &CGF,
+ FPOptions FPFeatures) {
+ auto NewRoundingBehavior = FPFeatures.getRoundingMode();
+ Builder.setDefaultConstrainedRounding(NewRoundingBehavior);
+ auto NewExceptionBehavior =
+ ToConstrainedExceptMD(FPFeatures.getExceptionMode());
+ Builder.setDefaultConstrainedExcept(NewExceptionBehavior);
+ auto FMF = Builder.getFastMathFlags();
+ updateFastMathFlags(FMF, FPFeatures);
+ Builder.setFastMathFlags(FMF);
+ assert((CGF.CurFuncDecl == nullptr || Builder.getIsFPConstrained() ||
+ isa<CXXConstructorDecl>(CGF.CurFuncDecl) ||
+ isa<CXXDestructorDecl>(CGF.CurFuncDecl) ||
+ (NewExceptionBehavior == llvm::fp::ebIgnore &&
+ NewRoundingBehavior == llvm::RoundingMode::NearestTiesToEven)) &&
+ "FPConstrained should be enabled on entire function");
+}

[mibintc]

I got a bug report about this patch, see https://bugs.llvm.org/show_bug.cgi?id=49479. I put a patch to fix it here, https://reviews.llvm.org/D98211

clang/docs/LanguageExtensions.rst

for(...) {
     d[i] += a;
   }
 
 
 The pragma can also be used with ``off`` which turns FP contraction off for a
 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)
 ===========================================================================
 
 The ``#pragma clang attribute`` directive can be used to apply an attribute to
 multiple declarations. The ``#pragma clang attribute push`` variation of the
 directive pushes a new "scope" of ``#pragma clang attribute`` that attributes
 can be added to. The ``#pragma clang attribute (...)`` variation adds an
 attribute to that scope, and the ``#pragma clang attribute pop`` variation pops

clang/include/clang/AST/Stmt.h

class BinaryOperatorBitfields {
     friend class BinaryOperator;
 
     unsigned : NumExprBits;
 
     unsigned Opc : 6;
 
     /// This is only meaningful for operations on floating point
     /// types and 0 otherwise.
-    unsigned FPFeatures : 8;
+    unsigned FPFeatures : 14;
 
     SourceLocation OpLoc;
   };
 
   class InitListExprBitfields {
     friend class InitListExpr;
 
     unsigned : NumExprBits;
class CXXOperatorCallExprBitfields {
 
     unsigned : NumCallExprBits;
 
     /// The kind of this overloaded operator. One of the enumerator
     /// value of OverloadedOperatorKind.
     unsigned OperatorKind : 6;
 
     // Only meaningful for floating point types.
-    unsigned FPFeatures : 8;
+    unsigned FPFeatures : 14;
   };
 
   class CXXRewrittenBinaryOperatorBitfields {
     friend class ASTStmtReader;
     friend class CXXRewrittenBinaryOperator;
 
     unsigned : NumCallExprBits;
 

 public:
   Stmt() = delete;
   Stmt(const Stmt &) = delete;
   Stmt(Stmt &&) = delete;
   Stmt &operator=(const Stmt &) = delete;
   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!");
     StmtBits.sClass = SC;
     StmtBits.IsOMPStructuredBlock = false;
     if (StatisticsEnabled) Stmt::addStmtClass(SC);
   }
 
   StmtClass getStmtClass() const {

clang/include/clang/Basic/DiagnosticParseKinds.td

 // - #pragma unused
 def warn_pragma_unused_expected_var : Warning<
   "expected '#pragma unused' argument to be a variable name">,
   InGroup<IgnoredPragmas>;
 // - #pragma init_seg
 def warn_pragma_init_seg_unsupported_target : Warning<
   "'#pragma init_seg' is only supported when targeting a "
   "Microsoft environment">,
   InGroup<IgnoredPragmas>;
-// - #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">,
    InGroup<UnknownPragmas>;
 def warn_stdc_fenv_access_not_supported :
    Warning<"pragma STDC FENV_ACCESS ON is not supported, ignoring pragma">,
    InGroup<UnknownPragmas>;
 // - #pragma comment
 def err_pragma_comment_malformed : Error<
   "pragma comment requires parenthesized identifier and optional string">;
 def err_pragma_comment_unknown_kind : Error<"unknown kind of pragma comment">;
 // PS4 recognizes only #pragma comment(lib)
 def warn_pragma_comment_ignored : Warning<"'#pragma comment %0' ignored">,
   InGroup<IgnoredPragmas>;
 // - #pragma detect_mismatch
 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"
   "'single_inheritance', 'multiple_inheritance', or 'virtual_inheritance'">;
 // - #pragma clang optimize on/off
 def err_pragma_optimize_invalid_argument : Error<
   "unexpected argument '%0' to '#pragma clang optimize'; "
   "expected 'on' or 'off'">;
def err_pragma_loop_invalid_option : Error<
   "vectorize_width, interleave, interleave_count, unroll, unroll_count, "
   "pipeline, pipeline_initiation_interval, vectorize_predicate, or distribute">;
 
 def err_pragma_fp_invalid_option : Error<
   "%select{invalid|missing}0 option%select{ %1|}0; expected contract">;
 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_invalid_keyword : Error<
   "invalid argument; expected 'enable'%select{|, 'full'}0%select{|, 'assume_safety'}1 or 'disable'">;
 def err_pragma_pipeline_invalid_keyword : Error<
     "invalid argument; expected 'disable'">;
 
 // Pragma unroll support.
 def warn_pragma_unroll_cuda_value_in_parens : Warning<
   "argument to '#pragma unroll' should not be in parentheses in CUDA C/C++">,

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def note_pragma_pack_pop_instead_reset : Note<
   "did you intend to use '#pragma pack (pop)' instead of '#pragma pack()'?">;
 // Follow the Microsoft implementation.
 def warn_pragma_pack_show : Warning<"value of #pragma pack(show) == %0">;
 def warn_pragma_pack_pop_identifier_and_alignment : Warning<
   "specifying both a name and alignment to 'pop' is undefined">;
 def warn_pragma_pop_failed : Warning<"#pragma %0(pop, ...) failed: %1">,
   InGroup<IgnoredPragmas>;
+def err_pragma_fc_pp_scope : Error<
+  "'#pragma float_control push/pop' can only appear at file scope">;
 def warn_cxx_ms_struct :
   Warning<"ms_struct may not produce Microsoft-compatible layouts for classes "
           "with base classes or virtual functions">,
   DefaultError, InGroup<IncompatibleMSStruct>;
 def err_section_conflict : Error<"%0 causes a section type conflict with %1">;
 def err_no_base_classes : Error<"invalid use of '__super', %0 has no base classes">;
 def err_invalid_super_scope : Error<"invalid use of '__super', "
   "this keyword can only be used inside class or member function scope">;
 def err_super_in_lambda_unsupported : Error<
   "use of '__super' inside a lambda is unsupported">;
 
 def warn_pragma_unused_undeclared_var : Warning<
   "undeclared variable %0 used as an argument for '#pragma unused'">,
   InGroup<IgnoredPragmas>;
 def warn_atl_uuid_deprecated : Warning<
   "specifying 'uuid' as an ATL attribute is deprecated; use __declspec instead">,

clang/include/clang/Basic/LangOptions.h

public:
 
   // TODO: merge FEnvAccessModeKind and FPContractModeKind
   enum FEnvAccessModeKind {
     FEA_Off,
 
     FEA_On
   };
 
   // Values of the following enumerations correspond to metadata arguments
   // specified for constrained floating-point intrinsics:
   // http://llvm.org/docs/LangRef.html#constrained-floating-point-intrinsics.
 
   /// Possible rounding modes.
   enum FPRoundingModeKind {
     /// Rounding to nearest, corresponds to "round.tonearest".
     FPR_ToNearest,
     /// Rounding toward -Inf, corresponds to "round.downward".
     FPR_Downward,
     /// Rounding toward +Inf, corresponds to "round.upward".
     FPR_Upward,
     /// Rounding toward zero, corresponds to "round.towardzero".
     FPR_TowardZero,
     /// Is determined by runtime environment, corresponds to "round.dynamic".
     FPR_Dynamic
   };
 
   /// Possible floating point exception behavior.
   enum FPExceptionModeKind {
     /// Assume that floating-point exceptions are masked.
     FPE_Ignore,
     /// Transformations do not cause new exceptions but may hide some.
     FPE_MayTrap,
     /// Strictly preserve the floating-point exception semantics.
public:
   /// Name of the IR file that contains the result of the OpenMP target
   /// host code generation.
   std::string OMPHostIRFile;
 
   /// Indicates whether the front-end is explicitly told that the
   /// input is a header file (i.e. -x c-header).
   bool IsHeaderFile = false;
 
   LangOptions();
 
   // Define accessors/mutators for language options of enumeration type.
 #define LANGOPT(Name, Bits, Default, Description)
 #define ENUM_LANGOPT(Name, Type, Bits, Default, Description) \
   Type get##Name() const { return static_cast<Type>(Name); } \
   void set##Name(Type Value) { Name = static_cast<unsigned>(Value); }
 #include "clang/Basic/LangOptions.def"
 

 };
 
 /// Floating point control options
 class FPOptions {
 public:
   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<LangOptions::FPContractModeKind>(I & 3)),
         fenv_access(static_cast<LangOptions::FEnvAccessModeKind>((I >> 2) & 1)),
         rounding(static_cast<LangOptions::FPRoundingModeKind>((I >> 3) & 7)),
-        exceptions(static_cast<LangOptions::FPExceptionModeKind>((I >> 6) & 3))
+        exceptions(static_cast<LangOptions::FPExceptionModeKind>((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;
   }
 
   bool allowFPContractAcrossStatement() const {
     return fp_contract == LangOptions::FPC_Fast;
   }
 
public:
   bool allowFEnvAccess() const {
     return fenv_access == LangOptions::FEA_On;
   }
 
   void setAllowFEnvAccess() {
     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 {
     return static_cast<LangOptions::FPRoundingModeKind>(rounding);
   }
 
   void setRoundingMode(LangOptions::FPRoundingModeKind RM) {
     rounding = RM;
   }
 
   LangOptions::FPExceptionModeKind getExceptionMode() const {
     return static_cast<LangOptions::FPExceptionModeKind>(exceptions);
   }
 
   void setExceptionMode(LangOptions::FPExceptionModeKind EM) {
     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 ||
            allowFEnvAccess();
   }
 
   /// 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<LangOptions::FPContractModeKind>(I & 3));
+    fenv_access = (static_cast<LangOptions::FEnvAccessModeKind>((I >> 2) & 1));
+    rounding = (static_cast<LangOptions::FPRoundingModeKind>((I >> 3) & 7));
+    exceptions = (static_cast<LangOptions::FPExceptionModeKind>((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:
   /// Adjust BinaryOperatorBitfields::FPFeatures and
   /// CXXOperatorCallExprBitfields::FPFeatures to match the total bit-field size
   /// of these fields.
   unsigned fp_contract : 2;
   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.
 enum TranslationUnitKind {
   /// The translation unit is a complete translation unit.
   TU_Complete,
 
   /// The translation unit is a prefix to a translation unit, and is

clang/include/clang/Basic/LangOptions.def

 LANGOPT(ROPI                         , 1, 0, "Read-only position independence")
 LANGOPT(RWPI                         , 1, 0, "Read-write position independence")
 COMPATIBLE_LANGOPT(GNUInline         , 1, 0, "GNU inline semantics")
 COMPATIBLE_LANGOPT(NoInlineDefine    , 1, 0, "__NO_INLINE__ predefined macro")
 COMPATIBLE_LANGOPT(Deprecated        , 1, 0, "__DEPRECATED predefined macro")
 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")
 
 BENIGN_LANGOPT(AccessControl     , 1, 1, "C++ access control")
 LANGOPT(CharIsSigned      , 1, 1, "signed char")
 LANGOPT(WCharSize         , 4, 0, "width of wchar_t")
 LANGOPT(WCharIsSigned        , 1, 0, "signed or unsigned wchar_t")
 ENUM_LANGOPT(MSPointerToMemberRepresentationMethod, PragmaMSPointersToMembersKind, 2, PPTMK_BestCase, "member-pointer representation method")

clang/include/clang/Basic/PragmaKinds.h

enum PragmaMSCommentKind {
   PCK_User      // #pragma comment(user, ...)
 };
 
 enum PragmaMSStructKind {
   PMSST_OFF, // #pragms ms_struct off
   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

clang/include/clang/Basic/TokenKinds.def

 // handles them.
 PRAGMA_ANNOTATION(pragma_fp_contract)
 
 // Annotation for #pragma STDC FENV_ACCESS
 // The lexer produces these so that they only take effect when the parser
 // 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.
 PRAGMA_ANNOTATION(pragma_ms_pointers_to_members)
 
 // Annotation for #pragma vtordisp...
 // The lexer produces these so that they only take effect when the parser
 // handles them.

clang/include/clang/Parse/Parser.h

class Parser : public CodeCompletionHandler {
   std::unique_ptr<PragmaHandler> WeakHandler;
   std::unique_ptr<PragmaHandler> RedefineExtnameHandler;
   std::unique_ptr<PragmaHandler> FPContractHandler;
   std::unique_ptr<PragmaHandler> OpenCLExtensionHandler;
   std::unique_ptr<PragmaHandler> OpenMPHandler;
   std::unique_ptr<PragmaHandler> PCSectionHandler;
   std::unique_ptr<PragmaHandler> MSCommentHandler;
   std::unique_ptr<PragmaHandler> MSDetectMismatchHandler;
+  std::unique_ptr<PragmaHandler> FloatControlHandler;
   std::unique_ptr<PragmaHandler> MSPointersToMembers;
   std::unique_ptr<PragmaHandler> MSVtorDisp;
   std::unique_ptr<PragmaHandler> MSInitSeg;
   std::unique_ptr<PragmaHandler> MSDataSeg;
   std::unique_ptr<PragmaHandler> MSBSSSeg;
   std::unique_ptr<PragmaHandler> MSConstSeg;
   std::unique_ptr<PragmaHandler> MSCodeSeg;
   std::unique_ptr<PragmaHandler> MSSection;
private:
   /// Handle the annotation token produced for
   /// #pragma STDC FP_CONTRACT...
   void HandlePragmaFPContract();
 
   /// Handle the annotation token produced for
   /// #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();
 
   /// Handle the annotation token produced for
   /// #pragma OPENCL EXTENSION...
   void HandlePragmaOpenCLExtension();
 

clang/include/clang/Sema/Sema.h

public:
   };
   SmallVector<PackIncludeState, 8> PackIncludeStack;
   // Segment #pragmas.
   PragmaStack<StringLiteral *> DataSegStack;
   PragmaStack<StringLiteral *> BSSSegStack;
   PragmaStack<StringLiteral *> ConstSegStack;
   PragmaStack<StringLiteral *> CodeSegStack;
 
+    // This stacks the current state of Sema.FPFeatures
+  PragmaStack<unsigned> 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 {
   public:
     PragmaStackSentinelRAII(Sema &S, StringRef SlotLabel, bool ShouldAct);
     ~PragmaStackSentinelRAII();
 
   private:

 
   /// Called on #pragma clang __debug dump II
   void ActOnPragmaDump(Scope *S, SourceLocation Loc, IdentifierInfo *II);
 
   /// ActOnPragmaDetectMismatch - Call on well-formed \#pragma detect_mismatch
   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,
                          SourceLocation PragmaLoc);
 
   /// ActOnPragmaVisibility - Called on well formed \#pragma GCC visibility... .
   void ActOnPragmaVisibility(const IdentifierInfo* VisType,
                              SourceLocation PragmaLoc);

clang/include/clang/Serialization/ASTBitCodes.h

enum ASTRecordTypes {
 
       /// Record code for \#pragma pack options.
       PACK_PRAGMA_OPTIONS = 61,
 
       /// The stack of open #ifs/#ifdefs recorded in a preamble.
       PP_CONDITIONAL_STACK = 62,
 
       /// A table of skipped ranges within the preprocessing record.
-      PPD_SKIPPED_RANGES = 63
+      PPD_SKIPPED_RANGES = 63,
+
+      /// Record code for \#pragma float_control options.
+      FLOAT_CONTROL_PRAGMA_OPTIONS = 64
     };
 
     /// Record types used within a source manager block.
     enum SourceManagerRecordTypes {
       /// Describes a source location entry (SLocEntry) for a
       /// file.
       SM_SLOC_FILE_ENTRY = 1,
 

clang/include/clang/Serialization/ASTReader.h

private:
 
   /// The PragmaMSStructKind pragma ms_struct state if set, or -1.
   int PragmaMSStructState = -1;
 
   /// The PragmaMSPointersToMembersKind pragma pointers_to_members state.
   int PragmaMSPointersToMembersState = -1;
   SourceLocation PointersToMembersPragmaLocation;
 
+  /// The pragma float_control state.
+  Optional<unsigned> FpPragmaCurrentValue;
+  SourceLocation FpPragmaCurrentLocation;
+  struct FpPragmaStackEntry {
+    unsigned Value;
+    SourceLocation Location;
+    SourceLocation PushLocation;
+    StringRef SlotLabel;
+  };
+  llvm::SmallVector<FpPragmaStackEntry, 2> FpPragmaStack;
+  llvm::SmallVector<std::string, 2> FpPragmaStrings;
+
   /// The pragma pack state.
   Optional<unsigned> PragmaPackCurrentValue;
   SourceLocation PragmaPackCurrentLocation;
   struct PragmaPackStackEntry {
     unsigned Value;
     SourceLocation Location;
     SourceLocation PushLocation;
     StringRef SlotLabel;

clang/include/clang/Serialization/ASTWriter.h

private:
   void WriteOpenCLExtensionDecls(Sema &SemaRef);
   void WriteCUDAPragmas(Sema &SemaRef);
   void WriteObjCCategories();
   void WriteLateParsedTemplates(Sema &SemaRef);
   void WriteOptimizePragmaOptions(Sema &SemaRef);
   void WriteMSStructPragmaOptions(Sema &SemaRef);
   void WriteMSPointersToMembersPragmaOptions(Sema &SemaRef);
   void WritePackPragmaOptions(Sema &SemaRef);
+  void WriteFloatControlPragmaOptions(Sema &SemaRef);
   void WriteModuleFileExtension(Sema &SemaRef,
                                 ModuleFileExtensionWriter &Writer);
 
   unsigned DeclParmVarAbbrev = 0;
   unsigned DeclContextLexicalAbbrev = 0;
   unsigned DeclContextVisibleLookupAbbrev = 0;
   unsigned UpdateVisibleAbbrev = 0;
   unsigned DeclRecordAbbrev = 0;

clang/lib/CodeGen/CGExprScalar.cpp

static bool CanElideOverflowCheck(const ASTContext &Ctx, const BinOpInfo &Op) {
   unsigned PromotedSize = Ctx.getTypeSize(Op.E->getType());
   return (2 * Ctx.getTypeSize(LHSTy)) < PromotedSize ||
          (2 * Ctx.getTypeSize(RHSTy)) < PromotedSize;
 }
 
 /// 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.
 static Value *propagateFMFlags(Value *V, const BinOpInfo &Op) {
   if (auto *I = dyn_cast<llvm::Instruction>(V)) {
     llvm::FastMathFlags FMF = I->getFastMathFlags();
     updateFastMathFlags(FMF, Op.FPFeatures);
     I->setFastMathFlags(FMF);
public:
   }
 
   //===--------------------------------------------------------------------===//
   //                            Visitor Methods
   //===--------------------------------------------------------------------===//
 
   Value *Visit(Expr *E) {
     ApplyDebugLocation DL(CGF, E);
+    if (BinaryOperator * BinOp = dyn_cast<BinaryOperator>(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<CXXConstructorDecl>(CGF.CurFuncDecl) ||
+             isa<CXXDestructorDecl>(CGF.CurFuncDecl) ||
+          (NewExceptionBehavior == llvm::fp::ebIgnore &&
+           NewRoundingBehavior == llvm::fp::rmToNearest)) &&
+          "FPConstrained should be enabled on entire function");
+
+      return StmtVisitor<ScalarExprEmitter, Value*>::Visit(E);
+    }
     return StmtVisitor<ScalarExprEmitter, Value*>::Visit(E);
   }
 
   Value *VisitStmt(Stmt *S) {
     S->dump(CGF.getContext().getSourceManager());
     llvm_unreachable("Stmt can't have complex result type!");
   }
   Value *VisitExpr(Expr *S);
Value *VisitExplicitCastExpr(ExplicitCastExpr *E) {
     return VisitCastExpr(E);
   }
   Value *VisitCastExpr(CastExpr *E);
 
   Value *VisitCallExpr(const CallExpr *E) {
     if (E->getCallReturnType(CGF.getContext())->isReferenceType())
       return EmitLoadOfLValue(E);
 
     Value *V = CGF.EmitCallExpr(E).getScalarVal();
 
     EmitLValueAlignmentAssumption(E, V);
     return V;
   }
 
   Value *VisitStmtExpr(const StmtExpr *E);
 
   // Unary Operators.
case Qualifiers::OCL_Weak:
     LHS = EmitCheckedLValue(E->getLHS(), CodeGenFunction::TCK_Store);
     RHS = CGF.EmitARCStoreWeak(LHS.getAddress(CGF), RHS, Ignore);
     break;
 
   case Qualifiers::OCL_None:
     // __block variables need to have the rhs evaluated first, plus
     // this should improve codegen just a little.
     RHS = Visit(E->getRHS());
     LHS = EmitCheckedLValue(E->getLHS(), CodeGenFunction::TCK_Store);
 
     // Store the value into the LHS.  Bit-fields are handled specially
     // because the result is altered by the store, i.e., [C99 6.5.16p1]
     // 'An assignment expression has the value of the left operand after
     // the assignment...'.
     if (LHS.isBitField()) {
       CGF.EmitStoreThroughBitfieldLValue(RValue::get(RHS), LHS, &RHS);
     } else {

clang/lib/CodeGen/CodeGenFunction.h

inline llvm::Value *DominatingLLVMValue::restore(CodeGenFunction &CGF,
   if (!value.getInt()) return value.getPointer();
 
   // Otherwise, it should be an alloca instruction, as set up in save().
   auto alloca = cast<llvm::AllocaInst>(value.getPointer());
   return CGF.Builder.CreateAlignedLoad(alloca, alloca->getAlign());
 }
 
 }  // 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

clang/lib/CodeGen/CodeGenFunction.cpp

CodeGenFunction::~CodeGenFunction() {
   // time of the CodeGenModule, because we have to ensure the IR has not yet
   // been "emitted" to the outside, thus, modifications are still sensible.
   if (llvm::OpenMPIRBuilder *OMPBuilder = CGM.getOpenMPIRBuilder())
     OMPBuilder->finalize();
 }
 
 // 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) {
   case LangOptions::FPR_ToNearest:  return llvm::fp::rmToNearest;
   case LangOptions::FPR_Downward:   return llvm::fp::rmDownward;
   case LangOptions::FPR_Upward:     return llvm::fp::rmUpward;
   case LangOptions::FPR_TowardZero: return llvm::fp::rmTowardZero;
   case LangOptions::FPR_Dynamic:    return llvm::fp::rmDynamic;
   }
   llvm_unreachable("Unsupported FP RoundingMode");
 }
 
 // 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) {
   case LangOptions::FPE_Ignore:  return llvm::fp::ebIgnore;
   case LangOptions::FPE_MayTrap: return llvm::fp::ebMayTrap;
   case LangOptions::FPE_Strict:  return llvm::fp::ebStrict;
   }
   llvm_unreachable("Unsupported FP Exception Behavior");
 }
 
 void CodeGenFunction::SetFPModel() {
   auto fpRoundingMode = ToConstrainedRoundingMD(
                           getLangOpts().getFPRoundingMode());
   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);
   }
 }
 
 CharUnits CodeGenFunction::getNaturalPointeeTypeAlignment(QualType T,
                                                     LValueBaseInfo *BaseInfo,
                                                     TBAAAccessInfo *TBAAInfo) {
   return getNaturalTypeAlignment(T->getPointeeType(), BaseInfo, TBAAInfo,
                                  /* forPointeeType= */ true);
#undef SANITIZER
   // If we're in C++ mode and the function name is "main", it is guaranteed
   // to be norecurse by the standard (3.6.1.3 "The function main shall not be
   // used within a program").
   if (getLangOpts().CPlusPlus)
     if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D))
       if (FD->isMain())
         Fn->addFnAttr(llvm::Attribute::NoRecurse);
 
-  if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D))
+  if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(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.
   if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D))
     if ((FD->isMain() || FD->isMSVCRTEntryPoint()) &&
         CGM.getCodeGenOpts().StackAlignment)
       Fn->addFnAttr("stackrealign");
 

clang/lib/Frontend/CompilerInvocation.cpp

case Language::Unknown:
     break;
   }
   llvm_unreachable("unknown input language");
 }
 
 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;
   if (const Arg *A = Args.getLastArg(OPT_std_EQ)) {
     LangStd = LangStandard::getLangKind(A->getValue());
     if (LangStd == LangStandard::lang_unspecified) {
       Diags.Report(diag::err_drv_invalid_value)
         << A->getAsString(Args) << A->getValue();
if (Arg *InlineArg = Args.getLastArg(
           options::OPT_finline_functions, options::OPT_finline_hint_functions,
           options::OPT_fno_inline_functions, options::OPT_fno_inline))
     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_cl_finite_math_only) ||
       Args.hasArg(OPT_cl_fast_relaxed_math);
   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();
     if (Val == "fast")
       Opts.setDefaultFPContractMode(LangOptions::FPC_Fast);
     else if (Val == "on")
       Opts.setDefaultFPContractMode(LangOptions::FPC_On);
     else if (Val == "off")
if (DashX.getFormat() == InputKind::Precompiled ||
     LangOpts.PICLevel = getLastArgIntValue(Args, OPT_pic_level, 0, Diags);
     LangOpts.PIE = Args.hasArg(OPT_pic_is_pie);
     parseSanitizerKinds("-fsanitize=", Args.getAllArgValues(OPT_fsanitize_EQ),
                         Diags, LangOpts.Sanitize);
   } else {
     // 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()) {
       // Supress the darwin-specific 'stdlibcxx-not-found' diagnostic for
       // preprocessed input as we don't expect it to be used with -std=libc++
       // anyway.
       Res.getDiagnosticOpts().Warnings.push_back("no-stdlibcxx-not-found");
     }

clang/lib/Parse/ParseDeclCXX.cpp

void Parser::ParseCXXMemberSpecification(SourceLocation RecordLoc,
   //   within function bodies, default arguments, exception-specifications, and
   //   brace-or-equal-initializers for non-static data members (including such
   //   things in nested classes).
   if (TagDecl && NonNestedClass) {
     // We are not inside a nested class. This class and its nested classes
     // 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());
     ParseLexedMethodDeclarations(getCurrentClass());
 
     // We've finished with all pending member declarations.
     Actions.ActOnFinishCXXMemberDecls();
 

clang/lib/Parse/ParsePragma.cpp

PragmaDetectMismatchHandler(Sema &Actions)
     : PragmaHandler("detect_mismatch"), Actions(Actions) {}
   void HandlePragma(Preprocessor &PP, PragmaIntroducer Introducer,
                     Token &FirstToken) override;
 
 private:
   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,
                     Token &FirstToken) override;
 };
 
 struct PragmaMSVtorDisp : public PragmaHandler {
   explicit PragmaMSVtorDisp() : PragmaHandler("vtordisp") {}
void Parser::initializePragmaHandlers() {
   PP.AddPragmaHandler(OpenMPHandler.get());
 
   if (getLangOpts().MicrosoftExt ||
       getTargetInfo().getTriple().isOSBinFormatELF()) {
     MSCommentHandler = std::make_unique<PragmaCommentHandler>(Actions);
     PP.AddPragmaHandler(MSCommentHandler.get());
   }
 
+  FloatControlHandler =
+      std::make_unique<PragmaFloatControlHandler>(Actions);
+  PP.AddPragmaHandler(FloatControlHandler.get());
   if (getLangOpts().MicrosoftExt) {
     MSDetectMismatchHandler =
         std::make_unique<PragmaDetectMismatchHandler>(Actions);
     PP.AddPragmaHandler(MSDetectMismatchHandler.get());
     MSPointersToMembers = std::make_unique<PragmaMSPointersToMembers>();
     PP.AddPragmaHandler(MSPointersToMembers.get());
     MSVtorDisp = std::make_unique<PragmaMSVtorDisp>();
     PP.AddPragmaHandler(MSVtorDisp.get());
if (getLangOpts().MicrosoftExt ||
       getTargetInfo().getTriple().isOSBinFormatELF()) {
     PP.RemovePragmaHandler(MSCommentHandler.get());
     MSCommentHandler.reset();
   }
 
   PP.RemovePragmaHandler("clang", PCSectionHandler.get());
   PCSectionHandler.reset();
 
+  PP.RemovePragmaHandler(FloatControlHandler.get());
+  FloatControlHandler.reset();
   if (getLangOpts().MicrosoftExt) {
     PP.RemovePragmaHandler(MSDetectMismatchHandler.get());
     MSDetectMismatchHandler.reset();
     PP.RemovePragmaHandler(MSPointersToMembers.get());
     MSPointersToMembers.reset();
     PP.RemovePragmaHandler(MSVtorDisp.get());
     MSVtorDisp.reset();
     PP.RemovePragmaHandler(MSInitSeg.get());
case tok::OOS_DEFAULT:
     FPC = getLangOpts().getDefaultFPContractMode();
     break;
   }
 
   Actions.ActOnPragmaFPContract(FPC);
   ConsumeAnnotationToken();
 }
 
+void Parser::HandlePragmaFloatControl() {
+  assert(Tok.is(tok::annot_pragma_float_control));
+
+  uintptr_t Value = reinterpret_cast<uintptr_t>(Tok.getAnnotationValue());
+  Sema::PragmaMsStackAction Action =
+      static_cast<Sema::PragmaMsStackAction>((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 =
     static_cast<tok::OnOffSwitch>(
     reinterpret_cast<uintptr_t>(Tok.getAnnotationValue()));
 
   LangOptions::FEnvAccessModeKind FPC;
   switch (OOS) {
void PragmaMSPragma::HandlePragma(Preprocessor &PP,
   std::copy(TokenVector.begin(), TokenVector.end(), TokenArray.get());
   auto Value = new (PP.getPreprocessorAllocator())
       std::pair<std::unique_ptr<Token[]>, size_t>(std::move(TokenArray),
                                                   TokenVector.size());
   AnnotTok.setAnnotationValue(Value);
   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<PragmaFloatControlKind>(
+      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 (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") {
+        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") {
+          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();
+  if (Tok.isNot(tok::eod)) {
+    PP.Diag(Tok.getLocation(), diag::warn_pragma_extra_tokens_at_eol)
+      << "float_control";
+    return;
+  }
+
+  // Note: there is no accomodation for PP callback for this pragma.
+
+  // Enter the annotation.
+  auto TokenArray = std::make_unique<Token[]>(1);
+  TokenArray[0].startToken();
+  TokenArray[0].setKind(tok::annot_pragma_float_control);
+  TokenArray[0].setLocation(FloatControlLoc);
+  TokenArray[0].setAnnotationEndLoc(EndLoc);
+  TokenArray[0].setAnnotationValue(reinterpret_cast<void *>(
+      static_cast<uintptr_t>((Action << 16) | (Kind & 0xFFFF))));
+  PP.EnterTokenStream(std::move(TokenArray), 1,
+                      /*DisableMacroExpansion=*/false, /*IsReinject=*/false);
+}
+
 /// Handle the Microsoft \#pragma detect_mismatch extension.
 ///
 /// The syntax is:
 /// \code
 ///   #pragma detect_mismatch("name", "value")
 /// \endcode
 /// Where 'name' and 'value' are quoted strings.  The values are embedded in
 /// the object file and passed along to the linker.  If the linker detects a

clang/lib/Parse/ParseStmt.cpp

Retry:
 
   case tok::annot_pragma_redefine_extname:
     ProhibitAttributes(Attrs);
     HandlePragmaRedefineExtname();
     return StmtEmpty();
 
   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();
 
   case tok::annot_pragma_fenv_access:
     ProhibitAttributes(Attrs);
     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();
     return StmtEmpty();
 
   case tok::annot_pragma_captured:
     ProhibitAttributes(Attrs);
     return HandlePragmaCaptured();
case tok::annot_pragma_fp_contract:
       HandlePragmaFPContract();
       break;
     case tok::annot_pragma_fp:
       HandlePragmaFP();
       break;
     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;
     case tok::annot_pragma_ms_pragma:
       HandlePragmaMSPragma();
       break;
     case tok::annot_pragma_ms_vtordisp:
       HandlePragmaMSVtorDisp();

clang/lib/Parse/Parser.cpp

case tok::annot_pragma_redefine_extname:
     HandlePragmaRedefineExtname();
     return nullptr;
   case tok::annot_pragma_fp_contract:
     HandlePragmaFPContract();
     return nullptr;
   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;
   case tok::annot_pragma_opencl_extension:
     HandlePragmaOpenCLExtension();
     return nullptr;
   case tok::annot_pragma_openmp: {
     AccessSpecifier AS = AS_none;

clang/lib/Sema/Sema.cpp

: ExternalSource(nullptr), isMultiplexExternalSource(false),
       Context(ctxt), Consumer(consumer), Diags(PP.getDiagnostics()),
       SourceMgr(PP.getSourceManager()), CollectStats(false),
       CodeCompleter(CodeCompleter), CurContext(nullptr),
       OriginalLexicalContext(nullptr), MSStructPragmaOn(false),
       MSPointerToMemberRepresentationMethod(
           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),
       StdCoroutineTraitsCache(nullptr), CXXTypeInfoDecl(nullptr),
       MSVCGuidDecl(nullptr), NSNumberDecl(nullptr), NSValueDecl(nullptr),
       NSStringDecl(nullptr), StringWithUTF8StringMethod(nullptr),
       ValueWithBytesObjCTypeMethod(nullptr), NSArrayDecl(nullptr),
       ArrayWithObjectsMethod(nullptr), NSDictionaryDecl(nullptr),

clang/lib/Sema/SemaAttr.cpp

 void Sema::ActOnPragmaDetectMismatch(SourceLocation Loc, StringRef Name,
                                      StringRef Value) {
   auto *PDMD = PragmaDetectMismatchDecl::Create(
       Context, Context.getTranslationUnitDecl(), Loc, Name, Value);
   Context.getTranslationUnitDecl()->addDecl(PDMD);
   Consumer.HandleTopLevelDecl(DeclGroupRef(PDMD));
 }
 
+void Sema::ActOnPragmaFloatControl(SourceLocation Loc,
+                              PragmaMsStackAction Action,
+                              PragmaFloatControlKind Value) {
+  auto NewValue = FpPragmaStack.CurrentValue;
+  if ((Action == PSK_Push_Set || Action == PSK_Push || Action == PSK_Pop) &&
+      !CurContext->isTranslationUnit()) {
+    // Push and pop can only occur at file scope.
+    Diag(Loc, diag::err_pragma_fc_pp_scope);
+    return;
+  }
+  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:
+  case PFC_Pop:
+    if (Value == PFC_Pop && FpPragmaStack.Stack.empty())
+      Diag(Loc, diag::warn_pragma_pop_failed) <<
+           "float_control" << "stack empty";
+    FpPragmaStack.Act(Loc, Action, StringRef(), NewValue);
+    if (Value == PFC_Pop) {
+      NewValue = FpPragmaStack.CurrentValue;
+      FPFeatures.Restore(NewValue);
+    }
+    break;
+  }
+}
+
 void Sema::ActOnPragmaMSPointersToMembers(
     LangOptions::PragmaMSPointersToMembersKind RepresentationMethod,
     SourceLocation PragmaLoc) {
   MSPointerToMemberRepresentationMethod = RepresentationMethod;
   ImplicitMSInheritanceAttrLoc = PragmaLoc;
 }
 
 void Sema::ActOnPragmaMSVtorDisp(PragmaMsStackAction Action,

 
 void Sema::setExceptionMode(LangOptions::FPExceptionModeKind FPE) {
   FPFeatures.setExceptionMode(FPE);
 }
 
 void Sema::ActOnPragmaFEnvAccess(LangOptions::FEnvAccessModeKind FPC) {
   switch (FPC) {
   case LangOptions::FEA_On:
     FPFeatures.setAllowFEnvAccess();
     break;
   case LangOptions::FEA_Off:
     FPFeatures.setDisallowFEnvAccess();
     break;
   }
 }
 
 

clang/lib/Sema/SemaExpr.cpp

if (getLangOpts().CPlusPlus && !RHS.isInvalid()) {
       VK = RHS.get()->getValueKind();
       OK = RHS.get()->getObjectKind();
     }
     break;
   }
   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<FunctionDecl>(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
   // arm64).
   assert(isVector(RHS.get()->getType(), Context.HalfTy) ==
          isVector(LHS.get()->getType(), Context.HalfTy) &&
          "both sides are half vectors or neither sides are");
   ConvertHalfVec = needsConversionOfHalfVec(ConvertHalfVec, Context,
                                             LHS.get()->getType());

clang/lib/Sema/SemaStmt.cpp

if (E->isGLValue() && E->getType().isVolatileQualified()) {
     Diag(Loc, diag::warn_unused_volatile) << R1 << R2;
     return;
   }
 
   DiagRuntimeBehavior(Loc, nullptr, PDiag(DiagID) << R1 << R2);
 }
 
 void Sema::ActOnStartOfCompoundStmt(bool IsStmtExpr) {
+  if (getFPOptions().isFPConstrained()) {
+    // Mark the current function as usng floating point constrained intrinsics
+    if (FunctionDecl *F = dyn_cast<FunctionDecl>(CurContext)) {
+      F->setUsesFPIntrin(true);
+    }
+  }
+
   PushCompoundScope(IsStmtExpr);
 }
 
 void Sema::ActOnFinishOfCompoundStmt() {
   PopCompoundScope();
 }
 
 sema::CompoundScopeInfo &Sema::getCurCompoundScope() const {
   return getCurFunction()->CompoundScopes.back();
 }
 
 StmtResult Sema::ActOnCompoundStmt(SourceLocation L, SourceLocation R,
                                    ArrayRef<Stmt *> Elts, bool isStmtExpr) {
   const unsigned NumElts = Elts.size();
 
   // If we're in C89 mode, check that we don't have any decls after stmts.  If
   // so, emit an extension diagnostic.
   if (!getLangOpts().C99 && !getLangOpts().CPlusPlus) {
     // Note that __extension__ can be around a decl.
     unsigned i = 0;
     // Skip over all declarations.
     for (; i != NumElts && isa<DeclStmt>(Elts[i]); ++i)
       /*empty*/;
 
     // We found the end of the list or a statement.  Scan for another declstmt.

clang/lib/Serialization/ASTReader.cpp

case PACK_PRAGMA_OPTIONS: {
         Entry.Location = ReadSourceLocation(F, Record[Idx++]);
         Entry.PushLocation = ReadSourceLocation(F, Record[Idx++]);
         PragmaPackStrings.push_back(ReadString(Record, Idx));
         Entry.SlotLabel = PragmaPackStrings.back();
         PragmaPackStack.push_back(Entry);
       }
       break;
     }
+
+    case FLOAT_CONTROL_PRAGMA_OPTIONS: {
+      if (Record.size() < 3) {
+        Error("invalid pragma pack record");
+        return Failure;
+      }
+      FpPragmaCurrentValue = Record[0];
+      FpPragmaCurrentLocation = ReadSourceLocation(F, Record[1]);
+      unsigned NumStackEntries = Record[2];
+      unsigned Idx = 3;
+      // Reset the stack when importing a new module.
+      FpPragmaStack.clear();
+      for (unsigned I = 0; I < NumStackEntries; ++I) {
+        FpPragmaStackEntry Entry;
+        Entry.Value = Record[Idx++];
+        Entry.Location = ReadSourceLocation(F, Record[Idx++]);
+        Entry.PushLocation = ReadSourceLocation(F, Record[Idx++]);
+        FpPragmaStrings.push_back(ReadString(Record, Idx));
+        Entry.SlotLabel = FpPragmaStrings.back();
+        FpPragmaStack.push_back(Entry);
+      }
+      break;
+    }
     }
   }
 }
 
 void ASTReader::ReadModuleOffsetMap(ModuleFile &F) const {
   assert(!F.ModuleOffsetMap.empty() && "no module offset map to read");
 
   // Additional remapping information.
if (PragmaPackCurrentLocation.isInvalid()) {
       assert(*PragmaPackCurrentValue == SemaObj->PackStack.DefaultValue &&
              "Expected a default alignment value");
       // Keep the current values.
     } else {
       SemaObj->PackStack.CurrentValue = *PragmaPackCurrentValue;
       SemaObj->PackStack.CurrentPragmaLocation = PragmaPackCurrentLocation;
     }
   }
+  if (FpPragmaCurrentValue) {
+    // The bottom of the stack might have a default value. It must be adjusted
+    // to the current value to ensure that fp-pragma state is preserved after
+    // popping entries that were included/imported from a PCH/module.
+    bool DropFirst = false;
+    if (!FpPragmaStack.empty() &&
+        FpPragmaStack.front().Location.isInvalid()) {
+      assert(FpPragmaStack.front().Value == SemaObj->FpPragmaStack.DefaultValue
+             && "Expected a default pragma float_control value");
+      SemaObj->FpPragmaStack.Stack.emplace_back(
+          FpPragmaStack.front().SlotLabel, SemaObj->FpPragmaStack.CurrentValue,
+          SemaObj->FpPragmaStack.CurrentPragmaLocation,
+          FpPragmaStack.front().PushLocation);
+      DropFirst = true;
+    }
+    for (const auto &Entry :
+         llvm::makeArrayRef(FpPragmaStack).drop_front(DropFirst ? 1 : 0))
+      SemaObj->FpPragmaStack.Stack.emplace_back(Entry.SlotLabel, Entry.Value,
+                                            Entry.Location, Entry.PushLocation);
+    if (FpPragmaCurrentLocation.isInvalid()) {
+      assert(*FpPragmaCurrentValue == SemaObj->FpPragmaStack.DefaultValue &&
+             "Expected a default pragma float_control value");
+      // Keep the current values.
+    } else {
+      SemaObj->FpPragmaStack.CurrentValue = *FpPragmaCurrentValue;
+      SemaObj->FpPragmaStack.CurrentPragmaLocation = FpPragmaCurrentLocation;
+    }
+  }
 }
 
 IdentifierInfo *ASTReader::get(StringRef Name) {
   // Note that we are loading an identifier.
   Deserializing AnIdentifier(this);
 
   IdentifierLookupVisitor Visitor(Name, /*PriorGeneration=*/0,
                                   NumIdentifierLookups,

clang/lib/Serialization/ASTWriter.cpp

for (const auto &StackEntry : SemaRef.PackStack.Stack) {
     Record.push_back(StackEntry.Value);
     AddSourceLocation(StackEntry.PragmaLocation, Record);
     AddSourceLocation(StackEntry.PragmaPushLocation, Record);
     AddString(StackEntry.StackSlotLabel, Record);
   }
   Stream.EmitRecord(PACK_PRAGMA_OPTIONS, Record);
 }
 
+/// Write the state of 'pragma float_control' at the end of the module.
+void ASTWriter::WriteFloatControlPragmaOptions(Sema &SemaRef) {
+  // Don't serialize pragma pack state for modules, since it should only take
+  // effect on a per-submodule basis.
+  if (WritingModule)
+    return;
+
+  RecordData Record;
+  Record.push_back(SemaRef.FpPragmaStack.CurrentValue);
+  AddSourceLocation(SemaRef.FpPragmaStack.CurrentPragmaLocation, Record);
+  Record.push_back(SemaRef.FpPragmaStack.Stack.size());
+  for (const auto &StackEntry : SemaRef.FpPragmaStack.Stack) {
+    Record.push_back(StackEntry.Value);
+    AddSourceLocation(StackEntry.PragmaLocation, Record);
+    AddSourceLocation(StackEntry.PragmaPushLocation, Record);
+    AddString(StackEntry.StackSlotLabel, Record);
+  }
+  Stream.EmitRecord(FLOAT_CONTROL_PRAGMA_OPTIONS, Record);
+}
+
 void ASTWriter::WriteModuleFileExtension(Sema &SemaRef,
                                          ModuleFileExtensionWriter &Writer) {
   // Enter the extension block.
   Stream.EnterSubblock(EXTENSION_BLOCK_ID, 4);
 
   // Emit the metadata record abbreviation.
   auto Abv = std::make_shared<llvm::BitCodeAbbrev>();
   Abv->Add(llvm::BitCodeAbbrevOp(EXTENSION_METADATA));
ASTFileSignature ASTWriter::WriteASTCore(Sema &SemaRef, StringRef isysroot,
 
   WriteObjCCategories();
   if(!WritingModule) {
     WriteOptimizePragmaOptions(SemaRef);
     WriteMSStructPragmaOptions(SemaRef);
     WriteMSPointersToMembersPragmaOptions(SemaRef);
   }
   WritePackPragmaOptions(SemaRef);
+  WriteFloatControlPragmaOptions(SemaRef);
 
   // Some simple statistics
   RecordData::value_type Record[] = {
       NumStatements, NumMacros, NumLexicalDeclContexts, NumVisibleDeclContexts};
   Stream.EmitRecord(STATISTICS, Record);
   Stream.ExitBlock();
 
   // Write the module file extension blocks.

clang/test/CodeGen/constrained-math-builtins.c

 };
 
 #pragma STDC FP_CONTRACT ON
 void bar(float f) {
   f * f + f;
   (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
 };

clang/test/CodeGen/fast-math.c

-// 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) {
   // CHECK-LABEL: define {{.*}}void @foo()
 
   // CHECK: fadd fast
   f0 = f1 + f2;
 
   // CHECK: ret
 }

clang/test/CodeGen/fp-contract-on-pragma.cpp

 // RUN: %clang_cc1 -O3 -triple %itanium_abi_triple -emit-llvm -o - %s | FileCheck %s
 
 // 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;
 }
 
 // Is FP_CONTRACT state cleared on exiting compound statements?
 float fp_contract_2(float a, float b, float c) {
   // CHECK: _Z13fp_contract_2fff
   // CHECK: %[[M:.+]] = fmul float %a, %b

 template <typename T>
 T template_muladd(T a, T b, T c) {
 #pragma clang fp contract(on)
   return a * b + c;
 }
 
 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<float>(a, b, c);
 }
 
 template <typename T>
 class fp_contract_4 {
   float method(float a, float b, float c) {
 #pragma clang fp contract(on)
     return a * b + c;
   }
 };
 
 template class fp_contract_4<int>;
 // 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;
 }
 
 #pragma clang fp contract(off)
 float fp_contract_6(float a, float b, float c) {
   // CHECK: _Z13fp_contract_6fff
   // CHECK: %[[M:.+]] = fmul float %a, %b
   // CHECK-NEXT: fadd float %[[M]], %c
   return a * b + c;
 }
 
 // If the multiply has multiple uses, don't produce fmuladd.
 // This used to assert (PR25719):
 // https://llvm.org/bugs/show_bug.cgi?id=25719
 
 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;
 }

clang/test/CodeGen/fp-contract-pragma.cpp

 // RUN: %clang_cc1 -O3 -triple %itanium_abi_triple -emit-llvm -o - %s | FileCheck %s
 
 // 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;
 }
 
 // Is FP_CONTRACT state cleared on exiting compound statements?
 float fp_contract_2(float a, float b, float c) {
 // CHECK: _Z13fp_contract_2fff
 // CHECK: %[[M:.+]] = fmul float %a, %b

 template <typename T>
 T template_muladd(T a, T b, T c) {
   #pragma STDC FP_CONTRACT ON
   return a * b + c;
 }
 
 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<float>(a, b, c);
 }
 
 template<typename T> class fp_contract_4 {
   float method(float a, float b, float c) {
     #pragma STDC FP_CONTRACT ON
     return a * b + c;
   }
 };
 
 template class fp_contract_4<int>;
 // 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;
 }
 
 #pragma STDC FP_CONTRACT OFF
 float fp_contract_6(float a, float b, float c) {
 // CHECK: _Z13fp_contract_6fff
 // CHECK: %[[M:.+]] = fmul float %a, %b
 // CHECK-NEXT: fadd float %[[M]], %c
   return a * b + c;
 }
 
 // If the multiply has multiple uses, don't produce fmuladd.
 // This used to assert (PR25719):
 // https://llvm.org/bugs/show_bug.cgi?id=25719
 
 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;
 }

clang/test/CodeGen/fp-floatcontrol-class.cpp

+// RUN: %clang -c -Xclang -emit-llvm -o - %s | FileCheck %s
+// XFAIL:*
+// 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;

clang/test/CodeGen/fp-floatcontrol-pragma.cpp

+// 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;
+}

clang/test/CodeGen/fp-floatcontrol-stack.cpp

+// 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
+// XFAIL:*
+
+#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;

clang/test/CodeGen/fpconstrained.c

 // 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
 float f0, f1, f2;
 
 void foo() {
   // CHECK-LABEL: define {{.*}}void @foo()
 
   // MAYTRAP: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.tonearest", metadata !"fpexcept.maytrap")
   // EXCEPT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.tonearest", metadata !"fpexcept.strict")
   // FPMODELSTRICT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.dynamic", metadata !"fpexcept.strict")
   // STRICTEXCEPT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.dynamic", metadata !"fpexcept.strict")
   // 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
 }

clang/test/CodeGen/fpconstrained.cpp

 // 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
 float f0, f1, f2;
 
   template <class>
   class aaaa {
    public:
float f0, f1, f2;
   } catch (...) {
   // MAYTRAP: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.tonearest", metadata !"fpexcept.maytrap")
   // EXCEPT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.tonearest", metadata !"fpexcept.strict")
   // FPMODELSTRICT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.dynamic", metadata !"fpexcept.strict")
   // STRICTEXCEPT: llvm.experimental.constrained.fadd.f32(float %{{.*}}, float %{{.*}}, metadata !"round.dynamic", metadata !"fpexcept.strict")
   // 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
   }
   }
   
   class d {
    public:

clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl

 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
 
 typedef unsigned int uint;
 typedef half __attribute__((ext_vector_type(2))) half2;
 typedef short __attribute__((ext_vector_type(2))) short2;
 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)
 
 // CHECK: call i32 @llvm.amdgcn.udot2(<2 x i16> %v2usA, <2 x i16> %v2usB, i32 %uiC, i1 false)
 // CHECK: call i32 @llvm.amdgcn.udot2(<2 x i16> %v2usA, <2 x i16> %v2usB, i32 %uiC, i1 true)
 
 // CHECK: call i32 @llvm.amdgcn.sdot4(i32 %siA, i32 %siB, i32 %siC, i1 false)

clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
 
 #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);
 }

clang/test/CodeGenOpenCL/builtins-amdgcn-interp.cl

void test_interp_f16(global half* out, float i, float j, int m0)
   half p2_0 = __builtin_amdgcn_interp_p2_f16(p1_0, j, 2, 3, false, m0);
   float p1_1 = __builtin_amdgcn_interp_p1_f16(i, 2, 3, true, m0);
   half p2_1 = __builtin_amdgcn_interp_p2_f16(p1_1, j, 2, 3, true, m0);
   *out = p2_0 + p2_1;
 }
 
 // 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);
   *out = __builtin_amdgcn_interp_p2(p1, j, 1, 4, m0);
 }
 
 // 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);
 }

clang/test/CodeGenOpenCL/builtins-amdgcn-mfma.cl

 typedef short  v2s   __attribute__((ext_vector_type(2)));
 typedef short  v4s   __attribute__((ext_vector_type(4)));
 typedef short  v16s  __attribute__((ext_vector_type(16)));
 typedef short  v32s  __attribute__((ext_vector_type(32)));
 typedef double v4d   __attribute__((ext_vector_type(4)));
 
 
 // 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);
 }
 
 // CHECK-LABEL: @test_mfma_i32_32x32x4i8
 // CHECK: call <32 x i32> @llvm.amdgcn.mfma.i32.32x32x4i8(i32 %a, i32 %b, <32 x i32> %c, i32 0, i32 0, i32 0)
 void test_mfma_i32_32x32x4i8(global v32i* out, int a, int b, v32i c)

 // CHECK-LABEL: @test_mfma_i32_16x16x16i8
 // CHECK: call <4 x i32> @llvm.amdgcn.mfma.i32.16x16x16i8(i32 %a, i32 %b, <4 x i32> %c, i32 0, i32 0, i32 0)
 void test_mfma_i32_16x16x16i8(global v4i* out, int a, int b, v4i c)
 {
   *out = __builtin_amdgcn_mfma_i32_16x16x16i8(a, b, c, 0, 0, 0);
 }
 
 // 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);
 }
 

clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 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);
 }
 
 // CHECK-LABEL: @test_frexp_exp_f16
 // CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16
 void test_frexp_exp_f16(global short* out, half a)
 {
   *out = __builtin_amdgcn_frexp_exph(a);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_class_f16
 // CHECK: call i1 @llvm.amdgcn.class.f16
 void test_class_f16(global half* out, half a, int b)

 // CHECK-LABEL: @test_update_dpp
 // CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
 void test_update_dpp(global int* out, int arg1, int arg2)
 {
   *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
 }
 
 // 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);
 }

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

 // CHECK: store i8 [[FLAGEXT]]
 void test_div_scale_f32_generic_ptr(global float* out, global int* flagout, float a, float b, global bool* flag_arg)
 {
   generic bool* flag = flag_arg;
   *out = __builtin_amdgcn_div_scalef(a, b, true, flag);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_frexp_exp_f32
 // CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32
 void test_frexp_exp_f32(global int* out, float a)
 {
   *out = __builtin_amdgcn_frexp_expf(a);
 }
 
 // CHECK-LABEL: @test_frexp_exp_f64
 // CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64
 void test_frexp_exp_f64(global int* out, double a)
 {
   *out = __builtin_amdgcn_frexp_exp(a);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_lerp
 // CHECK: call i32 @llvm.amdgcn.lerp
 void test_lerp(global int* out, int a, int b, int c)

 // CHECK: call void @llvm.amdgcn.s.decperflevel(i32 15)
 void test_s_decperflevel()
 {
   __builtin_amdgcn_s_decperflevel(1);
   __builtin_amdgcn_s_decperflevel(15);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_read_exec(
 // CHECK: call i64 @llvm.read_register.i64(metadata ![[$EXEC:[0-9]+]]) #[[$READ_EXEC_ATTRS:[0-9]+]]
 void test_read_exec(global ulong* out) {
   *out = __builtin_amdgcn_read_exec();
void test_get_local_id(int d, global int *out)
   case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
   case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
   case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
   default: *out = 0;
   }
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_s_getpc
 // CHECK: call i64 @llvm.amdgcn.s.getpc()
 void test_s_getpc(global ulong* out)

 
 // CHECK-LABEL: @test_sbfe(
 // CHECK: tail call i32 @llvm.amdgcn.sbfe.i32(i32 %src0, i32 %src1, i32 %src2)
 kernel void test_sbfe(global uint* out, uint src0, uint src1, uint src2) {
   *out = __builtin_amdgcn_sbfe(src0, src1, src2);
 }
 
 // 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);
 }
 
 // CHECK-LABEL: @test_cvt_pknorm_i16(
 // CHECK: tail call <2 x i16> @llvm.amdgcn.cvt.pknorm.i16(float %src0, float %src1)
 kernel void test_cvt_pknorm_i16(global short2* out, float src0, float src1) {
   *out = __builtin_amdgcn_cvt_pknorm_i16(src0, src1);

clang/test/CodeGenOpenCL/builtins-f16.cl

 // RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-darwin-apple %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 // CHECK-LABEL: define void @test_half_builtins
 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);
 }

clang/test/CodeGenOpenCL/builtins-r600.cl

 // REQUIRES: amdgpu-registered-target
 // 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);
 }
 
 #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);
 }
 #endif
 
 // CHECK-LABEL: @test_implicitarg_ptr
 // CHECK: call i8 addrspace(7)* @llvm.r600.implicitarg.ptr()

clang/test/CodeGenOpenCL/relaxed-fpmath.cl

 // 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-mad-enable -o - | FileCheck %s -check-prefix=MAD
 // RUN: %clang_cc1 %s -emit-llvm -cl-no-signed-zeros -o - | FileCheck %s -check-prefix=NOSIGNED
 
 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
 
 // NORMAL: "less-precise-fpmad"="false"
 // NORMAL: "no-infs-fp-math"="false"
 // NORMAL: "no-nans-fp-math"="false"
 // NORMAL: "no-signed-zeros-fp-math"="false"

clang/test/CodeGenOpenCL/single-precision-constant.cl

 // 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.;
 }

clang/test/PCH/pragma-floatcontrol.c

+// Test this without pch.
+// RUN: %clang_cc1 %s -include %s -verify -fsyntax-only -DSET
+// RUN: %clang_cc1 %s -include %s -verify -fsyntax-only -DPUSH
+// RUN: %clang_cc1 %s -include %s -verify -fsyntax-only -DPUSH_POP
+
+// Test with pch.
+// RUN: %clang_cc1 %s -DSET -emit-pch -o %t
+// RUN: %clang_cc1 %s -DSET -include-pch %t -emit-llvm -o - | FileCheck --check-prefix=CHECK-EBSTRICT %s
+// RUN: %clang_cc1 %s -DPUSH -emit-pch -o %t
+// RUN: %clang_cc1 %s -DPUSH -verify -include-pch %t
+// RUN: %clang_cc1 %s -DPUSH_POP -emit-pch -o %t
+// RUN: %clang_cc1 %s -DPUSH_POP -verify -include-pch %t
+
+#ifndef HEADER
+#define HEADER
+
+#ifdef SET
+#pragma float_control(except, on)
+#endif
+
+#ifdef PUSH
+#pragma float_control(precise, on)
+#pragma float_control (push)
+#pragma float_control(precise, off)
+#endif
+
+#ifdef PUSH_POP
+#pragma float_control (precise, on, push)
+#pragma float_control (push)
+#pragma float_control (pop)
+#endif
+#else
+
+#ifdef SET
+float fun(float a, float b) {
+// CHECK-LABEL: define float @fun{{.*}}
+//CHECK-EBSTRICT: llvm.experimental.constrained.fmul{{.*}}tonearest{{.*}}strict
+//CHECK-EBSTRICT: llvm.experimental.constrained.fadd{{.*}}tonearest{{.*}}strict
+  return a*b + 2;
+}
+#pragma float_control(pop) // expected-warning {{#pragma float_control(pop, ...) failed: stack empty}}
+#pragma float_control(pop) // expected-warning {{#pragma float_control(pop, ...) failed: stack empty}}
+#endif
+
+#ifdef PUSH
+#pragma float_control(pop)
+#pragma float_control(pop) // expected-warning {{#pragma float_control(pop, ...) failed: stack empty}}
+#endif
+
+#ifdef PUSH_POP
+#pragma float_control(pop)
+#pragma float_control(pop) // expected-warning {{#pragma float_control(pop, ...) failed: stack empty}}
+#endif
+
+#endif //ifndef HEADER

clang/test/Parser/fp-floatcontrol-syntax.cpp

+// RUN: %clang_cc1 -fsyntax-only -verify -DCHECK_ERROR %s
+
+float function_scope(float a) {
+#  pragma float_control(precise, on) junk // expected-warning {{extra tokens at end of '#pragma float_control' - ignored}}
+  return a;
+}
+
+#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}}
+#pragma float_control(except,on,push) // expected-error {{can only appear at file scope}}
+#pragma float_control(except,on,push,junk) // expected-error {{float_control is malformed}}
+  return;
+}
+#endif

llvm/include/llvm/IR/IRBuilder.h

public:
   }
 
   /// Get the floating point math metadata being used.
   MDNode *getDefaultFPMathTag() const { return DefaultFPMathTag; }
 
   /// 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(); }
 
   /// Set the floating point math metadata to be used.
   void setDefaultFPMathTag(MDNode *FPMathTag) { DefaultFPMathTag = FPMathTag; }
 
   /// Set the fast-math flags to be used with generated fp-math operators
   void setFastMathFlags(FastMathFlags NewFMF) { FMF = NewFMF; }
public:
   };
 
   // RAII object that stores the current fast math settings and restores
   // them when the object is destroyed.
   class FastMathFlagGuard {
     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;
 
     ~FastMathFlagGuard() {
       Builder.FMF = FMF;
       Builder.DefaultFPMathTag = FPMathTag;
+      Builder.IsFPConstrained = IsFPConstrained;
+      Builder.DefaultConstrainedExcept = DefaultConstrainedExcept;
+      Builder.DefaultConstrainedRounding = DefaultConstrainedRounding;
     }
   };
 
   //===--------------------------------------------------------------------===//
   // Miscellaneous creation methods.
   //===--------------------------------------------------------------------===//
 
   /// Make a new global variable with initializer type i8*