From 737394c490444e968a6f640b99a6614567ca7f28 Mon Sep 17 00:00:00 2001 From: Jeremy Morse Date: Thu, 5 Mar 2020 10:45:52 +0000 Subject: [PATCH] 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. --- clang/include/clang/Basic/CodeGenOptions.h | 6 +++--- clang/include/clang/Driver/ToolChain.h | 3 ++- clang/lib/Driver/ToolChains/Clang.cpp | 19 +++++-------------- clang/test/CodeGenCUDA/flush-denormals.cu | 9 +++++++-- clang/test/CodeGenCUDA/propagate-metadata.cu | 8 ++++---- clang/test/CodeGenOpenCL/amdgpu-features.cl | 14 +++++++------- clang/test/Driver/cuda-flush-denormals-to-zero.cu | 6 +----- clang/test/Driver/denormal-fp-math.c | 3 +-- clang/test/Driver/fp-model.c | 4 ---- 9 files changed, 30 insertions(+), 42 deletions(-) diff --git a/clang/include/clang/Basic/CodeGenOptions.h b/clang/include/clang/Basic/CodeGenOptions.h index d435beb..0a28ede 100644 --- a/clang/include/clang/Basic/CodeGenOptions.h +++ b/clang/include/clang/Basic/CodeGenOptions.h @@ -164,10 +164,10 @@ public: 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 88ce205..400ff9d 100644 --- a/clang/include/clang/Driver/ToolChain.h +++ b/clang/include/clang/Driver/ToolChain.h @@ -623,7 +623,8 @@ public: 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 c8da2d8..d387a1d 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 different. - 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 2753386..a372f3f 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 2514eeb..4b19769 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 b4ee2ed..d8eb2d2 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 84ef3587..c032732 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 63a2c3b..af18517 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 8e61b44..8bf53f6 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" -- 2.7.4