Commit 737394c4 authored by Jeremy Morse's avatar Jeremy Morse
Browse files

Revert "clang: Treat ieee mode as the default for denormal-fp-math" 

This reverts commit c64ca930.

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.
parent cada5b88

clang/include/clang/Basic/CodeGenOptions.h

+3 −3
Original line number Diff line number Diff line
@@ -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;

clang/include/clang/Driver/ToolChain.h

+2 −1
Original line number Diff line number Diff line
@@ -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();

  }

};

clang/lib/Driver/ToolChains/Clang.cpp

+5 −14
Original line number Diff line number Diff line
@@ -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;

clang/test/CodeGenCUDA/flush-denormals.cu

+7 −2
Original line number Diff line number Diff line 
// 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

clang/test/CodeGenCUDA/propagate-metadata.cu

+4 −4
Original line number Diff line number Diff line
@@ -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"