Consoldiate internal denormal flushing controls

[arsenm]

Summary

Currently there are 4 different mechanisms for controlling denormal
flushing behavior, and about as many equivalent frontend controls.

- AMDGPU uses the fp32-denormals and fp64-f16-denormals subtarget features
- NVPTX uses the nvptx-f32ftz attribute
- ARM directly uses the denormal-fp-math attribute
- Other targets indirectly use denormal-fp-math in one DAGCombine
- cl-denorms-are-zero has a corresponding denorms-are-zero attribute

AMDGPU wants a distinct control for f32 flushing from f16/f64, and as
far as I can tell the same is true for NVPTX (based on the attribute
name).

Work on consolidating these into the denormal-fp-math attribute, and a
new type specific denormal-fp-math-f32 variant. Only ARM seems to
support the two different flush modes, so this is overkill for the
other use cases. Ideally we would error on the unsupported
positive-zero mode on other targets from somewhere.

Move the logic for selecting the flush mode into the compiler driver,
instead of handling it in cc1. denormal-fp-math/denormal-fp-math-f32
are now both cc1 flags, but denormal-fp-math-f32 is not yet exposed as
a user flag.

-cl-denorms-are-zero, -fcuda-flush-denormals-to-zero and
-fno-cuda-flush-denormals-to-zero will be mapped to
-fp-denormal-math-f32=ieee or preserve-sign rather than the old
-attributes.

Stop emitting the denorms-are-zero attribute for the OpenCL flag. It
has no in-tree users. The meaning would also be target dependent, such
as the AMDGPU choice to treat this as only meaning allow flushing of
f32 and not f16 or f64. The naming is also potentially confusing,
since DAZ in other contexts refers to instructions implicitly treating
input denormals as zero, not necessarily flushing output denormals to
zero.

This also does not attempt to change the behavior for the current
attribute. The LangRef now states that the default is ieee behavior,
but this is inaccurate for the current implementation. The clang
handling is slightly hacky to avoid touching the existing
denormal-fp-math uses. Fixing this will be left for a future patch.

AMDGPU is still using the subtarget feature to control the denormal
mode, but the new attribute are now emitted. A future change will
switch this and remove the subtarget features.

[arsenm]

llvm/docs/LangRef.rst

1828–1831

On second thought I think this may be too permissive. I think based on the use in DAGCombiner, that flushing of outputs is compulsory.

[arsenm]

llvm/docs/LangRef.rst

1828–1831

It turns out the fast sqrt usage really cares about input denormals being implicitly treated as 0, not the output flushing (i.e. this only needs DAZ, not FTZ). I think being permissive on the output is OK, but if implicit input flushing is required then it's compulsory and a target is responsible for inserting a flush of some kind if the use instruction isn't known to follow this mode.

Because of this, I do think it's necessary to treat this as two separate modes. I'm thinking to comma separate output-mode,input-mode, and assume input-mode=output-mode if the second half isn't specified for compatibility with the existing attribute.

[arsenm]

Rename subnormal to denormal. Will defer splitting input and output setting into a future patch before switching default behavior

[Anastasia]

Stop emitting the denorms-are-zero attribute for the OpenCL flag. It
has no in-tree users. The meaning would also be target dependent, such
as the AMDGPU choice to treat this as only meaning allow flushing of
f32 and not f16 or f64. The naming is also potentially confusing,
since DAZ in other contexts refers to instructions implicitly treating
input denormals as zero, not necessarily flushing output denormals to
zero.

Would the targets supporting OpenCL need to define their own behavior in getDefaultDenormalModeForType?

clang/include/clang/Driver/ToolChain.h

619

Can you elaborate what has to be done in order to fix this?

[arsenm]

In D69878#1736865, @Anastasia wrote:

Stop emitting the denorms-are-zero attribute for the OpenCL flag. It
has no in-tree users. The meaning would also be target dependent, such
as the AMDGPU choice to treat this as only meaning allow flushing of
f32 and not f16 or f64. The naming is also potentially confusing,
since DAZ in other contexts refers to instructions implicitly treating
input denormals as zero, not necessarily flushing output denormals to
zero.

Would the targets supporting OpenCL need to define their own behavior in getDefaultDenormalModeForType?

Yes. The future ieee default should be conservatively correct though

clang/include/clang/Driver/ToolChain.h

619

The main problem is the current user assumes non-ieee by default. The main blocker is knowing what platforms should default to something different to avoid performance regressions. I have the patch almost ready to switch the default, it’s just missing toolchain overrides

[arsenm]

Fix name in documentation

[jlebar]

AMDGPU wants a distinct control for f32 flushing from f16/f64, and as far as I can tell the same is true for NVPTX (based on the attribute name).

I may be corrected, but I believe nvptx only supports ftz for f32.

Double-precision instructions support subnormal inputs and results. Single-precision instructions support subnormal inputs and results by default for sm_20 and subsequent targets, and flush subnormal inputs and results to sign-preserving zero for sm_1x targets. The optional .ftz modifier on single-precision instructions provides backward compatibility with sm_1x targets by flushing subnormal inputs and results to sign-preserving zero regardless of the target architecture.

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

1780

so where would denorms-are-zero be emitted now (in case some out of tree implementations rely on this)?

[arsenm]

clang/lib/CodeGen/CGCall.cpp

1780

Rely on in what sense? Do you have a concrete use of this?

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

1780

Since it has been emitted before in the module potentially some LLVM implementations could be using that attribute?

[arsenm]

clang/lib/CodeGen/CGCall.cpp

1780

I'm disinclined to leave things around just in case some unknown user might have been using them. We've dropped attributes like this before (I think the less-precise-fp-mad one for disuse). This also isn't needed for correctness, so it should be pretty safe to drop

[scanon]

llvm/docs/LangRef.rst

1834

Can you clarify this a little bit? I'd prefer something like "Same as `"denorm-fp-math"`, but only controls the behavior of the 32-bit float type.".

[arsenm]

Reword langref, fix name in langref

[Anastasia]

clang/lib/CodeGen/CGCall.cpp

1780

Ok, fair enough!

[arsenm]

ping

[cameron.mcinally]

This is looking pretty good to me, but I'm ignoring some of the target specific code that I'm not familiar with.

Is denormal-fp-math influenced by -Ofast? Or are there plans for that? Seems like -Ofast should imply DAZ and FTZ (if supported by target).

I think we discussed this before, but it's worth repeating. If denormal-fp-math isn't specified, we default to IEEE behavior, right? When this lands in master, there could be an unexpected performance hit for targets that aren't paying attention. E.g. I want to use denormal-fp-math to toggle whether a FSUB(-0.0,X) is converted to a FNEG(X) in SelectionDAGBuilder.

Apologies in advance if this has been discussed recently. I've been distracted with another project for the passed few months...

clang/lib/Driver/ToolChains/Clang.cpp

2311

Last line of comment was not removed.

Also, is it safe to remove TrappingMathPresent? Is that part of the work-in-progress to support ffp-exception-behavior?

[arsenm]

In D69878#1801508, @cameron.mcinally wrote:

This is looking pretty good to me, but I'm ignoring some of the target specific code that I'm not familiar with.

Is denormal-fp-math influenced by -Ofast? Or are there plans for that? Seems like -Ofast should imply DAZ and FTZ (if supported by target).

Yes, through the toolchain handling. I copied the logic for when crtfastmath is linked for the default mode for x86.

I think we discussed this before, but it's worth repeating. If denormal-fp-math isn't specified, we default to IEEE behavior, right? When this lands in master, there could be an unexpected performance hit for targets that aren't paying attention. E.g. I want to use denormal-fp-math to toggle whether a FSUB(-0.0,X) is converted to a FNEG(X) in SelectionDAGBuilder.

Apologies in advance if this has been discussed recently. I've been distracted with another project for the passed few months...

Yes, ieee should be the default. The dependent patches start adding the attribute by default for platforms with flushing enabled with fast math

clang/lib/Driver/ToolChains/Clang.cpp

2311

I think this is a rebase gone bad. The patch changing the strict math was revered and recommitted and I probably broke this

[arsenm]

In D69878#1805804, @arsenm wrote:

In D69878#1801508, @cameron.mcinally wrote:

This is looking pretty good to me, but I'm ignoring some of the target specific code that I'm not familiar with.

Is denormal-fp-math influenced by -Ofast? Or are there plans for that? Seems like -Ofast should imply DAZ and FTZ (if supported by target).

Yes, through the toolchain handling. I copied the logic for when crtfastmath is linked for the default mode for x86.

I think we discussed this before, but it's worth repeating. If denormal-fp-math isn't specified, we default to IEEE behavior, right? When this lands in master, there could be an unexpected performance hit for targets that aren't paying attention. E.g. I want to use denormal-fp-math to toggle whether a FSUB(-0.0,X) is converted to a FNEG(X) in SelectionDAGBuilder.

Apologies in advance if this has been discussed recently. I've been distracted with another project for the passed few months...

Yes, ieee should be the default. The dependent patches start adding the attribute by default for platforms with flushing enabled with fast math

To clarify this patch leaves the default and defers changing that to a later patch

[cameron.mcinally]

Ok, thanks for the clarifications. Looks good to me, but it would be good to have experts in OpenCL/Cuda/AMDGPU review the target specific changes.

