https://github.com/wenju-he updated https://github.com/llvm/llvm-project/pull/179019
>From 5ecb10c8a571b610dd6ba3413d99651b7c02bad8 Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Sat, 31 Jan 2026 07:37:02 +0100 Subject: [PATCH 1/2] [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]] +} >From 313161e4b8efbe57061ab4ff01c7e99f3b998c36 Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Sat, 31 Jan 2026 09:33:11 +0100 Subject: [PATCH 2/2] add a comment in dropUBImplyingAttrsAndUnknownMetadata --- llvm/lib/IR/Instruction.cpp | 1 + llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/llvm/lib/IR/Instruction.cpp b/llvm/lib/IR/Instruction.cpp index 8046585cdeccf..67fad4028df5f 100644 --- a/llvm/lib/IR/Instruction.cpp +++ b/llvm/lib/IR/Instruction.cpp @@ -572,6 +572,7 @@ void Instruction::dropUBImplyingAttrsAndUnknownMetadata( void Instruction::dropUBImplyingAttrsAndMetadata(ArrayRef<unsigned> Keep) { // !annotation and !prof metadata does not impact semantics. // !range, !nonnull and !align produce poison, so they are safe to speculate. + // !fpmath specifies floating-point precision and does not imply UB. // !noundef and various AA metadata must be dropped, as it generally produces // immediate undefined behavior. static const unsigned KnownIDs[] = { diff --git a/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll b/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll index 31e9556b4f80e..c986004f843ab 100644 --- a/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll +++ b/llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll @@ -72,7 +72,7 @@ EntryBlock: 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} + %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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
