[clang] [llvm] [IR] Optimize PHINode::removeIncomingValue() by swapping removed incoming value with the last incoming value. (PR #171963)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Dec 15 04:47:30 PST 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-transforms
Author: Mingjie Xu (Enna1)
<details>
<summary>Changes</summary>
Add an optional argument `KeepIncomingOrder`
- If `KeepIncomingOrder` is false, the new implementation simply moves the last incoming value and block into the position of the incoming value being removed.
- If `KeepIncomingOrder` is true, the relative order of incoming values remained is preserved, like previous behavior.
- `KeepIncomingOrder` defaults to false.
This improves compile-time for PHI nodes with many predecessors.
Depends:
https://github.com/llvm/llvm-project/pull/171955
https://github.com/llvm/llvm-project/pull/171956
https://github.com/llvm/llvm-project/pull/171960
https://github.com/llvm/llvm-project/pull/171962
---
Patch is 164.55 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/171963.diff
72 Files Affected:
- (modified) clang/test/Headers/__clang_hip_math.hip (+22-22)
- (modified) llvm/include/llvm/IR/Instructions.h (+4-2)
- (modified) llvm/lib/IR/Instructions.cpp (+13-9)
- (modified) llvm/lib/SandboxIR/Instruction.cpp (+2-1)
- (modified) llvm/test/Examples/IRTransforms/SimplifyCFG/tut-simplify-cfg5-del-phis-for-dead-block.ll (+1-1)
- (modified) llvm/test/Transforms/CodeGenPrepare/X86/sink-addrmode-base.ll (+1-1)
- (modified) llvm/test/Transforms/Coroutines/coro-catchswitch.ll (+1-1)
- (modified) llvm/test/Transforms/DFAJumpThreading/dfa-jump-threading-transform.ll (+9-11)
- (modified) llvm/test/Transforms/DFAJumpThreading/dfa-unfold-select.ll (+77-84)
- (modified) llvm/test/Transforms/DFAJumpThreading/equivalent-states.ll (+1-1)
- (modified) llvm/test/Transforms/IndVarSimplify/eliminate-backedge.ll (+2-2)
- (modified) llvm/test/Transforms/Inline/inline_invoke.ll (+2-2)
- (modified) llvm/test/Transforms/JumpThreading/ddt-crash.ll (+1-1)
- (modified) llvm/test/Transforms/JumpThreading/fold-not-thread.ll (+1-1)
- (modified) llvm/test/Transforms/JumpThreading/invalidate-lvi.ll (+1-1)
- (modified) llvm/test/Transforms/LoopSimplifyCFG/phi_with_duplicating_inputs.ll (+1-1)
- (modified) llvm/test/Transforms/LoopUnroll/AArch64/apple-unrolling-multi-exit.ll (+2-2)
- (modified) llvm/test/Transforms/LoopUnroll/AArch64/unrolling-multi-exit.ll (+2-2)
- (modified) llvm/test/Transforms/LoopUnroll/full-unroll-one-unpredictable-exit.ll (+2-2)
- (modified) llvm/test/Transforms/LoopUnroll/partial-unroll-non-latch-exit.ll (+1-1)
- (modified) llvm/test/Transforms/LoopUnroll/peel-loop.ll (+1-1)
- (modified) llvm/test/Transforms/LoopUnroll/runtime-loop-multiple-exits.ll (+10-10)
- (modified) llvm/test/Transforms/LoopUnroll/scevunroll.ll (+2-2)
- (modified) llvm/test/Transforms/LoopUnroll/unroll-header-exiting-with-phis-multiple-exiting-blocks.ll (+1-1)
- (modified) llvm/test/Transforms/LoopVectorize/AArch64/simple_early_exit.ll (+3-3)
- (modified) llvm/test/Transforms/LoopVectorize/single-early-exit-deref-assumptions.ll (+2-2)
- (modified) llvm/test/Transforms/LoopVectorize/single-early-exit-interleave.ll (+4-4)
- (modified) llvm/test/Transforms/LoopVectorize/single_early_exit.ll (+1-1)
- (modified) llvm/test/Transforms/LoopVectorize/single_early_exit_live_outs.ll (+18-18)
- (modified) llvm/test/Transforms/LoopVectorize/trip-count-expansion-may-introduce-ub.ll (+1-1)
- (modified) llvm/test/Transforms/LoopVectorize/vector-loop-backedge-elimination-early-exit.ll (+6-6)
- (modified) llvm/test/Transforms/LowerSwitch/2014-06-23-PHIlowering.ll (+1-1)
- (modified) llvm/test/Transforms/LowerSwitch/do-not-handle-impossible-values.ll (+1-1)
- (modified) llvm/test/Transforms/MergeICmps/X86/alias-merge-blocks.ll (+1-1)
- (modified) llvm/test/Transforms/MergeICmps/X86/entry-block-shuffled-2.ll (+1-1)
- (modified) llvm/test/Transforms/MergeICmps/X86/entry-block-shuffled.ll (+1-1)
- (modified) llvm/test/Transforms/MergeICmps/X86/pr59740.ll (+1-1)
- (modified) llvm/test/Transforms/PGOProfile/chr.ll (+11-11)
- (modified) llvm/test/Transforms/PhaseOrdering/AArch64/hoist-load-from-vector-loop.ll (+1-1)
- (modified) llvm/test/Transforms/PhaseOrdering/AArch64/std-find.ll (+1-1)
- (modified) llvm/test/Transforms/PhaseOrdering/X86/vector-reductions-logical.ll (+2-2)
- (modified) llvm/test/Transforms/PhaseOrdering/switch-sext.ll (+1-1)
- (modified) llvm/test/Transforms/SimpleLoopUnswitch/inject-invariant-conditions.ll (+10-10)
- (modified) llvm/test/Transforms/SimpleLoopUnswitch/nontrivial-unswitch-freeze.ll (+2-2)
- (modified) llvm/test/Transforms/SimpleLoopUnswitch/nontrivial-unswitch.ll (+2-2)
- (modified) llvm/test/Transforms/SimplifyCFG/EqualPHIEdgeBlockMerge.ll (+2-2)
- (modified) llvm/test/Transforms/SimplifyCFG/ForwardSwitchConditionToPHI.ll (+2-2)
- (modified) llvm/test/Transforms/SimplifyCFG/Hexagon/switch-to-lookup-table.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/HoistCode.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/RISCV/switch-of-powers-of-two.ll (+7-7)
- (modified) llvm/test/Transforms/SimplifyCFG/UnreachableEliminate.ll (+12-12)
- (modified) llvm/test/Transforms/SimplifyCFG/X86/debugloc-switch-powers-of-two.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/X86/empty-cleanuppad.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/X86/switch-of-powers-of-two.ll (+2-2)
- (modified) llvm/test/Transforms/SimplifyCFG/X86/switch-to-lookup-gep.ll (+2-2)
- (modified) llvm/test/Transforms/SimplifyCFG/X86/switch-to-lookup-globals.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/X86/switch_to_lookup_table.ll (+2-2)
- (modified) llvm/test/Transforms/SimplifyCFG/X86/switch_to_lookup_table_big.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/avoid-complex-phi.ll (+10-10)
- (modified) llvm/test/Transforms/SimplifyCFG/branch-fold-threshold.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/branch-fold.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/merge-phis-in-switch.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/multiple-phis.ll (+5-5)
- (modified) llvm/test/Transforms/SimplifyCFG/rangereduce.ll (+3-3)
- (modified) llvm/test/Transforms/SimplifyCFG/switch-dup-bbs.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/switch-on-const.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/switch-simplify-crash2.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/switch-to-select-two-case.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/switch-transformations-no-lut.ll (+3-3)
- (modified) llvm/test/Transforms/SimplifyCFG/switch_create-custom-dl.ll (+1-1)
- (modified) llvm/test/Transforms/SimplifyCFG/switch_create.ll (+1-1)
- (modified) llvm/test/Transforms/Util/lowerswitch.ll (+2-2)
``````````diff
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 426e5af319cbf..4a57e2d014dbb 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -179,7 +179,7 @@ extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
// CHECK-NEXT: [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP4]], 6
// CHECK-NEXT: br i1 [[OR_COND34_I]], label %[[IF_END31_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]]
// CHECK: [[IF_END31_I]]:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I]] ], [ -87, %[[IF_ELSE_I]] ], [ -55, %[[IF_ELSE17_I]] ]
+// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I]] ], [ -48, %[[WHILE_BODY_I]] ], [ -55, %[[IF_ELSE17_I]] ]
// CHECK-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I3]], 4
// CHECK-NEXT: [[CONV25_I:%.*]] = zext nneg i8 [[TMP1]] to i64
// CHECK-NEXT: [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]]
@@ -214,7 +214,7 @@ extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
// AMDGCNSPIRV-NEXT: [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP4]], 6
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I]], label %[[IF_END31_I]], label %[[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]]
// AMDGCNSPIRV: [[IF_END31_I]]:
-// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I]] ], [ -87, %[[IF_ELSE_I]] ], [ -55, %[[IF_ELSE17_I]] ]
+// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I]] ], [ -48, %[[WHILE_BODY_I]] ], [ -55, %[[IF_ELSE17_I]] ]
// AMDGCNSPIRV-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I3]], 4
// AMDGCNSPIRV-NEXT: [[CONV25_I:%.*]] = zext nneg i8 [[TMP1]] to i64
// AMDGCNSPIRV-NEXT: [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]]
@@ -269,7 +269,7 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// CHECK-NEXT: br i1 [[OR_COND34_I_I]], label %[[IF_END31_I_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]]
// CHECK: [[IF_END31_I_I]]:
-// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I31_I]] ], [ -87, %[[IF_ELSE_I_I]] ], [ -55, %[[IF_ELSE17_I_I]] ]
+// CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I]] ], [ -48, %[[WHILE_BODY_I31_I]] ], [ -55, %[[IF_ELSE17_I_I]] ]
// CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I29_I11]], 4
// CHECK-NEXT: [[CONV25_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// CHECK-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
@@ -311,7 +311,7 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// CHECK-NEXT: [[CMP_NOT_I17_I:%.*]] = icmp eq i8 [[TMP13]], 0
// CHECK-NEXT: br i1 [[CMP_NOT_I17_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], label %[[WHILE_BODY_I18_I]], !llvm.loop [[LOOP12]]
// CHECK: [[_ZL15__MAKE_MANTISSAPKC_EXIT]]:
-// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I]] ], [ 0, %[[WHILE_COND_I14_I_PREHEADER]] ], [ [[SUB_I_I]], %[[IF_THEN_I_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ 0, %[[IF_ELSE17_I_I]] ], [ [[SUB_I25_I]], %[[IF_THEN_I21_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ]
+// CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I]] ], [ 0, %[[WHILE_COND_I14_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_PREHEADER]] ], [ [[SUB_I_I]], %[[IF_THEN_I_I]] ], [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ 0, %[[IF_ELSE17_I_I]] ], [ [[SUB_I25_I]], %[[IF_THEN_I21_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ]
// CHECK-NEXT: ret i64 [[RETVAL_0_I]]
//
// AMDGCNSPIRV-LABEL: define spir_func i64 @test___make_mantissa(
@@ -347,7 +347,7 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I]], label %[[IF_END31_I_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]]
// AMDGCNSPIRV: [[IF_END31_I_I]]:
-// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I32_I]] ], [ -87, %[[IF_ELSE_I_I]] ], [ -55, %[[IF_ELSE17_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I]] ], [ -48, %[[WHILE_BODY_I32_I]] ], [ -55, %[[IF_ELSE17_I_I]] ]
// AMDGCNSPIRV-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I30_I7]], 4
// AMDGCNSPIRV-NEXT: [[CONV25_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// AMDGCNSPIRV-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
@@ -391,7 +391,7 @@ extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
// AMDGCNSPIRV-NEXT: [[__R_1_I26_I]] = select i1 [[OR_COND_I19_I]], i64 [[SUB_I23_I]], i64 [[__R_0_I16_I]]
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I19_I]], label %[[WHILE_COND_I14_I]], label %[[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP13]]
// AMDGCNSPIRV: [[_ZL15__MAKE_MANTISSAPKC_EXIT]]:
-// AMDGCNSPIRV-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ [[__R_0_I_I]], %[[WHILE_COND_I_I]] ], [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ 0, %[[IF_ELSE17_I_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ], [ [[__R_0_I16_I]], %[[WHILE_COND_I14_I]] ]
+// AMDGCNSPIRV-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I]] ], [ [[ADD28_I_I]], %[[IF_END31_I_I]] ], [ [[__R_0_I_I]], %[[WHILE_COND_I_I]] ], [ 0, %[[WHILE_BODY_I_I]] ], [ 0, %[[IF_ELSE17_I_I]] ], [ 0, %[[WHILE_BODY_I18_I]] ], [ [[__R_0_I16_I]], %[[WHILE_COND_I14_I]] ]
// AMDGCNSPIRV-NEXT: ret i64 [[RETVAL_0_I]]
//
extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
@@ -4357,7 +4357,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// DEFAULT: [[IF_END31_I_I_I]]:
-// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// DEFAULT-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
@@ -4399,7 +4399,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP12]]
// DEFAULT: [[_ZL4NANFPKC_EXIT]]:
-// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// DEFAULT-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// DEFAULT-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
@@ -4449,7 +4449,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// APPROX: [[IF_END31_I_I_I]]:
-// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// APPROX-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
@@ -4491,7 +4491,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP12]]
// APPROX: [[_ZL4NANFPKC_EXIT]]:
-// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// APPROX-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// APPROX-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// APPROX-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
@@ -4536,7 +4536,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// NCRDIV: [[IF_END31_I_I_I]]:
-// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// NCRDIV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
@@ -4578,7 +4578,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL4NANFPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP12]]
// NCRDIV: [[_ZL4NANFPKC_EXIT]]:
-// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
+// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// NCRDIV-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// NCRDIV-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// NCRDIV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
@@ -4618,7 +4618,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL4NANFPKC_EXIT]]
// AMDGCNSPIRV: [[IF_END31_I_I_I]]:
-// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I32_I_I]] ], [ -87, %[[IF_ELSE_I_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I32_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
// AMDGCNSPIRV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I30_I_I7]], 4
// AMDGCNSPIRV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// AMDGCNSPIRV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
@@ -4662,7 +4662,7 @@ extern "C" __device__ double test_modf(double x, double* y) {
// AMDGCNSPIRV-NEXT: [[__R_1_I26_I_I]] = select i1 [[OR_COND_I19_I_I]], i64 [[SUB_I23_I_I]], i64 [[__R_0_I16_I_I]]
// AMDGCNSPIRV-NEXT: br i1 [[OR_COND_I19_I_I]], label %[[WHILE_COND_I14_I_I]], label %[[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP13]]
// AMDGCNSPIRV: [[_ZL4NANFPKC_EXIT]]:
-// AMDGCNSPIRV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
+// AMDGCNSPIRV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], %[[WHILE_COND_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ], [ [[__R_0_I16_I_I]], %[[WHILE_COND_I14_I_I]] ]
// AMDGCNSPIRV-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
// AMDGCNSPIRV-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
// AMDGCNSPIRV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i32 [[BF_VALUE_I]], 2143289344
@@ -4711,7 +4711,7 @@ extern "C" __device__ float test_nanf(const char *tag) {
// DEFAULT-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// DEFAULT-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// DEFAULT: [[IF_END31_I_I_I]]:
-// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// DEFAULT-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
// DEFAULT-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
// DEFAULT-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// DEFAULT-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
@@ -4753,7 +4753,7 @@ extern "C" __device__ float test_nanf(const char *tag) {
// DEFAULT-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
// DEFAULT-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP12]]
// DEFAULT: [[_ZL3NANPKC_EXIT]]:
-// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
+// DEFAULT-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// DEFAULT-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// DEFAULT-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
// DEFAULT-NEXT: [[TMP14:%.*]] = bitcast i64 [[BF_SET9_I]] to double
@@ -4802,7 +4802,7 @@ extern "C" __device__ float test_nanf(const char *tag) {
// APPROX-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// APPROX-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// APPROX: [[IF_END31_I_I_I]]:
-// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// APPROX-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
// APPROX-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
// APPROX-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// APPROX-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
@@ -4844,7 +4844,7 @@ extern "C" __device__ float test_nanf(const char *tag) {
// APPROX-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
// APPROX-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP12]]
// APPROX: [[_ZL3NANPKC_EXIT]]:
-// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
+// APPROX-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// APPROX-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// APPROX-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
// APPROX-NEXT: [[TMP14:%.*]] = bitcast i64 [[BF_SET9_I]] to double
@@ -4888,7 +4888,7 @@ extern "C" __device__ float test_nanf(const char *tag) {
// NCRDIV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// NCRDIV-NEXT: br i1 [[OR_COND34_I_I_I]], label %[[IF_END31_I_I_I]], label %[[_ZL3NANPKC_EXIT]]
// NCRDIV: [[IF_END31_I_I_I]]:
-// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -87, %[[IF_ELSE_I_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
+// NCRDIV-NEXT: [[DOTSINK:%.*]] = phi i64 [ -87, %[[IF_ELSE_I_I_I]] ], [ -48, %[[WHILE_BODY_I31_I_I]] ], [ -55, %[[IF_ELSE17_I_I_I]] ]
// NCRDIV-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I29_I_I11]], 4
// NCRDIV-NEXT: [[CONV25_I_I_I:%.*]] = zext nneg i8 [[TMP3]] to i64
// NCRDIV-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
@@ -4930,7 +4930,7 @@ extern "C" __device__ float test_nanf(const char *tag) {
// NCRDIV-NEXT: [[CMP_NOT_I17_I_I:%.*]] = icmp eq i8 [[TMP13]], 0
// NCRDIV-NEXT: br i1 [[CMP_NOT_I17_I_I]], label %[[_ZL3NANPKC_EXIT]], label %[[WHILE_BODY_I18_I_I]], !llvm.loop [[LOOP12]]
// NCRDIV: [[_ZL3NANPKC_EXIT]]:
-// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
+// NCRDIV-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, %[[IF_THEN5_I_I]] ], [ 0, %[[WHILE_COND_I14_I_I_PREHEADER]] ], [ 0, %[[WHILE_COND_I_I_I_PREHEADER]] ], [ [[SUB_I_I_I]], %[[IF_THEN_I_I_I]] ], [ [[ADD28_I_I_I]], %[[IF_END31_I_I_I]] ], [ 0, %[[WHILE_BODY_I_I_I]] ], [ 0, %[[IF_ELSE17_I_I_I]] ], [ [[SUB_I25_I_I]], %[[IF_THEN_I21_I_I]] ], [ 0, %[[WHILE_BODY_I18_I_I]] ]
// NCRDIV-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
// NCRDIV-NEXT: [[BF_SET9_I:%.*]] = or disjoint i64 [[BF_VALUE_I]], 9221120237041090560
// NCRDIV-NEXT: [[TMP14:%.*]] = bitcast i64 [[BF_SET9_I]] to double
@@ -4969,7 +4969,7 @@ extern "C" __device__ float test_nanf(const char *tag) {
// AMDGCNSPIRV-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP6]], 6
// AMDGCNSPIRV-NEXT...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/171963
More information about the llvm-commits
mailing list