[andrew.w.kaylor]

llvm/docs/LangRef.rst

1837

Can you document which targets do support the option? What happens if I try to use the option on a target where it is not supported?

[arsenm]

llvm/docs/LangRef.rst

1837

I'm not sure where to document this, or if/how/where to diagnose it. I don't think the high level LangRef description is the right place to discuss specific target handling.

Currently it won't error or anything. Code checking the denorm mode will see the f32 specific mode, even if the target in the end isn't really going to respect this.

One problem is this potentially does require coordination with other toolchain components. For AMDGPU, the compiler can directly tell the driver what FP mode to set on each entry point, but for x86 it requires linking in crtfastmath to set the default mode bits. If another target had a similar runtime environment requirement, I don't think we can be sure the attribute is correct or not.

[andrew.w.kaylor]

llvm/docs/LangRef.rst

1837

There is precedent for describing target-specific behavior in LangRef. It just doesn't seem useful to say that not all targets support the attribute without saying which ones do. We should also say what is expected if a target doesn't support the attribute. It seems reasonable for the function attribute to be silently ignored.

One problem is this potentially does require coordination with other toolchain components. For AMDGPU, the compiler can directly tell the driver what FP mode to set on each entry point, but for x86 it requires linking in crtfastmath to set the default mode bits.

This is a point I'm interested in. I don't like the current crtfastmath.o handling. It feels almost accidental when FTZ works as expected. My understanding is we link crtfastmath.o if we find it but if not everything just goes about its business. The Intel compiler injects code into main() to explicitly set the FTZ/DAZ control modes. That obviously has problems too, but it's at least consistent and predictable. As I understand it, crtfastmath.o sets these modes from a static initializer, but I'm not sure anything is done to determine the order of that initializer relative to others.

How does the compiler identify entry points for AMDGPU? And does it emit code to set FTZ based on the function attribute here?

[arsenm]

llvm/docs/LangRef.rst

1837

The entry points are a specific calling convention. There's no real concept of main. Each kernel has an associated blob of metadata the driver uses to set up various config registers on dispatch.

I don't think specially recognizing main in the compiler is fundamentally different than having it done in a static constructor. It's still a construct not associated with any particular function or anything.

[andrew.w.kaylor]

llvm/docs/LangRef.rst

1837

The problem with having it done in a static constructor is that you have no certainty of when it will be done relative to other static constructors. If it's in main you can at least say that it's after all the static constructors (assuming main is your entry point).

[cameron.mcinally]

llvm/docs/LangRef.rst

1837

Yes and no. The linker should honor static constructor priorities. But, yeah, there's no guarantee that this constructor will run before other priority 101 constructors.

The performance penalty for setting denormal flushing in main could be significant (think C++). Also, there's precedent for using static constructors, like GCC's crtfastmath.o.

[andrew.w.kaylor]

llvm/docs/LangRef.rst

1837

Fair enough. I don't necessarily like how icc handles this. I don't have a problem with how gcc handles it. I just really don't like how LLVM does it. If we want to take the static constructor approach we should define our own, not depend on whether or not the GNU object file happens to be around.

Static initialization doesn't help for AMDGPU, and I suppose that's likely to be the case for any offload execution model. Since this patch is moving us toward a more consistent implementation I'm wondering if we can define some general rules for how this is supposed to work. Like when the function attribute will result in injected instructions setting the control flags and when it won't.

[arsenm]

llvm/docs/LangRef.rst

1837

I think the most we can expect of this attribute as informing codegen of the expected FP denormal handling mode, and not something responsible for ensuring the mode will really be set. AMDGPU conceptually could have a separate set of attributes for setting the denormal FP mode, but since it would look identical, this gets a bonus usage for setting it for kernels. This doesn't protect you from calling functions in modules compiled with different attributes, so similar problems outside the view of the compiler still exist

[cameron.mcinally]

llvm/docs/LangRef.rst

1837

If we want to take the static constructor approach we should define our own, not depend on whether or not the GNU object file happens to be around.

That's a good idea. There's subtle differences between targets in the GNU implementation. It would be good to standardize them.

[arsenm]

Mention support in langref

[arsenm]

ping

[andrew.w.kaylor]

clang/lib/Driver/ToolChains/Clang.cpp

2311

Looks like this is still wrong. You didn't intend to change either TrappingMath flag, did you?

2532

Shouldn't this also restore DenormalFP32Math to its default value?

[arsenm]

Forgot clang parts

[arsenm]

clang/lib/Driver/ToolChains/Clang.cpp

2572

I think this should just follow along with DenormalFPMath, but I'll put this off to the later patch since this one still is slightly awkward trying to avoid changing the meaning of the absence of the option

[andrew.w.kaylor]

I don't know if there were other reviewers who haven't commented on how you addressed their concerns, but this looks good to me.

Thanks for taking the time to improve our handling of this!

[cameron.mcinally]

LGTM too. Would be good if an expert reviewed the target-specific changes, but they seem safe enough either way.

[arsenm]

a4451d88ee456304c26d552749aea6a7f5154bde

clang/include/clang/Basic/CodeGenOptions.h

public:
   std::map<std::string, std::string> DebugPrefixMap;
 
   /// The ABI to use for passing floating point arguments.
   std::string FloatABI;
 
   /// The floating-point denormal mode to use.
   llvm::DenormalMode FPDenormalMode = llvm::DenormalMode::Invalid;
 
