Separately track input and output denormal mode

AMDGPU and x86 at least both have separate controls for whether
denormal results are flushed on output, and for whether denormals are
implicitly treated as 0 as an input. The current DAGCombiner use only
really cares about the input treatment of denormals.

clang/include/clang/Basic/CodeGenOptions.h

@@ -164,10 +164,10 @@ public:
   std::string FloatABI;
 
   /// The floating-point denormal mode to use.
-  llvm::DenormalMode FPDenormalMode = llvm::DenormalMode::Invalid;
+  llvm::DenormalMode FPDenormalMode;
 
   /// The floating-point subnormal mode to use, for float.
-  llvm::DenormalMode FP32DenormalMode = llvm::DenormalMode::Invalid;
+  llvm::DenormalMode FP32DenormalMode;
 
   /// The float precision limit to use, if non-empty.
   std::string LimitFloatPrecision;

clang/include/clang/Driver/ToolChain.h

@@ -617,7 +617,7 @@ public:
       Action::OffloadKind DeviceOffloadKind,
       const llvm::fltSemantics *FPType = nullptr) const {
     // FIXME: This should be IEEE when default handling is fixed.
-    return llvm::DenormalMode::Invalid;
+    return llvm::DenormalMode::getInvalid();
   }
 };
 

clang/lib/Basic/Targets/AMDGPU.cpp

@@ -247,7 +247,7 @@ void AMDGPUTargetInfo::adjustTargetOptions(const CodeGenOptions &CGOpts,
   if (!hasFP32Denormals)
     TargetOpts.Features.push_back(
       (Twine(hasFastFMAF() && hasFullRateDenormalsF32() &&
-             CGOpts.FP32DenormalMode == llvm::DenormalMode::IEEE
+             CGOpts.FP32DenormalMode.Output == llvm::DenormalMode::IEEE
              ? '+' : '-') + Twine("fp32-denormals"))
             .str());
   // Always do not flush fp64 or fp16 denorms.

clang/lib/CodeGen/CGCall.cpp

@@ -1749,14 +1749,14 @@ void CodeGenModule::ConstructDefaultFnAttrList(StringRef Name, bool HasOptnone,
       FuncAttrs.addAttribute("null-pointer-is-valid", "true");
 
     // TODO: Omit attribute when the default is IEEE.
-    if (CodeGenOpts.FPDenormalMode != llvm::DenormalMode::Invalid)
+    if (CodeGenOpts.FPDenormalMode.isValid())
       FuncAttrs.addAttribute("denormal-fp-math",
-                             llvm::denormalModeName(CodeGenOpts.FPDenormalMode));
+                             CodeGenOpts.FPDenormalMode.str());
-
+    if (CodeGenOpts.FP32DenormalMode.isValid()) {
-    if (CodeGenOpts.FP32DenormalMode != llvm::DenormalMode::Invalid)
       FuncAttrs.addAttribute(
           "denormal-fp-math-f32",
-          llvm::denormalModeName(CodeGenOpts.FP32DenormalMode));
+          CodeGenOpts.FP32DenormalMode.str());
+    }
 
     FuncAttrs.addAttribute("no-trapping-math",
                            llvm::toStringRef(CodeGenOpts.NoTrappingMath));

clang/lib/CodeGen/CodeGenModule.cpp

@@ -587,7 +587,7 @@ void CodeGenModule::Release() {
     // floating point values to 0.  (This corresponds to its "__CUDA_FTZ"
     // property.)
     getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-ftz",
-                              CodeGenOpts.FP32DenormalMode !=
+                              CodeGenOpts.FP32DenormalMode.Output !=
                                   llvm::DenormalMode::IEEE);
   }
 

clang/lib/Driver/ToolChains/AMDGPU.cpp

@@ -108,14 +108,14 @@ llvm::DenormalMode AMDGPUToolChain::getDefaultDenormalModeForType(
     const llvm::fltSemantics *FPType) const {
   // Denormals should always be enabled for f16 and f64.
   if (!FPType || FPType != &llvm::APFloat::IEEEsingle())
-    return llvm::DenormalMode::IEEE;
+    return llvm::DenormalMode::getIEEE();
 
   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;
+      return llvm::DenormalMode::getPreserveSign();
   }
 
   const StringRef GpuArch = DriverArgs.getLastArgValue(options::OPT_mcpu_EQ);
