Separately track input and output denormal mode

[arsenm]

Summary

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.

[andrew.w.kaylor]

llvm/docs/LangRef.rst

1822

I don't like the definition of this attribute. It's not reader-friendly. The comma-separated pair format has no indication which value refers to inputs and which refers to outputs. Also, while this predates your changes, I think the meanings of the current choices are unclear.

What would you think of a comma-separated list with the following possibilities?

allow-denormals (default)
inputs-are-zero (outputs not flushed)
inputs-are-zero, outputs-are-zero
inputs-are-zero, outputs-are-positive-zero
inputs-are-positivezero (outputs not flushed)
inputs-are-positivezero, outputs-are-zero
inputs-are-positivezero, outputs-are-positive-zero
denormal-outputs-are-zero (inputs are unchanged)
denormal-outputs-are-positive-zero (inputs are unchanged)

I'd also be open to abbreviations. I don't know if "daz" and "ftz" are readable to everyone, but I'm more comfortable with them. That would make the options something like this.

allow-denormals
daz
daz, ftz
daz, ftz+
daz+
daz+, ftz
daz+, ftz+
ftz
ftz+

[arsenm]

llvm/docs/LangRef.rst

1822

I'm trying to avoid needing to autoupgrade bitcode at this point, which leaving the names as-is accomplishes. I'm worried this could still end up not in the right place, and then we would need another level of auto upgrade to deal with it later. I think these are overly verbose (I'm also keeping in mind the fact that any use of these does a linear scan through all string attributes, and then needs to parse these).

I'm also unclear on what this weird ARM positive-zero really means. Does it mean inputs and outputs ignored the sign? Is there value in representing positive-zero on both sides?

[arsenm]

Rebase

[andrew.w.kaylor]

llvm/docs/LangRef.rst

1826

Is this saying that if a backend generates an instruction that doesn't handle the hardware daz mode then it must insert instructions to check for normals and convert them to zero? If so, do you intend this to apply to all such instructions or only instructions that aren't able to accept denormal inputs?

1827

Based on the changes below, if the second value is omitted the input mode will be assumed to be the same as the output mode. That should probably be documented. I guess you intend for that not to happen, but the documentation here leaves the result ambiguous if it does happen.

[arsenm]

llvm/docs/LangRef.rst

1826

Only in cases where denormal inputs are invalid or unhandled. The case I'm thinking of is the one user in DAGCombiner, where if denormals are not flushed the result ends up incorrect (see https://bugs.llvm.org/show_bug.cgi?id=34994)

1827

I've added a note for this

[arsenm]

Tweak langref

[andrew.w.kaylor]

lgtm

[arsenm]

a3c814d23497bc71b8ed53c35f773366aff02922

clang/include/clang/Basic/CodeGenOptions.h

public:
   std::string RecordCommandLine;
 
   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;
+  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;
 
   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

clang/include/clang/Driver/ToolChain.h

public:
   /// 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;
+    return llvm::DenormalMode::getInvalid();
   }
 };
 
 /// Set a ToolChain's effective triple. Reset it when the registration object
 /// is destroyed.
 class RegisterEffectiveTriple {
   const ToolChain &TC;
 

clang/lib/Basic/Targets/AMDGPU.cpp

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.FP32DenormalMode == llvm::DenormalMode::IEEE
+             CGOpts.FP32DenormalMode.Output == 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("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)
+    if (CodeGenOpts.FPDenormalMode.isValid())
       FuncAttrs.addAttribute("denormal-fp-math",
-                             llvm::subnormalModeName(CodeGenOpts.FPDenormalMode));
-    if (CodeGenOpts.FP32DenormalMode != llvm::DenormalMode::Invalid)
+                             CodeGenOpts.FPDenormalMode.str());
+    if (CodeGenOpts.FP32DenormalMode.isValid())
       FuncAttrs.addAttribute(
           "denormal-fp-math-f32",
-          llvm::subnormalModeName(CodeGenOpts.FP32DenormalMode));
+          CodeGenOpts.FP32DenormalMode.str());
 
     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");

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.FP32DenormalMode !=
+                              CodeGenOpts.FP32DenormalMode.Output !=
                                   llvm::DenormalMode::IEEE);
   }
 
   // Emit OpenCL specific module metadata: OpenCL/SPIR version.
   if (LangOpts.OpenCL) {
     EmitOpenCLMetadata();
     // Emit SPIR version.
     if (getTriple().isSPIR()) {

clang/lib/Driver/ToolChains/AMDGPU.cpp

AMDGPUToolChain::TranslateArgs(const DerivedArgList &Args, StringRef BoundArch,
   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;
+    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);
   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;
+  return DAZ ? llvm::DenormalMode::getPreserveSign() :
+               llvm::DenormalMode::getIEEE();
 }
 
 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

for (const Arg *A : Args) {
     case options::OPT_fno_reciprocal_math:  ReciprocalMath = false;   break;
     case options::OPT_fsigned_zeros:        SignedZeros = true;       break;
     case options::OPT_fno_signed_zeros:     SignedZeros = false;      break;
     case options::OPT_ftrapping_math:       TrappingMath = true;      break;
     case options::OPT_fno_trapping_math:    TrappingMath = false;     break;
 
     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();
       }
       break;
 
     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();
       }
       break;
 
     // Validate and pass through -fp-contract option.
     case options::OPT_ffp_contract: {
       StringRef Val = A->getValue();
static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
 
   if (ReciprocalMath)
     CmdArgs.push_back("-freciprocal-math");
 
   if (!TrappingMath)
     CmdArgs.push_back("-fno-trapping-math");
 
   // 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 (DenormalFPMath.isValid()) {
+    llvm::SmallString<64> DenormFlag;
+    llvm::raw_svector_ostream ArgStr(DenormFlag);
+    ArgStr << "-fdenormal-fp-math=" << DenormalFPMath;
+    CmdArgs.push_back(Args.MakeArgString(ArgStr.str()));
   }
 
