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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
