Skip to content

[clang][CodeGen] Remove "unsafe-fp-math" attribute support #162779

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Oct 10, 2025
Merged

[clang][CodeGen] Remove "unsafe-fp-math" attribute support #162779

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 0 additions & 7 deletions clang/lib/CodeGen/CGCall.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand Down
6 changes: 3 additions & 3 deletions clang/lib/CodeGen/CGCall.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
5 changes: 0 additions & 5 deletions clang/lib/CodeGen/CodeGenFunction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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() {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGen/backend-unsupported-error.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
Expand Down
9 changes: 4 additions & 5 deletions clang/test/CodeGen/fp-function-attrs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {{%.+}}, {{%.+}}

Expand All @@ -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"{{.*}} }
10 changes: 4 additions & 6 deletions clang/test/CodeGen/func-attr.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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;
Expand All @@ -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: }
13 changes: 3 additions & 10 deletions clang/test/CodeGenCUDA/propagate-attributes.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand Down Expand Up @@ -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]] = {

Expand All @@ -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"
6 changes: 0 additions & 6 deletions clang/test/CodeGenOpenCL/relaxed-fpmath.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down