-  if (DenormalFP32Math != llvm::DenormalMode::Invalid) {
-    CmdArgs.push_back(Args.MakeArgString(
-        "-fdenormal-fp-math-f32=" + llvm::subnormalModeName(DenormalFP32Math)));
+  if (DenormalFP32Math.isValid()) {
+    llvm::SmallString<64> DenormFlag;
+    llvm::raw_svector_ostream ArgStr(DenormFlag);
+    ArgStr << "-fdenormal-fp-math-f32=" << DenormalFP32Math;
+    CmdArgs.push_back(Args.MakeArgString(ArgStr.str()));
   }
 
   if (!FPContract.empty())
     CmdArgs.push_back(Args.MakeArgString("-ffp-contract=" + FPContract));
 
   ParseMRecip(D, Args, CmdArgs);
 
   // -ffast-math enables the __FAST_MATH__ preprocessor macro, but check for the

clang/lib/Driver/ToolChains/Cuda.cpp

 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;
+      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 {
   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) ||

clang/lib/Frontend/CompilerInvocation.cpp

if (Arg *A = Args.getLastArg(OPT_ftlsmodel_EQ)) {
     } else {
       Opts.setDefaultTLSModel(static_cast<CodeGenOptions::TLSModel>(Model));
     }
   }
 
   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;
   }
 
   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));

clang/test/CodeGen/denormalfpmode.c

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

clang/test/CodeGenCUDA/flush-denormals.cu

 // 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 = {{.*}} "denormal-fp-math-f32"="preserve-sign"
-// NOFTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="ieee"
+// FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// NOFTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="ieee,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

 // FIXME: These -NOT checks do not work as intended and do not check on the same
 // line.
 
 // Check the attribute list.
 // CHECK: attributes [[attr]] = {
 
 // FTZ-NOT: "denormal-fp-math"
 
-// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign"
-// NOFTZ-SAME: "denormal-fp-math-f32"="ieee"
+// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee"
 
 // CHECK-SAME: "no-trapping-math"="true"
 
 // FAST-SAME: "unsafe-fp-math"="true"
 // NOFAST-NOT: "unsafe-fp-math"="true"

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"
+// 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

 // 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"
