diff options
| author | Wenju He <wenju.he@intel.com> | 2026-02-02 09:58:58 +0800 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2026-02-02 09:58:58 +0800 |
| commit | dc152f0d2d085dcfb7542d0e71e19ebfa1aa3794 (patch) | |
| tree | 15e95ef10551b5eaf6222f3f33abd89612cc8131 | |
| parent | cf60af88b4d3791b5e225373e51186f4f86ad50f (diff) | |
| download | llvm-main.zip llvm-main.tar.gz llvm-main.tar.bz2 | |
`fpmath` is precision metadata rather than UB-implying metadata. This
avoids `fpmath` from being dropped in InstCombine FoldOpIntoSelect.
| -rw-r--r-- | clang/test/Headers/__clang_hip_math.hip | 14 | ||||
| -rw-r--r-- | llvm/lib/IR/Instruction.cpp | 4 | ||||
| -rw-r--r-- | llvm/test/Transforms/InstCombine/fold-fops-into-selects.ll | 9 |
3 files changed, 19 insertions, 8 deletions
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 6647d35..e91f572 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 3c35c65..67fad40 100644 --- a/llvm/lib/IR/Instruction.cpp +++ b/llvm/lib/IR/Instruction.cpp @@ -572,11 +572,13 @@ 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[] = { 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 22af7e3..c986004 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]] +} |
