[clang] 737394c - Revert "clang: Treat ieee mode as the default for denormal-fp-math"
Jeremy Morse via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 5 02:55:49 PST 2020
Author: Jeremy Morse
Date: 2020-03-05T10:55:24Z
New Revision: 737394c490444e968a6f640b99a6614567ca7f28
URL: https://github.com/llvm/llvm-project/commit/737394c490444e968a6f640b99a6614567ca7f28
DIFF: https://github.com/llvm/llvm-project/commit/737394c490444e968a6f640b99a6614567ca7f28.diff
LOG: Revert "clang: Treat ieee mode as the default for denormal-fp-math"
This reverts commit c64ca93053af235bac0ca4dcdcd21c8882478310.
This patch tripped a few build bots:
http://lab.llvm.org:8011/builders/clang-x86_64-debian-fast/builds/24703/
http://lab.llvm.org:8011/builders/clang-cmake-x86_64-avx2-linux/builds/13465/
http://lab.llvm.org:8011/builders/clang-with-lto-ubuntu/builds/15994/
Reverting to clear the bots.
Added:
Modified:
clang/include/clang/Basic/CodeGenOptions.h
clang/include/clang/Driver/ToolChain.h
clang/lib/Driver/ToolChains/Clang.cpp
clang/test/CodeGenCUDA/flush-denormals.cu
clang/test/CodeGenCUDA/propagate-metadata.cu
clang/test/CodeGenOpenCL/amdgpu-features.cl
clang/test/Driver/cuda-flush-denormals-to-zero.cu
clang/test/Driver/denormal-fp-math.c
clang/test/Driver/fp-model.c
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/CodeGenOptions.h b/clang/include/clang/Basic/CodeGenOptions.h
index d435bebcdcf9..0a28edefa1e6 100644
--- a/clang/include/clang/Basic/CodeGenOptions.h
+++ b/clang/include/clang/Basic/CodeGenOptions.h
@@ -164,10 +164,10 @@ class CodeGenOptions : public CodeGenOptionsBase {
std::string FloatABI;
/// The floating-point denormal mode to use.
- llvm::DenormalMode FPDenormalMode = llvm::DenormalMode::getIEEE();
+ llvm::DenormalMode FPDenormalMode;
- /// The floating-point denormal mode to use, for float.
- llvm::DenormalMode FP32DenormalMode = llvm::DenormalMode::getIEEE();
+ /// The floating-point subnormal mode to use, for float.
+ llvm::DenormalMode FP32DenormalMode;
/// The float precision limit to use, if non-empty.
std::string LimitFloatPrecision;
diff --git a/clang/include/clang/Driver/ToolChain.h b/clang/include/clang/Driver/ToolChain.h
index 88ce205228fd..400ff9d86664 100644
--- a/clang/include/clang/Driver/ToolChain.h
+++ b/clang/include/clang/Driver/ToolChain.h
@@ -623,7 +623,8 @@ class ToolChain {
const llvm::opt::ArgList &DriverArgs,
Action::OffloadKind DeviceOffloadKind,
const llvm::fltSemantics *FPType = nullptr) const {
- return llvm::DenormalMode::getIEEE();
+ // FIXME: This should be IEEE when default handling is fixed.
+ return llvm::DenormalMode::getInvalid();
}
};
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index c8da2d86490f..d387a1dc2079 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -2548,13 +2548,8 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
ReciprocalMath = false;
SignedZeros = true;
// -fno_fast_math restores default denormal and fpcontract handling
- FPContract = "";
DenormalFPMath = DefaultDenormalFPMath;
-
- // FIXME: The target may have picked a non-IEEE default mode here based on
- // -cl-denorms-are-zero. Should the target consider -fp-model interaction?
- DenormalFP32Math = DefaultDenormalFP32Math;
-
+ FPContract = "";
StringRef Val = A->getValue();
if (OFastEnabled && !Val.equals("fast")) {
// Only -ffp-model=fast is compatible with OFast, ignore.
@@ -2731,9 +2726,7 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
FPExceptionBehavior = "strict";
// -fno_unsafe_math_optimizations restores default denormal handling
DenormalFPMath = DefaultDenormalFPMath;
-
- // The target may have opted to flush just f32 by default, so force IEEE.
- DenormalFP32Math = llvm::DenormalMode::getIEEE();
+ DenormalFP32Math = DefaultDenormalFP32Math;
break;
case options::OPT_Ofast:
@@ -2774,12 +2767,11 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
if (StrictFPModel) {
// If -ffp-model=strict has been specified on command line but
// subsequent options conflict then emit warning diagnostic.
+ // TODO: How should this interact with DenormalFP32Math?
if (HonorINFs && HonorNaNs &&
!AssociativeMath && !ReciprocalMath &&
SignedZeros && TrappingMath && RoundingFPMath &&
- (FPContract.equals("off") || FPContract.empty()) &&
- DenormalFPMath == llvm::DenormalMode::getIEEE() &&
- DenormalFP32Math == llvm::DenormalMode::getIEEE())
+ (FPContract.equals("off") || FPContract.empty()))
// OK: Current Arg doesn't conflict with -ffp-model=strict
;
else {
@@ -2833,8 +2825,7 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
CmdArgs.push_back(Args.MakeArgString(ArgStr.str()));
}
- // Add f32 specific denormal mode flag if it's
diff erent.
- if (DenormalFP32Math != DenormalFPMath) {
+ if (DenormalFP32Math.isValid()) {
llvm::SmallString<64> DenormFlag;
llvm::raw_svector_ostream ArgStr(DenormFlag);
ArgStr << "-fdenormal-fp-math-f32=" << DenormalFP32Math;
diff --git a/clang/test/CodeGenCUDA/flush-denormals.cu b/clang/test/CodeGenCUDA/flush-denormals.cu
index 275338635bc6..a372f3faaf58 100644
--- a/clang/test/CodeGenCUDA/flush-denormals.cu
+++ b/clang/test/CodeGenCUDA/flush-denormals.cu
@@ -1,6 +1,6 @@
// RUN: %clang_cc1 -fcuda-is-device \
// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
-// RUN: FileCheck -check-prefix=NOFTZ %s
+// RUN: FileCheck -check-prefix=DEFAULT %s
// RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=ieee \
// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
@@ -10,9 +10,10 @@
// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
// RUN: FileCheck -check-prefix=FTZ %s
+// FIXME: Unspecified should default to ieee
// RUN: %clang_cc1 -fcuda-is-device -x hip \
// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
-// RUN: FileCheck -check-prefix=AMDNOFTZ %s
+// RUN: FileCheck -check-prefix=AMDFTZ %s
// RUN: %clang_cc1 -fcuda-is-device -x hip \
// RUN: -triple amdgcn-amd-amdhsa -target-cpu gfx900 -fdenormal-fp-math-f32=ieee -emit-llvm -o - %s | \
@@ -41,6 +42,10 @@ extern "C" __device__ void foo() {}
// FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
// NOFTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="ieee,ieee"
+
+// FIXME: This should be removed
+// DEFAULT-NOT: "denormal-fp-math-f32"
+
// AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals
// AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals
diff --git a/clang/test/CodeGenCUDA/propagate-metadata.cu b/clang/test/CodeGenCUDA/propagate-metadata.cu
index 2514eeb55d20..4b1976939993 100644
--- a/clang/test/CodeGenCUDA/propagate-metadata.cu
+++ b/clang/test/CodeGenCUDA/propagate-metadata.cu
@@ -15,7 +15,7 @@
// RUN: %s -o %t.bc -triple nvptx-unknown-unknown
// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc -o - \
-// RUN: -fno-trapping-math -fcuda-is-device -triple nvptx-unknown-unknown \
+// RUN: -fno-trapping-math -fcuda-is-device -fdenormal-fp-math-f32=ieee -triple nvptx-unknown-unknown \
// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ --check-prefix=NOFAST
// RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \
@@ -60,7 +60,8 @@ __global__ void kernel() { lib_fn(); }
// CHECK-SAME: convergent
// CHECK-SAME: norecurse
-// FTZ: "denormal-fp-math"="ieee,ieee"
+// FTZ-NOT: "denormal-fp-math"
+
// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
// NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee"
@@ -75,8 +76,7 @@ __global__ void kernel() { lib_fn(); }
// CHECK-SAME: convergent
// CHECK-NOT: norecurse
-// FTZ-SAME: "denormal-fp-math"="ieee,ieee"
-// NOFTZ-SAME: "denormal-fp-math"="ieee,ieee"
+// FTZ-NOT: "denormal-fp-math"
// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
// NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee"
diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl
index b4ee2edc5443..d8eb2d26b0e3 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-features.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl
@@ -14,13 +14,13 @@
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s
-// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime"
-// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime"
+// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime,-fp32-denormals"
+// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,-fp32-denormals"
+// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime,-fp32-denormals"
// GFX700: "target-features"="+ci-insts,+flat-address-space,+fp64-fp16-denormals,-fp32-denormals"
// GFX600: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
// GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
diff --git a/clang/test/Driver/cuda-flush-denormals-to-zero.cu b/clang/test/Driver/cuda-flush-denormals-to-zero.cu
index 84ef358758fd..c03273205414 100644
--- a/clang/test/Driver/cuda-flush-denormals-to-zero.cu
+++ b/clang/test/Driver/cuda-flush-denormals-to-zero.cu
@@ -9,9 +9,5 @@
// CPUFTZ-NOT: -fdenormal-fp-math
-// FTZ-NOT: -fdenormal-fp-math-f32=
// FTZ: "-fdenormal-fp-math-f32=preserve-sign,preserve-sign"
-
-// The default of ieee is omitted
-// NOFTZ-NOT: "-fdenormal-fp-math"
-// NOFTZ-NOT: "-fdenormal-fp-math-f32"
+// NOFTZ: "-fdenormal-fp-math=ieee,ieee"
diff --git a/clang/test/Driver/denormal-fp-math.c b/clang/test/Driver/denormal-fp-math.c
index 63a2c3b7e003..af18517a740a 100644
--- a/clang/test/Driver/denormal-fp-math.c
+++ b/clang/test/Driver/denormal-fp-math.c
@@ -8,8 +8,7 @@
// 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
-// TODO: ieee is the implied default, and the flag is not passed.
-// CHECK-IEEE: "-fdenormal-fp-math=ieee,ieee"
+// 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"
diff --git a/clang/test/Driver/fp-model.c b/clang/test/Driver/fp-model.c
index 8e61b4411cae..8bf53f6d997b 100644
--- a/clang/test/Driver/fp-model.c
+++ b/clang/test/Driver/fp-model.c
@@ -63,10 +63,6 @@
// RUN: | FileCheck --check-prefix=WARNf %s
// WARNf: warning: overriding '-ffp-model=strict' option with '-Ofast' [-Woverriding-t-option]
-// RUN: %clang -### -ffp-model=strict -fdenormal-fp-math=preserve-sign,preserve-sign -c %s 2>&1 \
-// RUN: | FileCheck --check-prefix=WARN10 %s
-// WARN10: warning: overriding '-ffp-model=strict' option with '-fdenormal-fp-math=preserve-sign,preserve-sign' [-Woverriding-t-option]
-
// RUN: %clang -### -c %s 2>&1 \
// RUN: | FileCheck --check-prefix=CHECK-NOROUND %s
// CHECK-NOROUND: "-cc1"
More information about the cfe-commits
mailing list