+  /// The floating-point subnormal mode to use, for float.
+  llvm::DenormalMode FP32DenormalMode = llvm::DenormalMode::Invalid;
+
   /// The float precision limit to use, if non-empty.
   std::string LimitFloatPrecision;
 
   struct BitcodeFileToLink {
     /// The filename of the bitcode file to link in.
     std::string Filename;
     /// If true, we set attributes functions in the bitcode library according to
     /// our CodeGenOptions, much as we set attrs on functions that we generate

clang/include/clang/Basic/CodeGenOptions.def

 CODEGENOPT(NoImplicitFloat   , 1, 0) ///< Set when -mno-implicit-float is enabled.
 CODEGENOPT(NoInfsFPMath      , 1, 0) ///< Assume FP arguments, results not +-Inf.
 CODEGENOPT(NoSignedZeros     , 1, 0) ///< Allow ignoring the signedness of FP zero
 CODEGENOPT(NullPointerIsValid , 1, 0) ///< Assume Null pointer deference is defined.
 CODEGENOPT(Reassociate       , 1, 0) ///< Allow reassociation of FP math ops
 CODEGENOPT(ReciprocalMath    , 1, 0) ///< Allow FP divisions to be reassociated.
 CODEGENOPT(NoTrappingMath    , 1, 0) ///< Set when -fno-trapping-math is enabled.
 CODEGENOPT(NoNaNsFPMath      , 1, 0) ///< Assume FP arguments, results not NaN.
-CODEGENOPT(FlushDenorm       , 1, 0) ///< Allow FP denorm numbers to be flushed to zero
 CODEGENOPT(CorrectlyRoundedDivSqrt, 1, 0) ///< -cl-fp32-correctly-rounded-divide-sqrt
 
 /// When false, this attempts to generate code as if the result of an
 /// overflowing conversion matches the overflowing behavior of a target's native
 /// float-to-int conversion instructions.
 CODEGENOPT(StrictFloatCastOverflow, 1, 1)
 
 CODEGENOPT(UniformWGSize     , 1, 0) ///< -cl-uniform-work-group-size

clang/include/clang/Driver/CC1Options.td

def msign_return_address_key_EQ : Joined<["-"], "msign-return-address-key=">,
     Values<"a_key,b_key">;
 def mbranch_target_enforce : Flag<["-"], "mbranch-target-enforce">;
 def fno_dllexport_inlines : Flag<["-"], "fno-dllexport-inlines">;
 def cfguard_no_checks : Flag<["-"], "cfguard-no-checks">,
     HelpText<"Emit Windows Control Flow Guard tables only (no checks)">;
 def cfguard : Flag<["-"], "cfguard">,
     HelpText<"Emit Windows Control Flow Guard tables and checks">;
 
+def fdenormal_fp_math_f32_EQ : Joined<["-"], "fdenormal-fp-math-f32=">,
+   Group<f_Group>;
+
 //===----------------------------------------------------------------------===//
 // Dependency Output Options
 //===----------------------------------------------------------------------===//
 
 def sys_header_deps : Flag<["-"], "sys-header-deps">,
   HelpText<"Include system headers in dependency output">;
 def module_file_deps : Flag<["-"], "module-file-deps">,
   HelpText<"Include module files in dependency output">;

clang/include/clang/Driver/Options.td

 def cl_fast_relaxed_math : Flag<["-"], "cl-fast-relaxed-math">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Sets -cl-finite-math-only and -cl-unsafe-math-optimizations, and defines __FAST_RELAXED_MATH__.">;
 def cl_mad_enable : Flag<["-"], "cl-mad-enable">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Allow use of less precise MAD computations in the generated binary.">;
 def cl_no_signed_zeros : Flag<["-"], "cl-no-signed-zeros">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Allow use of less precise no signed zeros computations in the generated binary.">;
 def cl_std_EQ : Joined<["-"], "cl-std=">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL language standard to compile for.">, Values<"cl,CL,cl1.1,CL1.1,cl1.2,CL1.2,cl2.0,CL2.0,clc++,CLC++">;
-def cl_denorms_are_zero : Flag<["-"], "cl-denorms-are-zero">, Group<opencl_Group>, Flags<[CC1Option]>,
+def cl_denorms_are_zero : Flag<["-"], "cl-denorms-are-zero">, Group<opencl_Group>,
   HelpText<"OpenCL only. Allow denormals to be flushed to zero.">;
 def cl_fp32_correctly_rounded_divide_sqrt : Flag<["-"], "cl-fp32-correctly-rounded-divide-sqrt">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Specify that single precision floating-point divide and sqrt used in the program source are correctly rounded.">;
 def cl_uniform_work_group_size : Flag<["-"], "cl-uniform-work-group-size">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel">;
 def client__name : JoinedOrSeparate<["-"], "client_name">;
 def combine : Flag<["-", "--"], "combine">, Flags<[DriverOption, Unsupported]>;
 def compatibility__version : JoinedOrSeparate<["-"], "compatibility_version">;

 def no_cuda_noopt_device_debug : Flag<["--"], "no-cuda-noopt-device-debug">;
 def cuda_path_EQ : Joined<["--"], "cuda-path=">, Group<i_Group>,
   HelpText<"CUDA installation path">;
 def cuda_path_ignore_env : Flag<["--"], "cuda-path-ignore-env">, Group<i_Group>,
   HelpText<"Ignore environment variables to detect CUDA installation">;
 def ptxas_path_EQ : Joined<["--"], "ptxas-path=">, Group<i_Group>,
   HelpText<"Path to ptxas (used for compiling CUDA code)">;
 def fcuda_flush_denormals_to_zero : Flag<["-"], "fcuda-flush-denormals-to-zero">,
-  Flags<[CC1Option]>, HelpText<"Flush denormal floating point values to zero in CUDA device mode.">;
+  HelpText<"Flush denormal floating point values to zero in CUDA device mode.">;
 def fno_cuda_flush_denormals_to_zero : Flag<["-"], "fno-cuda-flush-denormals-to-zero">;
 def fcuda_approx_transcendentals : Flag<["-"], "fcuda-approx-transcendentals">,
   Flags<[CC1Option]>, HelpText<"Use approximate transcendental functions">;
 def fno_cuda_approx_transcendentals : Flag<["-"], "fno-cuda-approx-transcendentals">;
 def fgpu_rdc : Flag<["-"], "fgpu-rdc">, Flags<[CC1Option]>,
   HelpText<"Generate relocatable device code, also known as separate compilation mode.">;
 def fno_gpu_rdc : Flag<["-"], "fno-gpu-rdc">;
 def : Flag<["-"], "fcuda-rdc">, Alias<fgpu_rdc>;

clang/include/clang/Driver/ToolChain.h

 
 #include "clang/Basic/DebugInfoOptions.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/Sanitizers.h"
 #include "clang/Driver/Action.h"
 #include "clang/Driver/Multilib.h"
 #include "clang/Driver/Types.h"
+#include "llvm/ADT/APFloat.h"
 #include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/FloatingPointMode.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/MC/MCTargetOptions.h"
 #include "llvm/Option/Option.h"
 #include "llvm/Support/VersionTuple.h"
 #include "llvm/Target/TargetOptions.h"
 #include <cassert>
public:
   /// Return sanitizers which are enabled by default.
   virtual SanitizerMask getDefaultSanitizers() const {
     return SanitizerMask();
   }
 
   /// Returns true when it's possible to split LTO unit to use whole
   /// program devirtualization and CFI santiizers.
   virtual bool canSplitThinLTOUnit() const { return true; }
+
+  /// Returns the output denormal handling type in the default floating point
+  /// environment for the given \p FPType if given. Otherwise, the default
+  /// assumed mode for any floating point type.
+  virtual llvm::DenormalMode getDefaultDenormalModeForType(
+      const llvm::opt::ArgList &DriverArgs,
+      Action::OffloadKind DeviceOffloadKind,
+      const llvm::fltSemantics *FPType = nullptr) const {
+    // FIXME: This should be IEEE when default handling is fixed.
+    return llvm::DenormalMode::Invalid;
+  }
 };
 
 /// Set a ToolChain's effective triple. Reset it when the registration object
 /// is destroyed.
 class RegisterEffectiveTriple {
   const ToolChain &TC;
 
 public:

clang/lib/Basic/Targets/AMDGPU.cpp

void AMDGPUTargetInfo::adjustTargetOptions(const CodeGenOptions &CGOpts,
   for (auto &I : TargetOpts.FeaturesAsWritten) {
     if (I == "+fp32-denormals" || I == "-fp32-denormals")
       hasFP32Denormals = true;
     if (I == "+fp64-fp16-denormals" || I == "-fp64-fp16-denormals")
       hasFP64Denormals = true;
   }
   if (!hasFP32Denormals)
     TargetOpts.Features.push_back(
-      (Twine(hasFastFMAF() && hasFullRateDenormalsF32() && !CGOpts.FlushDenorm
+      (Twine(hasFastFMAF() && hasFullRateDenormalsF32() &&
+             CGOpts.FP32DenormalMode == llvm::DenormalMode::IEEE
              ? '+' : '-') + Twine("fp32-denormals"))
             .str());
   // Always do not flush fp64 or fp16 denorms.
   if (!hasFP64Denormals && hasFP64())
     TargetOpts.Features.push_back("+fp64-fp16-denormals");
 }
 
 void AMDGPUTargetInfo::fillValidCPUList(

clang/lib/CodeGen/CGCall.cpp

if (AttrOnCallSite) {
     }
     FuncAttrs.addAttribute("frame-pointer", FpKind);
 
     FuncAttrs.addAttribute("less-precise-fpmad",
                            llvm::toStringRef(CodeGenOpts.LessPreciseFPMAD));
 
     if (CodeGenOpts.NullPointerIsValid)
       FuncAttrs.addAttribute("null-pointer-is-valid", "true");
+
+    // TODO: Omit attribute when the default is IEEE.
     if (CodeGenOpts.FPDenormalMode != llvm::DenormalMode::Invalid)
       FuncAttrs.addAttribute("denormal-fp-math",
                              llvm::denormalModeName(CodeGenOpts.FPDenormalMode));
 
+    if (CodeGenOpts.FP32DenormalMode != llvm::DenormalMode::Invalid)
+      FuncAttrs.addAttribute(
+          "denormal-fp-math-f32",
+          llvm::denormalModeName(CodeGenOpts.FP32DenormalMode));
+
     FuncAttrs.addAttribute("no-trapping-math",
                            llvm::toStringRef(CodeGenOpts.NoTrappingMath));
 
     // Strict (compliant) code is the default, so only add this attribute to
     // indicate that we are trying to workaround a problem case.
     if (!CodeGenOpts.StrictFloatCastOverflow)
       FuncAttrs.addAttribute("strict-float-cast-overflow", "false");
 
if (AttrOnCallSite) {
     FuncAttrs.addAttribute("stack-protector-buffer-size",
                            llvm::utostr(CodeGenOpts.SSPBufferSize));
     FuncAttrs.addAttribute("no-signed-zeros-fp-math",
                            llvm::toStringRef(CodeGenOpts.NoSignedZeros));
     FuncAttrs.addAttribute(
         "correctly-rounded-divide-sqrt-fp-math",
         llvm::toStringRef(CodeGenOpts.CorrectlyRoundedDivSqrt));
 
-    if (getLangOpts().OpenCL)
-      FuncAttrs.addAttribute("denorms-are-zero",
-                             llvm::toStringRef(CodeGenOpts.FlushDenorm));
-
     // TODO: Reciprocal estimate codegen options should apply to instructions?
     const std::vector<std::string> &Recips = CodeGenOpts.Reciprocals;
     if (!Recips.empty())
       FuncAttrs.addAttribute("reciprocal-estimates",
                              llvm::join(Recips, ","));
 
     if (!CodeGenOpts.PreferVectorWidth.empty() &&
         CodeGenOpts.PreferVectorWidth != "none")
if (getLangOpts().assumeFunctionsAreConvergent()) {
     // applied around them).  LLVM will remove this attribute where it safely
     // can.
     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
   }
 
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
     // Exceptions aren't supported in CUDA device code.
     FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
-
-    // Respect -fcuda-flush-denormals-to-zero.
-    if (CodeGenOpts.FlushDenorm)
-      FuncAttrs.addAttribute("nvptx-f32ftz", "true");
   }
 
   for (StringRef Attr : CodeGenOpts.DefaultFunctionAttrs) {
     StringRef Var, Value;
     std::tie(Var, Value) = Attr.split('=');
     FuncAttrs.addAttribute(Var, Value);
   }
 }

clang/lib/CodeGen/CodeGenModule.cpp

getModule().addModuleFlag(llvm::Module::Override, "cf-protection-branch",
                               1);
   }
 
   if (LangOpts.CUDAIsDevice && getTriple().isNVPTX()) {
     // Indicate whether __nvvm_reflect should be configured to flush denormal
     // floating point values to 0.  (This corresponds to its "__CUDA_FTZ"
     // property.)
     getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-ftz",
-                              CodeGenOpts.FlushDenorm ? 1 : 0);
+                              CodeGenOpts.FP32DenormalMode !=
+                                  llvm::DenormalMode::IEEE);
   }
 
   // Emit OpenCL specific module metadata: OpenCL/SPIR version.
   if (LangOpts.OpenCL) {
     EmitOpenCLMetadata();
     // Emit SPIR version.
     if (getTriple().isSPIR()) {
       // SPIR v2.0 s2.12 - The SPIR version used by the module is stored in the

clang/lib/Driver/ToolChains/AMDGPU.h

public:
 
   llvm::opt::DerivedArgList *
   TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch,
                 Action::OffloadKind DeviceOffloadKind) const override;
 
   void addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
                              llvm::opt::ArgStringList &CC1Args,
                              Action::OffloadKind DeviceOffloadKind) const override;
