diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index c5371e4c841cb..df28641904c04 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2012,13 +2012,6 @@ static void getTrivialDefaultFunctionAttributes( FuncAttrs.addAttribute("no-infs-fp-math", "true"); if (LangOpts.NoHonorNaNs) FuncAttrs.addAttribute("no-nans-fp-math", "true"); - if (LangOpts.AllowFPReassoc && LangOpts.AllowRecip && - LangOpts.NoSignedZero && LangOpts.ApproxFunc && - (LangOpts.getDefaultFPContractMode() == - LangOptions::FPModeKind::FPM_Fast || - LangOpts.getDefaultFPContractMode() == - LangOptions::FPModeKind::FPM_FastHonorPragmas)) - FuncAttrs.addAttribute("unsafe-fp-math", "true"); if (CodeGenOpts.SoftFloat) FuncAttrs.addAttribute("use-soft-float", "true"); FuncAttrs.addAttribute("stack-protector-buffer-size", diff --git a/clang/lib/CodeGen/CGCall.h b/clang/lib/CodeGen/CGCall.h index 935b5086f5983..1ef8a3f114573 100644 --- a/clang/lib/CodeGen/CGCall.h +++ b/clang/lib/CodeGen/CGCall.h @@ -410,10 +410,10 @@ class ReturnValueSlot { /// This is useful for adding attrs to bitcode modules that you want to link /// with but don't control, such as CUDA's libdevice. When linking with such /// a bitcode library, you might want to set e.g. its functions' -/// "unsafe-fp-math" attribute to match the attr of the functions you're +/// "denormal-fp-math" attribute to match the attr of the functions you're /// codegen'ing. Otherwise, LLVM will interpret the bitcode module's lack of -/// unsafe-fp-math attrs as tantamount to unsafe-fp-math=false, and then LLVM -/// will propagate unsafe-fp-math=false up to every transitive caller of a +/// denormal-fp-math attrs as tantamount to denormal-fp-math=ieee, and then LLVM +/// will propagate denormal-fp-math=ieee up to every transitive caller of a /// function in the bitcode library! /// /// With the exception of fast-math attrs, this will only make the attributes diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index acf8de4dee147..88628530cf66b 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -183,11 +183,6 @@ void CodeGenFunction::CGFPOptionsRAII::ConstructorHelper(FPOptions FPFeatures) { mergeFnAttrValue("no-infs-fp-math", FPFeatures.getNoHonorInfs()); mergeFnAttrValue("no-nans-fp-math", FPFeatures.getNoHonorNaNs()); mergeFnAttrValue("no-signed-zeros-fp-math", FPFeatures.getNoSignedZero()); - mergeFnAttrValue( - "unsafe-fp-math", - FPFeatures.getAllowFPReassociate() && FPFeatures.getAllowReciprocal() && - FPFeatures.getAllowApproxFunc() && FPFeatures.getNoSignedZero() && - FPFeatures.allowFPContractAcrossStatement()); } CodeGenFunction::CGFPOptionsRAII::~CGFPOptionsRAII() { diff --git a/clang/test/CodeGen/backend-unsupported-error.ll b/clang/test/CodeGen/backend-unsupported-error.ll index 2de2c87aa720d..47b397b6bb40b 100644 --- a/clang/test/CodeGen/backend-unsupported-error.ll +++ b/clang/test/CodeGen/backend-unsupported-error.ll @@ -21,7 +21,7 @@ entry: ret i32 %call, !dbg !15 } -attributes #0 = { nounwind noinline "disable-tail-calls"="false" "less-precise-fpmad"="false" "frame-pointer"="all" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #0 = { nounwind noinline "disable-tail-calls"="false" "less-precise-fpmad"="false" "frame-pointer"="all" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "use-soft-float"="false" } !llvm.dbg.cu = !{!0} !llvm.module.flags = !{!9, !10} diff --git a/clang/test/CodeGen/fp-function-attrs.cpp b/clang/test/CodeGen/fp-function-attrs.cpp index e92c26cdb75f7..3775bd5452d78 100644 --- a/clang/test/CodeGen/fp-function-attrs.cpp +++ b/clang/test/CodeGen/fp-function-attrs.cpp @@ -36,7 +36,7 @@ float test_reassociate_off_pragma(float a, float b, float c) { return tmp; } -// CHECK: define{{.*}} float @_Z27test_reassociate_off_pragmafff(float noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %b, float noundef nofpclass(nan inf) %c) [[NO_UNSAFE_ATTRS:#[0-9]+]] +// CHECK: define{{.*}} float @_Z27test_reassociate_off_pragmafff(float noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %b, float noundef nofpclass(nan inf) %c) // CHECK: fadd nnan ninf nsz arcp contract afn float {{%.+}}, {{%.+}} // CHECK: fadd fast float {{%.+}}, {{%.+}} @@ -49,10 +49,9 @@ float test_contract_on_pragma(float a, float b, float c) { return tmp; } -// CHECK: define{{.*}} float @_Z23test_contract_on_pragmafff(float noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %b, float noundef nofpclass(nan inf) %c) [[NO_UNSAFE_ATTRS:#[0-9]+]] +// CHECK: define{{.*}} float @_Z23test_contract_on_pragmafff(float noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %b, float noundef nofpclass(nan inf) %c) // CHECK: fmul fast float {{%.+}}, {{%.+}} // CHECK: fadd reassoc nnan ninf nsz arcp afn float {{%.+}}, {{%.+}} -// CHECK: attributes [[FAST_ATTRS]] = { {{.*}}"no-infs-fp-math"="true" {{.*}}"no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" {{.*}}"unsafe-fp-math"="true"{{.*}} } -// CHECK: attributes [[PRECISE_ATTRS]] = { {{.*}}"no-infs-fp-math"="false" {{.*}}"no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" {{.*}}"unsafe-fp-math"="false"{{.*}} } -// CHECK: attributes [[NO_UNSAFE_ATTRS]] = { {{.*}}"no-infs-fp-math"="true" {{.*}}"no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" {{.*}}"unsafe-fp-math"="false"{{.*}} } +// CHECK: attributes [[FAST_ATTRS]] = { {{.*}}"no-infs-fp-math"="true" {{.*}}"no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true"{{.*}} } +// CHECK: attributes [[PRECISE_ATTRS]] = { {{.*}}"no-infs-fp-math"="false" {{.*}}"no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false"{{.*}} } diff --git a/clang/test/CodeGen/func-attr.c b/clang/test/CodeGen/func-attr.c index 1b36c51b8ce9c..42b7f1c785a64 100644 --- a/clang/test/CodeGen/func-attr.c +++ b/clang/test/CodeGen/func-attr.c @@ -1,18 +1,18 @@ // RUN: %clang_cc1 -triple x86_64-linux-gnu -ffast-math \ // RUN: -ffp-contract=fast -emit-llvm -o - %s | \ -// RUN: FileCheck %s --check-prefixes=CHECK,CHECK-UNSAFE,FINITEONLY +// RUN: FileCheck %s --check-prefixes=CHECK,FINITEONLY // RUN: %clang_cc1 -triple x86_64-linux-gnu -funsafe-math-optimizations \ // RUN: -ffp-contract=fast -emit-llvm -o - %s | \ -// RUN: FileCheck %s --check-prefixes=CHECK,CHECK-UNSAFE,NOFINITEONLY +// RUN: FileCheck %s --check-prefixes=CHECK,NOFINITEONLY // RUN: %clang_cc1 -triple x86_64-linux-gnu -funsafe-math-optimizations \ // RUN: -ffp-contract=on -emit-llvm -o - %s | \ -// RUN: FileCheck %s --check-prefixes=CHECK,CHECK-NOUNSAFE,NOFINITEONLY +// RUN: FileCheck %s --check-prefixes=CHECK,NOFINITEONLY // RUN: %clang_cc1 -triple x86_64-linux-gnu -funsafe-math-optimizations \ // RUN: -ffp-contract=off -emit-llvm -o - %s | \ -// RUN: FileCheck %s --check-prefixes=CHECK,CHECK-NOUNSAFE,NOFINITEONLY +// RUN: FileCheck %s --check-prefixes=CHECK,NOFINITEONLY float foo(float a, float b) { return a+b; @@ -24,6 +24,4 @@ float foo(float a, float b) { // CHECK: attributes [[ATTRS]] = { // CHECK-SAME: "no-signed-zeros-fp-math"="true" // CHECK-SAME: "no-trapping-math"="true" -// CHECK-UNSAFE-SAME: "unsafe-fp-math"="true" -// CHECK-NOUNSAFE-NOT: "unsafe-fp-math"="true" // CHECK-SAME: } diff --git a/clang/test/CodeGenCUDA/propagate-attributes.cu b/clang/test/CodeGenCUDA/propagate-attributes.cu index 6dfd44487d1dc..a7e6b09ff97db 100644 --- a/clang/test/CodeGenCUDA/propagate-attributes.cu +++ b/clang/test/CodeGenCUDA/propagate-attributes.cu @@ -12,18 +12,17 @@ // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc -o - \ // RUN: -fcuda-is-device -triple nvptx-unknown-unknown \ -// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ --check-prefix=NOFAST +// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=NOFTZ // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \ // RUN: -fdenormal-fp-math-f32=preserve-sign -o - \ // RUN: -fcuda-is-device -triple nvptx-unknown-unknown \ -// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ \ -// RUN: --check-prefix=NOFAST +// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FTZ // RUN: %clang_cc1 -x cuda %s -emit-llvm -mlink-builtin-bitcode %t.bc \ // RUN: -fdenormal-fp-math-f32=preserve-sign -o - \ // RUN: -fcuda-is-device -funsafe-math-optimizations -triple nvptx-unknown-unknown \ -// RUN: | FileCheck %s --check-prefix=CHECK --check-prefix=FAST +// RUN: | FileCheck %s --check-prefix=CHECK #ifndef LIB #include "Inputs/cuda.h" @@ -65,9 +64,6 @@ __global__ void kernel() { lib_fn(); } // CHECK-SAME: "no-trapping-math"="true" -// FAST-SAME: "unsafe-fp-math"="true" -// NOFAST-NOT: "unsafe-fp-math"="true" - // Check the attribute list for lib_fn. // CHECK: attributes [[fattr]] = { @@ -81,6 +77,3 @@ __global__ void kernel() { lib_fn(); } // NOFTZ-NOT: "denormal-fp-math-f32" // CHECK-SAME: "no-trapping-math"="true" - -// FAST-SAME: "unsafe-fp-math"="true" -// NOFAST-NOT: "unsafe-fp-math"="true" diff --git a/clang/test/CodeGenOpenCL/relaxed-fpmath.cl b/clang/test/CodeGenOpenCL/relaxed-fpmath.cl index a5f0019dbc1e5..c113d23c3028d 100644 --- a/clang/test/CodeGenOpenCL/relaxed-fpmath.cl +++ b/clang/test/CodeGenOpenCL/relaxed-fpmath.cl @@ -33,37 +33,31 @@ float spscalardiv(float a, float b) { // NORMAL-NOT: "no-infs-fp-math" // NORMAL-NOT: "no-nans-fp-math" // NORMAL-NOT: "no-signed-zeros-fp-math" -// NORMAL-NOT: "unsafe-fp-math" // FAST: "less-precise-fpmad"="true" // FAST: "no-infs-fp-math"="true" // FAST: "no-nans-fp-math"="true" // FAST: "no-signed-zeros-fp-math"="true" -// FAST: "unsafe-fp-math"="true" // FINITE-NOT: "less-precise-fpmad" // FINITE: "no-infs-fp-math"="true" // FINITE: "no-nans-fp-math"="true" // FINITE-NOT: "no-signed-zeros-fp-math" -// FINITE-NOT: "unsafe-fp-math" // UNSAFE: "less-precise-fpmad"="true" // UNSAFE-NOT: "no-infs-fp-math" // UNSAFE-NOT: "no-nans-fp-math" // UNSAFE: "no-signed-zeros-fp-math"="true" -// UNSAFE: "unsafe-fp-math"="true" // MAD: "less-precise-fpmad"="true" // MAD-NOT: "no-infs-fp-math" // MAD-NOT: "no-nans-fp-math" // MAD-NOT: "no-signed-zeros-fp-math" -// MAD-NOT: "unsafe-fp-math" // NOSIGNED-NOT: "less-precise-fpmad" // NOSIGNED-NOT: "no-infs-fp-math" // NOSIGNED-NOT: "no-nans-fp-math" // NOSIGNED: "no-signed-zeros-fp-math"="true" -// NOSIGNED-NOT: "unsafe-fp-math" #else // Undefine this to avoid putting it in the PCH.