arsenm updated this revision to Diff 239196.
arsenm added a comment.
Herald added a subscriber: kerbowa.
Rebase
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D69978/new/
https://reviews.llvm.org/D69978
Files:
clang/include/clang/Basic/CodeGenOptions.h
clang/include/clang/Driver/ToolChain.h
clang/lib/Basic/Targets/AMDGPU.cpp
clang/lib/CodeGen/CGCall.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Driver/ToolChains/AMDGPU.cpp
clang/lib/Driver/ToolChains/Clang.cpp
clang/lib/Driver/ToolChains/Cuda.cpp
clang/lib/Frontend/CompilerInvocation.cpp
clang/test/CodeGen/denormalfpmode.c
clang/test/CodeGenCUDA/flush-denormals.cu
clang/test/CodeGenCUDA/propagate-metadata.cu
clang/test/Driver/cl-denorms-are-zero.cl
clang/test/Driver/cuda-flush-denormals-to-zero.cu
clang/test/Driver/denormal-fp-math.c
llvm/docs/LangRef.rst
llvm/include/llvm/ADT/FloatingPointMode.h
llvm/lib/CodeGen/MachineFunction.cpp
llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp
llvm/unittests/ADT/FloatingPointMode.cpp
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());
}
}
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/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/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<int64_t, 8> ByteOffsets(Width, INT64_MAX);
int64_t FirstOffset = INT64_MAX;
@@ -6632,15 +6632,15 @@
Value.getOpcode() == ISD::SRA) {
ConstantSDNode *ShiftOffset =
dyn_cast<ConstantSDNode>(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<bool> 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<LoadSDNode>(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/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/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<DenormalMode>(Str)
+ return StringSwitch<DenormalMode::DenormalModeKind>(Str)
.Cases("", "ieee", DenormalMode::IEEE)
.Case("preserve-sign", DenormalMode::PreserveSign)
.Case("positive-zero", DenormalMode::PositiveZero)
@@ -44,7 +113,7 @@
/// 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/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 <t_floating>` ``value`` to type ``ty2``.
Arguments:
@@ -15825,7 +15838,7 @@
The first argument to the '``llvm.experimental.constrained.fptosi``'
intrinsic must be :ref:`floating point <t_floating>` or :ref:`vector
-<t_vector>` of floating point values.
+<t_vector>` of floating point values.
The second argument specifies the exception behavior as described above.
@@ -15934,7 +15947,7 @@
<t_vector>` 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 <inttype>
@llvm.experimental.constrained.llround(<fptype> <op1>,
metadata <exception behavior>)
-
+
Overview:
"""""""""
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: 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/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/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/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/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/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/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/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/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/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/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/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/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/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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
