LangRef: Add "dynamic" option to "denormal-fp-math"

[arsenm]

Summary

This is stricter than the default "ieee", and should probably be the
default. This patch leaves the default alone. I can change this in a
future patch.

There are non-reversible transforms I would like to perform which are
legal under IEEE denormal handling, but illegal with flushing zero
behavior. Namely, conversions between llvm.is.fpclass and fcmp with
zeroes.

Under "ieee" handling, it is legal to translate between
llvm.is.fpclass(x, fcZero) and fcmp x, 0.

Under "preserve-sign" handling, it is legal to translate between
llvm.is.fpclass(x, fcSubnormal|fcZero) and fcmp x, 0.

I would like to compile and distribute some math library functions in
a mode where it's callable from code with and without denormals
enabled, which requires not changing the compares with denormals or
zeroes.

If an IEEE function transforms an llvm.is.fpclass call into an fcmp 0,
it is no longer possible to call the function from code with denormals
enabled, or write an optimization to move the function into a denormal
flushing mode. For the original function, if x was a denormal, the
class would evaluate to false. If the function compiled with denormal
handling was converted to or called from a preserve-sign function, the
fcmp now evaluates to true.

This could also be of use for strictfp handling, where code may be
changing the denormal mode.

Alternative name could be "unknown".

Replaces the old AMDGPU custom inlining logic with more conservative
logic which tries to permit inlining for callees with dynamic handling
and avoids inlining other mismatched modes.

[tra]

My $.02, mostly on the style and the mechanics of applying the new attribute. FP semantics aspects are above my pay grade.

clang/lib/CodeGen/CGCall.cpp

2059–2061

IIUIC, this changes denorm mode attributes on the functions with dynamic denorm mode that we link in.
Will that be a problem if the same function, when linked into different modules, would end up with different attributes? E.g. if a function is externally visible and is intended to be common'ed across multiple modules. Should dynamic denorm mode be restricted to the functions private to the module only? We do typically internalize linked bitcode for CUDA, but I don't think it's something we can always implicitly assume.

clang/test/CodeGenCUDA/Inputs/ocml-sample.cl

11

Cosmetic nit: order functions as f16/f32/f64?

clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu

15

Do we want to verify that the compiled samples have the correct function attributes?

20

Would it be useful to check the attributes the functions have w/o linking the sample bitcode? In these tests one can infer expected attributes from the flags and comments, so I'm fine not having explicit checks for that.

77

Nit: CHECK-LABEL ?

90–91

I assume these refer to linked in functions, not their calls. It may be useful to include match define/call to make it obvious.

94

I'm not sure whether it does what it's intended to.
AFAICT, at this point we will be past the call sites, so if it's intended to check the call sites in kernel_*, it will likely always succeed, even if we do litter call sites with unwanted attributes.

It's also possible that I have a wrong idea about what the expected IR looks like. If you could post it for reference, that would be helpful.

llvm/test/CodeGen/Generic/denormal-fp-math-cl-opt.ll

3

Edit: Check that the command line flag annotates the IR with the appropriate attributes.

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2059–2061

The whole point of -mlink-builtin-bitcode is to apply the attributes for the current compilation to what's linked in. The linked functions are always internalized. The only case where we might not want to internalize is for weak symbols (but it looks like we do internalize those today, but this is something I've thought about changing). I'll add a test with a weak library function

In the weak case the right thing to do is probably to not change from dynamic, simply because this linking process is outside of the user's control.

clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu

15

Maybe, but that's already tested separately. This test is a bit complex as it is (and could maybe use a few more combinations)

77

error: found 'CHECK-LABEL:' with variable definition or use

94

I can drop this, I later added the -implicit-check-not=denormal-fp-math to all the FileChecks

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2059–2061

It turns out we apply attributes prior to internalization. As a separate patch, we can either:

1. Skip functions that start as interposable, which has an observable change in the IR as it is
2. Move the link and internalize before setting attributes. This would be unobservable but would catch it if the internalization behavior ever changed

[arsenm]

Fix losing target-cpu

[arsenm]

Fix dropping target-cpu. Also skip interposable functions if we aren't internalizing (this seems to be a theoretical concern, since PropagateAttrs and Internalize are set as a pair)

[tra]

LGTM for the parts I've commented on.

clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu

77

Interesting.
In that case label and attribute checks could be separated into something like this:

CHECK-LABEL: name
CHECK-SAME: [[attribute]]

Up to you.

llvm/test/CodeGen/Generic/denormal-fp-math-cl-opt.ll

3

^^ The comment still needs to be edited.

[kpn]

We use "dynamic" for the constrained intrinsics. I'd stay consistent with our terminology and stick with "dynamic" here.

I like the amount of testing. You may have gotten every single combination of cases, but I didn't go far enough to check.

llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll

78

Are we changing the behavior in a way that may cause regressions? It looks like we've changed behavior in the absence of "dynamic".

[arsenm]

llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll

78

This case is broken to begin with, calling ieee from daz code. This makes the inlining more conservative / noticeable for debugging

[kpn]

llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll

78

Can I talk you into mentioning this in your commit message?

[arsenm]

Address comments

[jcranmer-intel]

Looking at the attribute logic here, there is conceptual room for both a dynamic and an unknown mode (i.e., you get a top and a bottom value), but I don't think there is value in distinguishing between them, so I'm fine with keeping just a dynamic.

I didn't bother to look at the clang changes, and my only comments are some minor ones around documentation:

llvm/docs/LangRef.rst

2166–2167

I feel like the description of this mode should mention that whether or not denormals are flushed is derived from the dynamic state of the FP environment.

2174–2178

This isn't your fault, but I noticed when reading the LangRef online that this paragraph has a slightly-different indentation that causes most of this attribute's documentation to gain an extra level of indentation.

[arsenm]

Documentation fixes

[andrew.w.kaylor]

In general, it seems like the denormal mode should be considered part of the floating point environment (though as far as I know the C standard, at least, doesn't document it as such). If it were considered part of the floating point environment, the LLVM rules would tell us we could assume the default setting, which I'd assume to be IEEE, and it would only be legal to change this mode in strict mode. However, for your use case preserving the behavior of fpclass seems like what users would want, even in fast-math modes. In this sense, this is a lot like the problem we have with preserving isnan() behavior when fast-math is enabled. Our rules allow it, but it's not what most people would want.

I think the new denormal mode is a good addition.

Do you need to do something with the inliner to handle the case where functions with different denormal modes are inlined into one another? We don't seem to handle that case correctly now (https://godbolt.org/z/PEsWaMEq6), but with the dynamic mode we could handle it without blocking inlining completely.

[arsenm]

In D142907#4119339, @andrew.w.kaylor wrote:

In general, it seems like the denormal mode should be considered part of the floating point environment (though as far as I know the C standard, at least, doesn't document it as such).

There’s no standardization of denormal flushing. OpenCL defines a flag for it but doesn’t really specify what it really means.

If it were considered part of the floating point environment, the LLVM rules would tell us we could assume the default setting, which I'd assume to be IEEE, and it would only be legal to change this mode in strict mode.

It is. The attribute is informative of the default mode. If we really wanted, a similar attribute could declare the assumed rounding mode if we really wanted for the function. It doesn’t imply you can change it dynamically, just that it should be in that mode before the function executes.

Do you need to do something with the inliner to handle the case where functions with different denormal modes are inlined into one another? We don't seem to handle that case correctly now (https://godbolt.org/z/PEsWaMEq6), but with the dynamic mode we could handle it without blocking inlining completely.

This patch fixed some bugs in the inlining to make it more conservative in undefined looking cases. The inlining of dynamic functions fully works

[jcranmer-intel]

Not entirely sure where the best place to effect this (I think somewhere in the clang driver code?), but on further reflection, it feels like strict fp-model in clang should set the denormal mode to dynamic.

[arsenm]

In D142907#4132318, @jcranmer-intel wrote:

Not entirely sure where the best place to effect this (I think somewhere in the clang driver code?), but on further reflection, it feels like strict fp-model in clang should set the denormal mode to dynamic.

I was thinking of changing the default in general to dynamic. I was going to at least change the strictfp default in a follow up

[kpn]

What's the plan for tying this to strictfp? Because I don't it should be tied to cases where we use the constrained intrinsics but the exceptions are ignored and the default rounding is in stated. Those instructions are supposed to behave the same as the non-constrained instructions. So keying off the presence of the strictfp attribute on the function definition, or the (equivalent) presence of constrained intrinsics, would be too simple.

I don't see an obvious connection between denormals and exception behavior, and the rounding mode has the same problem. It would be surprising if changing the rounding mode changed denormal handling or optimization even when the new rounding mode would have identical results. It would also be surprising if changing the rounding mode back to the default round-to-nearest changed how denormals are handled.

Would we get different denormal behavior with a clang flag vs using a #pragma at the top of a source file? That seems surprising as well.

In the abstract I can see lumping in denormal handing with the rest of the FP environment handling. But in the LLVM context I don't see how we can tie use of the constrained intrinsics to denormals.

[arsenm]

In D142907#4132543, @kpn wrote:

What's the plan for tying this to strictfp? Because I don't it should be tied to cases where we use the constrained intrinsics but the exceptions are ignored and the default rounding is in stated. Those instructions are supposed to behave the same as the non-constrained instructions. So keying off the presence of the strictfp attribute on the function definition, or the (equivalent) presence of constrained intrinsics, would be too simple.

The denormal mode is exactly parallel to the rounding mode, we just don't have a mirrored field in the constrained intrinsic metadata operands. If we defaulted to using the dynamic mode if you were to use strictfp, everything would be OK. You just couldn't optimize based on knowledge of the denormal mode. I don't really think it's worth putting in the same optimization effort as the rounding mode.

[jcranmer-intel]

In D142907#4132430, @arsenm wrote:

I was thinking of changing the default in general to dynamic. I was going to at least change the strictfp default in a follow up

I had the same thought too, but I reflected a little further that the default fp model implying that the environment being in the default state means we can assume the FTZ/DAZ are also in a default (IEEE) state.

In D142907#4132543, @kpn wrote:

What's the plan for tying this to strictfp? Because I don't it should be tied to cases where we use the constrained intrinsics but the exceptions are ignored and the default rounding is in stated. Those instructions are supposed to behave the same as the non-constrained instructions. So keying off the presence of the strictfp attribute on the function definition, or the (equivalent) presence of constrained intrinsics, would be too simple.

The way I see it, strictfp is an assertion that every FP instruction has a dependency on the FP environment, which is largely orthogonal to the denormal-mode attribute asserting that the FTZ/DAZ bits in the FP environment have a particular value. The constrained intrinsics also have the ability to assert some properties of the FP environment (specifically, rounding mode and exception behavior) on individual instructions. By not adding any metadata to constrained intrinsics at the same time, we don't get the ability to set the denormal-mode on a per-instruction basis-but I don't think there's much value to be gained by doing so (giving that we already have it at a per-function level).

Would we get different denormal behavior with a clang flag vs using a #pragma at the top of a source file? That seems surprising as well.

One of the consequences of having so many different ways of controlling compiler FP environment assumptions is that there's a crazy amount of interactions to consider. But I think there is ultimately a workable solution for the clang frontend to generate interactions that make sense.

[pengfei]

llvm/docs/LangRef.rst

2180

1. Does it mean users must specify dynamic when they change FTZ/DAZ in a function?
2. If 1) is true, is there a way to partially set on functions in its call stack, e.g.,

main
  f0
  f1
    f10
      setFtzDaz(true);
  f2
  f3

Ideally, users may want to tell compiler main, f1, f10 is dynamic, while f0 is ieee and f2, f3 is positive-zero, rather than dynamic for all.

3. If 2) is true, it looks silly to do it one by one manually. Should compiler help to deduce this information itself?

[arsenm]

llvm/docs/LangRef.rst

2180

This is really a question of how strictfp should interact with the default mode and shouldn't be a different policy from how strictfp functions treat the rounding mode. Arbitrary strictfp functions don't make assumptions based on the default rounding mode, assuming it's not the default. In that sense, denormal-fp-mode doesn't really matter for strictfp functions. They just can't make use of it to optimize. If we had a denormal annotation like the rounding mode, we could make use of it in the same way

This isn't an area for the backend to deduce, semantic meaning needs to be specific and explicit. I have no interest in making changing the denormal mode simple or easy. Turning on flushing isn't really semantically desirable and is basically obsolete on modern hardware.

[arsenm]

In D142907#4132836, @jcranmer-intel wrote:

In D142907#4132430, @arsenm wrote:

I was thinking of changing the default in general to dynamic. I was going to at least change the strictfp default in a follow up

I had the same thought too, but I reflected a little further that the default fp model implying that the environment being in the default state means we can assume the FTZ/DAZ are also in a default (IEEE) state.

In D142907#4132543, @kpn wrote:

What's the plan for tying this to strictfp? Because I don't it should be tied to cases where we use the constrained intrinsics but the exceptions are ignored and the default rounding is in stated. Those instructions are supposed to behave the same as the non-constrained instructions. So keying off the presence of the strictfp attribute on the function definition, or the (equivalent) presence of constrained intrinsics, would be too simple.