+
+  llvm::DenormalMode getDefaultDenormalModeForType(
+      const llvm::opt::ArgList &DriverArgs,
+      Action::OffloadKind DeviceOffloadKind,
+      const llvm::fltSemantics *FPType = nullptr) const override;
 };
 
 } // end namespace toolchains
 } // end namespace driver
 } // end namespace clang
 
 #endif // LLVM_CLANG_LIB_DRIVER_TOOLCHAINS_AMDGPU_H

clang/lib/Driver/ToolChains/AMDGPU.cpp

 //===--- AMDGPU.cpp - AMDGPU ToolChain Implementations ----------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "AMDGPU.h"
 #include "CommonArgs.h"
 #include "InputInfo.h"
 #include "clang/Driver/Compilation.h"
 #include "clang/Driver/DriverDiagnostic.h"
 #include "llvm/Option/ArgList.h"
+#include "llvm/Support/TargetParser.h"
 
 using namespace clang::driver;
 using namespace clang::driver::tools;
 using namespace clang::driver::toolchains;
 using namespace clang;
 using namespace llvm::opt;
 
 void amdgpu::Linker::ConstructJob(Compilation &C, const JobAction &JA,
if (!Args.hasArg(options::OPT_O, options::OPT_O0, options::OPT_O4,
                      options::OPT_Ofast))
       DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_O),
                         getOptionDefault(options::OPT_O));
   }
 
   return DAL;
 }
 
+llvm::DenormalMode AMDGPUToolChain::getDefaultDenormalModeForType(
+    const llvm::opt::ArgList &DriverArgs, Action::OffloadKind DeviceOffloadKind,
+    const llvm::fltSemantics *FPType) const {
+  // Denormals should always be enabled for f16 and f64.
+  if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
+    return llvm::DenormalMode::IEEE;
+
+  if (DeviceOffloadKind == Action::OFK_Cuda) {
+    if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
+        DriverArgs.hasFlag(options::OPT_fcuda_flush_denormals_to_zero,
+                           options::OPT_fno_cuda_flush_denormals_to_zero,
+                           false))
+      return llvm::DenormalMode::PreserveSign;
+  }
+
+  const StringRef GpuArch = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
+  auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
+
+  // Default to enabling f32 denormals by default on subtargets where fma is
+  // fast with denormals
+
+  const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind);
+  const bool DefaultDenormsAreZeroForTarget =
+    (ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) &&
+    (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32);
+
+  // TODO: There are way too many flags that change this. Do we need to check
+  // them all?
+  bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) ||
+             !DefaultDenormsAreZeroForTarget;
+  // Outputs are flushed to zero, preserving sign
+  return DAZ ? llvm::DenormalMode::PreserveSign : llvm::DenormalMode::IEEE;
+}
+
 void AMDGPUToolChain::addClangTargetOptions(
     const llvm::opt::ArgList &DriverArgs,
     llvm::opt::ArgStringList &CC1Args,
     Action::OffloadKind DeviceOffloadingKind) const {
   // Default to "hidden" visibility, as object level linking will not be
   // supported for the foreseeable future.
   if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
                          options::OPT_fvisibility_ms_compat)) {
     CC1Args.push_back("-fvisibility");
     CC1Args.push_back("hidden");
     CC1Args.push_back("-fapply-global-visibility-to-externs");
   }
 }

clang/lib/Driver/ToolChains/Clang.cpp

static void CollectArgsForIntegratedAssembler(Compilation &C,
   // forward -fembed-bitcode to assmebler
   if (C.getDriver().embedBitcodeEnabled() ||
       C.getDriver().embedBitcodeMarkerOnly())
     Args.AddLastArg(CmdArgs, options::OPT_fembed_bitcode_EQ);
 }
 
 static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
                                        bool OFastEnabled, const ArgList &Args,
