arsenm created this revision.
arsenm added reviewers: scanon, cameron.mcinally, spatel, andrew.w.kaylor,
mibintc, SjoerdMeijer.
Herald added subscribers: dexonsmith, hiraditya, tpr, nhaehnle, wdng, jvesely,
jholewinski.
Herald added a project: LLVM.
arsenm added parent revisions: D69878: Consoldiate internal denormal flushing
controls, D69598: Work on cleaning up denormal mode handling.
AMDGPU and x86 at least both have separate controls for whether
denormal results are flushed on output, and for whether denormals are
implicitly treated as 0 as an input. The current DAGCombiner use only
really cares about the input treatment of denormals.
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
@@ -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/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/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/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 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<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 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/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: 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
@@ -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/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
@@ -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/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/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/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
@@ -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/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/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/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
