[clang] c64ca93 - clang: Treat ieee mode as the default for denormal-fp-math
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 4 20:34:11 PST 2020
Author: Matt Arsenault
Date: 2020-03-04T23:34:02-05:00
New Revision: c64ca93053af235bac0ca4dcdcd21c8882478310
URL: https://github.com/llvm/llvm-project/commit/c64ca93053af235bac0ca4dcdcd21c8882478310
DIFF: https://github.com/llvm/llvm-project/commit/c64ca93053af235bac0ca4dcdcd21c8882478310.diff
LOG: clang: Treat ieee mode as the default for denormal-fp-math
The IR hasn't switched the default yet, so explicitly add the ieee
attributes.
I'm still not really sure how the target default denormal mode should
interact with -fno-unsafe-math-optimizations. The target may have
selected the default mode to be non-IEEE based on the flags or based
on its true behavior, but we don't know which is the case. Since the
only users of a non-IEEE mode without a flag still support IEEE mode,
just reset to IEEE.
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 0a28edefa1e6..d435bebcdcf9 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 FPDenormalMode = llvm::DenormalMode::getIEEE();
- /// The floating-point subnormal mode to use, for float.
- llvm::DenormalMode FP32DenormalMode;
+ /// The floating-point denormal mode to use, for float.
+ llvm::DenormalMode FP32DenormalMode = llvm::DenormalMode::getIEEE();
/// 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 400ff9d86664..88ce205228fd 100644
--- a/clang/include/clang/Driver/ToolChain.h
+++ b/clang/include/clang/Driver/ToolChain.h
@@ -623,8 +623,7 @@ class ToolChain {
const llvm::opt::ArgList &DriverArgs,
Action::OffloadKind DeviceOffloadKind,
const llvm::fltSemantics *FPType = nullptr) const {
- // FIXME: This should be IEEE when default handling is fixed.
- return llvm::DenormalMode::getInvalid();
+ return llvm::DenormalMode::getIEEE();
}
};
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index d387a1dc2079..c8da2d86490f 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -2548,8 +2548,13 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
ReciprocalMath = false;
SignedZeros = true;
// -fno_fast_math restores default denormal and fpcontract handling
- DenormalFPMath = DefaultDenormalFPMath;
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;
+
StringRef Val = A->getValue();
if (OFastEnabled && !Val.equals("fast")) {
// Only -ffp-model=fast is compatible with OFast, ignore.
@@ -2726,7 +2731,9 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
FPExceptionBehavior = "strict";
// -fno_unsafe_math_optimizations restores default denormal handling
DenormalFPMath = DefaultDenormalFPMath;
- DenormalFP32Math = DefaultDenormalFP32Math;
+
+ // The target may have opted to flush just f32 by default, so force IEEE.
+ DenormalFP32Math = llvm::DenormalMode::getIEEE();
break;
case options::OPT_Ofast:
@@ -2767,11 +2774,12 @@ 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()))
+ (FPContract.equals("off") || FPContract.empty()) &&
+ DenormalFPMath == llvm::DenormalMode::getIEEE() &&
+ DenormalFP32Math == llvm::DenormalMode::getIEEE())
// OK: Current Arg doesn't conflict with -ffp-model=strict
;
else {
@@ -2825,7 +2833,8 @@ static void RenderFloatingPointOptions(const ToolChain &TC, const Driver &D,
CmdArgs.push_back(Args.MakeArgString(ArgStr.str()));
}
- if (DenormalFP32Math.isValid()) {
+ // Add f32 specific denormal mode flag if it's
diff erent.
+ if (DenormalFP32Math != DenormalFPMath) {
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 a372f3faaf58..275338635bc6 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=DEFAULT %s
+// RUN: FileCheck -check-prefix=NOFTZ %s
// RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=ieee \
// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
@@ -10,10 +10,9 @@
// 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=AMDFTZ %s
+// RUN: FileCheck -check-prefix=AMDNOFTZ %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 | \
@@ -42,10 +41,6 @@ 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 4b1976939993..2514eeb55d20 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 -fdenormal-fp-math-f32=ieee -triple nvptx-unknown-unknown \
+// RUN: -fno-trapping-math -fcuda-is-device -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,8 +60,7 @@ __global__ void kernel() { lib_fn(); }
// CHECK-SAME: convergent
// CHECK-SAME: norecurse
-// FTZ-NOT: "denormal-fp-math"
-
+// FTZ: "denormal-fp-math"="ieee,ieee"
// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
// NOFTZ-SAME: "denormal-fp-math-f32"="ieee,ieee"
@@ -76,7 +75,8 @@ __global__ void kernel() { lib_fn(); }
// CHECK-SAME: convergent
// CHECK-NOT: norecurse
-// FTZ-NOT: "denormal-fp-math"
+// FTZ-SAME: "denormal-fp-math"="ieee,ieee"
+// NOFTZ-SAME: "denormal-fp-math"="ieee,ieee"
// 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 d8eb2d26b0e3..b4ee2edc5443 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,+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"
+// 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"
// 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 c03273205414..84ef358758fd 100644
--- a/clang/test/Driver/cuda-flush-denormals-to-zero.cu
+++ b/clang/test/Driver/cuda-flush-denormals-to-zero.cu
@@ -9,5 +9,9 @@
// CPUFTZ-NOT: -fdenormal-fp-math
+// FTZ-NOT: -fdenormal-fp-math-f32=
// FTZ: "-fdenormal-fp-math-f32=preserve-sign,preserve-sign"
-// NOFTZ: "-fdenormal-fp-math=ieee,ieee"
+
+// The default of ieee is omitted
+// NOFTZ-NOT: "-fdenormal-fp-math"
+// NOFTZ-NOT: "-fdenormal-fp-math-f32"
diff --git a/clang/test/Driver/denormal-fp-math.c b/clang/test/Driver/denormal-fp-math.c
index af18517a740a..63a2c3b7e003 100644
--- a/clang/test/Driver/denormal-fp-math.c
+++ b/clang/test/Driver/denormal-fp-math.c
@@ -8,7 +8,8 @@
// 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,ieee
+// TODO: ieee is the implied default, and the flag is not passed.
+// 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 8bf53f6d997b..8e61b4411cae 100644
--- a/clang/test/Driver/fp-model.c
+++ b/clang/test/Driver/fp-model.c
@@ -63,6 +63,10 @@
// 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