+// FTZ: "-fdenormal-fp-math-f32=preserve-sign,preserve-sign"
+// NOFTZ: "-fdenormal-fp-math=ieee,ieee"

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

 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee -v 2>&1 | FileCheck -check-prefix=CHECK-IEEE %s
 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=preserve-sign -v 2>&1 | FileCheck -check-prefix=CHECK-PS %s
 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=positive-zero -v 2>&1 | FileCheck -check-prefix=CHECK-PZ %s
 // RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=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-PS: "-fdenormal-fp-math=preserve-sign"
-// CHECK-PZ: "-fdenormal-fp-math=positive-zero"
+// CHECK-IEEE: -fdenormal-fp-math=ieee,ieee
+// CHECK-PS: "-fdenormal-fp-math=preserve-sign,preserve-sign"
+// 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

 ``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.
 
 ``"denorm-fp-mode"``
    This indicates the 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"``, subnormal outputs may
-   be flushed to zero by standard floating point operations. It is not
-   mandated that flushing to zero occurs, but if a subnormal output is
-   flushed to zero, it must respect the sign mode. Not all targets
-   support all modes.
+   default floating-point environment. This is a comma separated
+   pair. The elements may be one of ``"ieee"``, ``"preserve-sign"``,
+   or ``"positive-zero"``. The first entry indicates the flushing mode
+   for the result of floating point operations. The second indicates
+   the handling of denormal inputs to floating point instructions.
+
+   If this is attribute is not specified, the default is
+   ``"ieee,ieee"``.
+
+   If the output mode is ``"preserve-sign"``, or ``"positive-zero"``,
+   subnormal outputs may be flushed to zero by standard floating-point
+   operations. It is not mandated that flushing to zero occurs, but if
+   a subnormal output is flushed to zero, it must respect the sign
+   mode. Not all targets support all modes.
+
+   If the input mode is ``"preserve-sign"``, or ``"positive-zero"``, a
+   floating-point operation must treat any input denormal value as
+   zero. If an instruction does not respect this mode, the input
+   should be converted to 0 as if by ``@llvm.canonicalize`` during
+   lowering.
 
 ``"denorm-fp-mode-f32"``
    Same as ``"denorm-fp-mode-f32"``, except for float types. If both
    are present, this overrides ``"denorm-fp-mode"``.
 
 ``"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

llvm/include/llvm/ADT/FloatingPointMode.h

 //===- llvm/Support/FloatingPointMode.h -------------------------*- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Utilities for dealing with flags related to floating point mode controls.
 //
 //===----------------------------------------------------------------------===/
 
 #ifndef LLVM_FLOATINGPOINTMODE_H
 #define LLVM_FLOATINGPOINTMODE_H
 
 #include "llvm/ADT/StringSwitch.h"
+#include "llvm/Support/raw_ostream.h"
 
 namespace llvm {
 
+/// Represent denormal handling kind for floating point instruction inputs and
+/// outputs.
+struct DenormalMode {
   /// Represent handled modes for subnormal (aka denormal) modes in the floating
   /// point environment.
-enum class DenormalMode {
+  enum DenormalModeKind : char {
     Invalid = -1,
 
     /// IEEE-754 subnormal numbers preserved.
     IEEE,
 
     /// The sign of a flushed-to-zero number is preserved in the sign of 0
     PreserveSign,
 
     /// Denormals are flushed to positive zero.
     PositiveZero
   };
 
+  /// 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)
     .Default(DenormalMode::Invalid);
 }
 
 /// Return the name used for the subnormal 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";
   case DenormalMode::PreserveSign:
     return "preserve-sign";
   case DenormalMode::PositiveZero:
     return "positive-zero";
   default:
     return "";
   }
 }
 
+/// 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

DenormalMode MachineFunction::getDenormalMode(const fltSemantics &FPType) const {
   // 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();
   if (Val.empty())
-    return DenormalMode::Invalid;
+    return DenormalMode::getInvalid();
 
   return parseDenormalFPAttribute(Val);
 }
 
 /// Should we be emitting segmented stack stuff for the function
 bool MachineFunction::shouldSplitStack() const {
   return getFunction().hasFnAttribute("split-stack");
 }

llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp

if (Iterations) {
 
       if (!Reciprocal) {
         // The estimate is now completely wrong if the input was exactly 0.0 or
         // possibly a denormal. Force the answer to 0.0 for those cases.
         SDLoc DL(Op);
         EVT CCVT = getSetCCResultType(VT);
         ISD::NodeType SelOpcode = VT.isVector() ? ISD::VSELECT : ISD::SELECT;
         DenormalMode SubnormMode = DAG.getDenormalMode(VT);
-        if (SubnormMode == DenormalMode::IEEE) {
+        if (SubnormMode.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);
           SDValue NormC = DAG.getConstantFP(SmallestNorm, DL, VT);
           SDValue FPZero = DAG.getConstantFP(0.0, DL, VT);
           SDValue Fabs = DAG.getNode(ISD::FABS, DL, VT, Op);
           SDValue IsDenorm = DAG.getSetCC(DL, CCVT, Fabs, NormC, ISD::SETLT);
           Est = DAG.getNode(SelOpcode, DL, VT, IsDenorm, FPZero, Est);

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

 
 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;
   }
 
-  return MF.getDenormalMode(APFloat::IEEEsingle()) ==
+  return MF.getDenormalMode(APFloat::IEEEsingle()).Output ==
          DenormalMode::PreserveSign;
 }
 
 static bool IsPTXVectorType(MVT VT) {
   switch (VT.SimpleTy) {
   default:
     return false;
   case MVT::v2i1:

llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp

static Instruction *SimplifyNVVMIntrinsic(IntrinsicInst *II, InstCombiner &IC) {
   // 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) {
     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;
   }
 
   // Simplify to target-generic intrinsic.
   if (Action.IID) {
     SmallVector<Value *, 4> Args(II->arg_operands());

llvm/unittests/ADT/FloatingPointMode.cpp

 //===- llvm/unittest/ADT/FloatingPointMode.cpp ----------------------------===//
 //
 // 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 "llvm/ADT/FloatingPointMode.h"
 #include "gtest/gtest.h"
 
 using namespace llvm;
 
 namespace {
 
-TEST(FloatingPointModeTest, ParseDenormalFPAttribute) {
-  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttribute("ieee"));
-  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttribute(""));
+TEST(FloatingPointModeTest, ParseDenormalFPAttributeComponent) {
+  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttributeComponent("ieee"));
+  EXPECT_EQ(DenormalMode::IEEE, parseDenormalFPAttributeComponent(""));
   EXPECT_EQ(DenormalMode::PreserveSign,
-            parseDenormalFPAttribute("preserve-sign"));
+            parseDenormalFPAttributeComponent("preserve-sign"));
   EXPECT_EQ(DenormalMode::PositiveZero,
-            parseDenormalFPAttribute("positive-zero"));
-  EXPECT_EQ(DenormalMode::Invalid, parseDenormalFPAttribute("foo"));
+            parseDenormalFPAttributeComponent("positive-zero"));
+  EXPECT_EQ(DenormalMode::Invalid, parseDenormalFPAttributeComponent("foo"));
 }
 
 TEST(FloatingPointModeTest, DenormalAttributeName) {
-  EXPECT_EQ("ieee", denormalModeName(DenormalMode::IEEE));
-  EXPECT_EQ("preserve-sign", denormalModeName(DenormalMode::PreserveSign));
-  EXPECT_EQ("positive-zero", denormalModeName(DenormalMode::PositiveZero));
-  EXPECT_EQ("", denormalModeName(DenormalMode::Invalid));
+  EXPECT_EQ("ieee", denormalModeKindName(DenormalMode::IEEE));
+  EXPECT_EQ("preserve-sign", denormalModeKindName(DenormalMode::PreserveSign));
+  EXPECT_EQ("positive-zero", denormalModeKindName(DenormalMode::PositiveZero));
+  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());
 }
 
 }