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 @@ -242,7 +242,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 @@ -1749,14 +1749,14 @@ 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)); - - if (CodeGenOpts.FP32DenormalMode != llvm::DenormalMode::Invalid) + CodeGenOpts.FPDenormalMode.str()); + if (CodeGenOpts.FP32DenormalMode.isValid()) { FuncAttrs.addAttribute( "denormal-fp-math-f32", - llvm::denormalModeName(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 @@ -567,7 +567,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 @@ -2631,7 +2631,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(); } @@ -2639,7 +2639,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(); } @@ -2758,7 +2758,7 @@ 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 ; @@ -2806,14 +2806,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::denormalModeName(DenormalFPMath))); - } - - if (DenormalFP32Math != llvm::DenormalMode::Invalid) { - CmdArgs.push_back(Args.MakeArgString( - "-fdenormal-fp-math-f32=" + llvm::denormalModeName(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 @@ -723,11 +723,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 @@ -1272,14 +1272,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 @@ -61,8 +61,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,12 +1820,19 @@ not introduce any new floating-point instructions that may trap. ``"denormal-fp-math"`` - This indicates the denormal (subnormal) handling that may be assumed - for the default floating-point environment. This may be one of - ``"ieee"``, ``"preserve-sign"``, or ``"positive-zero"``. If this - is attribute is not specified, the default is ``"ieee"``. If the - mode is ``"preserve-sign"``, or ``"positive-zero"``, denormal - outputs may be flushed to zero by standard floating point + This indicates the denormal (subnormal) handling that may be + assumed for the default floating-point environment. This is a comma + separated pair. The elements may be one of ``"ieee"``, + ``"preserve-sign"``, or ``"positive-zero"``. The first entry + indicates the flushing mode for the result of floating point + operations. The second indicates the handling of denormal inputs 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"``, + 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 +1841,12 @@ 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. If an instruction does not respect this mode, the input + should be converted to 0 as if by ``@llvm.canonicalize`` during + lowering. + ``"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 @@ -15466,9 +15479,9 @@ data arguments and the return value are the same as the corresponding FP operation. -The rounding mode argument is a metadata string specifying what -assumptions, if any, the optimizer can make when transforming constant -values. Some constrained FP intrinsics omit this argument. If required +The rounding mode argument is a metadata string specifying what +assumptions, if any, the optimizer can make when transforming constant +values. Some constrained FP intrinsics omit this argument. If required by the intrinsic, this argument must be one of the following strings: :: @@ -15784,7 +15797,7 @@ Overview: """"""""" -The '``llvm.experimental.constrained.fptoui``' intrinsic converts a +The '``llvm.experimental.constrained.fptoui``' intrinsic converts a floating-point ``value`` to its unsigned integer equivalent of type ``ty2``. Arguments: @@ -15817,7 +15830,7 @@ Overview: """"""""" -The '``llvm.experimental.constrained.fptosi``' intrinsic converts +The '``llvm.experimental.constrained.fptosi``' intrinsic converts :ref:`floating-point ` ``value`` to type ``ty2``. Arguments: @@ -15825,7 +15838,7 @@ The first argument to the '``llvm.experimental.constrained.fptosi``' intrinsic must be :ref:`floating point ` or :ref:`vector -` of floating point values. +` of floating point values. The second argument specifies the exception behavior as described above. @@ -15934,7 +15947,7 @@ ` of floating point values. This argument must be larger in size than the result. -The second and third arguments specify the rounding mode and exception +The second and third arguments specify the rounding mode and exception behavior as described above. Semantics: @@ -15958,7 +15971,7 @@ Overview: """"""""" -The '``llvm.experimental.constrained.fpext``' intrinsic extends a +The '``llvm.experimental.constrained.fpext``' intrinsic extends a floating-point ``value`` to a larger floating-point value. Arguments: @@ -16912,7 +16925,7 @@ declare @llvm.experimental.constrained.llround( , metadata ) - + Overview: """"""""" 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 denormal (aka subnormal) modes in the floating -/// point environment. -enum class DenormalMode { - Invalid = -1, +/// Represent ssubnormal handling kind for floating point instruction inputs and +/// outputs. +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. - IEEE, + /// IEEE-754 denormal 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 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 @@ } } +/// 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 @@ -6610,9 +6610,9 @@ if (LegalOperations && !TLI.isOperationLegal(ISD::STORE, VT)) return SDValue(); - // Check if all the bytes of the combined value we are looking at are stored - // to the same base address. Collect bytes offsets from Base address into - // ByteOffsets. + // Check if all the bytes of the combined value we are looking at are stored + // to the same base address. Collect bytes offsets from Base address into + // ByteOffsets. SDValue CombinedValue; SmallVector ByteOffsets(Width, INT64_MAX); int64_t FirstOffset = INT64_MAX; @@ -6632,15 +6632,15 @@ Value.getOpcode() == ISD::SRA) { ConstantSDNode *ShiftOffset = dyn_cast(Value.getOperand(1)); - // Trying to match the following pattern. The shift offset must be + // Trying to match the following pattern. The shift offset must be // a constant and a multiple of 8. It is the byte offset in "y". - // + // // x = srl y, offset - // i8 z = trunc x + // i8 z = trunc x // store z, ... if (!ShiftOffset || (ShiftOffset->getSExtValue() % 8)) return SDValue(); - + Offset = ShiftOffset->getSExtValue()/8; Value = Value.getOperand(0); } @@ -6685,7 +6685,7 @@ assert(FirstOffset != INT64_MAX && "First byte offset must be set"); assert(FirstStore && "First store must be set"); - // Check if the bytes of the combined value we are looking at match with + // Check if the bytes of the combined value we are looking at match with // either big or little endian value store. Optional IsBigEndian = isBigEndian(ByteOffsets, FirstOffset); if (!IsBigEndian.hasValue()) @@ -8609,7 +8609,7 @@ // Create the actual or node if we can generate good code for it. if (!normalizeToSequence) { SDValue Or = DAG.getNode(ISD::OR, DL, N0.getValueType(), N0, N2_0); - return DAG.getNode(ISD::SELECT, DL, N1.getValueType(), Or, N1, + return DAG.getNode(ISD::SELECT, DL, N1.getValueType(), Or, N1, N2_2, Flags); } // Otherwise see if we can optimize to a better pattern. @@ -10480,7 +10480,7 @@ LoadSDNode *LN0 = cast(N0); // Reducing the width of a volatile load is illegal. For atomics, we may be - // able to reduce the width provided we never widen again. (see D66309) + // able to reduce the width provided we never widen again. (see D66309) if (!LN0->isSimple() || !isLegalNarrowLdSt(LN0, ExtType, ExtVT, ShAmt)) return SDValue(); @@ -20786,7 +20786,10 @@ 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); Index: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp =================================================================== --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -123,7 +123,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 @@ -1713,7 +1713,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()); } }