-                                       ArgStringList &CmdArgs) {
+                                       ArgStringList &CmdArgs,
+                                       Action::OffloadKind DeviceOffloadKind) {
   // Handle various floating point optimization flags, mapping them to the
   // appropriate LLVM code generation flags. This is complicated by several
   // "umbrella" flags, so we do this by stepping through the flags incrementally
   // adjusting what we think is enabled/disabled, then at the end setting the
   // LLVM flags based on the final state.
   bool HonorINFs = true;
   bool HonorNaNs = true;
   // -fmath-errno is the default on some platforms, e.g. BSD-derived OSes.
   bool MathErrno = TC.IsMathErrnoDefault();
   bool AssociativeMath = false;
   bool ReciprocalMath = false;
   bool SignedZeros = true;
-  bool TrappingMath = false; // Implemented via -ffp-exception-behavior
-  bool TrappingMathPresent = false; // Is trapping-math in args, and not
+  bool TrappingMath = true;
                                     // overriden by ffp-exception-behavior?
   bool RoundingFPMath = false;
   bool RoundingMathPresent = false; // Is rounding-math in args?
   // -ffp-model values: strict, fast, precise
   StringRef FPModel = "";
   // -ffp-exception-behavior options: strict, maytrap, ignore
   StringRef FPExceptionBehavior = "";
-  StringRef DenormalFPMath = "";
+  const llvm::DenormalMode DefaultDenormalFPMath =
+      TC.getDefaultDenormalModeForType(Args, DeviceOffloadKind);
+  llvm::DenormalMode DenormalFPMath = DefaultDenormalFPMath;
   StringRef FPContract = "";
   bool StrictFPModel = false;
 
+  llvm::DenormalMode DenormalFP32Math = TC.getDefaultDenormalModeForType(
+      Args, DeviceOffloadKind, &llvm::APFloat::IEEEsingle());
+
   if (const Arg *A = Args.getLastArg(options::OPT_flimited_precision_EQ)) {
     CmdArgs.push_back("-mlimit-float-precision");
     CmdArgs.push_back(A->getValue());
   }
 
   for (const Arg *A : Args) {
     auto optID = A->getOption().getID();
     bool PreciseFPModel = false;
case options::OPT_frounding_math:
       break;
 
     case options::OPT_fno_rounding_math:
       RoundingFPMath = false;
       RoundingMathPresent = false;
       break;
 
     case options::OPT_fdenormal_fp_math_EQ:
-      DenormalFPMath = A->getValue();
+      DenormalFPMath = llvm::parseDenormalFPAttribute(A->getValue());
+      if (DenormalFPMath == llvm::DenormalMode::Invalid) {
+        D.Diag(diag::err_drv_invalid_value)
+            << A->getAsString(Args) << A->getValue();
+      }
+      break;
+
+    case options::OPT_fdenormal_fp_math_f32_EQ:
+      DenormalFP32Math = llvm::parseDenormalFPAttribute(A->getValue());
+      if (DenormalFP32Math == llvm::DenormalMode::Invalid) {
+        D.Diag(diag::err_drv_invalid_value)
+            << A->getAsString(Args) << A->getValue();
+      }
       break;
 
     // Validate and pass through -ffp-contract option.
     case options::OPT_ffp_contract: {
       StringRef Val = A->getValue();
       if (PreciseFPModel) {
         // -ffp-model=precise enables ffp-contract=fast as a side effect
         // the FPContract value has already been set to a string literal
case options::OPT_funsafe_math_optimizations:
       break;
     case options::OPT_fno_unsafe_math_optimizations:
       AssociativeMath = false;
       ReciprocalMath = false;
       SignedZeros = true;
       TrappingMath = true;
       FPExceptionBehavior = "strict";
       // -fno_unsafe_math_optimizations restores default denormal handling
-      DenormalFPMath = "";
+      DenormalFPMath = DefaultDenormalFPMath;
       break;
 
     case options::OPT_Ofast:
       // If -Ofast is the optimization level, then -ffast-math should be enabled
       if (!OFastEnabled)
         continue;
       LLVM_FALLTHROUGH;
     case options::OPT_ffast_math:
case options::OPT_fno_fast_math:
       // toolchain default (which may be false).
       MathErrno = TC.IsMathErrnoDefault();
       AssociativeMath = false;
       ReciprocalMath = false;
       SignedZeros = true;
       TrappingMath = false;
       RoundingFPMath = false;
       // -fno_fast_math restores default denormal and fpcontract handling
-      DenormalFPMath = "";
+      DenormalFPMath = DefaultDenormalFPMath;
       FPContract = "";
       break;
     }
     if (StrictFPModel) {
       // If -ffp-model=strict has been specified on command line but
       // subsequent options conflict then emit warning diagnostic.
       if (HonorINFs && HonorNaNs &&
         !AssociativeMath && !ReciprocalMath &&
         SignedZeros && TrappingMath && RoundingFPMath &&
         DenormalFPMath.empty() && FPContract.empty())
         // OK: Current Arg doesn't conflict with -ffp-model=strict
         ;
       else {
         StrictFPModel = false;
         FPModel = "";
static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
 
   if (TrappingMath) {
     // FP Exception Behavior is also set to strict
     assert(FPExceptionBehavior.equals("strict"));
     CmdArgs.push_back("-ftrapping-math");
   } else if (TrappingMathPresent)
     CmdArgs.push_back("-fno-trapping-math");
 
-  if (!DenormalFPMath.empty())
-    CmdArgs.push_back(
-        Args.MakeArgString("-fdenormal-fp-math=" + DenormalFPMath));
+  // TODO: Omit flag for the default IEEE instead
+  if (DenormalFPMath != llvm::DenormalMode::Invalid) {
+    CmdArgs.push_back(Args.MakeArgString(
+        "-fdenormal-fp-math=" + llvm::subnormalModeName(DenormalFPMath)));
+  }
+
+  if (DenormalFP32Math != llvm::DenormalMode::Invalid) {
+    CmdArgs.push_back(Args.MakeArgString(
+        "-fdenormal-fp-math-f32=" + llvm::subnormalModeName(DenormalFP32Math)));
+  }
 
   if (!FPContract.empty())
     CmdArgs.push_back(Args.MakeArgString("-ffp-contract=" + FPContract));
 
   if (!RoundingFPMath)
     CmdArgs.push_back(Args.MakeArgString("-fno-rounding-math"));
 
   if (RoundingFPMath && RoundingMathPresent)
if (!TrivialAutoVarInit.empty()) {
     if (TrivialAutoVarInit == "zero" && !Args.hasArg(options::OPT_enable_trivial_var_init_zero))
       D.Diag(diag::err_drv_trivial_auto_var_init_zero_disabled);
     CmdArgs.push_back(
         Args.MakeArgString("-ftrivial-auto-var-init=" + TrivialAutoVarInit));
   }
 }
 
 static void RenderOpenCLOptions(const ArgList &Args, ArgStringList &CmdArgs) {
+  // cl-denorms-are-zero is not forwarded. It is translated into a generic flag
+  // for denormal flushing handling based on the target.
   const unsigned ForwardedArguments[] = {
       options::OPT_cl_opt_disable,
       options::OPT_cl_strict_aliasing,
       options::OPT_cl_single_precision_constant,
       options::OPT_cl_finite_math_only,
       options::OPT_cl_kernel_arg_info,
       options::OPT_cl_unsafe_math_optimizations,
       options::OPT_cl_fast_relaxed_math,
       options::OPT_cl_mad_enable,
       options::OPT_cl_no_signed_zeros,
-      options::OPT_cl_denorms_are_zero,
       options::OPT_cl_fp32_correctly_rounded_divide_sqrt,
       options::OPT_cl_uniform_work_group_size
   };
 
   if (Arg *A = Args.getLastArg(options::OPT_cl_std_EQ)) {
     std::string CLStdStr = std::string("-cl-std=") + A->getValue();
     CmdArgs.push_back(Args.MakeArgString(CLStdStr));
   }
for (const auto &A : Args)
         D.Diag(diag::err_drv_unsupported_embed_bitcode) << A->getSpelling();
 
     // Render the CodeGen options that need to be passed.
     if (!Args.hasFlag(options::OPT_foptimize_sibling_calls,
                       options::OPT_fno_optimize_sibling_calls))
       CmdArgs.push_back("-mdisable-tail-calls");
 
     RenderFloatingPointOptions(TC, D, isOptimizationLevelFast(Args), Args,
-                               CmdArgs);
+                               CmdArgs, JA.getOffloadingDeviceKind());
 
     // Render ABI arguments
     switch (TC.getArch()) {
     default: break;
     case llvm::Triple::arm:
     case llvm::Triple::armeb:
     case llvm::Triple::thumbeb:
       RenderARMABI(Triple, Args, CmdArgs);
#endif
 
   Args.AddLastArg(CmdArgs, options::OPT_ffine_grained_bitfield_accesses,
                   options::OPT_fno_fine_grained_bitfield_accesses);
 
   // Handle segmented stacks.
   if (Args.hasArg(options::OPT_fsplit_stack))
     CmdArgs.push_back("-split-stacks");
 
-  RenderFloatingPointOptions(TC, D, OFastEnabled, Args, CmdArgs);
+  RenderFloatingPointOptions(TC, D, OFastEnabled, Args, CmdArgs,
+                             JA.getOffloadingDeviceKind());
 
   if (Arg *A = Args.getLastArg(options::OPT_LongDouble_Group)) {
     if (TC.getArch() == llvm::Triple::x86 ||
         TC.getArch() == llvm::Triple::x86_64)
       A->render(Args, CmdArgs);
     else if ((TC.getArch() == llvm::Triple::ppc || TC.getTriple().isPPC64()) &&
              (A->getOption().getID() != options::OPT_mlong_double_80))
       A->render(Args, CmdArgs);

clang/lib/Driver/ToolChains/Cuda.h

public:
 
   llvm::opt::DerivedArgList *
   TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch,
                 Action::OffloadKind DeviceOffloadKind) const override;
   void addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
                              llvm::opt::ArgStringList &CC1Args,
                              Action::OffloadKind DeviceOffloadKind) const override;
 
+  llvm::DenormalMode getDefaultDenormalModeForType(
+      const llvm::opt::ArgList &DriverArgs,
+      Action::OffloadKind DeviceOffloadKind,
+      const llvm::fltSemantics *FPType = nullptr) const override;
+
   // Never try to use the integrated assembler with CUDA; always fork out to
   // ptxas.
   bool useIntegratedAs() const override { return false; }
   bool isCrossCompiling() const override { return true; }
   bool isPICDefault() const override { return false; }
   bool isPIEDefault() const override { return false; }
   bool isPICDefaultForced() const override { return false; }
   bool SupportsProfiling() const override { return false; }

clang/lib/Driver/ToolChains/Cuda.cpp

 #include "clang/Driver/Driver.h"
 #include "clang/Driver/DriverDiagnostic.h"
 #include "clang/Driver/Options.h"
 #include "llvm/Option/ArgList.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/Path.h"
 #include "llvm/Support/Process.h"
 #include "llvm/Support/Program.h"
+#include "llvm/Support/TargetParser.h"
 #include "llvm/Support/VirtualFileSystem.h"
 #include <system_error>
 
 using namespace clang::driver;
 using namespace clang::driver::toolchains;
 using namespace clang::driver::tools;
 using namespace clang;
 using namespace llvm::opt;
void CudaToolChain::addClangTargetOptions(
   assert(!GpuArch.empty() && "Must have an explicit GPU arch.");
   assert((DeviceOffloadingKind == Action::OFK_OpenMP ||
           DeviceOffloadingKind == Action::OFK_Cuda) &&
          "Only OpenMP or CUDA offloading kinds are supported for NVIDIA GPUs.");
 
   if (DeviceOffloadingKind == Action::OFK_Cuda) {
     CC1Args.push_back("-fcuda-is-device");
 
-    if (DriverArgs.hasFlag(options::OPT_fcuda_flush_denormals_to_zero,
-                           options::OPT_fno_cuda_flush_denormals_to_zero, false))
-      CC1Args.push_back("-fcuda-flush-denormals-to-zero");
-
     if (DriverArgs.hasFlag(options::OPT_fcuda_approx_transcendentals,
                            options::OPT_fno_cuda_approx_transcendentals, false))
       CC1Args.push_back("-fcuda-approx-transcendentals");
 
     if (DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
                            false))
       CC1Args.push_back("-fgpu-rdc");
   }
for (StringRef LibraryPath : LibraryPaths) {
       }
     }
     if (!FoundBCLibrary)
       getDriver().Diag(diag::warn_drv_omp_offload_target_missingbcruntime)
           << LibOmpTargetName;
   }
 }
 
+llvm::DenormalMode CudaToolChain::getDefaultDenormalModeForType(
+    const llvm::opt::ArgList &DriverArgs, Action::OffloadKind DeviceOffloadKind,
+    const llvm::fltSemantics *FPType) const {
+  if (DeviceOffloadKind == Action::OFK_Cuda) {
+    if (FPType && FPType == &llvm::APFloat::IEEEsingle() &&
+        DriverArgs.hasFlag(options::OPT_fcuda_flush_denormals_to_zero,
+                           options::OPT_fno_cuda_flush_denormals_to_zero,
+                           false))
+      return llvm::DenormalMode::PreserveSign;
+  }
+
+  assert(DeviceOffloadKind != Action::OFK_Host);
+  return llvm::DenormalMode::IEEE;
+}
+
 bool CudaToolChain::supportsDebugInfoOption(const llvm::opt::Arg *A) const {
   const Option &O = A->getOption();
   return (O.matches(options::OPT_gN_Group) &&
           !O.matches(options::OPT_gmodules)) ||
          O.matches(options::OPT_g_Flag) ||
          O.matches(options::OPT_ggdbN_Group) || O.matches(options::OPT_ggdb) ||
          O.matches(options::OPT_gdwarf) || O.matches(options::OPT_gdwarf_2) ||
          O.matches(options::OPT_gdwarf_3) || O.matches(options::OPT_gdwarf_4) ||

clang/lib/Driver/ToolChains/HIP.cpp

void HIPToolChain::addClangTargetOptions(
   (void) GpuArch;
   assert(DeviceOffloadingKind == Action::OFK_HIP &&
          "Only HIP offloading kinds are supported for GPUs.");
 
   CC1Args.push_back("-target-cpu");
   CC1Args.push_back(DriverArgs.MakeArgStringRef(GpuArch));
   CC1Args.push_back("-fcuda-is-device");
 
-  if (DriverArgs.hasFlag(options::OPT_fcuda_flush_denormals_to_zero,
-                         options::OPT_fno_cuda_flush_denormals_to_zero, false))
-    CC1Args.push_back("-fcuda-flush-denormals-to-zero");
-
   if (DriverArgs.hasFlag(options::OPT_fcuda_approx_transcendentals,
                          options::OPT_fno_cuda_approx_transcendentals, false))
     CC1Args.push_back("-fcuda-approx-transcendentals");
 
   if (DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
                          false))
     CC1Args.push_back("-fgpu-rdc");
 

clang/lib/Frontend/CompilerInvocation.cpp

Opts.NoNaNsFPMath = (Args.hasArg(OPT_menable_no_nans) ||
                        Args.hasArg(OPT_cl_unsafe_math_optimizations) ||
                        Args.hasArg(OPT_cl_finite_math_only) ||
                        Args.hasArg(OPT_cl_fast_relaxed_math));
   Opts.NoSignedZeros = (Args.hasArg(OPT_fno_signed_zeros) ||
                         Args.hasArg(OPT_cl_no_signed_zeros) ||
                         Args.hasArg(OPT_cl_unsafe_math_optimizations) ||
                         Args.hasArg(OPT_cl_fast_relaxed_math));
   Opts.Reassociate = Args.hasArg(OPT_mreassociate);
-  Opts.FlushDenorm = Args.hasArg(OPT_cl_denorms_are_zero) ||
-                     (Args.hasArg(OPT_fcuda_is_device) &&
-                      Args.hasArg(OPT_fcuda_flush_denormals_to_zero));
   Opts.CorrectlyRoundedDivSqrt =
       Args.hasArg(OPT_cl_fp32_correctly_rounded_divide_sqrt);
   Opts.UniformWGSize =
       Args.hasArg(OPT_cl_uniform_work_group_size);
   Opts.Reciprocals = Args.getAllArgValues(OPT_mrecip_EQ);
   Opts.ReciprocalMath = Args.hasArg(OPT_freciprocal_math);
   Opts.NoTrappingMath = Args.hasArg(OPT_fno_trapping_math);
   Opts.StrictFloatCastOverflow =
static bool ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, InputKind IK,
 
   if (Arg *A = Args.getLastArg(OPT_fdenormal_fp_math_EQ)) {
     StringRef Val = A->getValue();
     Opts.FPDenormalMode = llvm::parseDenormalFPAttribute(Val);
     if (Opts.FPDenormalMode == llvm::DenormalMode::Invalid)
       Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Val;
   }
 
+  if (Arg *A = Args.getLastArg(OPT_fdenormal_fp_math_f32_EQ)) {
+    StringRef Val = A->getValue();
+    Opts.FP32DenormalMode = llvm::parseDenormalFPAttribute(Val);
+    if (Opts.FP32DenormalMode == llvm::DenormalMode::Invalid)
+      Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Val;
+  }
+
   if (Arg *A = Args.getLastArg(OPT_fpcc_struct_return, OPT_freg_struct_return)) {
     if (A->getOption().matches(OPT_fpcc_struct_return)) {
       Opts.setStructReturnConvention(CodeGenOptions::SRCK_OnStack);
     } else {
       assert(A->getOption().matches(OPT_freg_struct_return));
       Opts.setStructReturnConvention(CodeGenOptions::SRCK_InRegs);
     }
   }

clang/test/CodeGenCUDA/flush-denormals.cu

 // RUN: %clang_cc1 -fcuda-is-device \
 // RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
-// RUN:   FileCheck %s -check-prefix CHECK -check-prefix NOFTZ
-// RUN: %clang_cc1 -fcuda-is-device -fcuda-flush-denormals-to-zero \
+// RUN:   FileCheck -check-prefix=DEFAULT %s
+
+// RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=ieee \
+// RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
+// RUN:   FileCheck -check-prefix=NOFTZ %s
+
+// RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=preserve-sign \
 // RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
-// RUN:   FileCheck %s -check-prefix CHECK -check-prefix FTZ
+// RUN:   FileCheck -check-prefix=FTZ %s
 
+// FIXME: Unspecified should default to ieee
 // RUN: %clang_cc1 -fcuda-is-device -x hip \
 // RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
-// RUN:   FileCheck %s -check-prefix CHECK -check-prefix AMDNOFTZ
-// RUN: %clang_cc1 -fcuda-is-device -x hip -fcuda-flush-denormals-to-zero \
+// RUN:   FileCheck -check-prefix=AMDFTZ %s
+
+// RUN: %clang_cc1 -fcuda-is-device -x hip \
+// RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -fdenormal-fp-math-f32=ieee -emit-llvm -o - %s | \
+// RUN:   FileCheck -check-prefix=AMDNOFTZ %s
+
+// RUN: %clang_cc1 -fcuda-is-device -x hip -fdenormal-fp-math-f32=preserve-sign \
 // RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
-// RUN:   FileCheck %s -check-prefix CHECK -check-prefix AMDFTZ
+// RUN:   FileCheck -check-prefix=AMDFTZ %s
 
 #include "Inputs/cuda.h"
 
-// Checks that device function calls get emitted with the "ntpvx-f32ftz"
-// attribute set to "true" when we compile CUDA device code with
-// -fcuda-flush-denormals-to-zero.  Further, check that we reflect the presence
-// or absence of -fcuda-flush-denormals-to-zero in a module flag.
+// Checks that device function calls get emitted with the "denormal-fp-math-f32"
+// attribute set when we compile CUDA device code with
+// -fdenormal-fp-math-f32. Further, check that we reflect the presence or
+// absence of -fcuda-flush-denormals-to-zero in a module flag.
 
 // AMDGCN targets always have +fp64-fp16-denormals.
 // AMDGCN targets without fast FMAF (e.g. gfx803) always have +fp32-denormals.
 // For AMDGCN target with fast FMAF (e.g. gfx900), it has +fp32-denormals
 // by default and -fp32-denormals when there is option
 // -fcuda-flush-denormals-to-zero.
 
 // CHECK-LABEL: define void @foo() #0
 extern "C" __device__ void foo() {}
 
-// FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true"
-// NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz"
+// FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign"
+// NOFTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="ieee"
+
+
+// FIXME: This should be removed
+// DEFAULT-NOT: "denormal-fp-math-f32"
+
 // AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals
 // AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals
 
 // FTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
 // FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
 
 // NOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
 // NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}

clang/test/CodeGenCUDA/propagate-metadata.cu

 // overrides the flag on the bitcode library.
 
 // Build the bitcode library.  This is not built in CUDA mode, otherwise it
 // might have incompatible attributes.  This mirrors how libdevice is built.
 // RUN: %clang_cc1 -x c++ -fconvergent-functions -emit-llvm-bc -ftrapping-math -DLIB \
 // RUN:   %s -o %t.bc -triple nvptx-unknown-unknown
 
 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc -o - \
-// RUN:   -fno-trapping-math -fcuda-is-device -triple nvptx-unknown-unknown \
+// RUN:   -fno-trapping-math -fcuda-is-device -fdenormal-fp-math-f32=ieee -triple nvptx-unknown-unknown \
 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ --check-prefix=NOFAST
 
 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
-// RUN:   -fno-trapping-math -fcuda-flush-denormals-to-zero -o - \
+// RUN:   -fno-trapping-math -fdenormal-fp-math-f32=preserve-sign -o - \
 // RUN:   -fcuda-is-device -triple nvptx-unknown-unknown \
 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ \
 // RUN:   --check-prefix=NOFAST
 
 // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
-// RUN:   -fno-trapping-math -fcuda-flush-denormals-to-zero -o - \
+// RUN:   -fno-trapping-math -fdenormal-fp-math-f32=preserve-sign -o - \
 // RUN:   -fcuda-is-device -menable-unsafe-fp-math -triple nvptx-unknown-unknown \
 // RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST
 
 // Wrap everything in extern "C" so we don't have to worry about name mangling
 // in the IR.
 extern "C" {
 #ifdef LIB
 

 
 #endif
 }
 
 // The kernel and lib function should have the same attributes.
 // CHECK: define void @kernel() [[attr:#[0-9]+]]
 // CHECK: define internal void @lib_fn() [[attr]]
 
+// FIXME: These -NOT checks do not work as intended and do not check on the same
+// line.
+
 // Check the attribute list.
 // CHECK: attributes [[attr]] = {
+
 // CHECK-SAME: convergent
 // CHECK-SAME: "no-trapping-math"="true"
 
-// FTZ-SAME: "nvptx-f32ftz"="true"
-// NOFTZ-NOT: "nvptx-f32ftz"="true"
+// FTZ-NOT: "denormal-fp-math"
+
+// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign"
+// NOFTZ-SAME: "denormal-fp-math-f32"="ieee"
+
+// CHECK-SAME: "no-trapping-math"="true"
 
 // FAST-SAME: "unsafe-fp-math"="true"
 // NOFAST-NOT: "unsafe-fp-math"="true"

clang/test/CodeGenOpenCL/amdgpu-features.cl

 // REQUIRES: amdgpu-registered-target
 
 // Check that appropriate features are defined for every supported AMDGPU
 // "-target" and "-mcpu" options.
 
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx904 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX904 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx908 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX908 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1010 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1011 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1012 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx801 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX801 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx700 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX700 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s
 
-// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime"
+// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime,-fp32-denormals"
 // GFX700: "target-features"="+ci-insts,+flat-address-space,+fp64-fp16-denormals,-fp32-denormals"
 // GFX600: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
 // GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
 
 kernel void test() {}

clang/test/CodeGenOpenCL/denorms-are-zero.cl

-// RUN: %clang_cc1 -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -emit-llvm -cl-denorms-are-zero -o - %s | FileCheck -check-prefix=DENORM-ZERO %s
-
-// Slow FMAF and slow f32 denormals
-// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa -target-cpu pitcairn %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
-// RUN: %clang_cc1 -emit-llvm -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu pitcairn %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH-OPT %s
-
-// Fast FMAF, but slow f32 denormals
-// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa -target-cpu tahiti %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
-// RUN: %clang_cc1 -emit-llvm -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu tahiti %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH-OPT %s
-
-// Fast F32 denormals, but slow FMAF
-// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa -target-cpu fiji %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
-// RUN: %clang_cc1 -emit-llvm -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu fiji %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH-OPT %s
-
-// Fast F32 denormals and fast FMAF
-// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa -target-cpu gfx900 %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-DENORM %s
-// RUN: %clang_cc1 -emit-llvm -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu gfx900 %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH-OPT %s
-
-// RUN: %clang_cc1 -emit-llvm -target-feature +fp32-denormals -target-feature -fp64-fp16-denormals -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu fiji %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FEATURE %s
-// RUN: %clang_cc1 -emit-llvm -target-feature +fp32-denormals -target-feature -fp64-fp16-denormals -cl-denorms-are-zero -o - -triple amdgcn--amdhsa -target-cpu pitcairn %s | FileCheck -check-prefixes=AMDGCN,AMDGCN-FEATURE %s
-
-
-
-// For all targets 'denorms-are-zero' attribute is set to 'true'
-// if '-cl-denorms-are-zero' was specified and  to 'false' otherwise.
-
-// CHECK-LABEL: define {{(dso_local )?}}void @f()
-// CHECK: attributes #{{[0-9]*}} = {{{[^}]*}} "denorms-are-zero"="false"
-//
-// DENORM-ZERO-LABEL: define {{(dso_local )?}}void @f()
-// DENORM-ZERO: attributes #{{[0-9]*}} = {{{[^}]*}} "denorms-are-zero"="true"
-
-// For amdgcn target cpu fiji, fp32 should be flushed since fiji does not support fp32 denormals, unless +fp32-denormals is
-// explicitly set. amdgcn target always do not flush fp64 denormals. The control for fp64 and fp16 denormals is the same.
-
-// AMDGCN-LABEL: define void @f()
-
-// AMDGCN-FLUSH: attributes #{{[0-9]*}} = {{{[^}]*}} "denorms-are-zero"="false" {{.*}} "target-features"="{{[^"]*}}+fp64-fp16-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}"
-// AMDGCN-FLUSH-OPT: attributes #{{[0-9]*}} = {{{[^}]*}} "denorms-are-zero"="true" {{.*}} "target-features"="{{[^"]*}}+fp64-fp16-denormals,{{[^"]*}}-fp32-denormals{{[^"]*}}"
-
-// AMDGCN-DENORM: attributes #{{[0-9]*}} = {{{[^}]*}} "denorms-are-zero"="false" {{.*}} "target-features"="{{[^"]*}}+fp32-denormals,{{[^"]*}}+fp64-fp16-denormals{{[^"]*}}"
-
-// AMDGCN-FEATURE: attributes #{{[0-9]*}} = {{{[^}]*}} "denorms-are-zero"="true" {{.*}} "target-features"="{{[^"]*}}+fp32-denormals,{{[^"]*}}-fp64-fp16-denormals{{[^"]*}}"
-void f() {}

clang/test/CodeGenOpenCL/gfx9-fp32-denorms.cl

-// REQUIRES: amdgpu-registered-target
-
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck --check-prefix=DEFAULT %s
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - -target-feature +fp32-denormals %s | FileCheck --check-prefix=FEATURE_FP32_DENORMALS_ON %s
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - -target-feature -fp32-denormals %s | FileCheck --check-prefix=FEATURE_FP32_DENORMALS_OFF %s
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - -cl-denorms-are-zero %s | FileCheck --check-prefix=OPT_DENORMS_ARE_ZERO %s
-
-// DEFAULT: +fp32-denormals
-// FEATURE_FP32_DENORMALS_ON: +fp32-denormals
-// FEATURE_FP32_DENORMALS_OFF: -fp32-denormals
-// OPT_DENORMS_ARE_ZERO: -fp32-denormals
-
-kernel void gfx9_fp32_denorms() {}

clang/test/Driver/cl-denorms-are-zero.cl

+// Slow FMAF and slow f32 denormals
+// RUN: %clang -### -target amdgcn--amdhsa -c -mcpu=pitcairn %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
+// RUN: %clang -### -cl-denorms-are-zero -o - -target amdgcn--amdhsa -c -mcpu=pitcairn %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
+
+// Fast FMAF, but slow f32 denormals
+// RUN: %clang -### -target amdgcn--amdhsa -c -mcpu=tahiti %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
+// RUN: %clang -### -cl-denorms-are-zero -o - -target amdgcn--amdhsa -c -mcpu=tahiti %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
+
+// Fast F32 denormals, but slow FMAF
+// RUN: %clang -### -target amdgcn--amdhsa -c -mcpu=fiji %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
+// RUN: %clang -### -cl-denorms-are-zero -o - -target amdgcn--amdhsa -c -mcpu=fiji %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
+
+// Fast F32 denormals and fast FMAF
+// RUN: %clang -### -target amdgcn--amdhsa -c -mcpu=gfx900 %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-DENORM %s
+// RUN: %clang -### -cl-denorms-are-zero -o - -target amdgcn--amdhsa -c -mcpu=gfx900 %s 2>&1 | FileCheck -check-prefixes=AMDGCN,AMDGCN-FLUSH %s
+
+// AMDGCN-FLUSH: "-fdenormal-fp-math-f32=preserve-sign"
+
+// This should be omitted and default to ieee
+// AMDGCN-DENORM-NOT: "-fdenormal-fp-math-f32"

clang/test/Driver/cuda-flush-denormals-to-zero.cu

+// Checks that cuda compilation does the right thing when passed
+// -fcuda-flush-denormals-to-zero. This should be translated to
+// -fdenormal-fp-math-f32=preserve-sign
+
+// RUN: %clang -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell--cuda-gpu-arch=sm_20 -fcuda-flush-denormals-to-zero -nocudainc -nocudalib %s 2>&1 | FileCheck -check-prefix=FTZ %s
+// RUN: %clang -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell--cuda-gpu-arch=sm_20 -fno-cuda-flush-denormals-to-zero -nocudainc -nocudalib %s 2>&1 | FileCheck -check-prefix=NOFTZ %s
+// RUN: %clang -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell--cuda-gpu-arch=sm_10 -fcuda-flush-denormals-to-zero -nocudainc -nocudalib %s 2>&1 | FileCheck -check-prefix=FTZ %s
+// RUN: %clang -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell--cuda-gpu-arch=sm_10 -fno-cuda-flush-denormals-to-zero -nocudainc -nocudalib %s 2>&1 | FileCheck -check-prefix=NOFTZ %s
+
+// CPUFTZ-NOT: -fdenormal-fp-math
+
+// FTZ: "-fdenormal-fp-math-f32=preserve-sign"
+// NOFTZ: "-fdenormal-fp-math=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=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-INVALID %s
 
-// CHECK-IEEE: "-fdenormal-fp-math=ieee"
+// CHECK-IEEE: -fdenormal-fp-math=ieee
 // CHECK-PS: "-fdenormal-fp-math=preserve-sign"
 // CHECK-PZ: "-fdenormal-fp-math=positive-zero"
 // CHECK-NO-UNSAFE-NOT: "-fdenormal-fp-math=ieee"
 // CHECK-INVALID: error: invalid value 'foo' in '-fdenormal-fp-math=foo'

clang/test/Driver/opencl.cl

 // CHECK-STRICT-ALIASING: "-cc1" {{.*}} "-cl-strict-aliasing"
 // CHECK-SINGLE-PRECISION-CONST: "-cc1" {{.*}} "-cl-single-precision-constant"
 // CHECK-FINITE-MATH-ONLY: "-cc1" {{.*}} "-cl-finite-math-only"
 // CHECK-KERNEL-ARG-INFO: "-cc1" {{.*}} "-cl-kernel-arg-info"
 // CHECK-UNSAFE-MATH-OPT: "-cc1" {{.*}} "-cl-unsafe-math-optimizations"
 // CHECK-FAST-RELAXED-MATH: "-cc1" {{.*}} "-cl-fast-relaxed-math"
 // CHECK-MAD-ENABLE: "-cc1" {{.*}} "-cl-mad-enable"
 // CHECK-NO-SIGNED-ZEROS: "-cc1" {{.*}} "-cl-no-signed-zeros"
-// CHECK-DENORMS-ARE-ZERO: "-cc1" {{.*}} "-cl-denorms-are-zero"
+
+// This is not forwarded
+// CHECK-DENORMS-ARE-ZERO-NOT: "-cl-denorms-are-zero"
+
 // CHECK-ROUND-DIV: "-cc1" {{.*}} "-cl-fp32-correctly-rounded-divide-sqrt"
 // CHECK-UNIFORM-WG: "-cc1" {{.*}} "-cl-uniform-work-group-size"
 // CHECK-C99: error: invalid value 'c99' in '-cl-std=c99'
 // CHECK-INVALID: error: invalid value 'invalid' in '-cl-std=invalid'
 
 kernel void func(void);

llvm/docs/LangRef.rst

``sspstrong``
     resulting function will have an ``sspstrong`` attribute.
 ``strictfp``
     This attribute indicates that the function was called from a scope that
     requires strict floating-point semantics.  LLVM will not attempt any
     optimizations that require assumptions about the floating-point rounding
     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 may be one of
+   ``"ieee"``, ``"preserve-sign"``, or ``"positive-zero"``.  If this
+   is attribute is not specified, the default is ``"ieee"``. If the
+   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.
+
+``"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.
+
 ``"thunk"``
     This attribute indicates that the function will delegate to some other
     function with a tail call. The prototype of a thunk should not be used for
     optimization purposes. The caller is expected to cast the thunk prototype to
     match the thunk target prototype.
 ``uwtable``
     This attribute indicates that the ABI being targeted requires that
     an unwind table entry be produced for this function even if we can

llvm/lib/CodeGen/MachineFunction.cpp

getOrCreateJumpTableInfo(unsigned EntryKind) {
   if (JumpTableInfo) return JumpTableInfo;
 
   JumpTableInfo = new (Allocator)
     MachineJumpTableInfo((MachineJumpTableInfo::JTEntryKind)EntryKind);
   return JumpTableInfo;
 }
 
 DenormalMode MachineFunction::getDenormalMode(const fltSemantics &FPType) const {
+  if (&FPType == &APFloat::IEEEsingle()) {
+    Attribute Attr = F.getFnAttribute("denormal-fp-math-f32");
+    StringRef Val = Attr.getValueAsString();
+    if (!Val.empty())
+      return parseDenormalFPAttribute(Val);
+
+    // If the f32 variant of the attribute isn't specified, try to use the
+    // generic one.
+  }
+
   // TODO: Should probably avoid the connection to the IR and store directly
   // in the MachineFunction.
   Attribute Attr = F.getFnAttribute("denormal-fp-math");
 
   // FIXME: This should assume IEEE behavior on an unspecified
   // attribute. However, the one current user incorrectly assumes a non-IEEE
   // target by default.
   StringRef Val = Attr.getValueAsString();

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

bool NVPTXTargetLowering::usePrecSqrtF32() const {
   }
 }
 
 bool NVPTXTargetLowering::useF32FTZ(const MachineFunction &MF) const {
   // TODO: Get rid of this flag; there can be only one way to do this.
   if (FtzEnabled.getNumOccurrences() > 0) {
     // If nvptx-f32ftz is used on the command-line, always honor it
     return FtzEnabled;
-  } else {
-    const Function &F = MF.getFunction();
-    // Otherwise, check for an nvptx-f32ftz attribute on the function
-    if (F.hasFnAttribute("nvptx-f32ftz"))
-      return F.getFnAttribute("nvptx-f32ftz").getValueAsString() == "true";
-    else
-      return false;
   }
+
+  return MF.getDenormalMode(APFloat::IEEEsingle()) ==
+         DenormalMode::PreserveSign;
 }
 
 static bool IsPTXVectorType(MVT VT) {
   switch (VT.SimpleTy) {
   default:
     return false;
   case MVT::v2i1:
   case MVT::v4i1:

llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "InstCombineInternal.h"
 #include "llvm/ADT/APFloat.h"
 #include "llvm/ADT/APInt.h"
 #include "llvm/ADT/APSInt.h"
 #include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/FloatingPointMode.h"
 #include "llvm/ADT/None.h"
 #include "llvm/ADT/Optional.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/Statistic.h"
 #include "llvm/ADT/Twine.h"
 #include "llvm/Analysis/AssumptionCache.h"
 #include "llvm/Analysis/InstructionSimplify.h"
const SimplifyAction Action = [II]() -> SimplifyAction {
     }
   }();
 
   // If Action.FtzRequirementTy is not satisfied by the module's ftz state, we
   // can bail out now.  (Notice that in the case that IID is not an NVVM
   // intrinsic, we don't have to look up any module metadata, as
   // FtzRequirementTy will be FTZ_Any.)
   if (Action.FtzRequirement != FTZ_Any) {
-    bool FtzEnabled =
-        II->getFunction()->getFnAttribute("nvptx-f32ftz").getValueAsString() ==
-        "true";
+    StringRef Attr = II->getFunction()
+                         ->getFnAttribute("denormal-fp-math-f32")
+                         .getValueAsString();
+    bool FtzEnabled = parseDenormalFPAttribute(Attr) != DenormalMode::IEEE;
 
     if (FtzEnabled != (Action.FtzRequirement == FTZ_MustBeOn))
       return nullptr;
   }
 
   // Simplify to target-generic intrinsic.
   if (Action.IID) {
     SmallVector<Value *, 4> Args(II->arg_operands());

llvm/test/CodeGen/NVPTX/fast-math.ll

 ; CHECK: mul.ftz.f32
   %x = fdiv float %a, %divisor
   %y = fdiv float %b, %divisor
   %z = select i1 %pred, float %x, float %y
   ret float %z
 }
 
 attributes #0 = { "unsafe-fp-math" = "true" }
-attributes #1 = { "nvptx-f32ftz" = "true" }
+attributes #1 = { "denormal-fp-math-f32" = "preserve-sign" }

llvm/test/CodeGen/NVPTX/math-intrins.ll

 ; CHECK-LABEL: @fma_double
 define double @fma_double(double %a, double %b, double %c) {
   ; CHECK: fma.rn.f64
   %x = call double @llvm.fma.f64(double %a, double %b, double %c)
   ret double %x
 }
 
 attributes #0 = { nounwind readnone }
-attributes #1 = { "nvptx-f32ftz" = "true" }
+attributes #1 = { "denormal-fp-math-f32" = "preserve-sign" }

llvm/test/CodeGen/NVPTX/sqrt-approx.ll

 ; CHECK-LABEL: test_sqrt64_refined_ftz
 define double @test_sqrt64_refined_ftz(double %a) #0 #1 #2 {
 ; CHECK: rsqrt.approx.f64
   %ret = tail call double @llvm.sqrt.f64(double %a)
   ret double %ret
 }
 
 attributes #0 = { "unsafe-fp-math" = "true" }
-attributes #1 = { "nvptx-f32ftz" = "true" }
+attributes #1 = { "denormal-fp-math-f32" = "preserve-sign" }
 attributes #2 = { "reciprocal-estimates" = "rsqrtf:1,rsqrtd:1,sqrtf:1,sqrtd:1" }

llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll

 ; Check that nvvm intrinsics get simplified to target-generic intrinsics where
 ; possible.
 ;
 ; We run this test twice; once with ftz on, and again with ftz off.  Behold the
 ; hackery:
 
 ; RUN: cat %s > %t.ftz
-; RUN: echo 'attributes #0 = { "nvptx-f32ftz" = "true" }' >> %t.ftz
+; RUN: echo 'attributes #0 = { "denormal-fp-math-f32" = "preserve-sign" }' >> %t.ftz
 ; RUN: opt < %t.ftz -instcombine -S | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ
 
 ; RUN: cat %s > %t.noftz
-; RUN: echo 'attributes #0 = { "nvptx-f32ftz" = "false" }' >> %t.noftz
+; RUN: echo 'attributes #0 = { "denormal-fp-math-f32" = "ieee" }' >> %t.noftz
 ; RUN: opt < %t.noftz -instcombine -S | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ
 
 ; We handle nvvm intrinsics with ftz variants as follows:
 ;  - If the module is in ftz mode, the ftz variant is transformed into the
 ;    regular llvm intrinsic, and the non-ftz variant is left alone.
 ;  - If the module is not in ftz mode, it's the reverse: Only the non-ftz
 ;    variant is transformed, and the ftz variant is left alone.