The way I see it, strictfp is an assertion that every FP instruction has a dependency on the FP environment, which is largely orthogonal to the denormal-mode attribute asserting that the FTZ/DAZ bits in the FP environment have a particular value. The constrained intrinsics also have the ability to assert some properties of the FP environment (specifically, rounding mode and exception behavior) on individual instructions. By not adding any metadata to constrained intrinsics at the same time, we don't get the ability to set the denormal-mode on a per-instruction basis-but I don't think there's much value to be gained by doing so (giving that we already have it at a per-function level)

I think this is not really useful for strictfp. We could make use of this if we were to inline non-strictfp functions into strictfp functions if we had a denormal mode annotation on the constrained intrinsic, which we don't. Without that, strictfp functions can't make assumptions based on the default rounding mode so this doesn't inform anything useful in the presence of a changeable mode

[arsenm]

ping

[jcranmer-intel]

I'm studiously ignoring the Clang and LLVM codegen changes here, but otherwise, I think the direction of this change is generally good.

llvm/lib/Analysis/ConstantFolding.cpp

1377–1379

You should change the doxygen documentation to indicate that this method returns nullptr if the denormal mode is dynamic.

Ditto for ConstantFoldFPInstOperands.

[arsenm]

Update doxygen comment

[arsenm]

ping

[arsenm]

ping

[arsenm]

ping

[arsenm]

ping

[arsenm]

ping

[efriedma]

I'm having trouble understanding the changes on the clang side.

If I'm following correctly; the "denormal-fp-math" setting is a promise from the user to the compiler: if the setting is not "dynamic", the user promises that the definition will only execute in the specified denormal mode. This is similar to, for example, the rounding mode pragmas: the user promises a specific rounding mode unless they explicitly request dynamic rounding.

Given that, I don't follow the whole "merging" thing... we should just be setting whatever mode is active. The attribute setting should not depend on whether the function is interposable. If you have a ODR function, all definitions must have a mode compatible with whatever mode will be used at runtime. If you have a non-ODR weak function, optimizations shouldn't propagate that mode from the callee to the caller.

[arsenm]

In D142907#4288247, @efriedma wrote:

Given that, I don't follow the whole "merging" thing... we should just be setting whatever mode is active. The attribute setting should not depend on whether the function is interposable. If you have a ODR function, all definitions must have a mode compatible with whatever mode will be used at runtime. If you have a non-ODR weak function, optimizations shouldn't propagate that mode from the callee to the caller.

The point is to have code that works with either mode. The intent is to avoid duplicating code for an extremely marginal difference. Having to duplicate every library function for denormal flushing and ieee handling defeats the purpose. The specialization will be realized after linking / inlining

[efriedma]

If you have a library function that's built with "denormal-fp-math"="dynamic,dynamic", you can link it into code built in any mode, and LTO should be able to propagate that mode from the caller to the callee. That doesn't require clang to do anything special; you can just specify -fdenormal-fp-math=dynamic while building the library, and the user specifies -fdenormal-fp-math=ieee while building their code.

I guess you're worried specifically about ODR inline functions, defined in headers? The user specifies a specific mode because they know their code honors it... but the user might not be aware of the effect on functions defined in library headers. Other libraries in the same binary might use the same header, but specify a different mode. So if the user specifies a denormal mode, we should ignore it for ODR inline functions, because they didn't actually mean to apply the denormal mode to those definitions?

I'm not sure about applying those semantics automatically; I don't think there's any precedent in clang for anything like this. The closest thing I can think of is -fvisibility-inlines-hidden. I'd prefer to RFC it separately from the rest of the patch, and loop in clang frontend owners, since the precedent we set here will apply to other sorts of attributes.

[arsenm]

In D142907#4288555, @efriedma wrote:

If you have a library function that's built with "denormal-fp-math"="dynamic,dynamic", you can link it into code built in any mode, and LTO should be able to propagate that mode from the caller to the callee. That doesn't require clang to do anything special; you can just specify -fdenormal-fp-math=dynamic while building the library, and the user specifies -fdenormal-fp-math=ieee while building their code.

That's essentially what this does. I think the part you are missing is the existing special treatment of the builtin device library functions. The default set of attributes for the current translation unit is forcibly set on functions in bitcode libraries linked in and internalized with -mlink-builtin-bitcode. We need to logically merge with the current translation unit's mode, or else we're potentially breaking the linked in function. The main reason I'm doing this in the first place is to move towards a model with less special treatment of these libraries.

I guess you're worried specifically about ODR inline functions, defined in headers? The user specifies a specific mode because they know their code honors it... but the user might not be aware of the effect on functions defined in library headers. Other libraries in the same binary might use the same header, but specify a different mode. So if the user specifies a denormal mode, we should ignore it for ODR inline functions, because they didn't actually mean to apply the denormal mode to those definitions?

No, the clang changes are to handle the headerless bitcode-only device libraries only. The user is supposed to be unaware the builtin libraries exist. They're an implementation detail managed by the clang driver (-mlink-builtin-bitcode is a cc1 only flag) and have a special contract with the compiler.

I'm not sure about applying those semantics automatically; I don't think there's any precedent in clang for anything like this. The closest thing I can think of is -fvisibility-inlines-hidden. I'd prefer to RFC it separately from the rest of the patch, and loop in clang frontend owners, since the precedent we set here will apply to other sorts of attributes.

This isn't new, isn't end user facing, and isn't general purpose. This is the minimum required update to the existing -mlink-builtin-bitcode handling.

[efriedma]

Oh, sorry, I missed that the new code specifically runs on functions imported using -mlink-builtin-bitcode. I somehow thought it was running on all functions.

LGTM

[arsenm]

bc37be1855773c1dcf8c6bf577a096a81fd58652

clang/lib/CodeGen/CGCall.cpp

static bool HasStrictReturn(const CodeGenModule &Module, QualType RetTy,
   // We don't want to be too aggressive with the return checking, unless
   // it's explicit in the code opts or we're using an appropriate sanitizer.
   // Try to respect what the programmer intended.
   return Module.getCodeGenOpts().StrictReturn ||
          !Module.MayDropFunctionReturn(Module.getContext(), RetTy) ||
          Module.getLangOpts().Sanitize.has(SanitizerKind::Return);
 }
 
-void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
-                                                 bool HasOptnone,
-                                                 bool AttrOnCallSite,
-                                               llvm::AttrBuilder &FuncAttrs) {
+/// Add denormal-fp-math and denormal-fp-math-f32 as appropriate for the
+/// requested denormal behavior, accounting for the overriding behavior of the
+/// -f32 case.
+static void addDenormalModeAttrs(llvm::DenormalMode FPDenormalMode,
+                                 llvm::DenormalMode FP32DenormalMode,
+                                 llvm::AttrBuilder &FuncAttrs) {
+  if (FPDenormalMode != llvm::DenormalMode::getDefault() &&
+      FPDenormalMode.isValid()) // xxx redundant check
+    FuncAttrs.addAttribute("denormal-fp-math", FPDenormalMode.str());
+
+  if (FP32DenormalMode != FPDenormalMode && FP32DenormalMode.isValid())
+    FuncAttrs.addAttribute("denormal-fp-math-f32", FP32DenormalMode.str());
+}
+
+/// Add default attributes to a function, which have merge semantics under
+/// -mlink-builtin-bitcode and should not simply overwrite any existing
+/// attributes in the linked library.
+static void
+addMergableDefaultFunctionAttributes(const CodeGenOptions &CodeGenOpts,
+                                     llvm::AttrBuilder &FuncAttrs) {
+  addDenormalModeAttrs(CodeGenOpts.FPDenormalMode, CodeGenOpts.FP32DenormalMode,
+                       FuncAttrs);
+}
+
+void CodeGenModule::getTrivialDefaultFunctionAttributes(
+    StringRef Name, bool HasOptnone, bool AttrOnCallSite,
   // OptimizeNoneAttr takes precedence over -Os or -Oz. No warning needed.
   if (!HasOptnone) {
     if (CodeGenOpts.OptimizeSize)
       FuncAttrs.addAttribute(llvm::Attribute::OptimizeForSize);
     if (CodeGenOpts.OptimizeSize == 2)
       FuncAttrs.addAttribute(llvm::Attribute::MinSize);
   }
 
if (AttrOnCallSite) {
     }
 
     if (CodeGenOpts.LessPreciseFPMAD)
       FuncAttrs.addAttribute("less-precise-fpmad", "true");
 
     if (CodeGenOpts.NullPointerIsValid)
       FuncAttrs.addAttribute(llvm::Attribute::NullPointerIsValid);
 
-    if (CodeGenOpts.FPDenormalMode != llvm::DenormalMode::getIEEE())
-      FuncAttrs.addAttribute("denormal-fp-math",
-                             CodeGenOpts.FPDenormalMode.str());
-    if (CodeGenOpts.FP32DenormalMode != CodeGenOpts.FPDenormalMode) {
-      FuncAttrs.addAttribute(
-          "denormal-fp-math-f32",
-          CodeGenOpts.FP32DenormalMode.str());
-    }
-
     if (LangOpts.getDefaultExceptionMode() == LangOptions::FPE_Ignore)
       FuncAttrs.addAttribute("no-trapping-math", "true");
 
     // TODO: Are these all needed?
     // unsafe/inf/nan/nsz are handled by instruction-level FastMathFlags.
     if (LangOpts.NoHonorInfs)
       FuncAttrs.addAttribute("no-infs-fp-math", "true");
     if (LangOpts.NoHonorNaNs)
void CodeGenModule::getTrivialDefaultFunctionAttributes(
 
   for (StringRef Attr : CodeGenOpts.DefaultFunctionAttrs) {
     StringRef Var, Value;
     std::tie(Var, Value) = Attr.split('=');
     FuncAttrs.addAttribute(Var, Value);
   }
 }
 
+void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
+                                                 bool HasOptnone,
+                                                 bool AttrOnCallSite,
+                                                 llvm::AttrBuilder &FuncAttrs) {
+  getTrivialDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite,
+                                      FuncAttrs);
+  if (!AttrOnCallSite) {
+    // If we're just getting the default, get the default values for mergeable
+    // attributes.
+    addMergableDefaultFunctionAttributes(CodeGenOpts, FuncAttrs);
+  }
+}
+
 void CodeGenModule::addDefaultFunctionDefinitionAttributes(llvm::Function &F) {
   llvm::AttrBuilder FuncAttrs(F.getContext());
   getDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
                                /* AttrOnCallSite = */ false, FuncAttrs);
   // TODO: call GetCPUAndFeaturesAttributes?
   F.addFnAttrs(FuncAttrs);
 }
 
