[clang] [llvm] [IR] Add `fpmath` to keep list of dropUBImplyingAttrsAndMetadata (PR #179019)
Wenju He via cfe-commits
cfe-commits at lists.llvm.org
Fri Jan 30 22:39:06 PST 2026
https://github.com/wenju-he created https://github.com/llvm/llvm-project/pull/179019
`fpmath` is precision metadata rather than UB-implying metadata. This avoids `fpmath` from dropped in InstCombine FoldOpIntoSelect.
>From 5ecb10c8a571b610dd6ba3413d99651b7c02bad8 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Sat, 31 Jan 2026 07:37:02 +0100
Subject: [PATCH] [IR] Add `fpmath` to keep list of
dropUBImplyingAttrsAndMetadata
`fpmath` is precision metadata rather than UB-implying metadata.
This avoids `fpmath` from dropped in InstCombine FoldOpIntoSelect.
---
clang/test/Headers/__clang_hip_math.hip | 14 +++++++-------
llvm/lib/IR/Instruction.cpp | 3 ++-
.../InstCombine/fold-fops-into-selects.ll | 9 +++++++++
3 files changed, 18 insertions(+), 8 deletions(-)
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 6647d35dc7d55..e91f5723a8460 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -5633,7 +5633,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL5NORMFIPKF_EXIT_LOOPEXIT:.*]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// NCRDIV: [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]]:
-// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract float @llvm.sqrt.f32(float [[ADD_I]])
+// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract float @llvm.sqrt.f32(float [[ADD_I]]), !fpmath [[META22:![0-9]+]]
// NCRDIV-NEXT: br label %[[_ZL5NORMFIPKF_EXIT]]
// NCRDIV: [[_ZL5NORMFIPKF_EXIT]]:
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, %[[ENTRY]] ], [ [[TMP1]], %[[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
@@ -5750,7 +5750,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
-// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL4NORMIPKD_EXIT_LOOPEXIT:.*]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL4NORMIPKD_EXIT_LOOPEXIT:.*]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
// NCRDIV: [[_ZL4NORMIPKD_EXIT_LOOPEXIT]]:
// NCRDIV-NEXT: [[TMP1:%.*]] = tail call contract double @llvm.sqrt.f64(double [[ADD_I]])
// NCRDIV-NEXT: br label %[[_ZL4NORMIPKD_EXIT]]
@@ -6391,7 +6391,7 @@ extern "C" __device__ double test_rint(double x) {
// NCRDIV-NEXT: [[ADD_I]] = fadd contract float [[__R_0_I4]], [[MUL_I]]
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 4
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
-// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL6RNORMFIPKF_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL6RNORMFIPKF_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
// NCRDIV: [[_ZL6RNORMFIPKF_EXIT]]:
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, %[[ENTRY]] ], [ [[ADD_I]], %[[WHILE_BODY_I]] ]
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_rsqrt_f32(float noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
@@ -6500,7 +6500,7 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
// NCRDIV-NEXT: [[ADD_I]] = fadd contract double [[__R_0_I4]], [[MUL_I]]
// NCRDIV-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds nuw i8, ptr [[__A_ADDR_0_I3]], i64 8
// NCRDIV-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
-// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL5RNORMIPKD_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// NCRDIV-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL5RNORMIPKD_EXIT]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
// NCRDIV: [[_ZL5RNORMIPKD_EXIT]]:
// NCRDIV-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, %[[ENTRY]] ], [ [[ADD_I]], %[[WHILE_BODY_I]] ]
// NCRDIV-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_rsqrt_f64(double noundef [[__R_0_I_LCSSA]]) #[[ATTR14]]
@@ -7459,7 +7459,7 @@ extern "C" __device__ double test_sinpi(double x) {
// NCRDIV-LABEL: define dso_local noundef float @test_sqrtf(
// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
-// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X]]), !fpmath [[META25:![0-9]+]]
+// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X]]), !fpmath [[META22]]
// NCRDIV-NEXT: ret float [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_sqrtf(
@@ -9421,10 +9421,10 @@ extern "C" __device__ int test_int_max(int x, int y) {
// NCRDIV: [[DOUBLE_TBAA19]] = !{[[META20:![0-9]+]], [[META20]], i64 0}
// NCRDIV: [[META20]] = !{!"double", [[META6]], i64 0}
// NCRDIV: [[LOOP21]] = distinct !{[[LOOP21]], [[META10]], [[META11]]}
-// NCRDIV: [[LOOP22]] = distinct !{[[LOOP22]], [[META10]], [[META11]]}
+// NCRDIV: [[META22]] = !{float 3.000000e+00}
// NCRDIV: [[LOOP23]] = distinct !{[[LOOP23]], [[META10]], [[META11]]}
// NCRDIV: [[LOOP24]] = distinct !{[[LOOP24]], [[META10]], [[META11]]}
-// NCRDIV: [[META25]] = !{float 3.000000e+00}
+// NCRDIV: [[LOOP25]] = distinct !{[[LOOP25]], [[META10]], [[META11]]}
// NCRDIV: [[LOOP26]] = distinct !{[[LOOP26]], [[META10]], [[META11]]}
// NCRDIV: [[LOOP27]] = distinct !{[[LOOP27]], [[META10]], [[META11]]}
//.
diff --git a/llvm/lib/IR/Instruction.cpp b/llvm/lib/IR/Instruction.cpp
index 3c35c656dca84..8046585cdeccf 100644
--- a/llvm/lib/IR/Instruction.cpp
+++ b/llvm/lib/IR/Instruction.cpp
@@ -576,7 +576,8 @@ void Instruction::dropUBImplyingAttrsAndMetadata(ArrayRef<unsigned> Keep) {
// immediate undefined behavior.
static const unsigned KnownIDs[] = {
LLVMContext::MD_annotation, LLVMContext::MD_range,
- LLVMContext::MD_nonnull, LLVMContext::MD_align, LLVMContext::MD_prof};
+ LLVMContext::MD_nonnull, LLVMContext::MD_align,
+ LLVMContext::MD_fpmath, LLVMContext::MD_prof};
SmallVector<unsigned> KeepIDs;
KeepIDs.reserve(Keep.size() + std::size(KnownIDs));
append_range(KeepIDs, (!ProfcheckDisableMetadataFixes ? KnownIDs
diff --git a/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll b/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll
index 22af7e3e75502..31e9556b4f80e 100644
--- a/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll
+++ b/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll
@@ -69,3 +69,12 @@ EntryBlock:
; CHECK: select i1 %A, float 0x3FD5555560000000, float [[OP]]
}
+define float @test8(i1 %A, float %B) {
+EntryBlock:
+ %cf = select i1 %A, float 1.000000e+00, float %B
+ %op= fdiv float 3.000000e+00, %cf, !fpmath !{float 2.5}
+ ret float %op
+; CHECK-LABEL: @test8(
+; CHECK: [[OP:%.*]] = fdiv float 3.000000e+00, %B, !fpmath
+; CHECK: select i1 %A, float 3.000000e+00, float [[OP]]
+}
More information about the cfe-commits
mailing list