Index: clang/include/clang/Basic/CodeGenOptions.h =================================================================== --- clang/include/clang/Basic/CodeGenOptions.h +++ clang/include/clang/Basic/CodeGenOptions.h @@ -164,10 +164,10 @@ 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; Index: clang/include/clang/Driver/ToolChain.h =================================================================== --- clang/include/clang/Driver/ToolChain.h +++ clang/include/clang/Driver/ToolChain.h @@ -617,7 +617,7 @@ 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(); } }; Index: clang/lib/Basic/Targets/AMDGPU.cpp =================================================================== --- clang/lib/Basic/Targets/AMDGPU.cpp +++ clang/lib/Basic/Targets/AMDGPU.cpp @@ -239,7 +239,7 @@ 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. Index: clang/lib/CodeGen/CGCall.cpp =================================================================== --- clang/lib/CodeGen/CGCall.cpp +++ clang/lib/CodeGen/CGCall.cpp @@ -1743,13 +1743,13 @@ 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)); Index: clang/lib/CodeGen/CodeGenModule.cpp =================================================================== --- clang/lib/CodeGen/CodeGenModule.cpp +++ clang/lib/CodeGen/CodeGenModule.cpp @@ -558,7 +558,7 @@ // 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); } Index: clang/lib/Driver/ToolChains/AMDGPU.cpp =================================================================== --- clang/lib/Driver/ToolChains/AMDGPU.cpp +++ clang/lib/Driver/ToolChains/AMDGPU.cpp @@ -108,14 +108,14 @@ 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 @@ 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( Index: clang/lib/Driver/ToolChains/Clang.cpp =================================================================== --- clang/lib/Driver/ToolChains/Clang.cpp +++ clang/lib/Driver/ToolChains/Clang.cpp @@ -2273,7 +2273,7 @@ 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(); } @@ -2281,7 +2281,7 @@ 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(); } @@ -2385,14 +2385,18 @@ 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 (DenormalFP32Math != llvm::DenormalMode::Invalid) { - CmdArgs.push_back(Args.MakeArgString( - "-fdenormal-fp-math-f32=" + llvm::subnormalModeName(DenormalFP32Math))); + 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.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()) Index: clang/lib/Driver/ToolChains/Cuda.cpp =================================================================== --- clang/lib/Driver/ToolChains/Cuda.cpp +++ clang/lib/Driver/ToolChains/Cuda.cpp @@ -722,11 +722,11 @@ 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 { Index: clang/lib/Frontend/CompilerInvocation.cpp =================================================================== --- clang/lib/Frontend/CompilerInvocation.cpp +++ clang/lib/Frontend/CompilerInvocation.cpp @@ -1264,14 +1264,14 @@ 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; } Index: clang/test/CodeGen/denormalfpmode.c =================================================================== --- clang/test/CodeGen/denormalfpmode.c +++ 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-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; Index: clang/test/CodeGenCUDA/flush-denormals.cu =================================================================== --- clang/test/CodeGenCUDA/flush-denormals.cu +++ 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" -// 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 Index: clang/test/CodeGenCUDA/propagate-metadata.cu =================================================================== --- clang/test/CodeGenCUDA/propagate-metadata.cu +++ clang/test/CodeGenCUDA/propagate-metadata.cu @@ -59,8 +59,8 @@ // 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" Index: clang/test/Driver/cl-denorms-are-zero.cl =================================================================== --- clang/test/Driver/cl-denorms-are-zero.cl +++ 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" Index: clang/test/Driver/cuda-flush-denormals-to-zero.cu =================================================================== --- clang/test/Driver/cuda-flush-denormals-to-zero.cu +++ 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" -// NOFTZ: "-fdenormal-fp-math=ieee" +// FTZ: "-fdenormal-fp-math-f32=preserve-sign,preserve-sign" +// NOFTZ: "-fdenormal-fp-math=ieee,ieee" Index: clang/test/Driver/denormal-fp-math.c =================================================================== --- clang/test/Driver/denormal-fp-math.c +++ 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-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' Index: llvm/docs/LangRef.rst =================================================================== --- llvm/docs/LangRef.rst +++ llvm/docs/LangRef.rst @@ -1820,15 +1820,27 @@ 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. + This indicates the subnormal handling that may be assumed for the + default floating-point environment. This is a comma separated + pair. The elements may be one of ``"ieee"``, ``"preserve-sign"``, + or ``"positive-zero"``. The first entry indicates the flushing mode + for the result of floating point operations. The second indicates + the handling of denormal inputs 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 Index: llvm/include/llvm/ADT/FloatingPointMode.h =================================================================== --- llvm/include/llvm/ADT/FloatingPointMode.h +++ llvm/include/llvm/ADT/FloatingPointMode.h @@ -14,28 +14,97 @@ #define LLVM_FLOATINGPOINTMODE_H #include "llvm/ADT/StringSwitch.h" +#include "llvm/Support/raw_ostream.h" namespace llvm { -/// Represent handled modes for subnormal (aka denormal) modes in the floating -/// point environment. -enum class DenormalMode { - Invalid = -1, +/// 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 DenormalModeKind : char { + Invalid = -1, - /// IEEE-754 subnormal numbers preserved. - IEEE, + /// IEEE-754 subnormal numbers preserved. + IEEE, - /// The sign of a flushed-to-zero number is preserved in the sign of 0 - PreserveSign, + /// The sign of a flushed-to-zero number is preserved in the sign of 0 + PreserveSign, - /// Denormals are flushed to positive zero. - PositiveZero + /// 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(Str) + return StringSwitch(Str) .Cases("", "ieee", DenormalMode::IEEE) .Case("preserve-sign", DenormalMode::PreserveSign) .Case("positive-zero", DenormalMode::PositiveZero) @@ -44,7 +113,7 @@ /// 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"; @@ -57,6 +126,26 @@ } } +/// 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 Index: llvm/lib/CodeGen/MachineFunction.cpp =================================================================== --- llvm/lib/CodeGen/MachineFunction.cpp +++ llvm/lib/CodeGen/MachineFunction.cpp @@ -290,7 +290,7 @@ // target by default. StringRef Val = Attr.getValueAsString(); if (Val.empty()) - return DenormalMode::Invalid; + return DenormalMode::getInvalid(); return parseDenormalFPAttribute(Val); } Index: llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp =================================================================== --- llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp +++ llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp @@ -20467,7 +20467,10 @@ 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); Index: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp =================================================================== --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -122,7 +122,7 @@ return FtzEnabled; } - return MF.getDenormalMode(APFloat::IEEEsingle()) == + return MF.getDenormalMode(APFloat::IEEEsingle()).Output == DenormalMode::PreserveSign; } Index: llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp =================================================================== --- llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp +++ llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp @@ -1707,7 +1707,8 @@ 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; Index: llvm/unittests/ADT/FloatingPointMode.cpp =================================================================== --- llvm/unittests/ADT/FloatingPointMode.cpp +++ llvm/unittests/ADT/FloatingPointMode.cpp @@ -13,21 +13,122 @@ 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()); } }