+/// Apply default attributes to \p F, accounting for merge semantics of
+/// attributes that should not overwrite existing attributes.
+void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
+    llvm::Function &F) {
+  llvm::AttrBuilder FuncAttrs(F.getContext());
+  getTrivialDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
+                                      /*AttrOnCallSite=*/false, FuncAttrs);
+
+  llvm::AttributeMask AttrsToRemove;
+
+  llvm::DenormalMode DenormModeToMerge = F.getDenormalModeRaw();
+  llvm::DenormalMode DenormModeToMergeF32 = F.getDenormalModeF32Raw();
+  llvm::DenormalMode Merged =
+      CodeGenOpts.FPDenormalMode.mergeCalleeMode(DenormModeToMerge);
+  llvm::DenormalMode MergedF32 = CodeGenOpts.FP32DenormalMode;
+
+  if (DenormModeToMergeF32.isValid()) {
+    MergedF32 =
+        CodeGenOpts.FP32DenormalMode.mergeCalleeMode(DenormModeToMergeF32);
+  }
+
+  if (Merged == llvm::DenormalMode::getDefault()) {
+    AttrsToRemove.addAttribute("denormal-fp-math");
+  } else if (Merged != DenormModeToMerge) {
+    // Overwrite existing attribute
+    FuncAttrs.addAttribute("denormal-fp-math",
+                           CodeGenOpts.FPDenormalMode.str());
+  }
+
+  if (MergedF32 == llvm::DenormalMode::getDefault()) {
+    AttrsToRemove.addAttribute("denormal-fp-math-f32");
+  } else if (MergedF32 != DenormModeToMergeF32) {
+    // Overwrite existing attribute
+    FuncAttrs.addAttribute("denormal-fp-math-f32",
+                           CodeGenOpts.FP32DenormalMode.str());
+  }
+
+  F.removeFnAttrs(AttrsToRemove);
+  addDenormalModeAttrs(Merged, MergedF32, FuncAttrs);
+  F.addFnAttrs(FuncAttrs);
+}
+
 void CodeGenModule::addDefaultFunctionDefinitionAttributes(
-                                                   llvm::AttrBuilder &attrs) {
   getDefaultFunctionAttributes(/*function name*/ "", /*optnone*/ false,
                                /*for call*/ false, attrs);
   GetCPUAndFeaturesAttributes(GlobalDecl(), attrs);
 }
 
 static void addNoBuiltinAttributes(llvm::AttrBuilder &FuncAttrs,
                                    const LangOptions &LangOpts,
                                    const NoBuiltinAttr *NBA = nullptr) {

clang/lib/CodeGen/CodeGenAction.cpp

public:
     bool LinkInModules() {
       for (auto &LM : LinkModules) {
         if (LM.PropagateAttrs)
           for (Function &F : *LM.Module) {
             // Skip intrinsics. Keep consistent with how intrinsics are created
             // in LLVM IR.
             if (F.isIntrinsic())
               continue;
-            Gen->CGM().addDefaultFunctionDefinitionAttributes(F);
+            Gen->CGM().mergeDefaultFunctionDefinitionAttributes(F);
           }
 
         CurLinkModule = LM.Module.get();
 
         bool Err;
         if (LM.Internalize) {
           Err = Linker::linkModules(
               *getModule(), std::move(LM.Module), LM.LinkFlags,

clang/lib/CodeGen/CodeGenModule.h

public:
   /// will propagate unsafe-fp-math=false up to every transitive caller of a
   /// function in the bitcode library!
   ///
   /// With the exception of fast-math attrs, this will only make the attributes
   /// on the function more conservative.  But it's unsafe to call this on a
   /// function which relies on particular fast-math attributes for correctness.
   /// It's up to you to ensure that this is safe.
   void addDefaultFunctionDefinitionAttributes(llvm::Function &F);
+  void mergeDefaultFunctionDefinitionAttributes(llvm::Function &F);
 
   /// Like the overload taking a `Function &`, but intended specifically
   /// for frontends that want to build on Clang's target-configuration logic.
   void addDefaultFunctionDefinitionAttributes(llvm::AttrBuilder &attrs);
 
   StringRef getMangledName(GlobalDecl GD);
   StringRef getBlockMangledName(GlobalDecl GD, const BlockDecl *BD);
   const GlobalDecl getMangledNameDecl(StringRef);
private:
   /// definitions whose linkage can change, e.g. implicit function instantions
   /// which may later be explicitly instantiated.
   bool MayBeEmittedEagerly(const ValueDecl *D);
 
   /// Check whether we can use a "simpler", more core exceptions personality
   /// function.
   void SimplifyPersonality();
 
+  /// Helper function for getDefaultFunctionAttributes. Builds a set of function
+  /// attributes which can be simply added to a function.
+  void getTrivialDefaultFunctionAttributes(StringRef Name, bool HasOptnone,
+                                           bool AttrOnCallSite,
+                                           llvm::AttrBuilder &FuncAttrs);
+
   /// Helper function for ConstructAttributeList and
   /// addDefaultFunctionDefinitionAttributes.  Builds a set of function
   /// attributes to add to a function with the given properties.
   void getDefaultFunctionAttributes(StringRef Name, bool HasOptnone,
                                     bool AttrOnCallSite,
                                     llvm::AttrBuilder &FuncAttrs);
 
   llvm::Metadata *CreateMetadataIdentifierImpl(QualType T, MetadataTypeMap &Map,
                                                StringRef Suffix);
 };
 
 }  // end namespace CodeGen
 }  // end namespace clang
 
 #endif // LLVM_CLANG_LIB_CODEGEN_CODEGENMODULE_H

clang/test/CodeGen/denormalfpmode-f32.c

 // RUN: %clang_cc1 -S %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-NONE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-NONE
+// RUN: %clang_cc1 -S -fdenormal-fp-math=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-NONE
 
 // RUN: %clang_cc1 -S -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-IEEE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-IEEE
+// RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-DYNAMIC
+
 
 // RUN: %clang_cc1 -S -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PS
 // RUN: %clang_cc1 -S -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PS
 // RUN: %clang_cc1 -S -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-NONE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-PS
+// RUN: %clang_cc1 -S -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-DYNAMIC
+
 
 // RUN: %clang_cc1 -S -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PZ
+// RUN: %clang_cc1 -S -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-DYNAMIC
 // RUN: %clang_cc1 -S -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PZ
+// RUN: %clang_cc1 -S -fdenormal-fp-math=dynamic -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-PZ
 // RUN: %clang_cc1 -S -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-PZ
 // RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-NONE
+// RUN: %clang_cc1 -S -fdenormal-fp-math=dynamic -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-NONE
+
 
 // CHECK-LABEL: main
 
 // CHECK-ATTR: attributes #0 =
 // CHECK-NONE-NOT:"denormal-fp-math"
 // CHECK-IEEE: "denormal-fp-math"="ieee,ieee"
 // CHECK-PS: "denormal-fp-math"="preserve-sign,preserve-sign"
 // CHECK-PZ: "denormal-fp-math"="positive-zero,positive-zero"
+// CHECK-DYNAMIC: "denormal-fp-math"="dynamic,dynamic"
+
 // CHECK-F32-NONE-NOT:"denormal-fp-math-f32"
 // CHECK-F32-IEEE: "denormal-fp-math-f32"="ieee,ieee"
 // CHECK-F32-PS: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
 // CHECK-F32-PZ: "denormal-fp-math-f32"="positive-zero,positive-zero"
 
+
+// CHECK-F32-DYNAMIC: "denormal-fp-math-f32"="dynamic,dynamic"
+
 int main(void) {
   return 0;
 }

clang/test/CodeGen/denormalfpmode.c

 // RUN: %clang_cc1 -S -fdenormal-fp-math=ieee %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-IEEE
 // RUN: %clang_cc1 -S -fdenormal-fp-math=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-PS
 // RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-PZ
+// RUN: %clang_cc1 -S -fdenormal-fp-math=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-DYNAMIC
 
 // CHECK-LABEL: main
 
 // The ieee,ieee is the default, so omit the attribute
 // CHECK-IEEE-NOT:"denormal-fp-math"
 // CHECK-PS: attributes #0 = {{.*}}"denormal-fp-math"="preserve-sign,preserve-sign"{{.*}}
 // CHECK-PZ: attributes #0 = {{.*}}"denormal-fp-math"="positive-zero,positive-zero"{{.*}}
+// CHECK-DYNAMIC: attributes #0 = {{.*}}"denormal-fp-math"="dynamic,dynamic"{{.*}}
 
 int main(void) {
   return 0;
 }

clang/test/CodeGenCUDA/Inputs/ocml-sample.cl

+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+float do_f32_stuff(float a, float b, float c) {
+  return __builtin_fmaf(a, b, c) + 4.0f;
+}
+
+double do_f64_stuff(double a, double b, double c) {
+  return __builtin_fma(a, b, c) + 4.0;
+}
+
+half do_f16_stuff(half a, half b, half c) {
+  return __builtin_fmaf16(a, b, c) + 4.0h;
+}

clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu

+// Verify the behavior of the denormal-fp-mode attributes in the way that
+// rocm-device-libs should be built with. The bitcode should be compiled with
+// denormal-fp-math-f32=dynamic, and should be replaced with the denormal mode
+// of the final TU.
+
+// Build the fake device library in the way rocm-device-libs should be built.
+//
+// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -fdenormal-fp-math-f32=dynamic \
+// RUN:   -mcode-object-version=none -emit-llvm-bc \
+// RUN:   %S/Inputs/ocml-sample.cl -o %t.dynamic.f32.bc
+//
+// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -fdenormal-fp-math=dynamic \
+// RUN:   -mcode-object-version=none -emit-llvm-bc \
+// RUN:   %S/Inputs/ocml-sample.cl -o %t.dynamic.full.bc
+
+
+
+// Check the default behavior with no denormal-fp-math arguments.
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
+// RUN:   -mlink-builtin-bitcode %t.dynamic.f32.bc \
+// RUN:   -emit-llvm %s -o - | FileCheck -implicit-check-not=denormal-fp-math %s --check-prefixes=CHECK,DEFAULT
+
+
+// Check an explicit full ieee request
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
+// RUN:    -fdenormal-fp-math=ieee \
+// RUN:   -mlink-builtin-bitcode %t.dynamic.f32.bc \
+// RUN:   -emit-llvm %s -o - | FileCheck -implicit-check-not=denormal-fp-math %s --check-prefixes=CHECK,DEFAULT
+
+
+// Check explicit f32-only flushing request
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN:   -fcuda-is-device -fdenormal-fp-math-f32=preserve-sign \
+// RUN:   -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,IEEEF64-PSZF32
+
+
+// Check explicit flush all request. Only the f32 component of the library is
+// dynamic, so the linked functions should use IEEE as the base mode and the new
+// functions preserve-sign.
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN:   -fcuda-is-device -fdenormal-fp-math=preserve-sign \
+// RUN:   -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,PSZ
+
+
+// Check explicit f32-only, ieee-other flushing request
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN:   -fcuda-is-device -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=preserve-sign \
+// RUN:   -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,IEEEF64-PSZF32
+
+
+// Check inverse of normal usage. Requesting IEEE f32, with flushed f16/f64
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN:   -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
+// RUN:   -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,IEEEF32-PSZF64-DYNF32
+
+
+// Check backwards from the normal usage where both library components can be
+// overridden.
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN:   -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
+// RUN:   -mlink-builtin-bitcode %t.dynamic.full.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,IEEEF32-PSZF64-DYNFULL
+
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+
+extern "C" {
+__device__ float do_f32_stuff(float a, float b, float c);
+__device__ double do_f64_stuff(double a, double b, double c);
+
+
+// CHECK: kernel_f32({{.*}}) #[[$KERNELATTR:[0-9]+]]
+__global__ void kernel_f32(float* out, float* a, float* b, float* c) {
+  int id = 0;
+  out[id] = do_f32_stuff(a[id], b[id], c[id]);
+}
+
+// CHECK: kernel_f64({{.*}}) #[[$KERNELATTR]]
+__global__ void kernel_f64(double* out, double* a, double* b, double* c) {
+  int id = 0;
+  out[id] = do_f64_stuff(a[id], b[id], c[id]);
+}
+}
+
+// CHECK: do_f32_stuff({{.*}}) #[[$FUNCATTR:[0-9]+]]
+// CHECK: do_f64_stuff({{.*}}) #[[$FUNCATTR]]
+
+// We should not be littering call sites with the attribute
+// CHECK-NOT: denormal-fp-math
+
+
+// Everything should use the default ieee with no explicit attribute
+// DEFAULT-NOT: "denormal-fp-math"
+// DEFAULT-NOT: "denormal-fp-math-f32"
+
+
+// PSZ: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" {{.*}} }
+
+// FIXME: Should check-not "denormal-fp-math" within the line
+
+// PSZ: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} }
+
+// FIXME: Should check-not "denormal-fp-math" within the line
+// IEEEF64-PSZF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} }
+// IEEEF64-PSZF32: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} }
+
+
+// IEEEF32-PSZF64-DYNF32-NOT: denormal-fp-math
+// IEEEF32-PSZF64-DYNF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"{{.*}} }
+// IEEEF32-PSZF64-DYNF32-NOT: denormal-fp-math
+
+// IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}}  "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} }
+// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}}  "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} }

clang/test/Driver/denormal-fp-math.c

 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee -v 2>&1 | FileCheck -check-prefix=CHECK-IEEE %s
 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=preserve-sign -v 2>&1 | FileCheck -check-prefix=CHECK-PS %s
 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=positive-zero -v 2>&1 | FileCheck -check-prefix=CHECK-PZ %s
+
+// RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=dynamic -v 2>&1 | FileCheck -check-prefix=CHECK-DYNAMIC %s
+
 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee -fno-fast-math -v 2>&1 | FileCheck -check-prefix=CHECK-NO-UNSAFE %s
 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee -fno-unsafe-math-optimizations -v 2>&1 | FileCheck -check-prefix=CHECK-NO-UNSAFE %s
 // RUN: not %clang -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=foo -v 2>&1 | FileCheck -check-prefix=CHECK-INVALID0 %s
 // RUN: not %clang -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee,foo -v 2>&1 | FileCheck -check-prefix=CHECK-INVALID1 %s
 // RUN: not %clang -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=foo,ieee -v 2>&1 | FileCheck -check-prefix=CHECK-INVALID2 %s
 // RUN: not %clang -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=foo,foo -v 2>&1 | FileCheck -check-prefix=CHECK-INVALID3 %s
 
 // IEEE is the implied default, and the flag is not passed.
 // CHECK-IEEE-NOT: -fdenormal-fp-math=
 // CHECK-PS: "-fdenormal-fp-math=preserve-sign,preserve-sign"
 // CHECK-PZ: "-fdenormal-fp-math=positive-zero,positive-zero"
+// CHECK-DYNAMIC: "-fdenormal-fp-math=dynamic,dynamic"
+
+
 // CHECK-NO-UNSAFE-NOT: "-fdenormal-fp-math=ieee"
 // CHECK-INVALID0: error: invalid value 'foo' in '-fdenormal-fp-math=foo'
 // CHECK-INVALID1: error: invalid value 'ieee,foo' in '-fdenormal-fp-math=ieee,foo'
 // CHECK-INVALID2: error: invalid value 'foo,ieee' in '-fdenormal-fp-math=foo,ieee'
 // CHECK-INVALID3: error: invalid value 'foo,foo' in '-fdenormal-fp-math=foo,foo'

llvm/docs/LangRef.rst

``strictfp``
     mode or that might alter the state of floating-point status flags that
     might otherwise be set or cleared by calling this function. LLVM will
     not introduce any new floating-point instructions that may trap.
 
 ``"denormal-fp-math"``
     This indicates the denormal (subnormal) handling that may be
     assumed for the default floating-point environment. This is a
     comma separated pair. The elements may be one of ``"ieee"``,
-    ``"preserve-sign"``, or ``"positive-zero"``. The first entry
-    indicates the flushing mode for the result of floating point
-    operations. The second indicates the handling of denormal inputs
+    ``"preserve-sign"``, ``"positive-zero"``, or ``"dynamic"``. The
+    first entry indicates the flushing mode for the result of floating
+    point operations. The second indicates the handling of denormal inputs
     to floating point instructions. For compatibility with older
     bitcode, if the second value is omitted, both input and output
     modes will assume the same mode.
 
-    If this is attribute is not specified, the default is
-    ``"ieee,ieee"``.
+    If this is attribute is not specified, the default is ``"ieee,ieee"``.
 
     If the output mode is ``"preserve-sign"``, or ``"positive-zero"``,
     denormal outputs may be flushed to zero by standard floating-point
     operations. It is not mandated that flushing to zero occurs, but if
     a denormal output is flushed to zero, it must respect the sign
-    mode. Not all targets support all modes. While this indicates the
-    expected floating point mode the function will be executed with,
-    this does not make any attempt to ensure the mode is
-    consistent. User or platform code is expected to set the floating
-    point mode appropriately before function entry.
+    mode. Not all targets support all modes.
+
+    If the mode is ``"dynamic"``, transformations which depend on the
+    behavior of denormal values should not be performed.
+
+    While this indicates the expected floating point mode the function will be executed
+    with, this does not make any attempt to ensure the mode is consistent. User or
+    platform code is expected to set the floating point mode appropriately before function
+    entry.
 
    If the input mode is ``"preserve-sign"``, or ``"positive-zero"``, a
    floating-point operation must treat any input denormal value as
    zero. In some situations, if an instruction does not respect this
    mode, the input may need to be converted to 0 as if by
    ``@llvm.canonicalize`` during lowering for correctness.
 
 ``"denormal-fp-math-f32"``
     Same as ``"denormal-fp-math"``, but only controls the behavior of
     the 32-bit float type (or vectors of 32-bit floats). If both are
     are present, this overrides ``"denormal-fp-math"``. Not all targets
     support separately setting the denormal mode per type, and no
     attempt is made to diagnose unsupported uses. Currently this
     attribute is respected by the AMDGPU and NVPTX backends.
 
 ``"thunk"``

llvm/include/llvm/ADT/FloatingPointMode.h

enum DenormalModeKind : int8_t {
 
     /// IEEE-754 denormal numbers preserved.
     IEEE,
 
     /// The sign of a flushed-to-zero number is preserved in the sign of 0
     PreserveSign,
 
     /// Denormals are flushed to positive zero.
-    PositiveZero
+    PositiveZero,
+
+    /// Denormals have unknown treatment.
+    Dynamic
   };
 
   /// Denormal flushing mode for floating point instruction results in the
   /// default floating point environment.
   DenormalModeKind Output = DenormalModeKind::Invalid;
 
   /// Denormal treatment kind for floating point instruction inputs in the
   /// default floating-point environment. If this is not DenormalModeKind::IEEE,
   /// floating-point instructions implicitly treat the input value as 0.
   DenormalModeKind Input = DenormalModeKind::Invalid;
 
   constexpr DenormalMode() = default;
   constexpr DenormalMode(DenormalModeKind Out, DenormalModeKind In) :
     Output(Out), Input(In) {}
 
 
   static constexpr DenormalMode getInvalid() {
     return DenormalMode(DenormalModeKind::Invalid, DenormalModeKind::Invalid);
   }
 
+  /// Return the assumed default mode for a function without denormal-fp-math.
+  static constexpr DenormalMode getDefault() {
+    return getIEEE();
+  }
+
   static constexpr DenormalMode getIEEE() {
     return DenormalMode(DenormalModeKind::IEEE, DenormalModeKind::IEEE);
   }
 
   static constexpr DenormalMode getPreserveSign() {
     return DenormalMode(DenormalModeKind::PreserveSign,
                         DenormalModeKind::PreserveSign);
   }
 
   static constexpr DenormalMode getPositiveZero() {
     return DenormalMode(DenormalModeKind::PositiveZero,
                         DenormalModeKind::PositiveZero);
   }
 
+  static constexpr DenormalMode getDynamic() {
+    return DenormalMode(DenormalModeKind::Dynamic, DenormalModeKind::Dynamic);
+  }
+
   bool operator==(DenormalMode Other) const {
     return Output == Other.Output && Input == Other.Input;
   }
 
   bool operator!=(DenormalMode Other) const {
     return !(*this == Other);
   }
 
   bool isSimple() const {
     return Input == Output;
   }
 
   bool isValid() const {
     return Output != DenormalModeKind::Invalid &&
            Input != DenormalModeKind::Invalid;
   }
 
+  /// Get the effective denormal mode if the mode if this caller calls into a
+  /// function with \p Callee. This promotes dynamic modes to the mode of the
+  /// caller.
+  DenormalMode mergeCalleeMode(DenormalMode Callee) const {
+    DenormalMode MergedMode = Callee;
+    if (Callee.Input == DenormalMode::Dynamic)
+      MergedMode.Input = Input;
+    if (Callee.Output == DenormalMode::Dynamic)
+      MergedMode.Output = Output;
+    return MergedMode;
+  }
+
   inline void print(raw_ostream &OS) const;
 
   inline std::string str() const {
     std::string storage;
     raw_string_ostream OS(storage);
     print(OS);
     return OS.str();
   }
 };
 
 inline raw_ostream& operator<<(raw_ostream &OS, DenormalMode Mode) {
   Mode.print(OS);
   return OS;
 }
 
 /// Parse the expected names from the denormal-fp-math attribute.
 inline DenormalMode::DenormalModeKind
 parseDenormalFPAttributeComponent(StringRef Str) {
   // Assume ieee on unspecified attribute.
   return StringSwitch<DenormalMode::DenormalModeKind>(Str)
       .Cases("", "ieee", DenormalMode::IEEE)
       .Case("preserve-sign", DenormalMode::PreserveSign)
       .Case("positive-zero", DenormalMode::PositiveZero)
+      .Case("dynamic", DenormalMode::Dynamic)
       .Default(DenormalMode::Invalid);
 }
 
 /// Return the name used for the denormal handling mode used by the the
 /// expected names from the denormal-fp-math attribute.
 inline StringRef denormalModeKindName(DenormalMode::DenormalModeKind Mode) {
   switch (Mode) {
   case DenormalMode::IEEE:
     return "ieee";
   case DenormalMode::PreserveSign:
     return "preserve-sign";
   case DenormalMode::PositiveZero:
     return "positive-zero";
+  case DenormalMode::Dynamic:
+    return "dynamic";
   default:
     return "";
   }
 }
 
 /// Returns the denormal mode to use for inputs and outputs.
 inline DenormalMode parseDenormalFPAttribute(StringRef Str) {
   StringRef OutputStr, InputStr;

llvm/include/llvm/IR/Attributes.td

 class IntAttr<string S, list<AttrProperty> P> : Attr<S, P>;
 
 /// Type attribute.
 class TypeAttr<string S, list<AttrProperty> P> : Attr<S, P>;
 
 /// StringBool attribute.
 class StrBoolAttr<string S> : Attr<S, []>;
 
+/// Arbitrary string attribute.
+class ComplexStrAttr<string S, list<AttrProperty> P> : Attr<S, P>;
+
 /// Target-independent enum attributes.
 
 /// Alignment of parameter (5 bits) stored as log2 of alignment with +1 bias.
 /// 0 means unaligned (different from align(1)).
 def Alignment : IntAttr<"align", [ParamAttr, RetAttr]>;
 
 /// Parameter of a function that tells us the alignment of an allocation, as in
 /// aligned_alloc and aligned ::operator::new.

 def ApproxFuncFPMath : StrBoolAttr<"approx-func-fp-math">;
 def NoSignedZerosFPMath : StrBoolAttr<"no-signed-zeros-fp-math">;
 def UnsafeFPMath : StrBoolAttr<"unsafe-fp-math">;
 def NoJumpTables : StrBoolAttr<"no-jump-tables">;
 def NoInlineLineTables : StrBoolAttr<"no-inline-line-tables">;
 def ProfileSampleAccurate : StrBoolAttr<"profile-sample-accurate">;
 def UseSampleProfile : StrBoolAttr<"use-sample-profile">;
 
+def DenormalFPMath : ComplexStrAttr<"denormal-fp-math", [FnAttr]>;
+def DenormalFPMathF32 : ComplexStrAttr<"denormal-fp-math-f32", [FnAttr]>;
+
 class CompatRule<string F> {
   // The name of the function called to check the attribute of the caller and
   // callee and decide whether inlining should be allowed. The function's
   // signature must match "bool(const Function&, const Function &)", where the
   // first parameter is the reference to the caller and the second parameter is
   // the reference to the callee. It must return false if the attributes of the
   // caller and callee are incompatible, and true otherwise.
   string CompatFunc = F;
 }
 
 def : CompatRule<"isEqual<SanitizeAddressAttr>">;
 def : CompatRule<"isEqual<SanitizeThreadAttr>">;
 def : CompatRule<"isEqual<SanitizeMemoryAttr>">;
 def : CompatRule<"isEqual<SanitizeHWAddressAttr>">;
 def : CompatRule<"isEqual<SanitizeMemTagAttr>">;
 def : CompatRule<"isEqual<SafeStackAttr>">;
 def : CompatRule<"isEqual<ShadowCallStackAttr>">;
 def : CompatRule<"isEqual<UseSampleProfileAttr>">;
 def : CompatRule<"isEqual<NoProfileAttr>">;
+def : CompatRule<"checkDenormMode">;
+
 
 class MergeRule<string F> {
   // The name of the function called to merge the attributes of the caller and
   // callee. The function's signature must match
   // "void(Function&, const Function &)", where the first parameter is the
   // reference to the caller and the second parameter is the reference to the
   // callee.
   string MergeFunc = F;

llvm/include/llvm/IR/Function.h

public:
   bool hasOptSize() const {
     return hasFnAttribute(Attribute::OptimizeForSize) || hasMinSize();
   }
 
   /// Returns the denormal handling type for the default rounding mode of the
   /// function.
   DenormalMode getDenormalMode(const fltSemantics &FPType) const;
 
+  /// Return the representational value of "denormal-fp-math". Code interested
+  /// in the semantics of the function should use getDenormalMode instead.
+  DenormalMode getDenormalModeRaw() const;
+
+  /// Return the representational value of "denormal-fp-math-f32". Code
+  /// interested in the semantics of the function should use getDenormalMode
+  /// instead.
+  DenormalMode getDenormalModeF32Raw() const;
+
   /// copyAttributesFrom - copy all additional attributes (those not needed to
   /// create a Function) from the Function Src to this one.
   void copyAttributesFrom(const Function *Src);
 
   /// deleteBody - This method deletes the body of the function, and converts
   /// the linkage to external.
   ///
   void deleteBody() {

llvm/lib/Analysis/ConstantFolding.cpp

if (auto *CE0 = dyn_cast<ConstantExpr>(Ops0)) {
     // operands and try again.
     Predicate = ICmpInst::getSwappedPredicate(Predicate);
     return ConstantFoldCompareInstOperands(Predicate, Ops1, Ops0, DL, TLI);
   }
 
   // Flush any denormal constant float input according to denormal handling
   // mode.
   Ops0 = FlushFPConstant(Ops0, I, /* IsOutput */ false);
+  if (!Ops0)
+    return nullptr;
   Ops1 = FlushFPConstant(Ops1, I, /* IsOutput */ false);
+  if (!Ops1)
+    return nullptr;
 
   return ConstantExpr::getCompare(Predicate, Ops0, Ops1);
 }
 
 Constant *llvm::ConstantFoldUnaryOpOperand(unsigned Opcode, Constant *Op,
                                            const DataLayout &DL) {
   assert(Instruction::isUnaryOp(Opcode));
 
Constant *llvm::FlushFPConstant(Constant *Operand, const Instruction *I,
   if (!I || !I->getParent() || !I->getFunction())
     return Operand;
 
   ConstantFP *CFP = dyn_cast<ConstantFP>(Operand);
   if (!CFP)
     return Operand;
 
   const APFloat &APF = CFP->getValueAPF();
+  // TODO: Should this canonicalize nans?
+  if (!APF.isDenormal())
+    return Operand;
+
   Type *Ty = CFP->getType();
   DenormalMode DenormMode =
       I->getFunction()->getDenormalMode(Ty->getFltSemantics());
   DenormalMode::DenormalModeKind Mode =
       IsOutput ? DenormMode.Output : DenormMode.Input;
   switch (Mode) {
   default:
     llvm_unreachable("unknown denormal mode");
-    return Operand;
+  case DenormalMode::Dynamic:
+    return nullptr;
   case DenormalMode::IEEE:
     return Operand;
   case DenormalMode::PreserveSign:
     if (APF.isDenormal()) {
       return ConstantFP::get(
           Ty->getContext(),
           APFloat::getZero(Ty->getFltSemantics(), APF.isNegative()));
     }
     return Operand;
   case DenormalMode::PositiveZero:
     if (APF.isDenormal()) {
       return ConstantFP::get(Ty->getContext(),
                              APFloat::getZero(Ty->getFltSemantics(), false));
     }
     return Operand;
   }
   return Operand;
 }
 
 Constant *llvm::ConstantFoldFPInstOperands(unsigned Opcode, Constant *LHS,
                                            Constant *RHS, const DataLayout &DL,
                                            const Instruction *I) {
   if (Instruction::isBinaryOp(Opcode)) {
     // Flush denormal inputs if needed.
     Constant *Op0 = FlushFPConstant(LHS, I, /* IsOutput */ false);
+    if (!Op0)
+      return nullptr;
     Constant *Op1 = FlushFPConstant(RHS, I, /* IsOutput */ false);
+    if (!Op1)
+      return nullptr;
 
     // Calculate constant result.
     Constant *C = ConstantFoldBinaryOpOperands(Opcode, Op0, Op1, DL);
     if (!C)
       return nullptr;
 
     // Flush denormal output if needed.
     return FlushFPConstant(C, I, /* IsOutput */ true);
static Constant *constantFoldCanonicalize(const Type *Ty, const CallBase *CI,
   // Denorms and nans may have special encodings, but it should be OK to fold a
   // totally average number.
   if (Src.isNormal() || Src.isInfinity())
     return ConstantFP::get(CI->getContext(), Src);
 
   if (Src.isDenormal() && CI->getParent() && CI->getFunction()) {
     DenormalMode DenormMode =
         CI->getFunction()->getDenormalMode(Src.getSemantics());
+
+    // TODO: Should allow folding for pure IEEE.
     if (DenormMode == DenormalMode::getIEEE())
       return nullptr;
 
+    if (DenormMode == DenormalMode::getDynamic())
+      return nullptr;
+
+    // If we know if either input or output is flushed, we can fold.
+    if ((DenormMode.Input == DenormalMode::Dynamic &&
+         DenormMode.Output == DenormalMode::IEEE) ||
+        (DenormMode.Input == DenormalMode::IEEE &&
+         DenormMode.Output == DenormalMode::Dynamic))
+      return nullptr;
+
     bool IsPositive =
         (!Src.isNegative() || DenormMode.Input == DenormalMode::PositiveZero ||
          (DenormMode.Output == DenormalMode::PositiveZero &&
           DenormMode.Input == DenormalMode::IEEE));
+
     return ConstantFP::get(CI->getContext(),
                            APFloat::getZero(Src.getSemantics(), !IsPositive));
   }
 
   return nullptr;
 }
 
 static Constant *ConstantFoldScalarCall1(StringRef Name,

llvm/lib/CodeGen/CommandFlags.cpp

#define CGBINDOPT(NAME)                                                        \
 
   static cl::opt<bool> EnableNoTrappingFPMath(
       "enable-no-trapping-fp-math",
       cl::desc("Enable setting the FP exceptions build "
                "attribute not to use exceptions"),
       cl::init(false));
   CGBINDOPT(EnableNoTrappingFPMath);
 
-  static const auto DenormFlagEnumOptions =
-  cl::values(clEnumValN(DenormalMode::IEEE, "ieee",
-                        "IEEE 754 denormal numbers"),
-             clEnumValN(DenormalMode::PreserveSign, "preserve-sign",
-                        "the sign of a  flushed-to-zero number is preserved "
-                        "in the sign of 0"),
-             clEnumValN(DenormalMode::PositiveZero, "positive-zero",
-                        "denormals are flushed to positive zero"));
+  static const auto DenormFlagEnumOptions = cl::values(
+      clEnumValN(DenormalMode::IEEE, "ieee", "IEEE 754 denormal numbers"),
+                 "denormals are flushed to positive zero"),
+      clEnumValN(DenormalMode::Dynamic, "dynamic",
+                 "denormals have unknown treatment"));
 
   // FIXME: Doesn't have way to specify separate input and output modes.
   static cl::opt<DenormalMode::DenormalModeKind> DenormalFPMath(
     "denormal-fp-math",
     cl::desc("Select which denormal numbers the code is permitted to require"),
     cl::init(DenormalMode::IEEE),
     DenormFlagEnumOptions);
   CGBINDOPT(DenormalFPMath);

llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp

 }
 
 SDValue TargetLowering::getSqrtInputTest(SDValue Op, SelectionDAG &DAG,
                                          const DenormalMode &Mode) const {
   SDLoc DL(Op);
   EVT VT = Op.getValueType();
   EVT CCVT = getSetCCResultType(DAG.getDataLayout(), *DAG.getContext(), VT);
   SDValue FPZero = DAG.getConstantFP(0.0, DL, VT);
-  // Testing it with denormal inputs to avoid wrong estimate.
-  if (Mode.Input == DenormalMode::IEEE) {
-    // This is specifically a check for the handling of denormal inputs,
-    // not the result.
 
-    // Test = fabs(X) < SmallestNormal
-    const fltSemantics &FltSem = DAG.EVTToAPFloatSemantics(VT);
-    APFloat SmallestNorm = APFloat::getSmallestNormalized(FltSem);
-    SDValue NormC = DAG.getConstantFP(SmallestNorm, DL, VT);
-    SDValue Fabs = DAG.getNode(ISD::FABS, DL, VT, Op);
-    return DAG.getSetCC(DL, CCVT, Fabs, NormC, ISD::SETLT);
-  }
-  // Test = X == 0.0
-  return DAG.getSetCC(DL, CCVT, Op, FPZero, ISD::SETEQ);
-}
+  // This is specifically a check for the handling of denormal inputs, not the
+  // result.
+  if (Mode.Input == DenormalMode::PreserveSign ||
+      Mode.Input == DenormalMode::PositiveZero) {
+    // Test = X == 0.0
+    return DAG.getSetCC(DL, CCVT, Op, FPZero, ISD::SETEQ);
+  }
+
+  // Testing it with denormal inputs to avoid wrong estimate.
+  //
 
 SDValue TargetLowering::getNegatedExpression(SDValue Op, SelectionDAG &DAG,
                                              bool LegalOps, bool OptForSize,
                                              NegatibleCost &Cost,
                                              unsigned Depth) const {
   // fneg is removable even if it has multiple uses.
   if (Op.getOpcode() == ISD::FNEG) {
     Cost = NegatibleCost::Cheaper;

llvm/lib/IR/Attributes.cpp

 AttributeMask AttributeFuncs::getUBImplyingAttributes() {
   AttributeMask AM;
   AM.addAttribute(Attribute::NoUndef);
   AM.addAttribute(Attribute::Dereferenceable);
   AM.addAttribute(Attribute::DereferenceableOrNull);
   return AM;
 }
 
+/// Callees with dynamic denormal modes are compatible with any caller mode.
+static bool denormModeCompatible(DenormalMode CallerMode,
+                                 DenormalMode CalleeMode) {
+  if (CallerMode == CalleeMode || CalleeMode == DenormalMode::getDynamic())
+    return true;
+
+  // If they don't exactly match, it's OK if the mismatched component is
+  // dynamic.
+  if (CalleeMode.Input == CallerMode.Input &&
+      CalleeMode.Output == DenormalMode::Dynamic)
+    return true;
+
+  if (CalleeMode.Output == CallerMode.Output &&
+      CalleeMode.Input == DenormalMode::Dynamic)
+    return true;
+  return false;
+}
+
+static bool checkDenormMode(const Function &Caller, const Function &Callee) {
+  DenormalMode CallerMode = Caller.getDenormalModeRaw();
+  DenormalMode CalleeMode = Callee.getDenormalModeRaw();
+
+  if (denormModeCompatible(CallerMode, CalleeMode)) {
+    DenormalMode CallerModeF32 = Caller.getDenormalModeF32Raw();
+    DenormalMode CalleeModeF32 = Callee.getDenormalModeF32Raw();
+    return denormModeCompatible(CallerModeF32, CalleeModeF32);
+  }
+
+  return false;
+}
+
 template<typename AttrClass>
 static bool isEqual(const Function &Caller, const Function &Callee) {
   return Caller.getFnAttribute(AttrClass::getKind()) ==
          Callee.getFnAttribute(AttrClass::getKind());
 }
 
 /// Compute the logical AND of the attributes of the caller and the
 /// callee.

llvm/lib/IR/Function.cpp

 void Function::addDereferenceableOrNullParamAttr(unsigned ArgNo,
                                                  uint64_t Bytes) {
   AttributeSets = AttributeSets.addDereferenceableOrNullParamAttr(getContext(),
                                                                   ArgNo, Bytes);
 }
 
 DenormalMode Function::getDenormalMode(const fltSemantics &FPType) const {
   if (&FPType == &APFloat::IEEEsingle()) {
-    Attribute Attr = getFnAttribute("denormal-fp-math-f32");
-    StringRef Val = Attr.getValueAsString();
-    if (!Val.empty())
-      return parseDenormalFPAttribute(Val);
-
+    DenormalMode Mode = getDenormalModeF32Raw();
     // If the f32 variant of the attribute isn't specified, try to use the
     // generic one.
+    if (Mode.isValid())
+      return Mode;
   }
 
+  return getDenormalModeRaw();
+}
+
+DenormalMode Function::getDenormalModeRaw() const {
   Attribute Attr = getFnAttribute("denormal-fp-math");
-  return parseDenormalFPAttribute(Attr.getValueAsString());
+  StringRef Val = Attr.getValueAsString();
+  return parseDenormalFPAttribute(Val);
+}
+
+DenormalMode Function::getDenormalModeF32Raw() const {
+  Attribute Attr = getFnAttribute("denormal-fp-math-f32");
+  if (Attr.isValid()) {
+    StringRef Val = Attr.getValueAsString();
+    return parseDenormalFPAttribute(Val);
+  }
+
+  return DenormalMode::getInvalid();
 }
 
 const std::string &Function::getGC() const {
   assert(hasGC() && "Function has no collector");
   return getContext().getGC(*this);
 }
 
 void Function::setGC(std::string Str) {

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h

if (FP64FP16Denormals == DenormalMode::getPreserveSign())
       return FP_DENORM_FLUSH_IN_FLUSH_OUT;
     if (FP64FP16Denormals.Output == DenormalMode::PreserveSign)
       return FP_DENORM_FLUSH_OUT;
     if (FP64FP16Denormals.Input == DenormalMode::PreserveSign)
       return FP_DENORM_FLUSH_IN;
     return FP_DENORM_FLUSH_NONE;
   }
 
-  /// Returns true if a flag is compatible if it's enabled in the callee, but
-  /// disabled in the caller.
-  static bool oneWayCompatible(bool CallerMode, bool CalleeMode) {
-    return CallerMode == CalleeMode || (!CallerMode && CalleeMode);
-  }
-
   // FIXME: Inlining should be OK for dx10-clamp, since the caller's mode should
   // be able to override.
   bool isInlineCompatible(SIModeRegisterDefaults CalleeMode) const {
-    if (DX10Clamp != CalleeMode.DX10Clamp)
-      return false;
-    if (IEEE != CalleeMode.IEEE)
-      return false;
-
-    // Allow inlining denormals enabled into denormals flushed functions.
-    return oneWayCompatible(FP64FP16Denormals.Input !=
-                                DenormalMode::PreserveSign,
-                            CalleeMode.FP64FP16Denormals.Input !=
-                                DenormalMode::PreserveSign) &&
-           oneWayCompatible(FP64FP16Denormals.Output !=
-                                DenormalMode::PreserveSign,
-                            CalleeMode.FP64FP16Denormals.Output !=
-                                DenormalMode::PreserveSign) &&
-           oneWayCompatible(FP32Denormals.Input != DenormalMode::PreserveSign,
-                            CalleeMode.FP32Denormals.Input !=
-                                DenormalMode::PreserveSign) &&
-           oneWayCompatible(FP32Denormals.Output != DenormalMode::PreserveSign,
-                            CalleeMode.FP32Denormals.Output !=
-                                DenormalMode::PreserveSign);
+    return DX10Clamp == CalleeMode.DX10Clamp && IEEE == CalleeMode.IEEE;
   }
 };
 
 } // end namespace AMDGPU
 
 raw_ostream &operator<<(raw_ostream &OS,
                         const AMDGPU::IsaInfo::TargetIDSetting S);
 
 } // end namespace llvm
 
 #endif // LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPUBASEINFO_H

llvm/test/CodeGen/Generic/denormal-fp-math-cl-opt.ll

+; RUN: llc -denormal-fp-math=dynamic --denormal-fp-math-f32=preserve-sign -stop-after=finalize-isel < %s | FileCheck %s
+
+; Check that we annotated the command line flag annotates the IR with the appropriate attributes
+
+; CHECK: attributes #0 = { "denormal-fp-math"="dynamic,dynamic" "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
+define float @foo(float %var) {
+  ret float %var
+}

llvm/test/CodeGen/X86/sqrt-fastmath.ll

 ; AVX-LABEL: sqrt_v4f32_check_denorms:
 ; AVX:       # %bb.0:
 ; AVX-NEXT:    vsqrtps %xmm0, %xmm0
 ; AVX-NEXT:    retq
   %call = tail call <4 x float> @llvm.sqrt.v4f32(<4 x float> %x) #2
   ret <4 x float> %call
 }
 
-define <4 x float> @sqrt_v4f32_check_denorms_ninf(<4 x float> %x) #3 {
-; SSE-LABEL: sqrt_v4f32_check_denorms_ninf:
+define <4 x float> @sqrt_v4f32_check_denorms_ieee_ninf(<4 x float> %x) #3 {
+; SSE-LABEL: sqrt_v4f32_check_denorms_ieee_ninf:
 ; SSE:       # %bb.0:
 ; SSE-NEXT:    rsqrtps %xmm0, %xmm1
 ; SSE-NEXT:    movaps %xmm0, %xmm2
 ; SSE-NEXT:    mulps %xmm1, %xmm2
 ; SSE-NEXT:    movaps {{.*#+}} xmm3 = [-5.0E-1,-5.0E-1,-5.0E-1,-5.0E-1]
 ; SSE-NEXT:    mulps %xmm2, %xmm3
 ; SSE-NEXT:    mulps %xmm1, %xmm2
 ; SSE-NEXT:    addps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm2
 ; SSE-NEXT:    mulps %xmm3, %xmm2
 ; SSE-NEXT:    andps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm0
 ; SSE-NEXT:    movaps {{.*#+}} xmm1 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
 ; SSE-NEXT:    cmpleps %xmm0, %xmm1
 ; SSE-NEXT:    andps %xmm2, %xmm1
 ; SSE-NEXT:    movaps %xmm1, %xmm0
 ; SSE-NEXT:    retq
 ;
-; AVX1-LABEL: sqrt_v4f32_check_denorms_ninf:
+; AVX1-LABEL: sqrt_v4f32_check_denorms_ieee_ninf:
 ; AVX1:       # %bb.0:
 ; AVX1-NEXT:    vrsqrtps %xmm0, %xmm1
 ; AVX1-NEXT:    vmulps %xmm1, %xmm0, %xmm2
 ; AVX1-NEXT:    vmulps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm2, %xmm3
 ; AVX1-NEXT:    vmulps %xmm1, %xmm2, %xmm1
 ; AVX1-NEXT:    vaddps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm1, %xmm1
 ; AVX1-NEXT:    vmulps %xmm1, %xmm3, %xmm1
 ; AVX1-NEXT:    vandps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm0, %xmm0
 ; AVX1-NEXT:    vmovaps {{.*#+}} xmm2 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
 ; AVX1-NEXT:    vcmpleps %xmm0, %xmm2, %xmm0
 ; AVX1-NEXT:    vandps %xmm1, %xmm0, %xmm0
 ; AVX1-NEXT:    retq
 ;
-; AVX512-LABEL: sqrt_v4f32_check_denorms_ninf:
+; AVX512-LABEL: sqrt_v4f32_check_denorms_ieee_ninf:
+; AVX512:       # %bb.0:
+; AVX512-NEXT:    vrsqrtps %xmm0, %xmm1
+; AVX512-NEXT:    vmulps %xmm1, %xmm0, %xmm2
+; AVX512-NEXT:    vbroadcastss {{.*#+}} xmm3 = [-3.0E+0,-3.0E+0,-3.0E+0,-3.0E+0]
+; AVX512-NEXT:    vfmadd231ps {{.*#+}} xmm3 = (xmm2 * xmm1) + xmm3
+; AVX512-NEXT:    vbroadcastss {{.*#+}} xmm1 = [-5.0E-1,-5.0E-1,-5.0E-1,-5.0E-1]
+; AVX512-NEXT:    vmulps %xmm1, %xmm2, %xmm1
+; AVX512-NEXT:    vmulps %xmm3, %xmm1, %xmm1
+; AVX512-NEXT:    vbroadcastss {{.*#+}} xmm2 = [NaN,NaN,NaN,NaN]
+; AVX512-NEXT:    vandps %xmm2, %xmm0, %xmm0
+; AVX512-NEXT:    vbroadcastss {{.*#+}} xmm2 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
+; AVX512-NEXT:    vcmpleps %xmm0, %xmm2, %xmm0
+; AVX512-NEXT:    vandps %xmm1, %xmm0, %xmm0
+; AVX512-NEXT:    retq
+  %call = tail call ninf afn <4 x float> @llvm.sqrt.v4f32(<4 x float> %x) #2
+  ret <4 x float> %call
+}
+
+define <4 x float> @sqrt_v4f32_check_denorms_dynamic_ninf(<4 x float> %x) #6 {
+; SSE-LABEL: sqrt_v4f32_check_denorms_dynamic_ninf:
+; SSE:       # %bb.0:
+; SSE-NEXT:    rsqrtps %xmm0, %xmm1
+; SSE-NEXT:    movaps %xmm0, %xmm2
+; SSE-NEXT:    mulps %xmm1, %xmm2
+; SSE-NEXT:    movaps {{.*#+}} xmm3 = [-5.0E-1,-5.0E-1,-5.0E-1,-5.0E-1]
+; SSE-NEXT:    mulps %xmm2, %xmm3
+; SSE-NEXT:    mulps %xmm1, %xmm2
+; SSE-NEXT:    addps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm2
+; SSE-NEXT:    mulps %xmm3, %xmm2
+; SSE-NEXT:    andps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm0
+; SSE-NEXT:    movaps {{.*#+}} xmm1 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
+; SSE-NEXT:    cmpleps %xmm0, %xmm1
+; SSE-NEXT:    andps %xmm2, %xmm1
+; SSE-NEXT:    movaps %xmm1, %xmm0
+; SSE-NEXT:    retq
+;
+; AVX1-LABEL: sqrt_v4f32_check_denorms_dynamic_ninf:
+; AVX1:       # %bb.0:
+; AVX1-NEXT:    vrsqrtps %xmm0, %xmm1
+; AVX1-NEXT:    vmulps %xmm1, %xmm0, %xmm2
+; AVX1-NEXT:    vmulps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm2, %xmm3
+; AVX1-NEXT:    vmulps %xmm1, %xmm2, %xmm1
+; AVX1-NEXT:    vaddps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm1, %xmm1
+; AVX1-NEXT:    vmulps %xmm1, %xmm3, %xmm1
+; AVX1-NEXT:    vandps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm0, %xmm0
+; AVX1-NEXT:    vmovaps {{.*#+}} xmm2 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
+; AVX1-NEXT:    vcmpleps %xmm0, %xmm2, %xmm0
+; AVX1-NEXT:    vandps %xmm1, %xmm0, %xmm0
+; AVX1-NEXT:    retq
+;
+; AVX512-LABEL: sqrt_v4f32_check_denorms_dynamic_ninf:
 ; AVX512:       # %bb.0:
 ; AVX512-NEXT:    vrsqrtps %xmm0, %xmm1
 ; AVX512-NEXT:    vmulps %xmm1, %xmm0, %xmm2
 ; AVX512-NEXT:    vbroadcastss {{.*#+}} xmm3 = [-3.0E+0,-3.0E+0,-3.0E+0,-3.0E+0]
 ; AVX512-NEXT:    vfmadd231ps {{.*#+}} xmm3 = (xmm2 * xmm1) + xmm3
 ; AVX512-NEXT:    vbroadcastss {{.*#+}} xmm1 = [-5.0E-1,-5.0E-1,-5.0E-1,-5.0E-1]
 ; AVX512-NEXT:    vmulps %xmm1, %xmm2, %xmm1
 ; AVX512-NEXT:    vmulps %xmm3, %xmm1, %xmm1
; AVX-NEXT:    retq
   %rsqrt = fdiv fast double 42.0, %sqrt
   store double %rsqrt, ptr %p, align 8
   ret double %sqrt_fast
 }
 
 attributes #0 = { "unsafe-fp-math"="true" "reciprocal-estimates"="!sqrtf,!vec-sqrtf,!divf,!vec-divf" }
 attributes #1 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" }
 attributes #2 = { nounwind readnone }
-attributes #3 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="ieee" }
+attributes #3 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="preserve-sign,ieee" }
 attributes #4 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="ieee,preserve-sign" }
 attributes #5 = { "unsafe-fp-math"="true" "reciprocal-estimates"="all:0" }
+attributes #6 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="preserve-sign,dynamic" }

llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll

 
 define i32 @func_ieee_psz() #3 {
 ; CHECK-LABEL: @func_ieee_psz(
 ; CHECK-NEXT:    ret i32 4
 ;
   ret i32 4
 }
 
+define i32 @func_dynamic_dynamic() #4 {
+; CHECK-LABEL: @func_dynamic_dynamic(
+; CHECK-NEXT:    ret i32 5
+;
+  ret i32 5
+}
+
+define i32 @func_dynamic_ieee() #5 {
+; CHECK-LABEL: @func_dynamic_ieee(
+; CHECK-NEXT:    ret i32 6
+;
+  ret i32 6
+}
+
+define i32 @func_ieee_dynamic() #6 {
+; CHECK-LABEL: @func_ieee_dynamic(
+; CHECK-NEXT:    ret i32 7
+;
+  ret i32 7
+}
+
+define i32 @func_psz_dynamic() #7 {
+; CHECK-LABEL: @func_psz_dynamic(
+; CHECK-NEXT:    ret i32 8
+;
+  ret i32 8
+}
+
+define i32 @func_dynamic_psz() #8 {
+; CHECK-LABEL: @func_dynamic_psz(
+; CHECK-NEXT:    ret i32 9
+;
+  ret i32 9
+}
+
 define i32 @call_default_from_psz_psz() #1 {
 ; CHECK-LABEL: @call_default_from_psz_psz(
-; CHECK-NEXT:    ret i32 0
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_default()
+; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_default()
   ret i32 %call
 }
 
 define i32 @call_ieee_ieee_from_ieee_ieee() #0 {
 ; CHECK-LABEL: @call_ieee_ieee_from_ieee_ieee(
 ; CHECK-NEXT:    ret i32 1
 ;
   %call = call i32 @func_ieee_ieee()
   ret i32 %call
 }
 
 define i32 @call_ieee_ieee_from_psz_psz() #1 {
 ; CHECK-LABEL: @call_ieee_ieee_from_psz_psz(
-; CHECK-NEXT:    ret i32 1
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_ieee_ieee()
   ret i32 %call
 }
 
 define i32 @call_psz_psz_from_ieee_ieee() #0 {
 ; CHECK-LABEL: @call_psz_psz_from_ieee_ieee(
 ; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_psz()
 ; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_psz_psz()
   ret i32 %call
 }
 
 define i32 @call_psz_psz_from_psz_psz() #1 {
 ; CHECK-LABEL: @call_psz_psz_from_psz_psz(
 ; CHECK-NEXT:    ret i32 2
 ;
   %call = call i32 @func_psz_psz()
   ret i32 %call
 }
 
 define i32 @call_psz_ieee_from_psz_psz() #1 {
 ; CHECK-LABEL: @call_psz_ieee_from_psz_psz(
-; CHECK-NEXT:    ret i32 3
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_psz_ieee()
   ret i32 %call
 }
 
 define i32 @call_ieee_psz_from_psz_psz() #1 {
 ; CHECK-LABEL: @call_ieee_psz_from_psz_psz(
-; CHECK-NEXT:    ret i32 4
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_ieee_psz()
   ret i32 %call
 }
 
 define i32 @call_psz_ieee_from_ieee_ieee() #0 {
 ; CHECK-LABEL: @call_psz_ieee_from_ieee_ieee(
 ; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()

 ; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_ieee_psz()
   ret i32 %call
 }
 
 define i32 @call_ieee_ieee_from_psz_ieee() #2 {
 ; CHECK-LABEL: @call_ieee_ieee_from_psz_ieee(
-; CHECK-NEXT:    ret i32 1
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_ieee_ieee()
   ret i32 %call
 }
 
 define i32 @call_psz_ieee_from_psz_ieee() #2 {
 ; CHECK-LABEL: @call_psz_ieee_from_psz_ieee(
 ; CHECK-NEXT:    ret i32 3

 ; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_psz_psz()
   ret i32 %call
 }
 
 define i32 @call_ieee_ieee_from_ieee_psz() #3 {
 ; CHECK-LABEL: @call_ieee_ieee_from_ieee_psz(
-; CHECK-NEXT:    ret i32 1
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_ieee_ieee()
   ret i32 %call
 }
 
 define i32 @call_psz_ieee_from_ieee_psz() #3 {
 ; CHECK-LABEL: @call_psz_ieee_from_ieee_psz(
 ; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()

 ; CHECK-LABEL: @call_psz_psz_from_ieee_psz(
 ; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_psz()
 ; CHECK-NEXT:    ret i32 [[CALL]]
 ;
   %call = call i32 @func_psz_psz()
   ret i32 %call
 }
 
+define i32 @call_dynamic_dynamic_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_ieee_ieee(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_dynamic_ieee_from_ieee_ieee(
+; CHECK-NEXT:    ret i32 6
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_ieee_dynamic_from_ieee_ieee(
+; CHECK-NEXT:    ret i32 7
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_dynamic_psz_from_ieee_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_psz_dynamic_from_ieee_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_psz_psz() #1 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_psz(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_psz_psz() #1 {
+; CHECK-LABEL: @call_ieee_dynamic_from_psz_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_psz_psz() #1 {
+; CHECK-LABEL: @call_dynamic_ieee_from_psz_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_psz_psz() #1 {
+; CHECK-LABEL: @call_psz_dynamic_from_psz_psz(
+; CHECK-NEXT:    ret i32 8
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_psz_psz() #1 {
+; CHECK-LABEL: @call_dynamic_psz_from_psz_psz(
+; CHECK-NEXT:    ret i32 9
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_ieee(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_dynamic_ieee_from_psz_ieee(
+; CHECK-NEXT:    ret i32 6
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_ieee_dynamic_from_psz_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_dynamic_psz_from_psz_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_psz_dynamic_from_psz_ieee(
+; CHECK-NEXT:    ret i32 8
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_ieee_psz(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_ieee_dynamic_from_ieee_psz(
+; CHECK-NEXT:    ret i32 7
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_dynamic_ieee_from_ieee_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_dynamic_psz_from_ieee_psz(
+; CHECK-NEXT:    ret i32 9
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_psz_dynamic_from_ieee_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_dynamic_dynamic(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_ieee_ieee_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_ieee_dynamic_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_dynamic_ieee_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_psz_dynamic_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_dynamic_psz_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_psz_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_psz_psz_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_psz_ieee_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_ieee_psz_from_dynamic_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_psz()
+  ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_ieee_ieee_from_dynamic_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_ieee_dynamic_from_dynamic_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_dynamic_ieee_from_dynamic_ieee(
+; CHECK-NEXT:    ret i32 6
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_psz_dynamic_from_dynamic_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_dynamic_psz_from_dynamic_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_psz_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_psz_psz_from_dynamic_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_psz_ieee_from_dynamic_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_ieee_psz_from_dynamic_ieee(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_psz()
+  ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_ieee_ieee_from_ieee_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_ieee_dynamic_from_ieee_dynamic(
+; CHECK-NEXT:    ret i32 7
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_dynamic_ieee_from_ieee_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_ieee_dynamic(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_dynamic_psz_from_ieee_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_psz_dynamic_from_ieee_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_psz_psz_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_psz_psz_from_ieee_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_psz_ieee_from_ieee_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_ieee_psz_from_ieee_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_psz()
+  ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_ieee_ieee_from_psz_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_ieee_dynamic_from_psz_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_dynamic_ieee_from_psz_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_dynamic_psz_from_psz_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_psz_dynamic_from_psz_dynamic(
+; CHECK-NEXT:    ret i32 8
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_dynamic(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_psz_ieee_from_psz_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_ieee_psz_from_psz_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_psz_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_psz_psz_from_psz_dynamic(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_psz()
+  ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_ieee_ieee_from_dynamic_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_ieee_psz_from_dynamic_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_psz_ieee_from_dynamic_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_ieee()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_dynamic_psz(
+; CHECK-NEXT:    ret i32 5
+;
+  %call = call i32 @func_dynamic_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_dynamic_psz_from_dynamic_psz(
+; CHECK-NEXT:    ret i32 9
+;
+  %call = call i32 @func_dynamic_psz()
+  ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_psz_dynamic_from_dynamic_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_dynamic_ieee_from_dynamic_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_dynamic_ieee()
+  ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_ieee_dynamic_from_dynamic_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_ieee_dynamic()
+  ret i32 %call
+}
+
+define i32 @call_psz_psz_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_psz_psz_from_dynamic_psz(
+; CHECK-NEXT:    [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT:    ret i32 [[CALL]]
+;
+  %call = call i32 @func_psz_psz()
+  ret i32 %call
+}
+
+; --------------------------------------------------------------------
+; denormal-fp-math-f32
+; --------------------------------------------------------------------
+
+define i32 @func_dynamic_dynamic_f32() #9 {
+; CHECK-LABEL: @func_dynamic_dynamic_f32(
+; CHECK-NEXT:    ret i32 10
+;
+  ret i32 10
+}
+
+define i32 @func_psz_psz_f32() #10 {
+; CHECK-LABEL: @func_psz_psz_f32(
+; CHECK-NEXT:    ret i32 11
+;
+  ret i32 11
+}
+
+define i32 @call_dynamic_dynamic_from_psz_psz_f32() #10 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_psz_f32(
+; CHECK-NEXT:    ret i32 10
+;
+  %result = call i32 @func_dynamic_dynamic_f32()
+  ret i32 %result
+}
+
+define i32 @call_psz_psz_from_psz_psz_f32() #10 {
+; CHECK-LABEL: @call_psz_psz_from_psz_psz_f32(
+; CHECK-NEXT:    ret i32 10
+;
+  %result = call i32 @func_dynamic_dynamic_f32()
+  ret i32 %result
+}
+
+define i32 @call_psz_psz_from_ieee_ieee_f32() #11 {
+; CHECK-LABEL: @call_psz_psz_from_ieee_ieee_f32(
+; CHECK-NEXT:    [[RESULT:%.*]] = call i32 @func_psz_psz_f32()
+; CHECK-NEXT:    ret i32 [[RESULT]]
+;
+  %result = call i32 @func_psz_psz_f32()
+  ret i32 %result
+}
+
 attributes #0 = { "denormal-fp-math"="ieee,ieee" }
 attributes #1 = { "denormal-fp-math"="preserve-sign,preserve-sign" }
 attributes #2 = { "denormal-fp-math"="preserve-sign,ieee" }
 attributes #3 = { "denormal-fp-math"="ieee,preserve-sign" }
+attributes #4 = { "denormal-fp-math"="dynamic,dynamic" }
+attributes #5 = { "denormal-fp-math"="dynamic,ieee" }
+attributes #6 = { "denormal-fp-math"="ieee,dynamic" }
+attributes #7 = { "denormal-fp-math"="preserve-sign,dynamic" }
+attributes #8 = { "denormal-fp-math"="dynamic,preserve-sign" }
+attributes #9 = { "denormal-fp-math-f32"="dynamic,dynamic" }
+attributes #10 = { "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
+attributes #11 = { "denormal-fp-math-f32"="ieee,ieee" "denormal-fp-math"="preserve-sign,preserve-sign" }

llvm/test/Transforms/InstSimplify/canonicalize.ll

 define float @canonicalize_neg_denorm_positive_zero_input() "denormal-fp-math"="ieee,positive-zero" {
 ; CHECK-LABEL: @canonicalize_neg_denorm_positive_zero_input(
 ; CHECK-NEXT:    ret float 0.000000e+00
 ;
   %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
   ret float %ret
 }
 
+define float @canonicalize_pos_denorm_dynamic_dynamic() "denormal-fp-math"="dynamic,dynamic" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_dynamic(
+; CHECK-NEXT:    [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0x380FFFFFC0000000)
+; CHECK-NEXT:    ret float [[RET]]
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+  ret float %ret
+}
+
+define float @canonicalize_neg_denorm_dynamic_dynamic() "denormal-fp-math"="dynamic,dynamic" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_dynamic(
+; CHECK-NEXT:    [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0xB80FFFFFC0000000)
+; CHECK-NEXT:    ret float [[RET]]
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+  ret float %ret
+}
+
+; Dynamic output - cannot flush
+define float @canonicalize_pos_denorm_dynamic_output() "denormal-fp-math"="dynamic,ieee" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_output(
+; CHECK-NEXT:    [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0x380FFFFFC0000000)
+; CHECK-NEXT:    ret float [[RET]]
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+  ret float %ret
+}
+
+; Dynamic output - cannot flush
+define float @canonicalize_neg_denorm_dynamic_output() "denormal-fp-math"="dynamic,ieee" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_output(
+; CHECK-NEXT:    [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0xB80FFFFFC0000000)
+; CHECK-NEXT:    ret float [[RET]]
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+  ret float %ret
+}
+
+; Dynamic input - cannot flush
+define float @canonicalize_pos_denorm_dynamic_input() "denormal-fp-math"="ieee,dynamic" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_input(
+; CHECK-NEXT:    [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0x380FFFFFC0000000)
+; CHECK-NEXT:    ret float [[RET]]
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+  ret float %ret
+}
+
+; Dynamic input - cannot flush
+define float @canonicalize_neg_denorm_dynamic_input() "denormal-fp-math"="ieee,dynamic" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_input(
+; CHECK-NEXT:    [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0xB80FFFFFC0000000)
+; CHECK-NEXT:    ret float [[RET]]
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+  ret float %ret
+}
+
+; Input is flushed, can fold
+define float @canonicalize_pos_denorm_dynamic_output_preserve_sign_input() "denormal-fp-math"="dynamic,preserve-sign" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_output_preserve_sign_input(
+; CHECK-NEXT:    ret float 0.000000e+00
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+  ret float %ret
+}
+
+; Input is flushed, can fold
+define float @canonicalize_neg_denorm_dynamic_output_preserve_sign_input() "denormal-fp-math"="dynamic,preserve-sign" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_output_preserve_sign_input(
+; CHECK-NEXT:    ret float -0.000000e+00
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+  ret float %ret
+}
+
+; Output is known flushed, can fold
+define float @canonicalize_pos_preserve_sign_output_denorm_dynamic_input() "denormal-fp-math"="preserve-sign,dynamic" {
+; CHECK-LABEL: @canonicalize_pos_preserve_sign_output_denorm_dynamic_input(
+; CHECK-NEXT:    ret float 0.000000e+00
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+  ret float %ret
+}
+
+; Output is known flushed, can fold
+define float @canonicalize_neg_denorm_preserve_sign_output_dynamic_input() "denormal-fp-math"="preserve-sign,dynamic" {
+; CHECK-LABEL: @canonicalize_neg_denorm_preserve_sign_output_dynamic_input(
+; CHECK-NEXT:    ret float -0.000000e+00
+;
+  %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+  ret float %ret
+}
+
 define float @canonicalize_inf() {
 ; CHECK-LABEL: @canonicalize_inf(
 ; CHECK-NEXT:    ret float 0x7FF0000000000000
 ;
   %ret = call float @llvm.canonicalize.f32(float 0x7FF0000000000000)
   ret float %ret
 }
 

llvm/test/Transforms/InstSimplify/constant-fold-fp-denormal.ll

 ; CHECK-NEXT:  entry:
 ; CHECK-NEXT:    ret i1 false
 ;
 entry:
   %cmp = fcmp uno double 0x0008000000000000, 0x1ff1000000000000
   ret i1 %cmp
 }
 
+; ============================================================================ ;
+; dynamic mode tests
+; ============================================================================ ;
+
+define float @test_float_fadd_dynamic_ieee() #9 {
+; CHECK-LABEL: @test_float_fadd_dynamic_ieee(
+; CHECK-NEXT:    [[RESULT:%.*]] = fadd float 0xB810000000000000, 0x3800000000000000
+; CHECK-NEXT:    ret float [[RESULT]]
+;
+  %result = fadd float 0xB810000000000000, 0x3800000000000000
+  ret float %result
+}
+
+define float @test_float_fadd_ieee_dynamic() #10 {
+; CHECK-LABEL: @test_float_fadd_ieee_dynamic(
+; CHECK-NEXT:    [[RESULT:%.*]] = fadd float 0xB810000000000000, 0x3800000000000000
+; CHECK-NEXT:    ret float [[RESULT]]
+;
+  %result = fadd float 0xB810000000000000, 0x3800000000000000
+  ret float %result
+}
+
+define float @test_float_fadd_dynamic_dynamic() #11 {
+; CHECK-LABEL: @test_float_fadd_dynamic_dynamic(
+; CHECK-NEXT:    [[RESULT:%.*]] = fadd float 0xB810000000000000, 0x3800000000000000
+; CHECK-NEXT:    ret float [[RESULT]]
+;
+  %result = fadd float 0xB810000000000000, 0x3800000000000000
+  ret float %result
+}
+
+; Check for failed to fold on each operand
+define float @test_float_fadd_dynamic_dynamic_commute() #11 {
+; CHECK-LABEL: @test_float_fadd_dynamic_dynamic_commute(
+; CHECK-NEXT:    [[RESULT:%.*]] = fadd float 0x3800000000000000, 0xB810000000000000
+; CHECK-NEXT:    ret float [[RESULT]]
+;
+  %result = fadd float 0x3800000000000000, 0xB810000000000000
+  ret float %result
+}
+
+define i1 @fcmp_double_dynamic_ieee() #9 {
+; CHECK-LABEL: @fcmp_double_dynamic_ieee(
+; CHECK-NEXT:    ret i1 true
+;
+  %cmp = fcmp une double 0x0008000000000000, 0x0
+  ret i1 %cmp
+}
+
+define i1 @fcmp_double_ieee_dynamic() #10 {
+; CHECK-LABEL: @fcmp_double_ieee_dynamic(
+; CHECK-NEXT:    [[CMP:%.*]] = fcmp une double 0x8000000000000, 0.000000e+00
+; CHECK-NEXT:    ret i1 [[CMP]]
+;
+  %cmp = fcmp une double 0x0008000000000000, 0x0
+  ret i1 %cmp
+}
+
+define i1 @fcmp_double_dynamic_dynamic() #11 {
+; CHECK-LABEL: @fcmp_double_dynamic_dynamic(
+; CHECK-NEXT:    [[CMP:%.*]] = fcmp une double 0x8000000000000, 0.000000e+00
+; CHECK-NEXT:    ret i1 [[CMP]]
+;
+  %cmp = fcmp une double 0x0008000000000000, 0x0
+  ret i1 %cmp
+}
+
+define i1 @fcmp_double_dynamic_dynamic_commute() #11 {
+; CHECK-LABEL: @fcmp_double_dynamic_dynamic_commute(
+; CHECK-NEXT:    [[CMP:%.*]] = fcmp une double 0.000000e+00, 0x8000000000000
+; CHECK-NEXT:    ret i1 [[CMP]]
+;
+  %cmp = fcmp une double 0x0, 0x0008000000000000
+  ret i1 %cmp
+}
+
+; Output doesn't matter.
+define i1 @fcmp_double_dynamic_psz() #12 {
+; CHECK-LABEL: @fcmp_double_dynamic_psz(
+; CHECK-NEXT:    ret i1 false
+;
+  %cmp = fcmp une double 0x0008000000000000, 0x0
+  ret i1 %cmp
+}
+
+; Non-denormal values should fold
+define float @test_float_fadd_dynamic_dynamic_normals() #11 {
+; CHECK-LABEL: @test_float_fadd_dynamic_dynamic_normals(
+; CHECK-NEXT:    ret float 3.000000e+00
+;
+  %result = fadd float 1.0, 2.0
+  ret float %result
+}
+
+; Non-denormal values should fold
+define i1 @fcmp_double_dynamic_dynamic_normals() #11 {
+; CHECK-LABEL: @fcmp_double_dynamic_dynamic_normals(
+; CHECK-NEXT:    ret i1 true
+;
+  %cmp = fcmp une double 1.0, 2.0
+  ret i1 %cmp
+}
+
 attributes #0 = { nounwind "denormal-fp-math"="ieee,ieee" }
 attributes #1 = { nounwind "denormal-fp-math"="positive-zero,ieee" }
 attributes #2 = { nounwind "denormal-fp-math"="preserve-sign,ieee" }
 attributes #3 = { nounwind "denormal-fp-math"="ieee,positive-zero" }
 attributes #4 = { nounwind "denormal-fp-math"="ieee,preserve-sign" }
 attributes #5 = { nounwind "denormal-fp-math"="ieee,ieee" "denormal-fp-math-f32"="positive-zero,ieee" }
 attributes #6 = { nounwind "denormal-fp-math"="positive-zero,positive-zero" }
 attributes #7 = { nounwind "denormal-fp-math"="preserve-sign,preserve-sign" }
 attributes #8 = { nounwind "denormal-fp-math"="ieee,ieee" "denormal-fp-math-f32"="positive-zero,positive-zero" }
+attributes #9 = { nounwind "denormal-fp-math"="dynamic,ieee" }
+attributes #10 = { nounwind "denormal-fp-math"="ieee,dynamic" }
+attributes #11 = { nounwind "denormal-fp-math"="dynamic,dynamic" }
+attributes #12 = { nounwind "denormal-fp-math"="dynamic,preserve-sign" }

llvm/unittests/ADT/FloatingPointMode.cpp

 
 TEST(FloatingPointModeTest, ParseDenormalFPAttributeComponent) {
   EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttributeComponent("ieee"));
   EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttributeComponent(""));
   EXPECT_EQ(DenormalMode::PreserveSign,
             parseDenormalFPAttributeComponent("preserve-sign"));
   EXPECT_EQ(DenormalMode::PositiveZero,
             parseDenormalFPAttributeComponent("positive-zero"));
+  EXPECT_EQ(DenormalMode::Dynamic,
+            parseDenormalFPAttributeComponent("dynamic"));
   EXPECT_EQ(DenormalMode::Invalid, parseDenormalFPAttributeComponent("foo"));
 }
 
 TEST(FloatingPointModeTest, DenormalAttributeName) {
   EXPECT_EQ("ieee", denormalModeKindName(DenormalMode::IEEE));
   EXPECT_EQ("preserve-sign", denormalModeKindName(DenormalMode::PreserveSign));
   EXPECT_EQ("positive-zero", denormalModeKindName(DenormalMode::PositiveZero));
+  EXPECT_EQ("dynamic", denormalModeKindName(DenormalMode::Dynamic));
   EXPECT_EQ("", denormalModeKindName(DenormalMode::Invalid));
 }
 
 TEST(FloatingPointModeTest, ParseDenormalFPAttribute) {
   EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE),
             parseDenormalFPAttribute("ieee"));
   EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE),
             parseDenormalFPAttribute("ieee,ieee"));
TEST(FloatingPointModeTest, ParseDenormalFPAttribute) {
   EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::PreserveSign),
             parseDenormalFPAttribute("preserve-sign,preserve-sign"));
 
   EXPECT_EQ(DenormalMode(DenormalMode::PositiveZero, DenormalMode::PositiveZero),
             parseDenormalFPAttribute("positive-zero"));
   EXPECT_EQ(DenormalMode(DenormalMode::PositiveZero, DenormalMode::PositiveZero),
             parseDenormalFPAttribute("positive-zero,positive-zero"));
 
+  EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic),
+            parseDenormalFPAttribute("dynamic"));
+  EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic),
+            parseDenormalFPAttribute("dynamic,dynamic"));
 
   EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::PositiveZero),
             parseDenormalFPAttribute("ieee,positive-zero"));
   EXPECT_EQ(DenormalMode(DenormalMode::PositiveZero, DenormalMode::IEEE),
             parseDenormalFPAttribute("positive-zero,ieee"));
 
   EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::IEEE),
             parseDenormalFPAttribute("preserve-sign,ieee"));
   EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::PreserveSign),
             parseDenormalFPAttribute("ieee,preserve-sign"));
 
+  EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::PreserveSign),
+            parseDenormalFPAttribute("dynamic,preserve-sign"));
+  EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::Dynamic),
+            parseDenormalFPAttribute("preserve-sign,dynamic"));
 
   EXPECT_EQ(DenormalMode(DenormalMode::Invalid, DenormalMode::Invalid),
             parseDenormalFPAttribute("foo"));
   EXPECT_EQ(DenormalMode(DenormalMode::Invalid, DenormalMode::Invalid),
             parseDenormalFPAttribute("foo,foo"));
   EXPECT_EQ(DenormalMode(DenormalMode::Invalid, DenormalMode::Invalid),
             parseDenormalFPAttribute("foo,bar"));
 }
TEST(FloatingPointModeTest, RenderDenormalFPAttribute) {
 
   EXPECT_EQ(
     "preserve-sign,ieee",
     DenormalMode(DenormalMode::PreserveSign, DenormalMode::IEEE).str());
 
   EXPECT_EQ(
     "preserve-sign,positive-zero",
     DenormalMode(DenormalMode::PreserveSign, DenormalMode::PositiveZero).str());
+
+  EXPECT_EQ("dynamic,dynamic",
+            DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic).str());
+  EXPECT_EQ("ieee,dynamic",
+            DenormalMode(DenormalMode::IEEE, DenormalMode::Dynamic).str());
+  EXPECT_EQ("dynamic,ieee",
+            DenormalMode(DenormalMode::Dynamic, DenormalMode::IEEE).str());
 }
 
 TEST(FloatingPointModeTest, DenormalModeIsSimple) {
   EXPECT_TRUE(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE).isSimple());
   EXPECT_FALSE(DenormalMode(DenormalMode::IEEE,
                             DenormalMode::Invalid).isSimple());
   EXPECT_FALSE(DenormalMode(DenormalMode::PreserveSign,
                             DenormalMode::PositiveZero).isSimple());
+  EXPECT_FALSE(DenormalMode(DenormalMode::PreserveSign, DenormalMode::Dynamic)
+                   .isSimple());
+  EXPECT_FALSE(DenormalMode(DenormalMode::Dynamic, DenormalMode::PreserveSign)
+                   .isSimple());
 }
 
 TEST(FloatingPointModeTest, DenormalModeIsValid) {
   EXPECT_TRUE(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE).isValid());
   EXPECT_FALSE(DenormalMode(DenormalMode::IEEE, DenormalMode::Invalid).isValid());
   EXPECT_FALSE(DenormalMode(DenormalMode::Invalid, DenormalMode::IEEE).isValid());
   EXPECT_FALSE(DenormalMode(DenormalMode::Invalid,
                             DenormalMode::Invalid).isValid());
 }
 
 TEST(FloatingPointModeTest, DenormalModeConstructor) {
   EXPECT_EQ(DenormalMode(DenormalMode::Invalid, DenormalMode::Invalid),
             DenormalMode::getInvalid());
   EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE),
             DenormalMode::getIEEE());
+  EXPECT_EQ(DenormalMode::getIEEE(), DenormalMode::getDefault());
+  EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic),
+            DenormalMode::getDynamic());
   EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::PreserveSign),
             DenormalMode::getPreserveSign());
   EXPECT_EQ(DenormalMode(DenormalMode::PositiveZero, DenormalMode::PositiveZero),
             DenormalMode::getPositiveZero());
 }
 
+TEST(FloatingPointModeTest, DenormalModeMerge) {
+  EXPECT_EQ(
+      DenormalMode::getInvalid(),
+      DenormalMode::getInvalid().mergeCalleeMode(DenormalMode::getInvalid()));
+  EXPECT_EQ(DenormalMode::getIEEE(), DenormalMode::getInvalid().mergeCalleeMode(
+                                         DenormalMode::getIEEE()));
+  EXPECT_EQ(DenormalMode::getInvalid(), DenormalMode::getIEEE().mergeCalleeMode(
+                                            DenormalMode::getInvalid()));
+
+  EXPECT_EQ(DenormalMode::getIEEE(), DenormalMode::getIEEE().mergeCalleeMode(
+                                         DenormalMode::getDynamic()));
+  EXPECT_EQ(DenormalMode::getPreserveSign(),
+            DenormalMode::getPreserveSign().mergeCalleeMode(
+                DenormalMode::getDynamic()));
+  EXPECT_EQ(DenormalMode::getPositiveZero(),
+            DenormalMode::getPositiveZero().mergeCalleeMode(
+                DenormalMode::getDynamic()));
+  EXPECT_EQ(
+      DenormalMode::getDynamic(),
+      DenormalMode::getDynamic().mergeCalleeMode(DenormalMode::getDynamic()));
+
+  EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::PreserveSign),
+            DenormalMode(DenormalMode::IEEE, DenormalMode::PreserveSign)
+                .mergeCalleeMode(
+                    DenormalMode(DenormalMode::IEEE, DenormalMode::Dynamic)));
+
+  EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::IEEE),
+            DenormalMode(DenormalMode::PreserveSign, DenormalMode::IEEE)
+                .mergeCalleeMode(
+                    DenormalMode(DenormalMode::Dynamic, DenormalMode::IEEE)));
+
+  EXPECT_EQ(
+      DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign),
+      DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign)
+          .mergeCalleeMode(
+              DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic)));
+
+  EXPECT_EQ(
+      DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign),
+      DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign)
+          .mergeCalleeMode(
+              DenormalMode(DenormalMode::PositiveZero, DenormalMode::Dynamic)));
+
+  EXPECT_EQ(
+      DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign),
+      DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign)
+          .mergeCalleeMode(
+              DenormalMode(DenormalMode::Dynamic, DenormalMode::PreserveSign)));
+
+  // Test some invalid / undefined behavior cases
+  EXPECT_EQ(
+      DenormalMode::getPreserveSign(),
+      DenormalMode::getIEEE().mergeCalleeMode(DenormalMode::getPreserveSign()));
+  EXPECT_EQ(
+      DenormalMode::getPreserveSign(),
+      DenormalMode::getIEEE().mergeCalleeMode(DenormalMode::getPreserveSign()));
+  EXPECT_EQ(
+      DenormalMode::getIEEE(),
+      DenormalMode::getPreserveSign().mergeCalleeMode(DenormalMode::getIEEE()));
+  EXPECT_EQ(
+      DenormalMode::getIEEE(),
+      DenormalMode::getPreserveSign().mergeCalleeMode(DenormalMode::getIEEE()));
+}
 }

llvm/utils/TableGen/Attributes.cpp

for (StringRef KindName : KindNames) {
       }
     }
     OS << "#undef " << MacroName << "\n\n";
   };
 
   // Emit attribute enums in the same order llvm::Attribute::operator< expects.
   Emit({"EnumAttr", "TypeAttr", "IntAttr"}, "ATTRIBUTE_ENUM");
   Emit({"StrBoolAttr"}, "ATTRIBUTE_STRBOOL");
+  Emit({"ComplexStrAttr"}, "ATTRIBUTE_COMPLEXSTR");
 
   OS << "#undef ATTRIBUTE_ALL\n";
   OS << "#endif\n\n";
 
   OS << "#ifdef GET_ATTR_ENUM\n";
   OS << "#undef GET_ATTR_ENUM\n";
   unsigned Value = 1; // Leave zero for AttrKind::None.
   for (StringRef KindName : {"EnumAttr", "TypeAttr", "IntAttr"}) {