@@ -134,7 +134,8 @@ llvm::DenormalMode AMDGPUToolChain::getDefaultDenormalModeForType(
   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;
+  return DAZ ? llvm::DenormalMode::getPreserveSign() :
+               llvm::DenormalMode::getIEEE();
 }
 
 void AMDGPUToolChain::addClangTargetOptions(

clang/lib/Driver/ToolChains/Clang.cpp

@@ -2641,7 +2641,7 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
 
     case options::OPT_fdenormal_fp_math_EQ:
       DenormalFPMath = llvm::parseDenormalFPAttribute(A->getValue());
-      if (DenormalFPMath == llvm::DenormalMode::Invalid) {
+      if (!DenormalFPMath.isValid()) {
         D.Diag(diag::err_drv_invalid_value)
             << A->getAsString(Args) << A->getValue();
       }
@@ -2649,7 +2649,7 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
 
     case options::OPT_fdenormal_fp_math_f32_EQ:
       DenormalFP32Math = llvm::parseDenormalFPAttribute(A->getValue());
-      if (DenormalFP32Math == llvm::DenormalMode::Invalid) {
+      if (!DenormalFP32Math.isValid()) {
         D.Diag(diag::err_drv_invalid_value)
             << A->getAsString(Args) << A->getValue();
       }
@@ -2768,7 +2768,7 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
       if (HonorINFs && HonorNaNs &&
         !AssociativeMath && !ReciprocalMath &&
         SignedZeros && TrappingMath && RoundingFPMath &&
-        DenormalFPMath != llvm::DenormalMode::IEEE &&
+        DenormalFPMath != llvm::DenormalMode::getIEEE() &&
         FPContract.empty())
         // OK: Current Arg doesn't conflict with -ffp-model=strict
         ;
@@ -2816,14 +2816,18 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
     CmdArgs.push_back("-fno-trapping-math");
 
   // TODO: Omit flag for the default IEEE instead
-  if (DenormalFPMath != llvm::DenormalMode::Invalid) {
+  if (DenormalFPMath.isValid()) {
-    CmdArgs.push_back(Args.MakeArgString(
+    llvm::SmallString<64> DenormFlag;
-        "-fdenormal-fp-math=" + llvm::denormalModeName(DenormalFPMath)));
+    llvm::raw_svector_ostream ArgStr(DenormFlag);
+    ArgStr << "-fdenormal-fp-math=" << DenormalFPMath;
+    CmdArgs.push_back(Args.MakeArgString(ArgStr.str()));
   }
 
-  if (DenormalFP32Math != llvm::DenormalMode::Invalid) {
+  if (DenormalFP32Math.isValid()) {
-    CmdArgs.push_back(Args.MakeArgString(
+    llvm::SmallString<64> DenormFlag;
-        "-fdenormal-fp-math-f32=" + llvm::denormalModeName(DenormalFP32Math)));
+    llvm::raw_svector_ostream ArgStr(DenormFlag);
+    ArgStr << "-fdenormal-fp-math-f32=" << DenormalFP32Math;
+    CmdArgs.push_back(Args.MakeArgString(ArgStr.str()));
   }
 
   if (!FPContract.empty())

clang/lib/Driver/ToolChains/Cuda.cpp

@@ -711,11 +711,11 @@ llvm::DenormalMode CudaToolChain::getDefaultDenormalModeForType(
         DriverArgs.hasFlag(options::OPT_fcuda_flush_denormals_to_zero,
                            options::OPT_fno_cuda_flush_denormals_to_zero,
                            false))
-      return llvm::DenormalMode::PreserveSign;
+      return llvm::DenormalMode::getPreserveSign();
   }
 
   assert(DeviceOffloadKind != Action::OFK_Host);
-  return llvm::DenormalMode::IEEE;
+  return llvm::DenormalMode::getIEEE();
 }
 
 bool CudaToolChain::supportsDebugInfoOption(const llvm::opt::Arg *A) const {

clang/lib/Frontend/CompilerInvocation.cpp

@@ -1286,14 +1286,14 @@ 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)
+    if (!Opts.FPDenormalMode.isValid())
       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)
+    if (!Opts.FP32DenormalMode.isValid())
       Diags.Report(diag::err_drv_invalid_value) << A->getAsString(Args) << Val;
   }
 

clang/test/CodeGen/denormalfpmode.c

@@ -3,9 +3,9 @@
 // RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-PZ
 
 // CHECK-LABEL: main
-// CHECK-IEEE: attributes #0 = {{.*}}"denormal-fp-math"="ieee"{{.*}}
+// CHECK-IEEE: attributes #0 = {{.*}}"denormal-fp-math"="ieee,ieee"{{.*}}
-// CHECK-PS: attributes #0 = {{.*}}"denormal-fp-math"="preserve-sign"{{.*}}
+// CHECK-PS: attributes #0 = {{.*}}"denormal-fp-math"="preserve-sign,preserve-sign"{{.*}}
-// CHECK-PZ: attributes #0 = {{.*}}"denormal-fp-math"="positive-zero"{{.*}}
+// CHECK-PZ: attributes #0 = {{.*}}"denormal-fp-math"="positive-zero,positive-zero"{{.*}}
 
 int main() {
   return 0;

clang/test/CodeGenCUDA/flush-denormals.cu

@@ -39,8 +39,8 @@
 // CHECK-LABEL: define void @foo() #0
 extern "C" __device__ void foo() {}
 
-// FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign"
+// FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
-// NOFTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="ieee"
+// NOFTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="ieee,ieee"
 
 
 // FIXME: This should be removed

clang/test/CodeGenCUDA/propagate-metadata.cu

@@ -61,8 +61,8 @@ __global__ void kernel() { lib_fn(); }
 
 // FTZ-NOT: "denormal-fp-math"
 
-// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign"
+// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
-// NOFTZ-SAME: "denormal-fp-math-f32"="ieee"
+// NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee"
 
 // CHECK-SAME: "no-trapping-math"="true"
 

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

@@ -14,7 +14,7 @@
 // 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"
+// AMDGCN-FLUSH: "-fdenormal-fp-math-f32=preserve-sign,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

@@ -9,5 +9,5 @@
 
 // CPUFTZ-NOT: -fdenormal-fp-math
 
-// FTZ: "-fdenormal-fp-math-f32=preserve-sign"
+// FTZ: "-fdenormal-fp-math-f32=preserve-sign,preserve-sign"
-// NOFTZ: "-fdenormal-fp-math=ieee"
+// NOFTZ: "-fdenormal-fp-math=ieee,ieee"

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

@@ -3,10 +3,16 @@
 // 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
+// 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
 
-// CHECK-IEEE: -fdenormal-fp-math=ieee
+// CHECK-IEEE: -fdenormal-fp-math=ieee,ieee
-// CHECK-PS: "-fdenormal-fp-math=preserve-sign"
+// CHECK-PS: "-fdenormal-fp-math=preserve-sign,preserve-sign"
-// CHECK-PZ: "-fdenormal-fp-math=positive-zero"
+// CHECK-PZ: "-fdenormal-fp-math=positive-zero,positive-zero"
 // CHECK-NO-UNSAFE-NOT: "-fdenormal-fp-math=ieee"
-// CHECK-INVALID: error: invalid value 'foo' in '-fdenormal-fp-math=foo'
+// 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

@@ -1820,12 +1820,21 @@ example:
     not introduce any new floating-point instructions that may trap.
 
 ``"denormal-fp-math"``
-    This indicates the denormal (subnormal) handling that may be assumed
+    This indicates the denormal (subnormal) handling that may be
-    for the default floating-point environment. This may be one of
+    assumed for the default floating-point environment. This is a
-    ``"ieee"``, ``"preserve-sign"``, or ``"positive-zero"``.  If this
+    comma separated pair. The elements may be one of ``"ieee"``,
-    is attribute is not specified, the default is ``"ieee"``. If the
+    ``"preserve-sign"``, or ``"positive-zero"``. The first entry
-    mode is ``"preserve-sign"``, or ``"positive-zero"``, denormal
+    indicates the flushing mode for the result of floating point
-    outputs may be flushed to zero by standard floating point
+    operations. The second indicates the handling of denormal inputs
+    to floating point instructions. For compatability 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 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
@@ -1834,6 +1843,12 @@ example:
     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

llvm/include/llvm/ADT/FloatingPointMode.h

@@ -14,12 +14,16 @@
 #define LLVM_FLOATINGPOINTMODE_H
 
 #include "llvm/ADT/StringSwitch.h"
+#include "llvm/Support/raw_ostream.h"
 
 namespace llvm {
 
-/// Represent handled modes for denormal (aka subnormal) modes in the floating
+/// Represent ssubnormal handling kind for floating point instruction inputs and
-/// point environment.
+/// outputs.
-enum class DenormalMode {
+struct DenormalMode {
+  /// Represent handled modes for denormal (aka subnormal) modes in the floating
+  /// point environment.
+  enum DenormalModeKind : char {
     Invalid = -1,
 
     /// IEEE-754 denormal numbers preserved.
@@ -30,12 +34,77 @@ enum class DenormalMode {
 
     /// Denormals are flushed to positive zero.
     PositiveZero
+  };
+
+  /// 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;
+
+  DenormalMode() = default;
+  DenormalMode(DenormalModeKind Out, DenormalModeKind In) :
+    Output(Out), Input(In) {}
+
+
+  static DenormalMode getInvalid() {
+    return DenormalMode(DenormalModeKind::Invalid, DenormalModeKind::Invalid);
+  }
+
+  static DenormalMode getIEEE() {
+    return DenormalMode(DenormalModeKind::IEEE, DenormalModeKind::IEEE);
+  }
+
+  static DenormalMode getPreserveSign() {
+    return DenormalMode(DenormalModeKind::PreserveSign,
+                        DenormalModeKind::PreserveSign);
+  }
+
+  static DenormalMode getPositiveZero() {
+    return DenormalMode(DenormalModeKind::PositiveZero,
+                        DenormalModeKind::PositiveZero);
+  }
+
+  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;
+  }
+
+  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 parseDenormalFPAttribute(StringRef Str) {
+inline DenormalMode::DenormalModeKind
+parseDenormalFPAttributeComponent(StringRef Str) {
   // Assume ieee on unspecified attribute.
-  return StringSwitch<DenormalMode>(Str)
+  return StringSwitch<DenormalMode::DenormalModeKind>(Str)
     .Cases("", "ieee", DenormalMode::IEEE)
     .Case("preserve-sign", DenormalMode::PreserveSign)
     .Case("positive-zero", DenormalMode::PositiveZero)
@@ -44,7 +113,7 @@ inline DenormalMode parseDenormalFPAttribute(StringRef Str) {
 
 /// Return the name used for the denormal handling mode used by the the
 /// expected names from the denormal-fp-math attribute.
-inline StringRef denormalModeName(DenormalMode Mode) {
+inline StringRef denormalModeKindName(DenormalMode::DenormalModeKind Mode) {
   switch (Mode) {
   case DenormalMode::IEEE:
     return "ieee";
@@ -57,6 +126,26 @@ inline StringRef denormalModeName(DenormalMode Mode) {
   }
 }
 
+/// Returns the denormal mode to use for inputs and outputs.
+inline DenormalMode parseDenormalFPAttribute(StringRef Str) {
+  StringRef OutputStr, InputStr;
+  std::tie(OutputStr, InputStr) = Str.split(',');
+
+  DenormalMode Mode;
+  Mode.Output = parseDenormalFPAttributeComponent(OutputStr);
+
+  // Maintain compatability with old form of the attribute which only specified
+  // one component.
+  Mode.Input = InputStr.empty() ? Mode.Output  :
+               parseDenormalFPAttributeComponent(InputStr);
+
+  return Mode;
+}
+
+void DenormalMode::print(raw_ostream &OS) const {
+  OS << denormalModeKindName(Output) << ',' << denormalModeKindName(Input);
+}
+
 }
 
 #endif // LLVM_FLOATINGPOINTMODE_H

llvm/lib/CodeGen/MachineFunction.cpp

@@ -290,7 +290,7 @@ DenormalMode MachineFunction::getDenormalMode(const fltSemantics &FPType) const
   // target by default.
   StringRef Val = Attr.getValueAsString();
   if (Val.empty())
-    return DenormalMode::Invalid;
+    return DenormalMode::getInvalid();
 
   return parseDenormalFPAttribute(Val);
 }

llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp

@@ -20820,7 +20820,10 @@ SDValue DAGCombiner::buildSqrtEstimateImpl(SDValue Op, SDNodeFlags Flags,
         EVT CCVT = getSetCCResultType(VT);
         ISD::NodeType SelOpcode = VT.isVector() ? ISD::VSELECT : ISD::SELECT;
         DenormalMode DenormMode = DAG.getDenormalMode(VT);
-        if (DenormMode == DenormalMode::IEEE) {
+        if (DenormMode.Input == DenormalMode::IEEE) {
+          // This is specifically a check for the handling of denormal inputs,
+          // not the result.
+
           // fabs(X) < SmallestNormal ? 0.0 : Est
           const fltSemantics &FltSem = DAG.EVTToAPFloatSemantics(VT);
           APFloat SmallestNorm = APFloat::getSmallestNormalized(FltSem);

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

@@ -123,7 +123,7 @@ bool NVPTXTargetLowering::useF32FTZ(const MachineFunction &MF) const {
     return FtzEnabled;
   }
 
-  return MF.getDenormalMode(APFloat::IEEEsingle()) ==
+  return MF.getDenormalMode(APFloat::IEEEsingle()).Output ==
          DenormalMode::PreserveSign;
 }
 

llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp

@@ -1706,7 +1706,8 @@ static Instruction *SimplifyNVVMIntrinsic(IntrinsicInst *II, InstCombiner &IC) {
     StringRef Attr = II->getFunction()
                          ->getFnAttribute("denormal-fp-math-f32")
                          .getValueAsString();
-    bool FtzEnabled = parseDenormalFPAttribute(Attr) != DenormalMode::IEEE;
+    DenormalMode Mode = parseDenormalFPAttribute(Attr);
+    bool FtzEnabled = Mode.Output != DenormalMode::IEEE;
 
     if (FtzEnabled != (Action.FtzRequirement == FTZ_MustBeOn))
       return nullptr;

llvm/unittests/ADT/FloatingPointMode.cpp

@@ -13,21 +13,122 @@ using namespace llvm;
 
 namespace {
 
-TEST(FloatingPointModeTest, ParseDenormalFPAttribute) {
+TEST(FloatingPointModeTest, ParseDenormalFPAttributeComponent) {
-  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttribute("ieee"));
+  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttributeComponent("ieee"));
-  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttribute(""));
+  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttributeComponent(""));
   EXPECT_EQ(DenormalMode::PreserveSign,
-            parseDenormalFPAttribute("preserve-sign"));
+            parseDenormalFPAttributeComponent("preserve-sign"));
   EXPECT_EQ(DenormalMode::PositiveZero,
-            parseDenormalFPAttribute("positive-zero"));
+            parseDenormalFPAttributeComponent("positive-zero"));
-  EXPECT_EQ(DenormalMode::Invalid, parseDenormalFPAttribute("foo"));
+  EXPECT_EQ(DenormalMode::Invalid, parseDenormalFPAttributeComponent("foo"));
 }
 
 TEST(FloatingPointModeTest, DenormalAttributeName) {
-  EXPECT_EQ("ieee", denormalModeName(DenormalMode::IEEE));
+  EXPECT_EQ("ieee", denormalModeKindName(DenormalMode::IEEE));
-  EXPECT_EQ("preserve-sign", denormalModeName(DenormalMode::PreserveSign));
+  EXPECT_EQ("preserve-sign", denormalModeKindName(DenormalMode::PreserveSign));
-  EXPECT_EQ("positive-zero", denormalModeName(DenormalMode::PositiveZero));
+  EXPECT_EQ("positive-zero", denormalModeKindName(DenormalMode::PositiveZero));
-  EXPECT_EQ("", denormalModeName(DenormalMode::Invalid));
+  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"));
+  EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE),
+            parseDenormalFPAttribute("ieee,"));
+  EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE),
+            parseDenormalFPAttribute(""));
+  EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE),
+            parseDenormalFPAttribute(","));
+
+  EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::PreserveSign),
+            parseDenormalFPAttribute("preserve-sign"));
+  EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::PreserveSign),
+            parseDenormalFPAttribute("preserve-sign,"));
+  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::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::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(DenormalMode(DenormalMode::Invalid, DenormalMode::Invalid),
+            parseDenormalFPAttribute("foo"));
+
+  EXPECT_EQ("ieee,ieee",
+            DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE).str());
+  EXPECT_EQ(",",
+            DenormalMode(DenormalMode::Invalid, DenormalMode::Invalid).str());
+
+  EXPECT_EQ(
+    "preserve-sign,preserve-sign",
+    DenormalMode(DenormalMode::PreserveSign, DenormalMode::PreserveSign).str());
+
+  EXPECT_EQ(
+    "positive-zero,positive-zero",
+    DenormalMode(DenormalMode::PositiveZero, DenormalMode::PositiveZero).str());
+
+  EXPECT_EQ(
+    "ieee,preserve-sign",
+    DenormalMode(DenormalMode::IEEE, DenormalMode::PreserveSign).str());
+
+  EXPECT_EQ(
+    "preserve-sign,ieee",
+    DenormalMode(DenormalMode::PreserveSign, DenormalMode::IEEE).str());
+
+  EXPECT_EQ(
+    "preserve-sign,positive-zero",
+    DenormalMode(DenormalMode::PreserveSign, DenormalMode::PositiveZero).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());
+}
+
+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(DenormalMode::PreserveSign, DenormalMode::PreserveSign),
+            DenormalMode::getPreserveSign());
+  EXPECT_EQ(DenormalMode(DenormalMode::PositiveZero, DenormalMode::PositiveZero),
+            DenormalMode::getPositiveZero());
 }
 
 }