mirrors/llvm-project
0
0 Fork 0
mirror of https://github.com/llvm/llvm-project.git synced 2026-10-26 13:53:31 +08:00
Code Activity

Revert "[MachineLICM] Use RegisterClassInfo::getRegPressureSetLimit (#119826)"

Browse Source 
This reverts commit b4e17d4a31.

This causes a large compile-time regression.
This commit is contained in:
Nikita Popov committed 2025-01-10 09:05:06 +01:00
1 parent e0f14e11c7
commit eeac0ffaf4
50 files changed
+32104 -32537
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

llvm/lib/CodeGen/MachineLICM.cpp
+1 -3
View File
@@ -124,7 +124,6 @@ namespace {
const TargetRegisterInfo *TRI = nullptr;
const MachineFrameInfo *MFI = nullptr;
MachineRegisterInfo *MRI = nullptr;
RegisterClassInfo RegClassInfo;
TargetSchedModel SchedModel;
bool PreRegAlloc = false;
bool HasProfileData = false;
@@ -393,7 +392,6 @@ bool MachineLICMImpl::run(MachineFunction &MF) {
MFI = &MF.getFrameInfo();
MRI = &MF.getRegInfo();
SchedModel.init(&ST);
RegClassInfo.runOnMachineFunction(MF);
HasProfileData = MF.getFunction().hasProfileData();
@@ -410,7 +408,7 @@ bool MachineLICMImpl::run(MachineFunction &MF) {
std::fill(RegPressure.begin(), RegPressure.end(), 0);
RegLimit.resize(NumRPS);
for (unsigned i = 0, e = NumRPS; i != e; ++i)
RegLimit[i] = RegClassInfo.getRegPressureSetLimit(i);
RegLimit[i] = TRI->getRegPressureSetLimit(MF, i);
}
if (HoistConstLoads)
llvm/test/CodeGen/AMDGPU/GlobalISel/atomicrmw_fmax.ll
+325 -345
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/GlobalISel/atomicrmw_fmin.ll
+325 -345
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll
+24 -20
View File
@@ -557,11 +557,11 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX908-NEXT: s_mul_hi_u32 s9, s0, s7
; GFX908-NEXT: s_mul_i32 s0, s0, s7
; GFX908-NEXT: s_add_i32 s1, s9, s1
; GFX908-NEXT: s_lshl_b64 s[0:1], s[0:1], 5
; GFX908-NEXT: s_lshl_b64 s[14:15], s[0:1], 5
; GFX908-NEXT: s_branch .LBB3_2
; GFX908-NEXT: .LBB3_1: ; %Flow20
; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
; GFX908-NEXT: s_andn2_b64 vcc, exec, s[14:15]
; GFX908-NEXT: s_andn2_b64 vcc, exec, s[0:1]
; GFX908-NEXT: s_cbranch_vccz .LBB3_12
; GFX908-NEXT: .LBB3_2: ; %bb9
; GFX908-NEXT: ; =>This Loop Header: Depth=1
@@ -571,15 +571,17 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX908-NEXT: ; %bb.3: ; %bb14
; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
; GFX908-NEXT: global_load_dwordx2 v[2:3], v[0:1], off
; GFX908-NEXT: v_cmp_gt_i64_e64 s[0:1], s[4:5], -1
; GFX908-NEXT: s_mov_b32 s7, s6
; GFX908-NEXT: v_cndmask_b32_e64 v6, 0, 1, s[0:1]
; GFX908-NEXT: v_mov_b32_e32 v4, s6
; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v6
; GFX908-NEXT: v_mov_b32_e32 v6, s6
; GFX908-NEXT: v_mov_b32_e32 v9, s7
; GFX908-NEXT: v_mov_b32_e32 v5, s7
; GFX908-NEXT: v_mov_b32_e32 v7, s7
; GFX908-NEXT: v_mov_b32_e32 v8, s6
; GFX908-NEXT: v_cmp_lt_i64_e64 s[14:15], s[4:5], 0
; GFX908-NEXT: v_cmp_gt_i64_e64 s[16:17], s[4:5], -1
; GFX908-NEXT: v_cmp_lt_i64_e64 s[16:17], s[4:5], 0
; GFX908-NEXT: v_mov_b32_e32 v11, v5
; GFX908-NEXT: s_mov_b64 s[18:19], s[10:11]
; GFX908-NEXT: v_mov_b32_e32 v10, v4
@@ -599,9 +601,9 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX908-NEXT: ; in Loop: Header=BB3_5 Depth=2
; GFX908-NEXT: v_add_co_u32_sdwa v2, vcc, v2, v16 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:WORD_0
; GFX908-NEXT: v_addc_co_u32_e32 v3, vcc, 0, v3, vcc
; GFX908-NEXT: s_add_u32 s18, s18, s0
; GFX908-NEXT: s_add_u32 s18, s18, s14
; GFX908-NEXT: v_cmp_lt_i64_e64 s[22:23], -1, v[2:3]
; GFX908-NEXT: s_addc_u32 s19, s19, s1
; GFX908-NEXT: s_addc_u32 s19, s19, s15
; GFX908-NEXT: s_mov_b64 s[20:21], 0
; GFX908-NEXT: s_andn2_b64 vcc, exec, s[22:23]
; GFX908-NEXT: s_cbranch_vccz .LBB3_9
@@ -620,7 +622,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX908-NEXT: s_waitcnt vmcnt(0)
; GFX908-NEXT: ds_read_b64 v[12:13], v19
; GFX908-NEXT: ds_read_b64 v[14:15], v0
; GFX908-NEXT: s_andn2_b64 vcc, exec, s[16:17]
; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1]
; GFX908-NEXT: s_waitcnt lgkmcnt(0)
; GFX908-NEXT: s_cbranch_vccnz .LBB3_7
; GFX908-NEXT: ; %bb.6: ; %bb51
@@ -648,7 +650,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX908-NEXT: s_mov_b64 s[20:21], -1
; GFX908-NEXT: s_branch .LBB3_4
; GFX908-NEXT: .LBB3_7: ; in Loop: Header=BB3_5 Depth=2
; GFX908-NEXT: s_mov_b64 s[20:21], s[14:15]
; GFX908-NEXT: s_mov_b64 s[20:21], s[16:17]
; GFX908-NEXT: s_andn2_b64 vcc, exec, s[20:21]
; GFX908-NEXT: s_cbranch_vccz .LBB3_4
; GFX908-NEXT: ; %bb.8: ; in Loop: Header=BB3_2 Depth=1
@@ -659,7 +661,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX908-NEXT: s_xor_b64 s[16:17], s[20:21], -1
; GFX908-NEXT: .LBB3_10: ; %Flow19
; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
; GFX908-NEXT: s_mov_b64 s[14:15], -1
; GFX908-NEXT: s_mov_b64 s[0:1], -1
; GFX908-NEXT: s_and_b64 vcc, exec, s[16:17]
; GFX908-NEXT: s_cbranch_vccz .LBB3_1
; GFX908-NEXT: ; %bb.11: ; %bb12
@@ -668,7 +670,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX908-NEXT: s_addc_u32 s5, s5, 0
; GFX908-NEXT: s_add_u32 s10, s10, s12
; GFX908-NEXT: s_addc_u32 s11, s11, s13
; GFX908-NEXT: s_mov_b64 s[14:15], 0
; GFX908-NEXT: s_mov_b64 s[0:1], 0
; GFX908-NEXT: s_branch .LBB3_1
; GFX908-NEXT: .LBB3_12: ; %DummyReturnBlock
; GFX908-NEXT: s_endpgm
@@ -718,11 +720,11 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX90A-NEXT: s_mul_hi_u32 s9, s0, s7
; GFX90A-NEXT: s_mul_i32 s0, s0, s7
; GFX90A-NEXT: s_add_i32 s1, s9, s1
; GFX90A-NEXT: s_lshl_b64 s[0:1], s[0:1], 5
; GFX90A-NEXT: s_lshl_b64 s[14:15], s[0:1], 5
; GFX90A-NEXT: s_branch .LBB3_2
; GFX90A-NEXT: .LBB3_1: ; %Flow20
; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[14:15]
; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[0:1]
; GFX90A-NEXT: s_cbranch_vccz .LBB3_12
; GFX90A-NEXT: .LBB3_2: ; %bb9
; GFX90A-NEXT: ; =>This Loop Header: Depth=1
@@ -732,12 +734,14 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX90A-NEXT: ; %bb.3: ; %bb14
; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
; GFX90A-NEXT: global_load_dwordx2 v[4:5], v[0:1], off
; GFX90A-NEXT: v_cmp_gt_i64_e64 s[0:1], s[4:5], -1
; GFX90A-NEXT: s_mov_b32 s7, s6
; GFX90A-NEXT: v_cndmask_b32_e64 v8, 0, 1, s[0:1]
; GFX90A-NEXT: v_pk_mov_b32 v[6:7], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v8
; GFX90A-NEXT: v_pk_mov_b32 v[8:9], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-NEXT: v_pk_mov_b32 v[10:11], s[6:7], s[6:7] op_sel:[0,1]
; GFX90A-NEXT: v_cmp_lt_i64_e64 s[14:15], s[4:5], 0
; GFX90A-NEXT: v_cmp_gt_i64_e64 s[16:17], s[4:5], -1
; GFX90A-NEXT: v_cmp_lt_i64_e64 s[16:17], s[4:5], 0
; GFX90A-NEXT: s_mov_b64 s[18:19], s[10:11]
; GFX90A-NEXT: v_pk_mov_b32 v[12:13], v[6:7], v[6:7] op_sel:[0,1]
; GFX90A-NEXT: s_waitcnt vmcnt(0)
@@ -756,8 +760,8 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX90A-NEXT: ; in Loop: Header=BB3_5 Depth=2
; GFX90A-NEXT: v_add_co_u32_sdwa v4, vcc, v4, v18 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:WORD_0
; GFX90A-NEXT: v_addc_co_u32_e32 v5, vcc, 0, v5, vcc
; GFX90A-NEXT: s_add_u32 s18, s18, s0
; GFX90A-NEXT: s_addc_u32 s19, s19, s1
; GFX90A-NEXT: s_add_u32 s18, s18, s14
; GFX90A-NEXT: s_addc_u32 s19, s19, s15
; GFX90A-NEXT: v_cmp_lt_i64_e64 s[22:23], -1, v[4:5]
; GFX90A-NEXT: s_mov_b64 s[20:21], 0
; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[22:23]
@@ -777,7 +781,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX90A-NEXT: s_waitcnt vmcnt(0)
; GFX90A-NEXT: ds_read_b64 v[14:15], v19
; GFX90A-NEXT: ds_read_b64 v[16:17], v0
; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[16:17]
; GFX90A-NEXT: s_and_b64 vcc, exec, s[0:1]
; GFX90A-NEXT: ; kill: killed $sgpr20 killed $sgpr21
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-NEXT: s_cbranch_vccnz .LBB3_7
@@ -798,7 +802,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX90A-NEXT: s_mov_b64 s[20:21], -1
; GFX90A-NEXT: s_branch .LBB3_4
; GFX90A-NEXT: .LBB3_7: ; in Loop: Header=BB3_5 Depth=2
; GFX90A-NEXT: s_mov_b64 s[20:21], s[14:15]
; GFX90A-NEXT: s_mov_b64 s[20:21], s[16:17]
; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[20:21]
; GFX90A-NEXT: s_cbranch_vccz .LBB3_4
; GFX90A-NEXT: ; %bb.8: ; in Loop: Header=BB3_2 Depth=1
@@ -809,7 +813,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX90A-NEXT: s_xor_b64 s[16:17], s[20:21], -1
; GFX90A-NEXT: .LBB3_10: ; %Flow19
; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
; GFX90A-NEXT: s_mov_b64 s[14:15], -1
; GFX90A-NEXT: s_mov_b64 s[0:1], -1
; GFX90A-NEXT: s_and_b64 vcc, exec, s[16:17]
; GFX90A-NEXT: s_cbranch_vccz .LBB3_1
; GFX90A-NEXT: ; %bb.11: ; %bb12
@@ -818,7 +822,7 @@ define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg
; GFX90A-NEXT: s_addc_u32 s5, s5, 0
; GFX90A-NEXT: s_add_u32 s10, s10, s12
; GFX90A-NEXT: s_addc_u32 s11, s11, s13
; GFX90A-NEXT: s_mov_b64 s[14:15], 0
; GFX90A-NEXT: s_mov_b64 s[0:1], 0
; GFX90A-NEXT: s_branch .LBB3_1
; GFX90A-NEXT: .LBB3_12: ; %DummyReturnBlock
; GFX90A-NEXT: s_endpgm
llvm/test/CodeGen/AMDGPU/buffer-fat-pointer-atomicrmw-fadd.ll
+3106 -3120
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/buffer-fat-pointer-atomicrmw-fmax.ll
+2225 -2241
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/buffer-fat-pointer-atomicrmw-fmin.ll
+2225 -2241
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/codegen-prepare-addrspacecast-non-null.ll
+4 -4
View File
@@ -245,14 +245,14 @@ define void @recursive_phis(i1 %cond, ptr addrspace(5) %ptr) {
; GISEL-ASM-NEXT: v_and_b32_e32 v0, 0xffff, v1
; GISEL-ASM-NEXT: ; %bb.2: ; %finallyendcf.split
; GISEL-ASM-NEXT: s_or_b64 exec, exec, s[6:7]
; GISEL-ASM-NEXT: s_mov_b64 s[6:7], 0
; GISEL-ASM-NEXT: s_mov_b64 s[8:9], src_private_base
; GISEL-ASM-NEXT: s_mov_b64 s[6:7], 0
; GISEL-ASM-NEXT: v_mov_b32_e32 v1, s9
; GISEL-ASM-NEXT: v_mov_b32_e32 v2, 7
; GISEL-ASM-NEXT: .LBB7_3: ; %finally
; GISEL-ASM-NEXT: ; =>This Inner Loop Header: Depth=1
; GISEL-ASM-NEXT: s_and_b64 s[10:11], exec, s[4:5]
; GISEL-ASM-NEXT: s_or_b64 s[6:7], s[10:11], s[6:7]
; GISEL-ASM-NEXT: v_mov_b32_e32 v1, s9
; GISEL-ASM-NEXT: s_and_b64 s[8:9], exec, s[4:5]
; GISEL-ASM-NEXT: s_or_b64 s[6:7], s[8:9], s[6:7]
; GISEL-ASM-NEXT: flat_store_dword v[0:1], v2
; GISEL-ASM-NEXT: s_waitcnt vmcnt(0)
; GISEL-ASM-NEXT: s_andn2_b64 exec, exec, s[6:7]
llvm/test/CodeGen/AMDGPU/exec-mask-opt-cannot-create-empty-or-backward-segment.ll
+52 -47
View File
@@ -6,90 +6,95 @@ define amdgpu_kernel void @cannot_create_empty_or_backwards_segment(i1 %arg, i1
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_mov_b64 s[26:27], s[2:3]
; CHECK-NEXT: s_mov_b64 s[24:25], s[0:1]
; CHECK-NEXT: s_load_dword s4, s[8:9], 0x0
; CHECK-NEXT: s_load_dword s2, s[8:9], 0x0
; CHECK-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
; CHECK-NEXT: s_load_dword s12, s[8:9], 0x4
; CHECK-NEXT: s_load_dword s6, s[8:9], 0x4
; CHECK-NEXT: s_add_u32 s24, s24, s15
; CHECK-NEXT: s_addc_u32 s25, s25, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_bitcmp1_b32 s4, 0
; CHECK-NEXT: s_cselect_b64 s[2:3], -1, 0
; CHECK-NEXT: s_bitcmp1_b32 s4, 8
; CHECK-NEXT: s_cselect_b64 s[6:7], -1, 0
; CHECK-NEXT: s_bitcmp1_b32 s4, 16
; CHECK-NEXT: s_cselect_b64 s[8:9], -1, 0
; CHECK-NEXT: s_bitcmp1_b32 s0, 24
; CHECK-NEXT: s_cselect_b64 s[4:5], -1, 0
; CHECK-NEXT: s_xor_b64 s[14:15], s[4:5], -1
; CHECK-NEXT: s_bitcmp1_b32 s1, 0
; CHECK-NEXT: s_bitcmp1_b32 s2, 0
; CHECK-NEXT: s_cselect_b64 s[16:17], -1, 0
; CHECK-NEXT: s_bitcmp1_b32 s2, 8
; CHECK-NEXT: s_cselect_b64 s[10:11], -1, 0
; CHECK-NEXT: s_bitcmp1_b32 s12, 8
; CHECK-NEXT: s_bitcmp1_b32 s2, 16
; CHECK-NEXT: s_cselect_b64 s[2:3], -1, 0
; CHECK-NEXT: s_bitcmp1_b32 s0, 24
; CHECK-NEXT: s_cselect_b64 s[8:9], -1, 0
; CHECK-NEXT: s_xor_b64 s[4:5], s[8:9], -1
; CHECK-NEXT: s_bitcmp1_b32 s1, 0
; CHECK-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[2:3]
; CHECK-NEXT: s_cselect_b64 s[12:13], -1, 0
; CHECK-NEXT: s_and_b64 s[0:1], exec, s[14:15]
; CHECK-NEXT: s_bitcmp1_b32 s6, 8
; CHECK-NEXT: v_cmp_ne_u32_e64 s[2:3], 1, v0
; CHECK-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[16:17]
; CHECK-NEXT: s_cselect_b64 s[14:15], -1, 0
; CHECK-NEXT: s_and_b64 s[4:5], exec, s[4:5]
; CHECK-NEXT: s_and_b64 s[6:7], exec, s[10:11]
; CHECK-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
; CHECK-NEXT: v_mov_b32_e32 v0, 0
; CHECK-NEXT: s_branch .LBB0_3
; CHECK-NEXT: .LBB0_1: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[16:17], 0
; CHECK-NEXT: s_mov_b64 s[18:19], -1
; CHECK-NEXT: s_mov_b64 s[14:15], -1
; CHECK-NEXT: s_mov_b64 s[18:19], 0
; CHECK-NEXT: s_mov_b64 s[20:21], -1
; CHECK-NEXT: s_mov_b64 s[16:17], -1
; CHECK-NEXT: s_mov_b64 s[22:23], -1
; CHECK-NEXT: .LBB0_2: ; %Flow7
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_and_b64 vcc, exec, s[20:21]
; CHECK-NEXT: s_and_b64 vcc, exec, s[22:23]
; CHECK-NEXT: s_cbranch_vccnz .LBB0_12
; CHECK-NEXT: .LBB0_3: ; %bb7
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[8:9]
; CHECK-NEXT: s_and_b64 vcc, exec, s[2:3]
; CHECK-NEXT: s_cbranch_vccnz .LBB0_1
; CHECK-NEXT: ; %bb.4: ; %bb8
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 vcc, s[0:1]
; CHECK-NEXT: s_mov_b64 vcc, s[4:5]
; CHECK-NEXT: s_cbranch_vccz .LBB0_6
; CHECK-NEXT: ; %bb.5: ; %bb9
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[14:15], 0
; CHECK-NEXT: s_mov_b64 s[16:17], -1
; CHECK-NEXT: s_mov_b64 s[20:21], s[6:7]
; CHECK-NEXT: s_mov_b64 s[16:17], 0
; CHECK-NEXT: s_mov_b64 s[18:19], -1
; CHECK-NEXT: s_mov_b64 s[22:23], s[10:11]
; CHECK-NEXT: s_cbranch_execz .LBB0_7
; CHECK-NEXT: s_branch .LBB0_8
; CHECK-NEXT: .LBB0_6: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[14:15], -1
; CHECK-NEXT: s_mov_b64 s[16:17], 0
; CHECK-NEXT: s_mov_b64 s[20:21], 0
; CHECK-NEXT: s_mov_b64 s[16:17], -1
; CHECK-NEXT: s_mov_b64 s[18:19], 0
; CHECK-NEXT: s_mov_b64 s[22:23], 0
; CHECK-NEXT: .LBB0_7: ; %bb10
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[16:17], -1
; CHECK-NEXT: s_mov_b64 s[14:15], 0
; CHECK-NEXT: s_mov_b64 s[20:21], s[12:13]
; CHECK-NEXT: s_mov_b64 s[18:19], -1
; CHECK-NEXT: s_mov_b64 s[16:17], 0
; CHECK-NEXT: s_mov_b64 s[22:23], s[14:15]
; CHECK-NEXT: .LBB0_8: ; %Flow9
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[18:19], -1
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[20:21]
; CHECK-NEXT: s_mov_b64 s[20:21], -1
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[22:23]
; CHECK-NEXT: s_mov_b64 s[22:23], -1
; CHECK-NEXT: s_cbranch_vccnz .LBB0_2
; CHECK-NEXT: ; %bb.9: ; %bb13
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[14:15], 0
; CHECK-NEXT: s_and_b64 vcc, exec, s[6:7]
; CHECK-NEXT: s_mov_b64 s[18:19], 0
; CHECK-NEXT: s_mov_b64 s[16:17], 0
; CHECK-NEXT: s_mov_b64 s[20:21], 0
; CHECK-NEXT: s_mov_b64 vcc, s[6:7]
; CHECK-NEXT: s_cbranch_vccz .LBB0_11
; CHECK-NEXT: ; %bb.10: ; %bb16
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[18:19], -1
; CHECK-NEXT: s_mov_b64 s[20:21], s[10:11]
; CHECK-NEXT: s_mov_b64 s[20:21], -1
; CHECK-NEXT: s_mov_b64 s[22:23], s[12:13]
; CHECK-NEXT: .LBB0_11: ; %Flow11
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[16:17], 0
; CHECK-NEXT: s_mov_b64 s[18:19], 0
; CHECK-NEXT: s_branch .LBB0_2
; CHECK-NEXT: .LBB0_12: ; %loop.exit.guard6
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_xor_b64 s[20:21], s[18:19], -1
; CHECK-NEXT: s_mov_b64 s[18:19], -1
; CHECK-NEXT: s_and_b64 vcc, exec, s[20:21]
; CHECK-NEXT: s_xor_b64 s[22:23], s[20:21], -1
; CHECK-NEXT: s_mov_b64 s[20:21], -1
; CHECK-NEXT: s_and_b64 vcc, exec, s[22:23]
; CHECK-NEXT: s_cbranch_vccz .LBB0_16
; CHECK-NEXT: ; %bb.13: ; %bb14
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[2:3]
; CHECK-NEXT: s_and_b64 vcc, exec, s[0:1]
; CHECK-NEXT: s_cbranch_vccnz .LBB0_15
; CHECK-NEXT: ; %bb.14: ; %bb15
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
@@ -97,22 +102,22 @@ define amdgpu_kernel void @cannot_create_empty_or_backwards_segment(i1 %arg, i1
; CHECK-NEXT: buffer_store_dword v0, off, s[24:27], 0
; CHECK-NEXT: .LBB0_15: ; %Flow
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_mov_b64 s[18:19], 0
; CHECK-NEXT: s_mov_b64 s[20:21], 0
; CHECK-NEXT: .LBB0_16: ; %Flow13
; CHECK-NEXT: ; in Loop: Header=BB0_3 Depth=1
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[18:19]
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[20:21]
; CHECK-NEXT: s_cbranch_vccnz .LBB0_3
; CHECK-NEXT: ; %bb.17: ; %loop.exit.guard
; CHECK-NEXT: s_and_b64 vcc, exec, s[14:15]
; CHECK-NEXT: s_and_b64 vcc, exec, s[16:17]
; CHECK-NEXT: s_cbranch_vccnz .LBB0_22
; CHECK-NEXT: ; %bb.18: ; %loop.exit.guard5
; CHECK-NEXT: s_and_b64 vcc, exec, s[16:17]
; CHECK-NEXT: s_and_b64 vcc, exec, s[18:19]
; CHECK-NEXT: s_cbranch_vccnz .LBB0_23
; CHECK-NEXT: ; %bb.19: ; %bb17
; CHECK-NEXT: s_and_b64 vcc, exec, s[4:5]
; CHECK-NEXT: s_and_b64 vcc, exec, s[8:9]
; CHECK-NEXT: s_cbranch_vccz .LBB0_21
; CHECK-NEXT: ; %bb.20: ; %bb19
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[2:3]
; CHECK-NEXT: s_and_b64 vcc, exec, s[0:1]
; CHECK-NEXT: s_cbranch_vccz .LBB0_22
; CHECK-NEXT: .LBB0_21: ; %bb18
; CHECK-NEXT: s_endpgm
llvm/test/CodeGen/AMDGPU/flat-atomicrmw-fadd.ll
+1269 -1257
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/flat-atomicrmw-fmax.ll
+2828 -2846
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/flat-atomicrmw-fmin.ll
+2828 -2846
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/flat-atomicrmw-fsub.ll
+1191 -1175
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/flat_atomics_i32_system.ll
+780 -926
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/flat_atomics_i64_system.ll
+741 -927
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/flat_atomics_i64_system_noprivate.ll
+964 -1064
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/fp64-atomics-gfx90a.ll
+6 -6
View File
@@ -1426,23 +1426,23 @@ main_body:
define amdgpu_kernel void @flat_atomic_fadd_f64_noret_pat_agent_safe(ptr %ptr) {
; GFX90A-LABEL: flat_atomic_fadd_f64_noret_pat_agent_safe:
; GFX90A: ; %bb.0: ; %main_body
; GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
; GFX90A-NEXT: s_mov_b64 s[2:3], 0
; GFX90A-NEXT: s_load_dwordx2 s[2:3], s[4:5], 0x24
; GFX90A-NEXT: s_mov_b64 s[0:1], 0
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]
; GFX90A-NEXT: flat_load_dwordx2 v[2:3], v[0:1]
; GFX90A-NEXT: v_pk_mov_b32 v[4:5], s[2:3], s[2:3] op_sel:[0,1]
; GFX90A-NEXT: .LBB50_1: ; %atomicrmw.start
; GFX90A-NEXT: ; =>This Inner Loop Header: Depth=1
; GFX90A-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GFX90A-NEXT: v_add_f64 v[0:1], v[2:3], 4.0
; GFX90A-NEXT: v_pk_mov_b32 v[4:5], s[0:1], s[0:1] op_sel:[0,1]
; GFX90A-NEXT: flat_atomic_cmpswap_x2 v[0:1], v[4:5], v[0:3] glc
; GFX90A-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GFX90A-NEXT: buffer_wbinvl1_vol
; GFX90A-NEXT: v_cmp_eq_u64_e32 vcc, v[0:1], v[2:3]
; GFX90A-NEXT: s_or_b64 s[2:3], vcc, s[2:3]
; GFX90A-NEXT: s_or_b64 s[0:1], vcc, s[0:1]
; GFX90A-NEXT: v_pk_mov_b32 v[2:3], v[0:1], v[0:1] op_sel:[0,1]
; GFX90A-NEXT: s_andn2_b64 exec, exec, s[2:3]
; GFX90A-NEXT: s_andn2_b64 exec, exec, s[0:1]
; GFX90A-NEXT: s_cbranch_execnz .LBB50_1
; GFX90A-NEXT: ; %bb.2: ; %atomicrmw.end
; GFX90A-NEXT: s_endpgm
llvm/test/CodeGen/AMDGPU/global-atomicrmw-fadd.ll
+1510 -1496
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/global-atomicrmw-fmax.ll
+2499 -2486
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/global-atomicrmw-fmin.ll
+2499 -2486
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/global-atomicrmw-fsub.ll
+1200 -1184
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/global-load-saddr-to-vaddr.ll
+7 -7
View File
@@ -56,19 +56,19 @@ define amdgpu_kernel void @test_move_load_address_to_vgpr_d16_hi(ptr addrspace(1
; GCN: ; %bb.0: ; %bb
; GCN-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
; GCN-NEXT: v_mov_b32_e32 v1, 0
; GCN-NEXT: s_movk_i32 s2, 0x100
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: global_load_ushort v0, v1, s[0:1] glc
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: v_mov_b32_e32 v2, s1
; GCN-NEXT: s_movk_i32 s1, 0x100
; GCN-NEXT: .LBB1_1: ; %bb3
; GCN-NEXT: ; =>This Inner Loop Header: Depth=1
; GCN-NEXT: v_lshlrev_b64 v[2:3], 1, v[0:1]
; GCN-NEXT: v_mov_b32_e32 v0, s1
; GCN-NEXT: v_add_co_u32_e32 v2, vcc, s0, v2
; GCN-NEXT: v_addc_co_u32_e32 v3, vcc, v0, v3, vcc
; GCN-NEXT: global_load_short_d16_hi v0, v[2:3], off glc
; GCN-NEXT: v_lshlrev_b64 v[3:4], 1, v[0:1]
; GCN-NEXT: v_add_co_u32_e32 v3, vcc, s0, v3
; GCN-NEXT: v_addc_co_u32_e32 v4, vcc, v2, v4, vcc
; GCN-NEXT: global_load_short_d16_hi v0, v[3:4], off glc
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: v_cmp_eq_u32_e32 vcc, s2, v0
; GCN-NEXT: v_cmp_eq_u32_e32 vcc, s1, v0
; GCN-NEXT: s_cbranch_vccz .LBB1_1
; GCN-NEXT: ; %bb.2: ; %bb2
; GCN-NEXT: s_endpgm
llvm/test/CodeGen/AMDGPU/global_atomics_i32_system.ll
+338 -368
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/global_atomics_i64_system.ll
+962 -982
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/global_atomics_scan_fmax.ll
+328 -327
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/global_atomics_scan_fmin.ll
+328 -327
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/insert-delay-alu-bug.ll
+34 -36
View File
@@ -54,18 +54,18 @@ define amdgpu_kernel void @f2(i32 %arg, i32 %arg1, i32 %arg2, i1 %arg3, i32 %arg
; GFX11: ; %bb.0: ; %bb
; GFX11-NEXT: s_mov_b64 s[16:17], s[4:5]
; GFX11-NEXT: v_mov_b32_e32 v31, v0
; GFX11-NEXT: s_load_b32 s24, s[16:17], 0x24
; GFX11-NEXT: s_load_b32 s19, s[16:17], 0x24
; GFX11-NEXT: s_mov_b32 s12, s13
; GFX11-NEXT: s_mov_b64 s[10:11], s[6:7]
; GFX11-NEXT: s_mov_b64 s[6:7], s[2:3]
; GFX11-NEXT: v_and_b32_e32 v0, 0x3ff, v31
; GFX11-NEXT: s_mov_b64 s[4:5], s[0:1]
; GFX11-NEXT: s_mov_b32 s19, 0
; GFX11-NEXT: s_mov_b32 s20, 0
; GFX11-NEXT: s_mov_b32 s0, -1
; GFX11-NEXT: s_mov_b32 s3, exec_lo
; GFX11-NEXT: s_mov_b32 s32, 0
; GFX11-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-NEXT: v_mul_lo_u32 v0, s24, v0
; GFX11-NEXT: v_mul_lo_u32 v0, s19, v0
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-NEXT: v_cmpx_eq_u32_e32 0, v0
; GFX11-NEXT: s_cbranch_execz .LBB2_13
@@ -74,7 +74,7 @@ define amdgpu_kernel void @f2(i32 %arg, i32 %arg1, i32 %arg2, i1 %arg3, i32 %arg
; GFX11-NEXT: s_mov_b32 s18, 0
; GFX11-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-NEXT: s_bitcmp1_b32 s21, 0
; GFX11-NEXT: s_cselect_b32 s19, -1, 0
; GFX11-NEXT: s_cselect_b32 s24, -1, 0
; GFX11-NEXT: s_bitcmp0_b32 s21, 0
; GFX11-NEXT: s_cbranch_scc0 .LBB2_3
; GFX11-NEXT: ; %bb.2: ; %bb15
@@ -110,60 +110,58 @@ define amdgpu_kernel void @f2(i32 %arg, i32 %arg1, i32 %arg2, i1 %arg3, i32 %arg
; GFX11-NEXT: s_cbranch_scc0 .LBB2_8
; GFX11-NEXT: ; %bb.5: ; %bb18.preheader
; GFX11-NEXT: s_load_b128 s[28:31], s[16:17], 0x44
; GFX11-NEXT: s_mov_b32 vcc_lo, 0
; GFX11-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-NEXT: s_mul_hi_u32 s0, s29, s28
; GFX11-NEXT: s_mul_i32 s1, s29, s28
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
; GFX11-NEXT: v_alignbit_b32 v0, s0, s1, 1
; GFX11-NEXT: s_mov_b32 s1, 0
; GFX11-NEXT: s_mov_b32 s13, s1
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
; GFX11-NEXT: v_readfirstlane_b32 s0, v0
; GFX11-NEXT: v_mov_b32_e32 v0, 0
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(SALU_CYCLE_1)
; GFX11-NEXT: s_or_b32 s0, s0, 1
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
; GFX11-NEXT: s_lshr_b32 s0, s0, s30
; GFX11-NEXT: s_mul_i32 s0, s0, s22
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
; GFX11-NEXT: s_mul_i32 s0, s0, s22
; GFX11-NEXT: s_mul_i32 s0, s0, s20
; GFX11-NEXT: s_or_b32 s0, s24, s0
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
; GFX11-NEXT: s_or_b32 s0, s19, s0
; GFX11-NEXT: s_lshl_b64 s[20:21], s[0:1], 1
; GFX11-NEXT: global_load_u16 v0, v0, s[20:21]
; GFX11-NEXT: s_mov_b32 s0, s1
; GFX11-NEXT: global_load_u16 v1, v0, s[20:21]
; GFX11-NEXT: v_cndmask_b32_e64 v0, 0, 1, s24
; GFX11-NEXT: s_waitcnt vmcnt(0)
; GFX11-NEXT: v_cmp_eq_u32_e64 s0, 0, v0
; GFX11-NEXT: v_cmp_eq_u32_e32 vcc_lo, 0, v1
; GFX11-NEXT: v_cndmask_b32_e64 v1, 0, 1, vcc_lo
; GFX11-NEXT: s_mov_b32 vcc_lo, 0
; GFX11-NEXT: .p2align 6
; GFX11-NEXT: .LBB2_6: ; %bb18
; GFX11-NEXT: ; =>This Inner Loop Header: Depth=1
; GFX11-NEXT: s_and_b32 s1, 0xffff, s1
; GFX11-NEXT: v_cndmask_b32_e64 v0, 0, 1, s19
; GFX11-NEXT: v_readfirstlane_b32 s13, v0
; GFX11-NEXT: s_cmp_lg_u32 s1, 0
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
; GFX11-NEXT: v_cndmask_b32_e64 v1, 0, 1, s0
; GFX11-NEXT: s_cselect_b32 s1, -1, 0
; GFX11-NEXT: v_readfirstlane_b32 s20, v0
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
; GFX11-NEXT: v_cndmask_b32_e64 v2, 0, 1, s1
; GFX11-NEXT: s_and_b32 s1, s8, s1
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
; GFX11-NEXT: s_and_b32 s1, s1, exec_lo
; GFX11-NEXT: v_readfirstlane_b32 s21, v2
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1)
; GFX11-NEXT: s_cselect_b32 s1, s21, s20
; GFX11-NEXT: s_and_b32 s20, 0xffff, s13
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
; GFX11-NEXT: v_readfirstlane_b32 s19, v2
; GFX11-NEXT: s_cselect_b32 s1, s19, s13
; GFX11-NEXT: s_and_b32 s13, 0xffff, s0
; GFX11-NEXT: s_and_b32 s1, s1, 1
; GFX11-NEXT: s_cmp_lg_u32 s20, 0
; GFX11-NEXT: s_cselect_b32 s20, -1, 0
; GFX11-NEXT: s_and_b32 s22, s9, exec_lo
; GFX11-NEXT: v_cndmask_b32_e64 v0, 0, 1, s20
; GFX11-NEXT: v_readfirstlane_b32 s20, v1
; GFX11-NEXT: s_cmp_lg_u32 s13, 0
; GFX11-NEXT: s_cselect_b32 s13, -1, 0
; GFX11-NEXT: s_and_b32 s20, s9, exec_lo
; GFX11-NEXT: v_cndmask_b32_e64 v2, 0, 1, s13
; GFX11-NEXT: v_readfirstlane_b32 s13, v1
; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
; GFX11-NEXT: v_readfirstlane_b32 s21, v0
; GFX11-NEXT: s_cselect_b32 s20, s21, s20
; GFX11-NEXT: v_readfirstlane_b32 s19, v2
; GFX11-NEXT: s_cselect_b32 s13, s19, s13
; GFX11-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
; GFX11-NEXT: s_bitcmp1_b32 s20, 0
; GFX11-NEXT: s_cselect_b32 s20, 0x100, 0
; GFX11-NEXT: s_or_b32 s13, s20, s13
; GFX11-NEXT: s_bitcmp1_b32 s13, 0
; GFX11-NEXT: s_cselect_b32 s13, 0x100, 0
; GFX11-NEXT: s_or_b32 s0, s13, s0
; GFX11-NEXT: s_cbranch_vccz .LBB2_6
; GFX11-NEXT: ; %bb.7: ; %Flow
; GFX11-NEXT: s_mov_b32 s0, 0
@@ -181,7 +179,7 @@ define amdgpu_kernel void @f2(i32 %arg, i32 %arg1, i32 %arg2, i1 %arg3, i32 %arg
; GFX11-NEXT: ; %bb.11: ; %Flow6
; GFX11-NEXT: s_mov_b32 s18, -1
; GFX11-NEXT: .LBB2_12: ; %Flow11
; GFX11-NEXT: s_and_b32 s19, s2, exec_lo
; GFX11-NEXT: s_and_b32 s20, s2, exec_lo
; GFX11-NEXT: s_or_not1_b32 s0, s18, exec_lo
; GFX11-NEXT: .LBB2_13: ; %Flow9
; GFX11-NEXT: s_or_b32 exec_lo, exec_lo, s3
@@ -198,10 +196,10 @@ define amdgpu_kernel void @f2(i32 %arg, i32 %arg1, i32 %arg2, i1 %arg3, i32 %arg
; GFX11-NEXT: s_mov_b32 s14, s15
; GFX11-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-NEXT: s_swappc_b64 s[30:31], s[0:1]
; GFX11-NEXT: s_or_b32 s19, s19, exec_lo
; GFX11-NEXT: s_or_b32 s20, s20, exec_lo
; GFX11-NEXT: .LBB2_15: ; %Flow14
; GFX11-NEXT: s_or_b32 exec_lo, exec_lo, s3
; GFX11-NEXT: s_and_saveexec_b32 s0, s19
; GFX11-NEXT: s_and_saveexec_b32 s0, s20
; GFX11-NEXT: ; %bb.16: ; %UnifiedUnreachableBlock
; GFX11-NEXT: ; divergent unreachable
; GFX11-NEXT: ; %bb.17: ; %UnifiedReturnBlock
llvm/test/CodeGen/AMDGPU/licm-regpressure.mir
+26 -26
View File
@@ -32,38 +32,38 @@ body: |
; GCN-NEXT: [[COPY15:%[0-9]+]]:vgpr_32 = COPY $vgpr0
; GCN-NEXT: [[COPY16:%[0-9]+]]:vgpr_32 = COPY $vgpr0
; GCN-NEXT: [[COPY17:%[0-9]+]]:vgpr_32 = COPY $vgpr0
; GCN-NEXT: [[V_CVT_F64_I32_e32_:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_1:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY1]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_2:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY2]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_3:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY3]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_4:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY4]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_5:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY5]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_6:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY6]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_7:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY7]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_8:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY8]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_9:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY9]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_10:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY10]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_11:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY11]], implicit $mode, implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_12:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY12]], implicit $mode, implicit $exec
; GCN-NEXT: {{ $}}
; GCN-NEXT: bb.1:
; GCN-NEXT: successors: %bb.2(0x04000000), %bb.1(0x7c000000)
; GCN-NEXT: liveins: $vcc
; GCN-NEXT: {{ $}}
; GCN-NEXT: $vcc = S_AND_B64 $exec, $vcc, implicit-def $scc
; GCN-NEXT: [[V_CVT_F64_I32_e32_:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_1:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY1]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_1]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_2:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY2]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_2]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_3:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY3]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_3]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_4:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY4]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_4]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_5:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY5]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_5]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_6:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY6]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_6]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_7:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY7]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_7]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_8:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY8]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_8]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_9:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY9]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_9]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_10:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY10]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_10]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_11:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY11]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_11]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_12:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY12]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_12]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_1]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_2]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_3]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_4]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_5]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_6]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_7]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_8]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_9]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_10]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_11]], implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, [[V_CVT_F64_I32_e32_12]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_13:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY13]], implicit $mode, implicit $exec
; GCN-NEXT: $vcc = V_CMP_EQ_U64_e64 $vcc, killed [[V_CVT_F64_I32_e32_13]], implicit $exec
; GCN-NEXT: [[V_CVT_F64_I32_e32_14:%[0-9]+]]:vreg_64 = nofpexcept V_CVT_F64_I32_e32 [[COPY14]], implicit $mode, implicit $exec
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.struct.atomic.buffer.load.ll
+75 -75
View File
@@ -6,20 +6,20 @@ define amdgpu_kernel void @struct_atomic_buffer_load_i32(<4 x i32> %addr, i32 %i
; CHECK-LABEL: struct_atomic_buffer_load_i32:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB0_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB0_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -67,20 +67,20 @@ define amdgpu_kernel void @struct_atomic_buffer_load_i32_off(<4 x i32> %addr, i3
; CHECK-LABEL: struct_atomic_buffer_load_i32_off:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB2_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB2_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -99,20 +99,20 @@ define amdgpu_kernel void @struct_atomic_buffer_load_i32_soff(<4 x i32> %addr, i
; CHECK-LABEL: struct_atomic_buffer_load_i32_soff:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB3_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 4 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 4 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB3_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -130,20 +130,20 @@ define amdgpu_kernel void @struct_atomic_buffer_load_i32_dlc(<4 x i32> %addr, i3
; CHECK-LABEL: struct_atomic_buffer_load_i32_dlc:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB4_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen offset:4 dlc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen offset:4 dlc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB4_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -194,20 +194,20 @@ define amdgpu_kernel void @struct_atomic_buffer_load_i64(<4 x i32> %addr, i32 %i
; CHECK-LABEL: struct_atomic_buffer_load_i64:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_and_b32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v2, s6
; CHECK-NEXT: .LBB6_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v2, s4
; CHECK-NEXT: buffer_load_b64 v[2:3], v2, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b64 v[3:4], v2, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u64_e32 vcc_lo, v[2:3], v[0:1]
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u64_e32 vcc_lo, v[3:4], v[0:1]
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB6_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -227,20 +227,20 @@ define amdgpu_kernel void @struct_atomic_buffer_load_v2i16(<4 x i32> %addr, i32
; CHECK-LABEL: struct_atomic_buffer_load_v2i16:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB7_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB7_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -260,23 +260,23 @@ define amdgpu_kernel void @struct_atomic_buffer_load_v4i16(<4 x i32> %addr, i32
; CHECK-LABEL: struct_atomic_buffer_load_v4i16:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB8_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b64 v[1:2], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b64 v[2:3], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_and_b32_e32 v1, 0xffff, v1
; CHECK-NEXT: v_and_b32_e32 v2, 0xffff, v2
; CHECK-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
; CHECK-NEXT: v_lshl_or_b32 v1, v2, 16, v1
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_lshl_or_b32 v2, v3, 16, v2
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB8_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -297,20 +297,20 @@ define amdgpu_kernel void @struct_atomic_buffer_load_v4i32(<4 x i32> %addr, i32
; CHECK-LABEL: struct_atomic_buffer_load_v4i32:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB9_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b128 v[1:4], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b128 v[2:5], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v4, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v5, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB9_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -330,22 +330,22 @@ define amdgpu_kernel void @struct_atomic_buffer_load_ptr(<4 x i32> %addr, i32 %i
; CHECK-LABEL: struct_atomic_buffer_load_ptr:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB10_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b64 v[1:2], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b64 v[2:3], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: flat_load_b32 v1, v[1:2]
; CHECK-NEXT: flat_load_b32 v2, v[2:3]
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB10_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.struct.ptr.atomic.buffer.load.ll
+75 -75
View File
@@ -6,20 +6,20 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_i32(ptr addrspace(8) %p
; CHECK-LABEL: struct_ptr_atomic_buffer_load_i32:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB0_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB0_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -67,20 +67,20 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_i32_off(ptr addrspace(8
; CHECK-LABEL: struct_ptr_atomic_buffer_load_i32_off:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB2_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB2_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -99,20 +99,20 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_i32_soff(ptr addrspace(
; CHECK-LABEL: struct_ptr_atomic_buffer_load_i32_soff:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB3_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 4 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 4 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB3_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -130,20 +130,20 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_i32_dlc(ptr addrspace(8
; CHECK-LABEL: struct_ptr_atomic_buffer_load_i32_dlc:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB4_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen offset:4 dlc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen offset:4 dlc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB4_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -194,20 +194,20 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_i64(ptr addrspace(8) %p
; CHECK-LABEL: struct_ptr_atomic_buffer_load_i64:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_and_b32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v2, s6
; CHECK-NEXT: .LBB6_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v2, s4
; CHECK-NEXT: buffer_load_b64 v[2:3], v2, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b64 v[3:4], v2, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u64_e32 vcc_lo, v[2:3], v[0:1]
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u64_e32 vcc_lo, v[3:4], v[0:1]
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB6_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -227,20 +227,20 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_v2i16(ptr addrspace(8)
; CHECK-LABEL: struct_ptr_atomic_buffer_load_v2i16:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB7_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b32 v1, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: buffer_load_b32 v2, v1, s[0:3], 0 idxen glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB7_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -260,23 +260,23 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_v4i16(ptr addrspace(8)
; CHECK-LABEL: struct_ptr_atomic_buffer_load_v4i16:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB8_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b64 v[1:2], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b64 v[2:3], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_and_b32_e32 v1, 0xffff, v1
; CHECK-NEXT: v_and_b32_e32 v2, 0xffff, v2
; CHECK-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
; CHECK-NEXT: v_lshl_or_b32 v1, v2, 16, v1
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_lshl_or_b32 v2, v3, 16, v2
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB8_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -297,20 +297,20 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_v4i32(ptr addrspace(8)
; CHECK-LABEL: struct_ptr_atomic_buffer_load_v4i32:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB9_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b128 v[1:4], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b128 v[2:5], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v4, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v5, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB9_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
@@ -330,22 +330,22 @@ define amdgpu_kernel void @struct_ptr_atomic_buffer_load_ptr(ptr addrspace(8) %p
; CHECK-LABEL: struct_ptr_atomic_buffer_load_ptr:
; CHECK: ; %bb.0: ; %bb
; CHECK-NEXT: s_clause 0x1
; CHECK-NEXT: s_load_b32 s6, s[4:5], 0x34
; CHECK-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
; CHECK-NEXT: s_load_b32 s4, s[4:5], 0x34
; CHECK-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; CHECK-NEXT: s_mov_b32 s5, 0
; CHECK-NEXT: s_mov_b32 s4, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s6
; CHECK-NEXT: .LBB10_1: ; %bb1
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_mov_b32_e32 v1, s4
; CHECK-NEXT: buffer_load_b64 v[1:2], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: buffer_load_b64 v[2:3], v1, s[0:3], 0 idxen offset:4 glc
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: flat_load_b32 v1, v[1:2]
; CHECK-NEXT: flat_load_b32 v2, v[2:3]
; CHECK-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v1, v0
; CHECK-NEXT: s_or_b32 s5, vcc_lo, s5
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc_lo, v2, v0
; CHECK-NEXT: s_or_b32 s4, vcc_lo, s4
; CHECK-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s5
; CHECK-NEXT: s_and_not1_b32 exec_lo, exec_lo, s4
; CHECK-NEXT: s_cbranch_execnz .LBB10_1
; CHECK-NEXT: ; %bb.2: ; %bb2
; CHECK-NEXT: s_endpgm
llvm/test/CodeGen/AMDGPU/local-atomicrmw-fadd.ll
+512 -522
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/local-atomicrmw-fmax.ll
+692 -688
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/local-atomicrmw-fmin.ll
+692 -688
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/local-atomicrmw-fsub.ll
+562 -558
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/AMDGPU/memcpy-crash-issue63986.ll
+5 -5
View File
@@ -161,13 +161,13 @@ define void @issue63986_reduced_expanded(i64 %idxprom) {
; CHECK-NEXT: v_lshlrev_b64 v[0:1], 1, v[0:1]
; CHECK-NEXT: s_cbranch_execnz .LBB1_8
; CHECK-NEXT: .LBB1_5: ; %loop-memcpy-residual.preheader
; CHECK-NEXT: s_mov_b64 s[6:7], 0
; CHECK-NEXT: .LBB1_6: ; %loop-memcpy-residual
; CHECK-NEXT: s_add_u32 s6, s6, 1
; CHECK-NEXT: v_mov_b32_e32 v0, s4
; CHECK-NEXT: s_mov_b64 s[6:7], 0
; CHECK-NEXT: v_mov_b32_e32 v1, s5
; CHECK-NEXT: s_addc_u32 s7, s7, 0
; CHECK-NEXT: v_cmp_lt_u64_e32 vcc, s[6:7], v[0:1]
; CHECK-NEXT: .LBB1_6: ; %loop-memcpy-residual
; CHECK-NEXT: s_add_u32 s4, s6, 1
; CHECK-NEXT: s_addc_u32 s5, s7, 0
; CHECK-NEXT: v_cmp_lt_u64_e32 vcc, s[4:5], v[0:1]
; CHECK-NEXT: s_mov_b64 s[6:7], 1
; CHECK-NEXT: s_cbranch_vccnz .LBB1_6
; CHECK-NEXT: ; %bb.7: ; %Flow
llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll
+29 -25
View File
@@ -5,21 +5,23 @@
define amdgpu_kernel void @matmul_kernel(i32 %a0, i32 %a1) {
; GFX942-LABEL: matmul_kernel:
; GFX942: ; %bb.0: ; %entry
; GFX942-NEXT: s_load_dwordx2 s[2:3], s[4:5], 0x0
; GFX942-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
; GFX942-NEXT: v_mov_b32_e32 v1, 0
; GFX942-NEXT: s_mov_b32 s0, 0
; GFX942-NEXT: s_mov_b32 s2, 0
; GFX942-NEXT: v_accvgpr_write_b32 a0, v1
; GFX942-NEXT: s_mov_b32 s1, 0
; GFX942-NEXT: s_mov_b32 s3, 0
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
; GFX942-NEXT: s_cmp_lg_u32 s2, 0
; GFX942-NEXT: s_cselect_b64 s[2:3], -1, 0
; GFX942-NEXT: s_cmp_lg_u32 s0, 0
; GFX942-NEXT: s_cselect_b64 s[0:1], -1, 0
; GFX942-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
; GFX942-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
; GFX942-NEXT: s_branch .LBB0_2
; GFX942-NEXT: .LBB0_1: ; %bb2
; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1
; GFX942-NEXT: s_or_b32 s4, s1, 1
; GFX942-NEXT: s_ashr_i32 s5, s1, 31
; GFX942-NEXT: s_mov_b32 s1, s0
; GFX942-NEXT: v_mov_b64_e32 v[4:5], s[0:1]
; GFX942-NEXT: s_or_b32 s4, s3, 1
; GFX942-NEXT: s_ashr_i32 s5, s3, 31
; GFX942-NEXT: s_mov_b32 s3, s2
; GFX942-NEXT: v_mov_b64_e32 v[4:5], s[2:3]
; GFX942-NEXT: v_accvgpr_read_b32 v0, a0
; GFX942-NEXT: v_mov_b32_e32 v2, v1
; GFX942-NEXT: v_mov_b32_e32 v3, v1
@@ -27,54 +29,56 @@ define amdgpu_kernel void @matmul_kernel(i32 %a0, i32 %a1) {
; GFX942-NEXT: v_accvgpr_write_b32 a1, v1
; GFX942-NEXT: v_accvgpr_write_b32 a2, v2
; GFX942-NEXT: v_accvgpr_write_b32 a3, v3
; GFX942-NEXT: s_and_b32 s1, s5, s4
; GFX942-NEXT: s_and_b32 s3, s5, s4
; GFX942-NEXT: s_nop 0
; GFX942-NEXT: v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[4:5], a[0:3]
; GFX942-NEXT: s_cbranch_execz .LBB0_4
; GFX942-NEXT: .LBB0_2: ; %bb
; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1
; GFX942-NEXT: s_andn2_b64 vcc, exec, s[2:3]
; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1]
; GFX942-NEXT: s_cbranch_vccz .LBB0_1
; GFX942-NEXT: ; %bb.3:
; GFX942-NEXT: ; implicit-def: $sgpr1
; GFX942-NEXT: ; implicit-def: $sgpr3
; GFX942-NEXT: .LBB0_4: ; %common.ret
; GFX942-NEXT: s_endpgm
;
; GFX908-LABEL: matmul_kernel:
; GFX908: ; %bb.0: ; %entry
; GFX908-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x0
; GFX908-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
; GFX908-NEXT: v_mov_b32_e32 v1, 0
; GFX908-NEXT: s_mov_b32 s0, 0
; GFX908-NEXT: s_mov_b32 s1, 0
; GFX908-NEXT: s_mov_b32 s2, 0
; GFX908-NEXT: s_mov_b32 s3, 0
; GFX908-NEXT: v_accvgpr_write_b32 a0, v1
; GFX908-NEXT: s_waitcnt lgkmcnt(0)
; GFX908-NEXT: s_cmp_lg_u32 s2, 0
; GFX908-NEXT: s_cselect_b64 s[2:3], -1, 0
; GFX908-NEXT: s_cmp_lg_u32 s0, 0
; GFX908-NEXT: s_cselect_b64 s[0:1], -1, 0
; GFX908-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
; GFX908-NEXT: s_branch .LBB0_2
; GFX908-NEXT: .LBB0_1: ; %bb2
; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1
; GFX908-NEXT: s_or_b32 s4, s1, 1
; GFX908-NEXT: s_ashr_i32 s5, s1, 31
; GFX908-NEXT: s_mov_b32 s1, s0
; GFX908-NEXT: s_or_b32 s4, s3, 1
; GFX908-NEXT: s_ashr_i32 s5, s3, 31
; GFX908-NEXT: s_mov_b32 s3, s2
; GFX908-NEXT: s_nop 3
; GFX908-NEXT: v_accvgpr_read_b32 v0, a0
; GFX908-NEXT: v_mov_b32_e32 v5, s1
; GFX908-NEXT: v_mov_b32_e32 v4, s0
; GFX908-NEXT: v_mov_b32_e32 v5, s3
; GFX908-NEXT: v_mov_b32_e32 v4, s2
; GFX908-NEXT: v_mov_b32_e32 v2, v1
; GFX908-NEXT: v_mov_b32_e32 v3, v1
; GFX908-NEXT: v_accvgpr_write_b32 a0, v0
; GFX908-NEXT: v_accvgpr_write_b32 a1, v1
; GFX908-NEXT: v_accvgpr_write_b32 a2, v2
; GFX908-NEXT: v_accvgpr_write_b32 a3, v3
; GFX908-NEXT: s_and_b32 s1, s5, s4
; GFX908-NEXT: s_and_b32 s3, s5, s4
; GFX908-NEXT: v_mfma_f32_16x16x16f16 a[0:3], v[4:5], v[4:5], a[0:3]
; GFX908-NEXT: s_cbranch_execz .LBB0_4
; GFX908-NEXT: .LBB0_2: ; %bb
; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1
; GFX908-NEXT: s_andn2_b64 vcc, exec, s[2:3]
; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1]
; GFX908-NEXT: s_cbranch_vccz .LBB0_1
; GFX908-NEXT: ; %bb.3:
; GFX908-NEXT: ; implicit-def: $sgpr1
; GFX908-NEXT: ; implicit-def: $sgpr3
; GFX908-NEXT: .LBB0_4: ; %common.ret
; GFX908-NEXT: s_endpgm
entry:
llvm/test/CodeGen/AMDGPU/optimize-negated-cond.ll
+4 -2
View File
@@ -23,8 +23,10 @@ define amdgpu_kernel void @negated_cond(ptr addrspace(1) %arg1) {
; GCN-NEXT: ; Child Loop BB0_4 Depth 2
; GCN-NEXT: buffer_load_dword v1, off, s[8:11], 0
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: v_cmp_eq_u32_e64 s[0:1], 0, v1
; GCN-NEXT: v_cmp_ne_u32_e64 s[2:3], 0, v1
; GCN-NEXT: v_cmp_eq_u32_e32 vcc, 0, v1
; GCN-NEXT: v_cndmask_b32_e64 v1, 0, 1, vcc
; GCN-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v1
; GCN-NEXT: s_mov_b32 s12, s6
; GCN-NEXT: s_branch .LBB0_4
; GCN-NEXT: .LBB0_3: ; %Flow1
@@ -34,7 +36,7 @@ define amdgpu_kernel void @negated_cond(ptr addrspace(1) %arg1) {
; GCN-NEXT: .LBB0_4: ; %bb2
; GCN-NEXT: ; Parent Loop BB0_2 Depth=1
; GCN-NEXT: ; => This Inner Loop Header: Depth=2
; GCN-NEXT: s_andn2_b64 vcc, exec, s[0:1]
; GCN-NEXT: s_and_b64 vcc, exec, s[0:1]
; GCN-NEXT: s_lshl_b32 s12, s12, 5
; GCN-NEXT: s_cbranch_vccz .LBB0_6
; GCN-NEXT: ; %bb.5: ; in Loop: Header=BB0_4 Depth=2
llvm/test/CodeGen/AMDGPU/sdiv64.ll
+1 -1
View File
@@ -1744,8 +1744,8 @@ define i64 @v_test_sdiv_pow2_k_den_i64(i64 %x) {
; GCN-IR-NEXT: v_addc_u32_e64 v5, s[4:5], 0, -1, vcc
; GCN-IR-NEXT: s_mov_b64 s[10:11], 0
; GCN-IR-NEXT: v_mov_b32_e32 v9, 0
; GCN-IR-NEXT: v_mov_b32_e32 v3, 0
; GCN-IR-NEXT: s_movk_i32 s12, 0x7fff
; GCN-IR-NEXT: v_mov_b32_e32 v3, 0
; GCN-IR-NEXT: .LBB13_3: ; %udiv-do-while
; GCN-IR-NEXT: ; =>This Inner Loop Header: Depth=1
; GCN-IR-NEXT: v_lshl_b64 v[6:7], v[6:7], 1
llvm/test/CodeGen/AMDGPU/srem64.ll
+1 -1
View File
@@ -1865,8 +1865,8 @@ define i64 @v_test_srem_pow2_k_den_i64(i64 %x) {
; GCN-IR-NEXT: v_addc_u32_e64 v7, s[4:5], 0, -1, vcc
; GCN-IR-NEXT: s_mov_b64 s[10:11], 0
; GCN-IR-NEXT: v_mov_b32_e32 v11, 0
; GCN-IR-NEXT: v_mov_b32_e32 v5, 0
; GCN-IR-NEXT: s_movk_i32 s12, 0x7fff
; GCN-IR-NEXT: v_mov_b32_e32 v5, 0
; GCN-IR-NEXT: .LBB13_3: ; %udiv-do-while
; GCN-IR-NEXT: ; =>This Inner Loop Header: Depth=1
; GCN-IR-NEXT: v_lshl_b64 v[8:9], v[8:9], 1
llvm/test/CodeGen/AMDGPU/tuple-allocation-failure.ll
+168 -124
View File
@@ -32,51 +32,71 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-LABEL: kernel:
; GLOBALNESS1: ; %bb.0: ; %bb
; GLOBALNESS1-NEXT: s_mov_b64 s[36:37], s[6:7]
; GLOBALNESS1-NEXT: s_load_dwordx4 s[64:67], s[8:9], 0x0
; GLOBALNESS1-NEXT: s_load_dwordx4 s[76:79], s[8:9], 0x0
; GLOBALNESS1-NEXT: s_load_dword s6, s[8:9], 0x14
; GLOBALNESS1-NEXT: v_mov_b32_e32 v41, v0
; GLOBALNESS1-NEXT: v_mov_b32_e32 v42, 0
; GLOBALNESS1-NEXT: v_pk_mov_b32 v[0:1], 0, 0
; GLOBALNESS1-NEXT: global_store_dword v[0:1], v42, off
; GLOBALNESS1-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS1-NEXT: global_load_dword v2, v42, s[64:65]
; GLOBALNESS1-NEXT: s_add_u32 flat_scratch_lo, s12, s17
; GLOBALNESS1-NEXT: s_addc_u32 flat_scratch_hi, s13, 0
; GLOBALNESS1-NEXT: s_add_u32 s0, s0, s17
; GLOBALNESS1-NEXT: s_addc_u32 s1, s1, 0
; GLOBALNESS1-NEXT: global_load_dword v2, v42, s[76:77]
; GLOBALNESS1-NEXT: s_mov_b64 s[40:41], s[4:5]
; GLOBALNESS1-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x18
; GLOBALNESS1-NEXT: s_load_dword s7, s[8:9], 0x20
; GLOBALNESS1-NEXT: s_bitcmp1_b32 s66, 0
; GLOBALNESS1-NEXT: s_cselect_b64 s[68:69], -1, 0
; GLOBALNESS1-NEXT: s_add_u32 flat_scratch_lo, s12, s17
; GLOBALNESS1-NEXT: s_addc_u32 flat_scratch_hi, s13, 0
; GLOBALNESS1-NEXT: s_add_u32 s0, s0, s17
; GLOBALNESS1-NEXT: v_mov_b32_e32 v0, 0
; GLOBALNESS1-NEXT: s_xor_b64 s[70:71], s[68:69], -1
; GLOBALNESS1-NEXT: s_addc_u32 s1, s1, 0
; GLOBALNESS1-NEXT: v_mov_b32_e32 v1, 0x40994400
; GLOBALNESS1-NEXT: s_bitcmp1_b32 s6, 0
; GLOBALNESS1-NEXT: s_bitcmp1_b32 s78, 0
; GLOBALNESS1-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS1-NEXT: v_cmp_ngt_f64_e64 s[42:43], s[4:5], v[0:1]
; GLOBALNESS1-NEXT: v_cmp_ngt_f64_e64 s[64:65], s[4:5], 0
; GLOBALNESS1-NEXT: v_cmp_ngt_f64_e32 vcc, s[4:5], v[0:1]
; GLOBALNESS1-NEXT: v_cmp_ngt_f64_e64 s[4:5], s[4:5], 0
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v1, 0, 1, s[4:5]
; GLOBALNESS1-NEXT: s_cselect_b64 s[4:5], -1, 0
; GLOBALNESS1-NEXT: s_xor_b64 s[72:73], s[4:5], -1
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v3, 0, 1, s[4:5]
; GLOBALNESS1-NEXT: s_xor_b64 s[4:5], s[4:5], -1
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc
; GLOBALNESS1-NEXT: s_bitcmp1_b32 s6, 0
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[42:43], 1, v0
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[4:5]
; GLOBALNESS1-NEXT: s_cselect_b64 s[4:5], -1, 0
; GLOBALNESS1-NEXT: s_xor_b64 s[4:5], s[4:5], -1
; GLOBALNESS1-NEXT: s_bitcmp1_b32 s7, 0
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[48:49], 1, v0
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[4:5]
; GLOBALNESS1-NEXT: s_cselect_b64 s[4:5], -1, 0
; GLOBALNESS1-NEXT: s_mov_b32 s60, s16
; GLOBALNESS1-NEXT: s_xor_b64 s[4:5], s[4:5], -1
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[50:51], 1, v0
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[4:5]
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[52:53], 1, v0
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[44:45], 1, v1
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[46:47], 1, v3
; GLOBALNESS1-NEXT: s_mov_b32 s70, s16
; GLOBALNESS1-NEXT: s_mov_b64 s[38:39], s[8:9]
; GLOBALNESS1-NEXT: s_mov_b32 s61, s15
; GLOBALNESS1-NEXT: s_mov_b32 s62, s14
; GLOBALNESS1-NEXT: s_mov_b32 s71, s15
; GLOBALNESS1-NEXT: s_mov_b32 s72, s14
; GLOBALNESS1-NEXT: s_mov_b64 s[34:35], s[10:11]
; GLOBALNESS1-NEXT: s_xor_b64 s[74:75], s[4:5], -1
; GLOBALNESS1-NEXT: s_mov_b32 s32, 0
; GLOBALNESS1-NEXT: ; implicit-def: $vgpr44_vgpr45
; GLOBALNESS1-NEXT: s_waitcnt vmcnt(0)
; GLOBALNESS1-NEXT: v_cmp_gt_i32_e64 s[44:45], 0, v2
; GLOBALNESS1-NEXT: v_cmp_gt_i32_e64 s[46:47], 1, v2
; GLOBALNESS1-NEXT: v_cmp_eq_u32_e64 s[48:49], 1, v2
; GLOBALNESS1-NEXT: v_cmp_eq_u32_e64 s[50:51], 0, v2
; GLOBALNESS1-NEXT: v_cmp_gt_i32_e32 vcc, 0, v2
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc
; GLOBALNESS1-NEXT: v_cmp_gt_i32_e32 vcc, 1, v2
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v1, 0, 1, vcc
; GLOBALNESS1-NEXT: v_cmp_eq_u32_e32 vcc, 1, v2
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v3, 0, 1, vcc
; GLOBALNESS1-NEXT: v_cmp_eq_u32_e32 vcc, 0, v2
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v2, 0, 1, vcc
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[54:55], 1, v0
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[56:57], 1, v1
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[58:59], 1, v3
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[60:61], 1, v2
; GLOBALNESS1-NEXT: s_branch .LBB1_4
; GLOBALNESS1-NEXT: .LBB1_1: ; %bb70.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[50:51]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[60:61]
; GLOBALNESS1-NEXT: s_cbranch_vccz .LBB1_29
; GLOBALNESS1-NEXT: .LBB1_2: ; %Flow15
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
@@ -105,23 +125,23 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS1-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS1-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS1-NEXT: s_mov_b32 s12, s62
; GLOBALNESS1-NEXT: s_mov_b32 s13, s61
; GLOBALNESS1-NEXT: s_mov_b32 s14, s60
; GLOBALNESS1-NEXT: s_mov_b32 s12, s72
; GLOBALNESS1-NEXT: s_mov_b32 s13, s71
; GLOBALNESS1-NEXT: s_mov_b32 s14, s70
; GLOBALNESS1-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS1-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS1-NEXT: s_swappc_b64 s[30:31], s[16:17]
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[68:69]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[46:47]
; GLOBALNESS1-NEXT: s_mov_b64 s[6:7], -1
; GLOBALNESS1-NEXT: ; implicit-def: $sgpr4_sgpr5
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_9
; GLOBALNESS1-NEXT: ; %bb.5: ; %NodeBlock
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS1-NEXT: s_cmp_lt_i32 s67, 1
; GLOBALNESS1-NEXT: s_cmp_lt_i32 s79, 1
; GLOBALNESS1-NEXT: s_cbranch_scc1 .LBB1_7
; GLOBALNESS1-NEXT: ; %bb.6: ; %LeafBlock12
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS1-NEXT: s_cmp_lg_u32 s67, 1
; GLOBALNESS1-NEXT: s_cmp_lg_u32 s79, 1
; GLOBALNESS1-NEXT: s_mov_b64 s[4:5], -1
; GLOBALNESS1-NEXT: s_cselect_b64 s[6:7], -1, 0
; GLOBALNESS1-NEXT: s_cbranch_execz .LBB1_8
@@ -131,7 +151,7 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-NEXT: ; implicit-def: $sgpr4_sgpr5
; GLOBALNESS1-NEXT: .LBB1_8: ; %LeafBlock
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS1-NEXT: s_cmp_lg_u32 s67, 0
; GLOBALNESS1-NEXT: s_cmp_lg_u32 s79, 0
; GLOBALNESS1-NEXT: s_mov_b64 s[4:5], 0
; GLOBALNESS1-NEXT: s_cselect_b64 s[6:7], -1, 0
; GLOBALNESS1-NEXT: .LBB1_9: ; %Flow25
@@ -143,15 +163,15 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-NEXT: v_pk_mov_b32 v[2:3], 0, 0
; GLOBALNESS1-NEXT: flat_load_dword v0, v[2:3]
; GLOBALNESS1-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GLOBALNESS1-NEXT: v_cmp_gt_i32_e64 s[52:53], 0, v0
; GLOBALNESS1-NEXT: v_cmp_gt_i32_e64 s[62:63], 0, v0
; GLOBALNESS1-NEXT: v_mov_b32_e32 v0, 0
; GLOBALNESS1-NEXT: v_mov_b32_e32 v1, 0x3ff00000
; GLOBALNESS1-NEXT: s_and_saveexec_b64 s[76:77], s[52:53]
; GLOBALNESS1-NEXT: s_and_saveexec_b64 s[74:75], s[62:63]
; GLOBALNESS1-NEXT: s_cbranch_execz .LBB1_26
; GLOBALNESS1-NEXT: ; %bb.11: ; %bb33.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS1-NEXT: global_load_dwordx2 v[0:1], v[2:3], off
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[44:45]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[54:55]
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_13
; GLOBALNESS1-NEXT: ; %bb.12: ; %bb39.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
@@ -163,70 +183,72 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e32 vcc, 0, v46
; GLOBALNESS1-NEXT: v_cndmask_b32_e32 v2, 0, v40, vcc
; GLOBALNESS1-NEXT: s_waitcnt vmcnt(0)
; GLOBALNESS1-NEXT: v_cmp_nlt_f64_e64 s[54:55], 0, v[0:1]
; GLOBALNESS1-NEXT: v_cmp_eq_u32_e64 s[56:57], 0, v2
; GLOBALNESS1-NEXT: v_cmp_nlt_f64_e32 vcc, 0, v[0:1]
; GLOBALNESS1-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc
; GLOBALNESS1-NEXT: v_cmp_eq_u32_e64 s[64:65], 0, v2
; GLOBALNESS1-NEXT: v_cmp_ne_u32_e64 s[66:67], 1, v0
; GLOBALNESS1-NEXT: s_branch .LBB1_16
; GLOBALNESS1-NEXT: .LBB1_14: ; %Flow16
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_or_b64 exec, exec, s[4:5]
; GLOBALNESS1-NEXT: .LBB1_15: ; %bb63.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[74:75]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[52:53]
; GLOBALNESS1-NEXT: s_cbranch_vccz .LBB1_25
; GLOBALNESS1-NEXT: .LBB1_16: ; %bb44.i
; GLOBALNESS1-NEXT: ; Parent Loop BB1_4 Depth=1
; GLOBALNESS1-NEXT: ; => This Inner Loop Header: Depth=2
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[70:71]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[48:49]
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_15
; GLOBALNESS1-NEXT: ; %bb.17: ; %bb46.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[72:73]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[50:51]
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_15
; GLOBALNESS1-NEXT: ; %bb.18: ; %bb50.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[42:43]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[42:43]
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_21
; GLOBALNESS1-NEXT: ; %bb.19: ; %bb3.i.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[64:65]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[44:45]
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_21
; GLOBALNESS1-NEXT: ; %bb.20: ; %bb6.i.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[54:55]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[66:67]
; GLOBALNESS1-NEXT: .LBB1_21: ; %spam.exit.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[46:47]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[56:57]
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_15
; GLOBALNESS1-NEXT: ; %bb.22: ; %bb55.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS1-NEXT: s_add_u32 s58, s38, 40
; GLOBALNESS1-NEXT: s_addc_u32 s59, s39, 0
; GLOBALNESS1-NEXT: s_add_u32 s68, s38, 40
; GLOBALNESS1-NEXT: s_addc_u32 s69, s39, 0
; GLOBALNESS1-NEXT: s_getpc_b64 s[4:5]
; GLOBALNESS1-NEXT: s_add_u32 s4, s4, wobble@gotpcrel32@lo+4
; GLOBALNESS1-NEXT: s_addc_u32 s5, s5, wobble@gotpcrel32@hi+12
; GLOBALNESS1-NEXT: s_load_dwordx2 s[78:79], s[4:5], 0x0
; GLOBALNESS1-NEXT: s_load_dwordx2 s[76:77], s[4:5], 0x0
; GLOBALNESS1-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS1-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS1-NEXT: s_mov_b64 s[8:9], s[58:59]
; GLOBALNESS1-NEXT: s_mov_b64 s[8:9], s[68:69]
; GLOBALNESS1-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS1-NEXT: s_mov_b32 s12, s62
; GLOBALNESS1-NEXT: s_mov_b32 s13, s61
; GLOBALNESS1-NEXT: s_mov_b32 s14, s60
; GLOBALNESS1-NEXT: s_mov_b32 s12, s72
; GLOBALNESS1-NEXT: s_mov_b32 s13, s71
; GLOBALNESS1-NEXT: s_mov_b32 s14, s70
; GLOBALNESS1-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS1-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS1-NEXT: s_swappc_b64 s[30:31], s[78:79]
; GLOBALNESS1-NEXT: s_swappc_b64 s[30:31], s[76:77]
; GLOBALNESS1-NEXT: v_pk_mov_b32 v[46:47], 0, 0
; GLOBALNESS1-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS1-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS1-NEXT: s_mov_b64 s[8:9], s[58:59]
; GLOBALNESS1-NEXT: s_mov_b64 s[8:9], s[68:69]
; GLOBALNESS1-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS1-NEXT: s_mov_b32 s12, s62
; GLOBALNESS1-NEXT: s_mov_b32 s13, s61
; GLOBALNESS1-NEXT: s_mov_b32 s14, s60
; GLOBALNESS1-NEXT: s_mov_b32 s12, s72
; GLOBALNESS1-NEXT: s_mov_b32 s13, s71
; GLOBALNESS1-NEXT: s_mov_b32 s14, s70
; GLOBALNESS1-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS1-NEXT: global_store_dwordx2 v[46:47], v[44:45], off
; GLOBALNESS1-NEXT: s_swappc_b64 s[30:31], s[78:79]
; GLOBALNESS1-NEXT: s_and_saveexec_b64 s[4:5], s[56:57]
; GLOBALNESS1-NEXT: s_swappc_b64 s[30:31], s[76:77]
; GLOBALNESS1-NEXT: s_and_saveexec_b64 s[4:5], s[64:65]
; GLOBALNESS1-NEXT: s_cbranch_execz .LBB1_14
; GLOBALNESS1-NEXT: ; %bb.23: ; %bb62.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_16 Depth=2
@@ -242,12 +264,12 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-NEXT: v_pk_mov_b32 v[0:1], 0, 0
; GLOBALNESS1-NEXT: .LBB1_26: ; %Flow24
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS1-NEXT: s_or_b64 exec, exec, s[76:77]
; GLOBALNESS1-NEXT: s_and_saveexec_b64 s[4:5], s[52:53]
; GLOBALNESS1-NEXT: s_or_b64 exec, exec, s[74:75]
; GLOBALNESS1-NEXT: s_and_saveexec_b64 s[4:5], s[62:63]
; GLOBALNESS1-NEXT: s_cbranch_execz .LBB1_2
; GLOBALNESS1-NEXT: ; %bb.27: ; %bb67.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS1-NEXT: s_andn2_b64 vcc, exec, s[48:49]
; GLOBALNESS1-NEXT: s_and_b64 vcc, exec, s[58:59]
; GLOBALNESS1-NEXT: s_cbranch_vccnz .LBB1_1
; GLOBALNESS1-NEXT: ; %bb.28: ; %bb69.i
; GLOBALNESS1-NEXT: ; in Loop: Header=BB1_4 Depth=1
@@ -271,9 +293,9 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS1-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS1-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS1-NEXT: s_mov_b32 s12, s62
; GLOBALNESS1-NEXT: s_mov_b32 s13, s61
; GLOBALNESS1-NEXT: s_mov_b32 s14, s60
; GLOBALNESS1-NEXT: s_mov_b32 s12, s72
; GLOBALNESS1-NEXT: s_mov_b32 s13, s71
; GLOBALNESS1-NEXT: s_mov_b32 s14, s70
; GLOBALNESS1-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS1-NEXT: s_getpc_b64 s[16:17]
; GLOBALNESS1-NEXT: s_add_u32 s16, s16, widget@rel32@lo+4
@@ -289,9 +311,9 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS1-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS1-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS1-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS1-NEXT: s_mov_b32 s12, s62
; GLOBALNESS1-NEXT: s_mov_b32 s13, s61
; GLOBALNESS1-NEXT: s_mov_b32 s14, s60
; GLOBALNESS1-NEXT: s_mov_b32 s12, s72
; GLOBALNESS1-NEXT: s_mov_b32 s13, s71
; GLOBALNESS1-NEXT: s_mov_b32 s14, s70
; GLOBALNESS1-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS1-NEXT: s_getpc_b64 s[16:17]
; GLOBALNESS1-NEXT: s_add_u32 s16, s16, widget@rel32@lo+4
@@ -302,51 +324,71 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-LABEL: kernel:
; GLOBALNESS0: ; %bb.0: ; %bb
; GLOBALNESS0-NEXT: s_mov_b64 s[36:37], s[6:7]
; GLOBALNESS0-NEXT: s_load_dwordx4 s[64:67], s[8:9], 0x0
; GLOBALNESS0-NEXT: s_load_dwordx4 s[72:75], s[8:9], 0x0
; GLOBALNESS0-NEXT: s_load_dword s6, s[8:9], 0x14
; GLOBALNESS0-NEXT: v_mov_b32_e32 v41, v0
; GLOBALNESS0-NEXT: v_mov_b32_e32 v42, 0
; GLOBALNESS0-NEXT: v_pk_mov_b32 v[0:1], 0, 0
; GLOBALNESS0-NEXT: global_store_dword v[0:1], v42, off
; GLOBALNESS0-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS0-NEXT: global_load_dword v2, v42, s[64:65]
; GLOBALNESS0-NEXT: s_add_u32 flat_scratch_lo, s12, s17
; GLOBALNESS0-NEXT: s_addc_u32 flat_scratch_hi, s13, 0
; GLOBALNESS0-NEXT: s_add_u32 s0, s0, s17
; GLOBALNESS0-NEXT: s_addc_u32 s1, s1, 0
; GLOBALNESS0-NEXT: global_load_dword v2, v42, s[72:73]
; GLOBALNESS0-NEXT: s_mov_b64 s[40:41], s[4:5]
; GLOBALNESS0-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x18
; GLOBALNESS0-NEXT: s_load_dword s7, s[8:9], 0x20
; GLOBALNESS0-NEXT: s_bitcmp1_b32 s66, 0
; GLOBALNESS0-NEXT: s_cselect_b64 s[68:69], -1, 0
; GLOBALNESS0-NEXT: s_add_u32 flat_scratch_lo, s12, s17
; GLOBALNESS0-NEXT: s_addc_u32 flat_scratch_hi, s13, 0
; GLOBALNESS0-NEXT: s_add_u32 s0, s0, s17
; GLOBALNESS0-NEXT: v_mov_b32_e32 v0, 0
; GLOBALNESS0-NEXT: s_xor_b64 s[70:71], s[68:69], -1
; GLOBALNESS0-NEXT: s_addc_u32 s1, s1, 0
; GLOBALNESS0-NEXT: v_mov_b32_e32 v1, 0x40994400
; GLOBALNESS0-NEXT: s_bitcmp1_b32 s6, 0
; GLOBALNESS0-NEXT: s_bitcmp1_b32 s74, 0
; GLOBALNESS0-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS0-NEXT: v_cmp_ngt_f64_e64 s[42:43], s[4:5], v[0:1]
; GLOBALNESS0-NEXT: v_cmp_ngt_f64_e64 s[64:65], s[4:5], 0
; GLOBALNESS0-NEXT: v_cmp_ngt_f64_e32 vcc, s[4:5], v[0:1]
; GLOBALNESS0-NEXT: v_cmp_ngt_f64_e64 s[4:5], s[4:5], 0
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v1, 0, 1, s[4:5]
; GLOBALNESS0-NEXT: s_cselect_b64 s[4:5], -1, 0
; GLOBALNESS0-NEXT: s_xor_b64 s[72:73], s[4:5], -1
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v3, 0, 1, s[4:5]
; GLOBALNESS0-NEXT: s_xor_b64 s[4:5], s[4:5], -1
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc
; GLOBALNESS0-NEXT: s_bitcmp1_b32 s6, 0
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[42:43], 1, v0
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[4:5]
; GLOBALNESS0-NEXT: s_cselect_b64 s[4:5], -1, 0
; GLOBALNESS0-NEXT: s_xor_b64 s[4:5], s[4:5], -1
; GLOBALNESS0-NEXT: s_bitcmp1_b32 s7, 0
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[48:49], 1, v0
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[4:5]
; GLOBALNESS0-NEXT: s_cselect_b64 s[4:5], -1, 0
; GLOBALNESS0-NEXT: s_mov_b32 s58, s16
; GLOBALNESS0-NEXT: s_xor_b64 s[4:5], s[4:5], -1
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[50:51], 1, v0
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[4:5]
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[52:53], 1, v0
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[44:45], 1, v1
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[46:47], 1, v3
; GLOBALNESS0-NEXT: s_mov_b32 s68, s16
; GLOBALNESS0-NEXT: s_mov_b64 s[38:39], s[8:9]
; GLOBALNESS0-NEXT: s_mov_b32 s59, s15
; GLOBALNESS0-NEXT: s_mov_b32 s60, s14
; GLOBALNESS0-NEXT: s_mov_b32 s69, s15
; GLOBALNESS0-NEXT: s_mov_b32 s70, s14
; GLOBALNESS0-NEXT: s_mov_b64 s[34:35], s[10:11]
; GLOBALNESS0-NEXT: s_xor_b64 s[74:75], s[4:5], -1
; GLOBALNESS0-NEXT: s_mov_b32 s32, 0
; GLOBALNESS0-NEXT: ; implicit-def: $vgpr44_vgpr45
; GLOBALNESS0-NEXT: s_waitcnt vmcnt(0)
; GLOBALNESS0-NEXT: v_cmp_gt_i32_e64 s[44:45], 0, v2
; GLOBALNESS0-NEXT: v_cmp_gt_i32_e64 s[46:47], 1, v2
; GLOBALNESS0-NEXT: v_cmp_eq_u32_e64 s[48:49], 1, v2
; GLOBALNESS0-NEXT: v_cmp_eq_u32_e64 s[50:51], 0, v2
; GLOBALNESS0-NEXT: v_cmp_gt_i32_e32 vcc, 0, v2
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc
; GLOBALNESS0-NEXT: v_cmp_gt_i32_e32 vcc, 1, v2
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v1, 0, 1, vcc
; GLOBALNESS0-NEXT: v_cmp_eq_u32_e32 vcc, 1, v2
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v3, 0, 1, vcc
; GLOBALNESS0-NEXT: v_cmp_eq_u32_e32 vcc, 0, v2
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v2, 0, 1, vcc
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[54:55], 1, v0
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[56:57], 1, v1
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[58:59], 1, v3
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[60:61], 1, v2
; GLOBALNESS0-NEXT: s_branch .LBB1_4
; GLOBALNESS0-NEXT: .LBB1_1: ; %bb70.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[50:51]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[60:61]
; GLOBALNESS0-NEXT: s_cbranch_vccz .LBB1_29
; GLOBALNESS0-NEXT: .LBB1_2: ; %Flow15
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
@@ -375,23 +417,23 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS0-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS0-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS0-NEXT: s_mov_b32 s12, s60
; GLOBALNESS0-NEXT: s_mov_b32 s13, s59
; GLOBALNESS0-NEXT: s_mov_b32 s14, s58
; GLOBALNESS0-NEXT: s_mov_b32 s12, s70
; GLOBALNESS0-NEXT: s_mov_b32 s13, s69
; GLOBALNESS0-NEXT: s_mov_b32 s14, s68
; GLOBALNESS0-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS0-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS0-NEXT: s_swappc_b64 s[30:31], s[16:17]
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[68:69]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[46:47]
; GLOBALNESS0-NEXT: s_mov_b64 s[6:7], -1
; GLOBALNESS0-NEXT: ; implicit-def: $sgpr4_sgpr5
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_9
; GLOBALNESS0-NEXT: ; %bb.5: ; %NodeBlock
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS0-NEXT: s_cmp_lt_i32 s67, 1
; GLOBALNESS0-NEXT: s_cmp_lt_i32 s75, 1
; GLOBALNESS0-NEXT: s_cbranch_scc1 .LBB1_7
; GLOBALNESS0-NEXT: ; %bb.6: ; %LeafBlock12
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS0-NEXT: s_cmp_lg_u32 s67, 1
; GLOBALNESS0-NEXT: s_cmp_lg_u32 s75, 1
; GLOBALNESS0-NEXT: s_mov_b64 s[4:5], -1
; GLOBALNESS0-NEXT: s_cselect_b64 s[6:7], -1, 0
; GLOBALNESS0-NEXT: s_cbranch_execz .LBB1_8
@@ -401,7 +443,7 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-NEXT: ; implicit-def: $sgpr4_sgpr5
; GLOBALNESS0-NEXT: .LBB1_8: ; %LeafBlock
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS0-NEXT: s_cmp_lg_u32 s67, 0
; GLOBALNESS0-NEXT: s_cmp_lg_u32 s75, 0
; GLOBALNESS0-NEXT: s_mov_b64 s[4:5], 0
; GLOBALNESS0-NEXT: s_cselect_b64 s[6:7], -1, 0
; GLOBALNESS0-NEXT: .LBB1_9: ; %Flow25
@@ -413,15 +455,15 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-NEXT: v_pk_mov_b32 v[2:3], 0, 0
; GLOBALNESS0-NEXT: flat_load_dword v0, v[2:3]
; GLOBALNESS0-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GLOBALNESS0-NEXT: v_cmp_gt_i32_e64 s[52:53], 0, v0
; GLOBALNESS0-NEXT: v_cmp_gt_i32_e64 s[62:63], 0, v0
; GLOBALNESS0-NEXT: v_mov_b32_e32 v0, 0
; GLOBALNESS0-NEXT: v_mov_b32_e32 v1, 0x3ff00000
; GLOBALNESS0-NEXT: s_and_saveexec_b64 s[76:77], s[52:53]
; GLOBALNESS0-NEXT: s_and_saveexec_b64 s[76:77], s[62:63]
; GLOBALNESS0-NEXT: s_cbranch_execz .LBB1_26
; GLOBALNESS0-NEXT: ; %bb.11: ; %bb33.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS0-NEXT: global_load_dwordx2 v[0:1], v[2:3], off
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[44:45]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[54:55]
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_13
; GLOBALNESS0-NEXT: ; %bb.12: ; %bb39.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
@@ -433,70 +475,72 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e32 vcc, 0, v46
; GLOBALNESS0-NEXT: v_cndmask_b32_e32 v2, 0, v40, vcc
; GLOBALNESS0-NEXT: s_waitcnt vmcnt(0)
; GLOBALNESS0-NEXT: v_cmp_nlt_f64_e64 s[54:55], 0, v[0:1]
; GLOBALNESS0-NEXT: v_cmp_eq_u32_e64 s[56:57], 0, v2
; GLOBALNESS0-NEXT: v_cmp_nlt_f64_e32 vcc, 0, v[0:1]
; GLOBALNESS0-NEXT: v_cndmask_b32_e64 v0, 0, 1, vcc
; GLOBALNESS0-NEXT: v_cmp_eq_u32_e64 s[64:65], 0, v2
; GLOBALNESS0-NEXT: v_cmp_ne_u32_e64 s[66:67], 1, v0
; GLOBALNESS0-NEXT: s_branch .LBB1_16
; GLOBALNESS0-NEXT: .LBB1_14: ; %Flow16
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_or_b64 exec, exec, s[4:5]
; GLOBALNESS0-NEXT: .LBB1_15: ; %bb63.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[74:75]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[52:53]
; GLOBALNESS0-NEXT: s_cbranch_vccz .LBB1_25
; GLOBALNESS0-NEXT: .LBB1_16: ; %bb44.i
; GLOBALNESS0-NEXT: ; Parent Loop BB1_4 Depth=1
; GLOBALNESS0-NEXT: ; => This Inner Loop Header: Depth=2
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[70:71]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[48:49]
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_15
; GLOBALNESS0-NEXT: ; %bb.17: ; %bb46.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[72:73]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[50:51]
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_15
; GLOBALNESS0-NEXT: ; %bb.18: ; %bb50.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[42:43]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[42:43]
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_21
; GLOBALNESS0-NEXT: ; %bb.19: ; %bb3.i.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[64:65]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[44:45]
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_21
; GLOBALNESS0-NEXT: ; %bb.20: ; %bb6.i.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[54:55]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[66:67]
; GLOBALNESS0-NEXT: .LBB1_21: ; %spam.exit.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[46:47]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[56:57]
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_15
; GLOBALNESS0-NEXT: ; %bb.22: ; %bb55.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
; GLOBALNESS0-NEXT: s_add_u32 s62, s38, 40
; GLOBALNESS0-NEXT: s_addc_u32 s63, s39, 0
; GLOBALNESS0-NEXT: s_add_u32 s72, s38, 40
; GLOBALNESS0-NEXT: s_addc_u32 s73, s39, 0
; GLOBALNESS0-NEXT: s_getpc_b64 s[4:5]
; GLOBALNESS0-NEXT: s_add_u32 s4, s4, wobble@gotpcrel32@lo+4
; GLOBALNESS0-NEXT: s_addc_u32 s5, s5, wobble@gotpcrel32@hi+12
; GLOBALNESS0-NEXT: s_load_dwordx2 s[78:79], s[4:5], 0x0
; GLOBALNESS0-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS0-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS0-NEXT: s_mov_b64 s[8:9], s[62:63]
; GLOBALNESS0-NEXT: s_mov_b64 s[8:9], s[72:73]
; GLOBALNESS0-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS0-NEXT: s_mov_b32 s12, s60
; GLOBALNESS0-NEXT: s_mov_b32 s13, s59
; GLOBALNESS0-NEXT: s_mov_b32 s14, s58
; GLOBALNESS0-NEXT: s_mov_b32 s12, s70
; GLOBALNESS0-NEXT: s_mov_b32 s13, s69
; GLOBALNESS0-NEXT: s_mov_b32 s14, s68
; GLOBALNESS0-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS0-NEXT: s_waitcnt lgkmcnt(0)
; GLOBALNESS0-NEXT: s_swappc_b64 s[30:31], s[78:79]
; GLOBALNESS0-NEXT: v_pk_mov_b32 v[46:47], 0, 0
; GLOBALNESS0-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS0-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS0-NEXT: s_mov_b64 s[8:9], s[62:63]
; GLOBALNESS0-NEXT: s_mov_b64 s[8:9], s[72:73]
; GLOBALNESS0-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS0-NEXT: s_mov_b32 s12, s60
; GLOBALNESS0-NEXT: s_mov_b32 s13, s59
; GLOBALNESS0-NEXT: s_mov_b32 s14, s58
; GLOBALNESS0-NEXT: s_mov_b32 s12, s70
; GLOBALNESS0-NEXT: s_mov_b32 s13, s69
; GLOBALNESS0-NEXT: s_mov_b32 s14, s68
; GLOBALNESS0-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS0-NEXT: global_store_dwordx2 v[46:47], v[44:45], off
; GLOBALNESS0-NEXT: s_swappc_b64 s[30:31], s[78:79]
; GLOBALNESS0-NEXT: s_and_saveexec_b64 s[4:5], s[56:57]
; GLOBALNESS0-NEXT: s_and_saveexec_b64 s[4:5], s[64:65]
; GLOBALNESS0-NEXT: s_cbranch_execz .LBB1_14
; GLOBALNESS0-NEXT: ; %bb.23: ; %bb62.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_16 Depth=2
@@ -513,11 +557,11 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-NEXT: .LBB1_26: ; %Flow24
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS0-NEXT: s_or_b64 exec, exec, s[76:77]
; GLOBALNESS0-NEXT: s_and_saveexec_b64 s[4:5], s[52:53]
; GLOBALNESS0-NEXT: s_and_saveexec_b64 s[4:5], s[62:63]
; GLOBALNESS0-NEXT: s_cbranch_execz .LBB1_2
; GLOBALNESS0-NEXT: ; %bb.27: ; %bb67.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
; GLOBALNESS0-NEXT: s_andn2_b64 vcc, exec, s[48:49]
; GLOBALNESS0-NEXT: s_and_b64 vcc, exec, s[58:59]
; GLOBALNESS0-NEXT: s_cbranch_vccnz .LBB1_1
; GLOBALNESS0-NEXT: ; %bb.28: ; %bb69.i
; GLOBALNESS0-NEXT: ; in Loop: Header=BB1_4 Depth=1
@@ -541,9 +585,9 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS0-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS0-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS0-NEXT: s_mov_b32 s12, s60
; GLOBALNESS0-NEXT: s_mov_b32 s13, s59
; GLOBALNESS0-NEXT: s_mov_b32 s14, s58
; GLOBALNESS0-NEXT: s_mov_b32 s12, s70
; GLOBALNESS0-NEXT: s_mov_b32 s13, s69
; GLOBALNESS0-NEXT: s_mov_b32 s14, s68
; GLOBALNESS0-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS0-NEXT: s_getpc_b64 s[16:17]
; GLOBALNESS0-NEXT: s_add_u32 s16, s16, widget@rel32@lo+4
@@ -559,9 +603,9 @@ define amdgpu_kernel void @kernel(ptr addrspace(1) %arg1.global, i1 %tmp3.i.i, i
; GLOBALNESS0-NEXT: s_mov_b64 s[4:5], s[40:41]
; GLOBALNESS0-NEXT: s_mov_b64 s[6:7], s[36:37]
; GLOBALNESS0-NEXT: s_mov_b64 s[10:11], s[34:35]
; GLOBALNESS0-NEXT: s_mov_b32 s12, s60
; GLOBALNESS0-NEXT: s_mov_b32 s13, s59
; GLOBALNESS0-NEXT: s_mov_b32 s14, s58
; GLOBALNESS0-NEXT: s_mov_b32 s12, s70
; GLOBALNESS0-NEXT: s_mov_b32 s13, s69
; GLOBALNESS0-NEXT: s_mov_b32 s14, s68
; GLOBALNESS0-NEXT: v_mov_b32_e32 v31, v41
; GLOBALNESS0-NEXT: s_getpc_b64 s[16:17]
; GLOBALNESS0-NEXT: s_add_u32 s16, s16, widget@rel32@lo+4
llvm/test/CodeGen/AMDGPU/udiv64.ll
+1 -1
View File
@@ -1187,8 +1187,8 @@ define i64 @v_test_udiv_pow2_k_den_i64(i64 %x) {
; GCN-IR-NEXT: v_addc_u32_e64 v1, s[4:5], 0, -1, vcc
; GCN-IR-NEXT: s_mov_b64 s[10:11], 0
; GCN-IR-NEXT: v_mov_b32_e32 v9, 0
; GCN-IR-NEXT: v_mov_b32_e32 v5, 0
; GCN-IR-NEXT: s_movk_i32 s12, 0x7fff
; GCN-IR-NEXT: v_mov_b32_e32 v5, 0
; GCN-IR-NEXT: .LBB10_3: ; %udiv-do-while
; GCN-IR-NEXT: ; =>This Inner Loop Header: Depth=1
; GCN-IR-NEXT: v_lshl_b64 v[6:7], v[6:7], 1
llvm/test/CodeGen/AMDGPU/undefined-subreg-liverange.ll
+5 -2
View File
@@ -52,10 +52,13 @@ define amdgpu_ps float @valley_partially_undef_copy() #0 {
; CHECK-NEXT: v_mov_b32_e32 v2, 0x7fc00000
; CHECK-NEXT: buffer_store_dwordx4 v[0:3], off, s[0:3], 0
; CHECK-NEXT: buffer_store_dword v2, off, s[0:3], 0
; CHECK-NEXT: v_cmp_ne_u32_e64 s[0:1], 0, v1
; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v1
; CHECK-NEXT: s_waitcnt expcnt(1)
; CHECK-NEXT: v_cndmask_b32_e64 v1, 0, 1, vcc
; CHECK-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v1
; CHECK-NEXT: .LBB1_1: ; %bb9
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[0:1]
; CHECK-NEXT: s_and_b64 vcc, exec, s[0:1]
; CHECK-NEXT: s_cbranch_vccnz .LBB1_1
; CHECK-NEXT: ; %bb.2: ; %bb11
; CHECK-NEXT: s_mov_b32 s3, 0xf000
llvm/test/CodeGen/AMDGPU/urem64.ll
+1 -1
View File
@@ -1294,8 +1294,8 @@ define i64 @v_test_urem_pow2_k_den_i64(i64 %x) {
; GCN-IR-NEXT: v_addc_u32_e64 v7, s[4:5], 0, -1, vcc
; GCN-IR-NEXT: s_mov_b64 s[10:11], 0
; GCN-IR-NEXT: v_mov_b32_e32 v11, 0
; GCN-IR-NEXT: v_mov_b32_e32 v5, 0
; GCN-IR-NEXT: s_movk_i32 s12, 0x7fff
; GCN-IR-NEXT: v_mov_b32_e32 v5, 0
; GCN-IR-NEXT: .LBB9_3: ; %udiv-do-while
; GCN-IR-NEXT: ; =>This Inner Loop Header: Depth=1
; GCN-IR-NEXT: v_lshl_b64 v[8:9], v[8:9], 1
llvm/test/CodeGen/AMDGPU/vgpr-descriptor-waterfall-loop-idom-update.ll
+3 -3
View File
@@ -6,15 +6,15 @@ define void @vgpr_descriptor_waterfall_loop_idom_update(ptr %arg) #0 {
; GCN-LABEL: vgpr_descriptor_waterfall_loop_idom_update:
; GCN: ; %bb.0: ; %entry
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GCN-NEXT: v_add_co_u32 v6, vcc_lo, v0, 8
; GCN-NEXT: v_add_co_ci_u32_e32 v7, vcc_lo, 0, v1, vcc_lo
; GCN-NEXT: .LBB0_1: ; %bb0
; GCN-NEXT: ; =>This Loop Header: Depth=1
; GCN-NEXT: ; Child Loop BB0_2 Depth 2
; GCN-NEXT: v_add_co_u32 v6, vcc_lo, v0, 8
; GCN-NEXT: v_add_co_ci_u32_e32 v7, vcc_lo, 0, v1, vcc_lo
; GCN-NEXT: s_mov_b32 s5, exec_lo
; GCN-NEXT: s_clause 0x1
; GCN-NEXT: flat_load_dwordx2 v[4:5], v[6:7]
; GCN-NEXT: flat_load_dwordx2 v[2:3], v[0:1]
; GCN-NEXT: s_mov_b32 s5, exec_lo
; GCN-NEXT: .LBB0_2: ; Parent Loop BB0_1 Depth=1
; GCN-NEXT: ; => This Inner Loop Header: Depth=2
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
llvm/test/CodeGen/LoongArch/jr-without-ra.ll
+56 -56
View File
@@ -20,101 +20,101 @@ define void @jr_without_ra(ptr %rtwdev, ptr %chan, ptr %h2c, i8 %.pre, i1 %cmp.i
; CHECK-NEXT: st.d $s6, $sp, 24 # 8-byte Folded Spill
; CHECK-NEXT: st.d $s7, $sp, 16 # 8-byte Folded Spill
; CHECK-NEXT: st.d $s8, $sp, 8 # 8-byte Folded Spill
; CHECK-NEXT: move $s6, $zero
; CHECK-NEXT: move $s1, $zero
; CHECK-NEXT: move $s7, $zero
; CHECK-NEXT: move $s0, $zero
; CHECK-NEXT: ld.d $t0, $sp, 184
; CHECK-NEXT: ld.d $t1, $sp, 176
; CHECK-NEXT: ld.d $s2, $sp, 168
; CHECK-NEXT: ld.d $t2, $sp, 160
; CHECK-NEXT: ld.d $t3, $sp, 152
; CHECK-NEXT: ld.d $t4, $sp, 144
; CHECK-NEXT: ld.d $t5, $sp, 136
; CHECK-NEXT: ld.d $t6, $sp, 128
; CHECK-NEXT: ld.d $t7, $sp, 120
; CHECK-NEXT: ld.d $t8, $sp, 112
; CHECK-NEXT: ld.d $fp, $sp, 104
; CHECK-NEXT: ld.d $s0, $sp, 96
; CHECK-NEXT: ld.d $s2, $sp, 176
; CHECK-NEXT: ld.d $s1, $sp, 168
; CHECK-NEXT: ld.d $t1, $sp, 160
; CHECK-NEXT: ld.d $t2, $sp, 152
; CHECK-NEXT: ld.d $t3, $sp, 144
; CHECK-NEXT: ld.d $t4, $sp, 136
; CHECK-NEXT: ld.d $t5, $sp, 128
; CHECK-NEXT: ld.d $t6, $sp, 120
; CHECK-NEXT: ld.d $t7, $sp, 112
; CHECK-NEXT: ld.d $t8, $sp, 104
; CHECK-NEXT: ld.d $fp, $sp, 96
; CHECK-NEXT: andi $a4, $a4, 1
; CHECK-NEXT: alsl.d $a6, $a6, $s2, 4
; CHECK-NEXT: pcalau12i $s2, %pc_hi20(.LJTI0_0)
; CHECK-NEXT: addi.d $s2, $s2, %pc_lo12(.LJTI0_0)
; CHECK-NEXT: alsl.d $a6, $a6, $s1, 4
; CHECK-NEXT: pcalau12i $s1, %pc_hi20(.LJTI0_0)
; CHECK-NEXT: addi.d $s1, $s1, %pc_lo12(.LJTI0_0)
; CHECK-NEXT: slli.d $s3, $s2, 2
; CHECK-NEXT: alsl.d $s2, $s2, $s3, 1
; CHECK-NEXT: add.d $s2, $t5, $s2
; CHECK-NEXT: addi.w $s4, $zero, -41
; CHECK-NEXT: ori $s3, $zero, 1
; CHECK-NEXT: ori $s4, $zero, 50
; CHECK-NEXT: ori $s5, $zero, 3
; CHECK-NEXT: lu32i.d $s5, 262144
; CHECK-NEXT: slli.d $s4, $s4, 3
; CHECK-NEXT: ori $s6, $zero, 3
; CHECK-NEXT: lu32i.d $s6, 262144
; CHECK-NEXT: b .LBB0_4
; CHECK-NEXT: .p2align 4, , 16
; CHECK-NEXT: .LBB0_1: # %sw.bb27.i.i
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: ori $s7, $zero, 1
; CHECK-NEXT: ori $s8, $zero, 1
; CHECK-NEXT: .LBB0_2: # %if.else.i106
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: alsl.d $s8, $s1, $s1, 3
; CHECK-NEXT: alsl.d $s1, $s8, $s1, 1
; CHECK-NEXT: add.d $s1, $t0, $s1
; CHECK-NEXT: ldx.bu $s7, $s1, $s7
; CHECK-NEXT: alsl.d $s5, $s0, $s0, 3
; CHECK-NEXT: alsl.d $s0, $s5, $s0, 1
; CHECK-NEXT: add.d $s0, $t0, $s0
; CHECK-NEXT: ldx.bu $s8, $s0, $s8
; CHECK-NEXT: .LBB0_3: # %phy_tssi_get_ofdm_de.exit
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: st.b $zero, $t6, 0
; CHECK-NEXT: st.b $s6, $t4, 0
; CHECK-NEXT: st.b $zero, $fp, 0
; CHECK-NEXT: st.b $zero, $t2, 0
; CHECK-NEXT: st.b $zero, $t5, 0
; CHECK-NEXT: st.b $s7, $t3, 0
; CHECK-NEXT: st.b $zero, $t8, 0
; CHECK-NEXT: st.b $zero, $t1, 0
; CHECK-NEXT: st.b $zero, $a1, 0
; CHECK-NEXT: st.b $zero, $t3, 0
; CHECK-NEXT: st.b $s7, $a5, 0
; CHECK-NEXT: ori $s1, $zero, 1
; CHECK-NEXT: move $s6, $a3
; CHECK-NEXT: st.b $zero, $t2, 0
; CHECK-NEXT: st.b $s8, $a5, 0
; CHECK-NEXT: ori $s0, $zero, 1
; CHECK-NEXT: move $s7, $a3
; CHECK-NEXT: .LBB0_4: # %for.body
; CHECK-NEXT: # =>This Inner Loop Header: Depth=1
; CHECK-NEXT: beqz $a4, .LBB0_9
; CHECK-NEXT: # %bb.5: # %calc_6g.i
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: move $s6, $zero
; CHECK-NEXT: move $s7, $zero
; CHECK-NEXT: bnez $zero, .LBB0_8
; CHECK-NEXT: # %bb.6: # %calc_6g.i
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: slli.d $s7, $zero, 3
; CHECK-NEXT: ldx.d $s7, $s7, $s2
; CHECK-NEXT: jr $s7
; CHECK-NEXT: slli.d $s8, $zero, 3
; CHECK-NEXT: ldx.d $s8, $s8, $s1
; CHECK-NEXT: jr $s8
; CHECK-NEXT: .LBB0_7: # %sw.bb12.i.i
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: ori $s6, $zero, 1
; CHECK-NEXT: ori $s7, $zero, 1
; CHECK-NEXT: .LBB0_8: # %if.else58.i
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: ldx.bu $s6, $a6, $s6
; CHECK-NEXT: ldx.bu $s7, $a6, $s7
; CHECK-NEXT: b .LBB0_11
; CHECK-NEXT: .p2align 4, , 16
; CHECK-NEXT: .LBB0_9: # %if.end.i
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: andi $s6, $s6, 255
; CHECK-NEXT: bltu $s4, $s6, .LBB0_15
; CHECK-NEXT: andi $s7, $s7, 255
; CHECK-NEXT: ori $s5, $zero, 50
; CHECK-NEXT: bltu $s5, $s7, .LBB0_15
; CHECK-NEXT: # %bb.10: # %if.end.i
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: sll.d $s6, $s3, $s6
; CHECK-NEXT: and $s7, $s6, $s5
; CHECK-NEXT: move $s6, $s0
; CHECK-NEXT: beqz $s7, .LBB0_15
; CHECK-NEXT: sll.d $s7, $s3, $s7
; CHECK-NEXT: and $s8, $s7, $s6
; CHECK-NEXT: move $s7, $fp
; CHECK-NEXT: beqz $s8, .LBB0_15
; CHECK-NEXT: .LBB0_11: # %phy_tssi_get_ofdm_trim_de.exit
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: move $s7, $zero
; CHECK-NEXT: st.b $zero, $t8, 0
; CHECK-NEXT: slli.d $s8, $t1, 2
; CHECK-NEXT: alsl.d $s8, $t1, $s8, 1
; CHECK-NEXT: add.d $s8, $t6, $s8
; CHECK-NEXT: ldx.b $s8, $s8, $t5
; CHECK-NEXT: move $s8, $zero
; CHECK-NEXT: st.b $zero, $t7, 0
; CHECK-NEXT: ldx.b $ra, $s2, $t4
; CHECK-NEXT: st.b $zero, $a2, 0
; CHECK-NEXT: st.b $zero, $a7, 0
; CHECK-NEXT: st.b $zero, $t7, 0
; CHECK-NEXT: st.b $s8, $a0, 0
; CHECK-NEXT: st.b $zero, $t6, 0
; CHECK-NEXT: st.b $ra, $a0, 0
; CHECK-NEXT: bnez $s3, .LBB0_13
; CHECK-NEXT: # %bb.12: # %phy_tssi_get_ofdm_trim_de.exit
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: addi.w $s8, $zero, -41
; CHECK-NEXT: slli.d $s8, $s8, 3
; CHECK-NEXT: pcalau12i $ra, %pc_hi20(.LJTI0_1)
; CHECK-NEXT: addi.d $ra, $ra, %pc_lo12(.LJTI0_1)
; CHECK-NEXT: ldx.d $s8, $s8, $ra
; CHECK-NEXT: jr $s8
; CHECK-NEXT: ldx.d $s5, $s4, $ra
; CHECK-NEXT: jr $s5
; CHECK-NEXT: .LBB0_13: # %phy_tssi_get_ofdm_trim_de.exit
; CHECK-NEXT: # in Loop: Header=BB0_4 Depth=1
; CHECK-NEXT: bnez $s3, .LBB0_1
llvm/test/CodeGen/RISCV/rvv/vxrm-insert-out-of-loop.ll
+2 -3
View File
@@ -489,9 +489,8 @@ define void @test1(ptr nocapture noundef writeonly %dst, i32 noundef signext %i_
; RV64-NEXT: j .LBB0_11
; RV64-NEXT: .LBB0_8: # %vector.ph
; RV64-NEXT: # in Loop: Header=BB0_6 Depth=1
; RV64-NEXT: slli t6, t0, 1
; RV64-NEXT: slli s0, t0, 28
; RV64-NEXT: sub t6, s0, t6
; RV64-NEXT: slli t6, t0, 28
; RV64-NEXT: sub t6, t6, t1
; RV64-NEXT: and t6, t6, a6
; RV64-NEXT: csrwi vxrm, 0
; RV64-NEXT: mv s0, a2
llvm/test/CodeGen/Thumb2/mve-blockplacement.ll
+65 -63
View File
@@ -353,8 +353,8 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: push.w {r4, r5, r6, r7, r8, r9, r10, r11, lr}
; CHECK-NEXT: .pad #4
; CHECK-NEXT: sub sp, #4
; CHECK-NEXT: .vsave {d8, d9, d10, d11, d12, d13}
; CHECK-NEXT: vpush {d8, d9, d10, d11, d12, d13}
; CHECK-NEXT: .vsave {d8, d9, d10, d11, d12, d13, d14, d15}
; CHECK-NEXT: vpush {d8, d9, d10, d11, d12, d13, d14, d15}
; CHECK-NEXT: .pad #16
; CHECK-NEXT: sub sp, #16
; CHECK-NEXT: mov lr, r0
@@ -364,48 +364,50 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: @ %bb.1: @ %for.cond2.preheader.lr.ph
; CHECK-NEXT: movs r0, #1
; CHECK-NEXT: cmp r2, #1
; CHECK-NEXT: csel r3, r2, r0, lt
; CHECK-NEXT: csel r7, r2, r0, lt
; CHECK-NEXT: mov r12, r1
; CHECK-NEXT: mov r1, r3
; CHECK-NEXT: cmp r3, #3
; CHECK-NEXT: mov r1, r7
; CHECK-NEXT: cmp r7, #3
; CHECK-NEXT: it ls
; CHECK-NEXT: movls r1, #3
; CHECK-NEXT: mov r4, r2
; CHECK-NEXT: subs r1, r1, r3
; CHECK-NEXT: subs r1, r1, r7
; CHECK-NEXT: movw r2, #43691
; CHECK-NEXT: adds r1, #2
; CHECK-NEXT: movt r2, #43690
; CHECK-NEXT: ldr r6, [sp, #112]
; CHECK-NEXT: movw r9, :lower16:c
; CHECK-NEXT: ldr r6, [sp, #128]
; CHECK-NEXT: movw r8, :lower16:c
; CHECK-NEXT: umull r1, r2, r1, r2
; CHECK-NEXT: adr.w r8, .LCPI1_1
; CHECK-NEXT: movt r8, :upper16:c
; CHECK-NEXT: movs r1, #4
; CHECK-NEXT: @ implicit-def: $r10
; CHECK-NEXT: @ implicit-def: $r5
; CHECK-NEXT: @ implicit-def: $r11
; CHECK-NEXT: @ implicit-def: $r7
; CHECK-NEXT: movt r9, :upper16:c
; CHECK-NEXT: mov.w r10, #12
; CHECK-NEXT: str r3, [sp, #12] @ 4-byte Spill
; CHECK-NEXT: mov.w r9, #12
; CHECK-NEXT: str r4, [sp, #12] @ 4-byte Spill
; CHECK-NEXT: add.w r1, r1, r2, lsr #1
; CHECK-NEXT: add.w r0, r0, r2, lsr #1
; CHECK-NEXT: bic r2, r1, #3
; CHECK-NEXT: bic r3, r1, #3
; CHECK-NEXT: adr r1, .LCPI1_0
; CHECK-NEXT: vldrw.u32 q0, [r1]
; CHECK-NEXT: vdup.32 q5, r0
; CHECK-NEXT: adr r1, .LCPI1_1
; CHECK-NEXT: vldrw.u32 q5, [r1]
; CHECK-NEXT: vdup.32 q6, r0
; CHECK-NEXT: strd r2, r4, [sp, #4] @ 8-byte Folded Spill
; CHECK-NEXT: vadd.i32 q4, q0, r3
; CHECK-NEXT: vadd.i32 q4, q0, r7
; CHECK-NEXT: vdup.32 q7, r0
; CHECK-NEXT: strd r3, r7, [sp, #4] @ 8-byte Folded Spill
; CHECK-NEXT: b .LBB1_6
; CHECK-NEXT: .LBB1_2: @ %for.body6.preheader
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: mov r0, r7
; CHECK-NEXT: cmn.w r7, #4
; CHECK-NEXT: mov r0, r11
; CHECK-NEXT: cmn.w r11, #4
; CHECK-NEXT: it le
; CHECK-NEXT: mvnle r0, #3
; CHECK-NEXT: movw r2, #18725
; CHECK-NEXT: adds r0, #6
; CHECK-NEXT: movt r2, #9362
; CHECK-NEXT: subs r1, r0, r7
; CHECK-NEXT: sub.w r1, r0, r11
; CHECK-NEXT: mov r10, r3
; CHECK-NEXT: umull r2, r3, r1, r2
; CHECK-NEXT: subs r2, r1, r3
; CHECK-NEXT: add.w r2, r3, r2, lsr #1
@@ -413,18 +415,19 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: lsls r3, r3, #3
; CHECK-NEXT: sub.w r2, r3, r2, lsr #2
; CHECK-NEXT: subs r1, r2, r1
; CHECK-NEXT: mov r3, r10
; CHECK-NEXT: add r0, r1
; CHECK-NEXT: .LBB1_3: @ %for.cond.cleanup5.loopexit134.split.loop.exit139
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: adds r7, r0, #7
; CHECK-NEXT: add.w r11, r0, #7
; CHECK-NEXT: .LBB1_4: @ %for.cond.cleanup5
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: movs r5, #0
; CHECK-NEXT: mov.w r10, #0
; CHECK-NEXT: .LBB1_5: @ %for.cond.cleanup5
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: add.w r11, r11, #2
; CHECK-NEXT: subs.w r1, r11, lr
; CHECK-NEXT: asr.w r0, r11, #31
; CHECK-NEXT: adds r5, #2
; CHECK-NEXT: subs.w r1, r5, lr
; CHECK-NEXT: asr.w r0, r5, #31
; CHECK-NEXT: sbcs.w r0, r0, r12
; CHECK-NEXT: bge.w .LBB1_28
; CHECK-NEXT: .LBB1_6: @ %for.cond2.preheader
@@ -433,35 +436,36 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: @ Child Loop BB1_10 Depth 2
; CHECK-NEXT: @ Child Loop BB1_12 Depth 3
; CHECK-NEXT: @ Child Loop BB1_14 Depth 3
; CHECK-NEXT: cmp r7, #2
; CHECK-NEXT: cmp.w r11, #2
; CHECK-NEXT: bgt .LBB1_5
; CHECK-NEXT: @ %bb.7: @ %for.body6.lr.ph
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: ldr r0, [sp, #12] @ 4-byte Reload
; CHECK-NEXT: cmp r0, #5
; CHECK-NEXT: cmp r7, #5
; CHECK-NEXT: bhi .LBB1_17
; CHECK-NEXT: @ %bb.8: @ %for.body6.us.preheader
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: ldrd r2, r3, [sp, #104]
; CHECK-NEXT: ldrd r2, r3, [sp, #120]
; CHECK-NEXT: movs r0, #32
; CHECK-NEXT: movs r1, #0
; CHECK-NEXT: mov r6, r12
; CHECK-NEXT: mov r4, lr
; CHECK-NEXT: mov r4, r6
; CHECK-NEXT: mov r7, r12
; CHECK-NEXT: mov r6, lr
; CHECK-NEXT: bl __aeabi_ldivmod
; CHECK-NEXT: mov lr, r4
; CHECK-NEXT: mov r12, r6
; CHECK-NEXT: ldr r4, [sp, #8] @ 4-byte Reload
; CHECK-NEXT: vdup.32 q0, r2
; CHECK-NEXT: ldr r6, [sp, #112]
; CHECK-NEXT: mov r0, r7
; CHECK-NEXT: mov lr, r6
; CHECK-NEXT: mov r6, r4
; CHECK-NEXT: mov r12, r7
; CHECK-NEXT: ldr r3, [sp, #4] @ 4-byte Reload
; CHECK-NEXT: ldr r4, [sp, #12] @ 4-byte Reload
; CHECK-NEXT: vdup.32 q0, r2
; CHECK-NEXT: ldr r7, [sp, #8] @ 4-byte Reload
; CHECK-NEXT: mov r0, r11
; CHECK-NEXT: b .LBB1_10
; CHECK-NEXT: .LBB1_9: @ %for.cond.cleanup17.us
; CHECK-NEXT: @ in Loop: Header=BB1_10 Depth=2
; CHECK-NEXT: adds r7, r0, #7
; CHECK-NEXT: add.w r11, r0, #7
; CHECK-NEXT: cmn.w r0, #4
; CHECK-NEXT: mov.w r5, #0
; CHECK-NEXT: mov r0, r7
; CHECK-NEXT: mov.w r10, #0
; CHECK-NEXT: mov r0, r11
; CHECK-NEXT: bge .LBB1_5
; CHECK-NEXT: .LBB1_10: @ %for.body6.us
; CHECK-NEXT: @ Parent Loop BB1_6 Depth=1
@@ -484,14 +488,13 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: @ Parent Loop BB1_6 Depth=1
; CHECK-NEXT: @ Parent Loop BB1_10 Depth=2
; CHECK-NEXT: @ => This Inner Loop Header: Depth=3
; CHECK-NEXT: vldrw.u32 q2, [r8]
; CHECK-NEXT: subs r2, #4
; CHECK-NEXT: vqadd.u32 q2, q2, r1
; CHECK-NEXT: add.w r1, r1, #4
; CHECK-NEXT: vcmp.u32 hi, q6, q2
; CHECK-NEXT: vqadd.u32 q2, q5, r1
; CHECK-NEXT: adds r1, #4
; CHECK-NEXT: vcmp.u32 hi, q7, q2
; CHECK-NEXT: vshl.i32 q2, q1, #2
; CHECK-NEXT: vadd.i32 q2, q2, r9
; CHECK-NEXT: vadd.i32 q1, q1, r10
; CHECK-NEXT: subs r2, #4
; CHECK-NEXT: vadd.i32 q2, q2, r8
; CHECK-NEXT: vadd.i32 q1, q1, r9
; CHECK-NEXT: vpst
; CHECK-NEXT: vstrwt.32 q0, [q2]
; CHECK-NEXT: bne .LBB1_12
@@ -504,14 +507,13 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: @ Parent Loop BB1_6 Depth=1
; CHECK-NEXT: @ Parent Loop BB1_10 Depth=2
; CHECK-NEXT: @ => This Inner Loop Header: Depth=3
; CHECK-NEXT: vldrw.u32 q2, [r8]
; CHECK-NEXT: subs r2, #4
; CHECK-NEXT: vqadd.u32 q2, q2, r1
; CHECK-NEXT: add.w r1, r1, #4
; CHECK-NEXT: vcmp.u32 hi, q5, q2
; CHECK-NEXT: vqadd.u32 q2, q5, r1
; CHECK-NEXT: adds r1, #4
; CHECK-NEXT: vcmp.u32 hi, q6, q2
; CHECK-NEXT: vshl.i32 q2, q1, #2
; CHECK-NEXT: vadd.i32 q2, q2, r9
; CHECK-NEXT: vadd.i32 q1, q1, r10
; CHECK-NEXT: subs r2, #4
; CHECK-NEXT: vadd.i32 q2, q2, r8
; CHECK-NEXT: vadd.i32 q1, q1, r9
; CHECK-NEXT: vpst
; CHECK-NEXT: vstrwt.32 q0, [q2]
; CHECK-NEXT: bne .LBB1_14
@@ -521,7 +523,7 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: beq .LBB1_9
; CHECK-NEXT: @ %bb.16: @ %for.cond9.for.cond15.preheader_crit_edge.us
; CHECK-NEXT: @ in Loop: Header=BB1_10 Depth=2
; CHECK-NEXT: eor r1, r5, #1
; CHECK-NEXT: eor r1, r10, #1
; CHECK-NEXT: lsls r1, r1, #31
; CHECK-NEXT: bne .LBB1_9
; CHECK-NEXT: b .LBB1_26
@@ -530,11 +532,11 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: cmp r6, #0
; CHECK-NEXT: beq.w .LBB1_2
; CHECK-NEXT: @ %bb.18: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: mov r0, r7
; CHECK-NEXT: mov r0, r11
; CHECK-NEXT: .LBB1_19: @ %for.body6.us60
; CHECK-NEXT: @ Parent Loop BB1_6 Depth=1
; CHECK-NEXT: @ => This Inner Loop Header: Depth=2
; CHECK-NEXT: lsls r1, r5, #31
; CHECK-NEXT: lsls.w r1, r10, #31
; CHECK-NEXT: bne .LBB1_27
; CHECK-NEXT: @ %bb.20: @ %for.cond.cleanup17.us63
; CHECK-NEXT: @ in Loop: Header=BB1_19 Depth=2
@@ -550,19 +552,19 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: bgt .LBB1_25
; CHECK-NEXT: @ %bb.23: @ %for.cond.cleanup17.us63.3
; CHECK-NEXT: @ in Loop: Header=BB1_19 Depth=2
; CHECK-NEXT: add.w r7, r0, #28
; CHECK-NEXT: add.w r11, r0, #28
; CHECK-NEXT: cmn.w r0, #25
; CHECK-NEXT: mov.w r5, #0
; CHECK-NEXT: mov r0, r7
; CHECK-NEXT: mov.w r10, #0
; CHECK-NEXT: mov r0, r11
; CHECK-NEXT: blt .LBB1_19
; CHECK-NEXT: b .LBB1_5
; CHECK-NEXT: .LBB1_24: @ %for.cond.cleanup5.loopexit134.split.loop.exit137
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: add.w r7, r0, #14
; CHECK-NEXT: add.w r11, r0, #14
; CHECK-NEXT: b .LBB1_4
; CHECK-NEXT: .LBB1_25: @ %for.cond.cleanup5.loopexit134.split.loop.exit135
; CHECK-NEXT: @ in Loop: Header=BB1_6 Depth=1
; CHECK-NEXT: add.w r7, r0, #21
; CHECK-NEXT: add.w r11, r0, #21
; CHECK-NEXT: b .LBB1_4
; CHECK-NEXT: .LBB1_26: @ %for.inc19.us
; CHECK-NEXT: @ =>This Inner Loop Header: Depth=1
@@ -572,7 +574,7 @@ define i32 @d(i64 %e, i32 %f, i64 %g, i32 %h) {
; CHECK-NEXT: b .LBB1_27
; CHECK-NEXT: .LBB1_28: @ %for.cond.cleanup
; CHECK-NEXT: add sp, #16
; CHECK-NEXT: vpop {d8, d9, d10, d11, d12, d13}
; CHECK-NEXT: vpop {d8, d9, d10, d11, d12, d13, d14, d15}
; CHECK-NEXT: add sp, #4
; CHECK-NEXT: pop.w {r4, r5, r6, r7, r8, r9, r10, r11, pc}
; CHECK-NEXT: .p2align 4
llvm/test/CodeGen/Thumb2/mve-gather-increment.ll
+445 -423
View File
File diff suppressed because it is too large. Load diff
llvm/test/CodeGen/Thumb2/mve-gather-scatter-optimisation.ll
+70 -70
View File
@@ -602,57 +602,60 @@ define dso_local void @arm_mat_mult_q15(ptr noalias nocapture readonly %A, ptr n
; CHECK-NEXT: push.w {r4, r5, r6, r7, r8, r9, r10, r11, lr}
; CHECK-NEXT: .pad #4
; CHECK-NEXT: sub sp, #4
; CHECK-NEXT: .vsave {d8, d9, d10, d11, d12, d13}
; CHECK-NEXT: vpush {d8, d9, d10, d11, d12, d13}
; CHECK-NEXT: .pad #24
; CHECK-NEXT: sub sp, #24
; CHECK-NEXT: str r0, [sp, #20] @ 4-byte Spill
; CHECK-NEXT: .vsave {d8, d9, d10, d11, d12, d13, d14, d15}
; CHECK-NEXT: vpush {d8, d9, d10, d11, d12, d13, d14, d15}
; CHECK-NEXT: .pad #32
; CHECK-NEXT: sub sp, #32
; CHECK-NEXT: strd r0, r2, [sp, #24] @ 8-byte Folded Spill
; CHECK-NEXT: cmp r3, #0
; CHECK-NEXT: str r3, [sp, #4] @ 4-byte Spill
; CHECK-NEXT: str r3, [sp, #8] @ 4-byte Spill
; CHECK-NEXT: itt ne
; CHECK-NEXT: ldrne r0, [sp, #112]
; CHECK-NEXT: ldrne r0, [sp, #136]
; CHECK-NEXT: cmpne r0, #0
; CHECK-NEXT: bne .LBB10_2
; CHECK-NEXT: .LBB10_1: @ %for.cond.cleanup
; CHECK-NEXT: add sp, #24
; CHECK-NEXT: vpop {d8, d9, d10, d11, d12, d13}
; CHECK-NEXT: add sp, #32
; CHECK-NEXT: vpop {d8, d9, d10, d11, d12, d13, d14, d15}
; CHECK-NEXT: add sp, #4
; CHECK-NEXT: pop.w {r4, r5, r6, r7, r8, r9, r10, r11, pc}
; CHECK-NEXT: .LBB10_2: @ %for.cond1.preheader.us.preheader
; CHECK-NEXT: ldr.w r12, [sp, #116]
; CHECK-NEXT: ldr.w r12, [sp, #140]
; CHECK-NEXT: movs r7, #1
; CHECK-NEXT: mov.w r11, #0
; CHECK-NEXT: vmov.i32 q0, #0x0
; CHECK-NEXT: movs r5, #0
; CHECK-NEXT: bic r0, r12, #3
; CHECK-NEXT: subs r3, r0, #4
; CHECK-NEXT: add.w r3, r7, r3, lsr #2
; CHECK-NEXT: str r3, [sp, #12] @ 4-byte Spill
; CHECK-NEXT: ldr r3, [sp, #112]
; CHECK-NEXT: lsl.w r7, r12, #1
; CHECK-NEXT: str r7, [sp] @ 4-byte Spill
; CHECK-NEXT: movs r7, #0
; CHECK-NEXT: vdup.32 q1, r3
; CHECK-NEXT: lsls r6, r3, #1
; CHECK-NEXT: vshl.i32 q2, q1, #2
; CHECK-NEXT: ldr r3, [sp, #20] @ 4-byte Reload
; CHECK-NEXT: str r3, [sp, #16] @ 4-byte Spill
; CHECK-NEXT: bic r2, r12, #3
; CHECK-NEXT: subs r3, r2, #4
; CHECK-NEXT: add.w r0, r7, r3, lsr #2
; CHECK-NEXT: ldr r7, [sp, #136]
; CHECK-NEXT: adr r3, .LCPI10_0
; CHECK-NEXT: str r0, [sp, #16] @ 4-byte Spill
; CHECK-NEXT: lsl.w r0, r12, #1
; CHECK-NEXT: vdup.32 q1, r7
; CHECK-NEXT: vldrw.u32 q2, [r3]
; CHECK-NEXT: str r0, [sp, #4] @ 4-byte Spill
; CHECK-NEXT: ldr r0, [sp, #24] @ 4-byte Reload
; CHECK-NEXT: lsls r6, r7, #1
; CHECK-NEXT: vshl.i32 q3, q1, #2
; CHECK-NEXT: movs r3, #0
; CHECK-NEXT: str r0, [sp, #20] @ 4-byte Spill
; CHECK-NEXT: b .LBB10_5
; CHECK-NEXT: .LBB10_3: @ %for.cond5.preheader.us73.preheader
; CHECK-NEXT: @ in Loop: Header=BB10_5 Depth=1
; CHECK-NEXT: add.w r3, r2, r8, lsl #1
; CHECK-NEXT: ldr r0, [sp, #28] @ 4-byte Reload
; CHECK-NEXT: add.w r3, r0, r5, lsl #1
; CHECK-NEXT: wlstp.8 lr, r6, .LBB10_4
; CHECK-NEXT: b .LBB10_15
; CHECK-NEXT: .LBB10_4: @ %for.cond1.for.cond.cleanup3_crit_edge.us
; CHECK-NEXT: @ in Loop: Header=BB10_5 Depth=1
; CHECK-NEXT: ldr r3, [sp] @ 4-byte Reload
; CHECK-NEXT: add r7, r12
; CHECK-NEXT: ldr r5, [sp, #16] @ 4-byte Reload
; CHECK-NEXT: add r5, r3
; CHECK-NEXT: str r5, [sp, #16] @ 4-byte Spill
; CHECK-NEXT: ldr r5, [sp, #8] @ 4-byte Reload
; CHECK-NEXT: ldr r3, [sp, #4] @ 4-byte Reload
; CHECK-NEXT: adds r5, #1
; CHECK-NEXT: cmp r5, r3
; CHECK-NEXT: ldr r0, [sp, #4] @ 4-byte Reload
; CHECK-NEXT: add r11, r12
; CHECK-NEXT: ldr r3, [sp, #20] @ 4-byte Reload
; CHECK-NEXT: add r3, r0
; CHECK-NEXT: str r3, [sp, #20] @ 4-byte Spill
; CHECK-NEXT: ldr r3, [sp, #12] @ 4-byte Reload
; CHECK-NEXT: ldr r0, [sp, #8] @ 4-byte Reload
; CHECK-NEXT: adds r3, #1
; CHECK-NEXT: cmp r3, r0
; CHECK-NEXT: beq .LBB10_1
; CHECK-NEXT: .LBB10_5: @ %for.cond1.preheader.us
; CHECK-NEXT: @ =>This Loop Header: Depth=1
@@ -660,22 +663,21 @@ define dso_local void @arm_mat_mult_q15(ptr noalias nocapture readonly %A, ptr n
; CHECK-NEXT: @ Child Loop BB10_11 Depth 3
; CHECK-NEXT: @ Child Loop BB10_14 Depth 3
; CHECK-NEXT: @ Child Loop BB10_15 Depth 2
; CHECK-NEXT: ldr r3, [sp, #112]
; CHECK-NEXT: mul r5, r3, r7
; CHECK-NEXT: cmp.w r12, #0
; CHECK-NEXT: str r5, [sp, #8] @ 4-byte Spill
; CHECK-NEXT: mul r8, r5, r3
; CHECK-NEXT: str r3, [sp, #12] @ 4-byte Spill
; CHECK-NEXT: beq .LBB10_3
; CHECK-NEXT: @ %bb.6: @ %for.cond5.preheader.us.us.preheader
; CHECK-NEXT: @ in Loop: Header=BB10_5 Depth=1
; CHECK-NEXT: mov.w r9, #0
; CHECK-NEXT: mov.w r8, #0
; CHECK-NEXT: b .LBB10_8
; CHECK-NEXT: .LBB10_7: @ %for.cond5.for.cond.cleanup7_crit_edge.us.us
; CHECK-NEXT: @ in Loop: Header=BB10_8 Depth=2
; CHECK-NEXT: add.w r3, r9, r8
; CHECK-NEXT: add.w r9, r9, #1
; CHECK-NEXT: strh.w r10, [r2, r3, lsl #1]
; CHECK-NEXT: ldr r3, [sp, #112]
; CHECK-NEXT: cmp r9, r3
; CHECK-NEXT: ldr r3, [sp, #28] @ 4-byte Reload
; CHECK-NEXT: add.w r0, r8, r5
; CHECK-NEXT: add.w r8, r8, #1
; CHECK-NEXT: cmp r8, r7
; CHECK-NEXT: strh.w r10, [r3, r0, lsl #1]
; CHECK-NEXT: beq .LBB10_4
; CHECK-NEXT: .LBB10_8: @ %for.cond5.preheader.us.us
; CHECK-NEXT: @ Parent Loop BB10_5 Depth=1
@@ -690,48 +692,46 @@ define dso_local void @arm_mat_mult_q15(ptr noalias nocapture readonly %A, ptr n
; CHECK-NEXT: b .LBB10_13
; CHECK-NEXT: .LBB10_10: @ %vector.ph
; CHECK-NEXT: @ in Loop: Header=BB10_8 Depth=2
; CHECK-NEXT: adr r3, .LCPI10_0
; CHECK-NEXT: vmov q4, q1
; CHECK-NEXT: vldrw.u32 q5, [r3]
; CHECK-NEXT: ldr r3, [sp, #12] @ 4-byte Reload
; CHECK-NEXT: vmov.i32 q3, #0x0
; CHECK-NEXT: dls lr, r3
; CHECK-NEXT: vmlas.i32 q4, q5, r9
; CHECK-NEXT: ldr r3, [sp, #16] @ 4-byte Reload
; CHECK-NEXT: ldr r0, [sp, #16] @ 4-byte Reload
; CHECK-NEXT: vmov q5, q1
; CHECK-NEXT: vmov.i32 q4, #0x0
; CHECK-NEXT: vmlas.i32 q5, q2, r8
; CHECK-NEXT: dls lr, r0
; CHECK-NEXT: ldr r3, [sp, #20] @ 4-byte Reload
; CHECK-NEXT: .LBB10_11: @ %vector.body
; CHECK-NEXT: @ Parent Loop BB10_5 Depth=1
; CHECK-NEXT: @ Parent Loop BB10_8 Depth=2
; CHECK-NEXT: @ => This Inner Loop Header: Depth=3
; CHECK-NEXT: vadd.i32 q5, q4, q2
; CHECK-NEXT: vldrh.s32 q6, [r1, q4, uxtw #1]
; CHECK-NEXT: vldrh.s32 q4, [r3], #8
; CHECK-NEXT: vmul.i32 q4, q6, q4
; CHECK-NEXT: vadd.i32 q3, q4, q3
; CHECK-NEXT: vmov q4, q5
; CHECK-NEXT: vadd.i32 q6, q5, q3
; CHECK-NEXT: vldrh.s32 q7, [r1, q5, uxtw #1]
; CHECK-NEXT: vldrh.s32 q5, [r3], #8
; CHECK-NEXT: vmul.i32 q5, q7, q5
; CHECK-NEXT: vadd.i32 q4, q5, q4
; CHECK-NEXT: vmov q5, q6
; CHECK-NEXT: le lr, .LBB10_11
; CHECK-NEXT: @ %bb.12: @ %middle.block
; CHECK-NEXT: @ in Loop: Header=BB10_8 Depth=2
; CHECK-NEXT: vaddv.u32 r10, q3
; CHECK-NEXT: cmp r0, r12
; CHECK-NEXT: mov r4, r0
; CHECK-NEXT: vaddv.u32 r10, q4
; CHECK-NEXT: cmp r2, r12
; CHECK-NEXT: mov r4, r2
; CHECK-NEXT: beq .LBB10_7
; CHECK-NEXT: .LBB10_13: @ %for.body8.us.us.preheader
; CHECK-NEXT: @ in Loop: Header=BB10_8 Depth=2
; CHECK-NEXT: ldr r3, [sp, #112]
; CHECK-NEXT: mla r3, r7, r4, r8
; CHECK-NEXT: add.w r0, r11, r4
; CHECK-NEXT: ldr r7, [sp, #24] @ 4-byte Reload
; CHECK-NEXT: sub.w lr, r12, r4
; CHECK-NEXT: ldr r5, [sp, #20] @ 4-byte Reload
; CHECK-NEXT: mla r3, r3, r4, r9
; CHECK-NEXT: add.w r11, r1, r3, lsl #1
; CHECK-NEXT: adds r3, r7, r4
; CHECK-NEXT: add.w r3, r5, r3, lsl #1
; CHECK-NEXT: add.w r9, r7, r0, lsl #1
; CHECK-NEXT: ldr r7, [sp, #136]
; CHECK-NEXT: add.w r3, r1, r3, lsl #1
; CHECK-NEXT: .LBB10_14: @ %for.body8.us.us
; CHECK-NEXT: @ Parent Loop BB10_5 Depth=1
; CHECK-NEXT: @ Parent Loop BB10_8 Depth=2
; CHECK-NEXT: @ => This Inner Loop Header: Depth=3
; CHECK-NEXT: ldrsh.w r5, [r11]
; CHECK-NEXT: add r11, r6
; CHECK-NEXT: ldrsh r4, [r3], #2
; CHECK-NEXT: smlabb r10, r5, r4, r10
; CHECK-NEXT: ldrsh.w r4, [r3]
; CHECK-NEXT: add r3, r6
; CHECK-NEXT: ldrsh r0, [r9], #2
; CHECK-NEXT: smlabb r10, r4, r0, r10
; CHECK-NEXT: le lr, .LBB10_14
; CHECK-NEXT: b .LBB10_7
; CHECK-NEXT: .LBB10_15: @ Parent Loop BB10_5 Depth=1
llvm/test/Transforms/InferAddressSpaces/AMDGPU/flat_atomic.ll
+15 -13
View File
@@ -121,29 +121,31 @@ define protected amdgpu_kernel void @InferPHI(i32 %a, ptr addrspace(1) %b, doubl
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: s_ashr_i32 s7, s6, 31
; CHECK-NEXT: s_lshl_b64 s[4:5], s[6:7], 3
; CHECK-NEXT: s_add_u32 s4, s0, s4
; CHECK-NEXT: s_addc_u32 s5, s1, s5
; CHECK-NEXT: s_add_u32 s0, s4, -8
; CHECK-NEXT: s_addc_u32 s1, s5, -1
; CHECK-NEXT: s_cmp_eq_u64 s[4:5], 9
; CHECK-NEXT: s_cselect_b64 s[4:5], -1, 0
; CHECK-NEXT: s_add_u32 s0, s0, s4
; CHECK-NEXT: s_addc_u32 s1, s1, s5
; CHECK-NEXT: s_add_u32 s4, s0, -8
; CHECK-NEXT: s_addc_u32 s5, s1, -1
; CHECK-NEXT: s_cmp_eq_u64 s[0:1], 9
; CHECK-NEXT: s_cselect_b64 s[0:1], -1, 0
; CHECK-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
; CHECK-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
; CHECK-NEXT: .LBB3_1: ; %bb0
; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
; CHECK-NEXT: s_andn2_b64 vcc, exec, s[4:5]
; CHECK-NEXT: s_and_b64 vcc, exec, s[0:1]
; CHECK-NEXT: s_cbranch_vccnz .LBB3_1
; CHECK-NEXT: ; %bb.2: ; %bb1
; CHECK-NEXT: s_mov_b64 s[4:5], exec
; CHECK-NEXT: v_mbcnt_lo_u32_b32 v0, s4, 0
; CHECK-NEXT: v_mbcnt_hi_u32_b32 v0, s5, v0
; CHECK-NEXT: s_mov_b64 s[0:1], exec
; CHECK-NEXT: v_mbcnt_lo_u32_b32 v0, s0, 0
; CHECK-NEXT: v_mbcnt_hi_u32_b32 v0, s1, v0
; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v0
; CHECK-NEXT: s_and_saveexec_b64 s[6:7], vcc
; CHECK-NEXT: s_cbranch_execz .LBB3_4
; CHECK-NEXT: ; %bb.3:
; CHECK-NEXT: s_bcnt1_i32_b64 s4, s[4:5]
; CHECK-NEXT: v_cvt_f64_u32_e32 v[0:1], s4
; CHECK-NEXT: s_bcnt1_i32_b64 s0, s[0:1]
; CHECK-NEXT: v_cvt_f64_u32_e32 v[0:1], s0
; CHECK-NEXT: v_mul_f64 v[0:1], s[2:3], v[0:1]
; CHECK-NEXT: v_mov_b32_e32 v2, 0
; CHECK-NEXT: global_atomic_add_f64 v2, v[0:1], s[0:1]
; CHECK-NEXT: global_atomic_add_f64 v2, v[0:1], s[4:5]
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: buffer_wbinvl1_vol
; CHECK-NEXT: .LBB3_4: