Skip to content

Commit 4d72ab2

Browse files
bcahoonbcahoon
committed
Add end-to-end test
1 parent 32a1a9a commit 4d72ab2

File tree

1 file changed

+124
-0
lines changed

1 file changed

+124
-0
lines changed

‎llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll‎

Lines changed: 124 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,124 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc -O3 -mtriple=amdgcn--amdhsa -mcpu=gfx942 %s -o - | FileCheck %s --check-prefixes=GFX942
3+
; RUN: llc -O3 -mtriple=amdgcn--amdhsa -mcpu=gfx908 %s -o - | FileCheck %s --check-prefixes=GFX908
4+
5+
define amdgpu_kernel void @test(i1 %arg, i1 %arg1) {
6+
; GFX942-LABEL: test:
7+
; GFX942: ; %bb.0: ; %bb
8+
; GFX942-NEXT: s_load_dword s0, s[4:5], 0x0
9+
; GFX942-NEXT: v_mov_b32_e32 v1, 0
10+
; GFX942-NEXT: s_mov_b32 s2, 0
11+
; GFX942-NEXT: v_accvgpr_write_b32 a0, v1
12+
; GFX942-NEXT: s_mov_b32 s3, 0
13+
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
14+
; GFX942-NEXT: s_bitcmp1_b32 s0, 0
15+
; GFX942-NEXT: s_cselect_b64 s[0:1], -1, 0
16+
; GFX942-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
17+
; GFX942-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
18+
; GFX942-NEXT: s_branch .LBB0_2
19+
; GFX942-NEXT: .LBB0_1: ; %Flow
20+
; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1
21+
; GFX942-NEXT: s_andn2_b64 vcc, exec, s[4:5]
22+
; GFX942-NEXT: s_cbranch_vccz .LBB0_4
23+
; GFX942-NEXT: .LBB0_2: ; %bb2
24+
; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1
25+
; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1]
26+
; GFX942-NEXT: s_mov_b64 s[4:5], -1
27+
; GFX942-NEXT: s_cbranch_vccnz .LBB0_1
28+
; GFX942-NEXT: ; %bb.3: ; %bb4
29+
; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1
30+
; GFX942-NEXT: s_or_b32 s4, s3, 1
31+
; GFX942-NEXT: s_ashr_i32 s5, s3, 31
32+
; GFX942-NEXT: s_mov_b32 s3, s2
33+
; GFX942-NEXT: v_mov_b64_e32 v[4:5], s[2:3]
34+
; GFX942-NEXT: v_accvgpr_read_b32 v0, a0
35+
; GFX942-NEXT: v_mov_b32_e32 v2, v1
36+
; GFX942-NEXT: v_mov_b32_e32 v3, v1
37+
; GFX942-NEXT: v_accvgpr_write_b32 a0, v0
38+
; GFX942-NEXT: v_accvgpr_write_b32 a1, v1
39+
; GFX942-NEXT: v_accvgpr_write_b32 a2, v2
40+
; GFX942-NEXT: v_accvgpr_write_b32 a3, v3
41+
; GFX942-NEXT: s_and_b32 s3, s5, s4
42+
; GFX942-NEXT: s_mov_b64 s[4:5], 0
43+
; GFX942-NEXT: v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[4:5], a[0:3]
44+
; GFX942-NEXT: s_branch .LBB0_1
45+
; GFX942-NEXT: .LBB0_4: ; %common.ret
46+
; GFX942-NEXT: s_endpgm
47+
;
48+
; GFX908-LABEL: test:
49+
; GFX908: ; %bb.0: ; %bb
50+
; GFX908-NEXT: s_load_dword s0, s[8:9], 0x0
51+
; GFX908-NEXT: v_mov_b32_e32 v1, 0
52+
; GFX908-NEXT: s_mov_b32 s2, 0
53+
; GFX908-NEXT: s_mov_b32 s3, 0
54+
; GFX908-NEXT: v_accvgpr_write_b32 a0, v1
55+
; GFX908-NEXT: s_waitcnt lgkmcnt(0)
56+
; GFX908-NEXT: s_bitcmp1_b32 s0, 0
57+
; GFX908-NEXT: s_cselect_b64 s[0:1], -1, 0
58+
; GFX908-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
59+
; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
60+
; GFX908-NEXT: s_branch .LBB0_2
61+
; GFX908-NEXT: .LBB0_1: ; %Flow
62+
; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1
63+
; GFX908-NEXT: s_andn2_b64 vcc, exec, s[4:5]
64+
; GFX908-NEXT: s_cbranch_vccz .LBB0_4
65+
; GFX908-NEXT: .LBB0_2: ; %bb2
66+
; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1
67+
; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1]
68+
; GFX908-NEXT: s_mov_b64 s[4:5], -1
69+
; GFX908-NEXT: s_cbranch_vccnz .LBB0_1
70+
; GFX908-NEXT: ; %bb.3: ; %bb4
71+
; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1
72+
; GFX908-NEXT: s_or_b32 s4, s3, 1
73+
; GFX908-NEXT: s_ashr_i32 s5, s3, 31
74+
; GFX908-NEXT: s_mov_b32 s3, s2
75+
; GFX908-NEXT: v_accvgpr_read_b32 v0, a0
76+
; GFX908-NEXT: v_mov_b32_e32 v5, s3
77+
; GFX908-NEXT: v_mov_b32_e32 v4, s2
78+
; GFX908-NEXT: v_mov_b32_e32 v2, v1
79+
; GFX908-NEXT: v_mov_b32_e32 v3, v1
80+
; GFX908-NEXT: v_accvgpr_write_b32 a0, v0
81+
; GFX908-NEXT: v_accvgpr_write_b32 a1, v1
82+
; GFX908-NEXT: v_accvgpr_write_b32 a2, v2
83+
; GFX908-NEXT: v_accvgpr_write_b32 a3, v3
84+
; GFX908-NEXT: s_and_b32 s3, s5, s4
85+
; GFX908-NEXT: v_mfma_f32_16x16x16f16 a[0:3], v[4:5], v[4:5], a[0:3]
86+
; GFX908-NEXT: s_mov_b64 s[4:5], 0
87+
; GFX908-NEXT: s_branch .LBB0_1
88+
; GFX908-NEXT: .LBB0_4: ; %common.ret
89+
; GFX908-NEXT: s_endpgm
90+
bb:
91+
br label %bb2
92+
93+
bb2:
94+
%i = phi { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } [ %i12, %bb4 ], [ zeroinitializer, %bb ]
95+
%i3 = phi i32 [ %i7, %bb4 ], [ 0, %bb ]
96+
br i1 %arg, label %bb4, label %bb13
97+
98+
bb4:
99+
%i5 = or i32 %i3, 1
100+
%i6 = icmp slt i32 %i3, 0
101+
%i7 = select i1 %i6, i32 %i5, i32 0
102+
%i8 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 123
103+
%i9 = insertelement <4 x float> zeroinitializer, float %i8, i32 0
104+
%i10 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i9, i32 0, i32 0, i32 0)
105+
%i11 = extractelement <4 x float> %i10, i32 0
106+
%i12 = insertvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } zeroinitializer, float %i11, 123
107+
br label %bb2
108+
109+
bb13:
110+
br i1 %arg1, label %bb14, label %common.ret
111+
112+
common.ret:
113+
ret void
114+
115+
bb14:
116+
%i15 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 0
117+
%i16 = insertelement <4 x float> zeroinitializer, float %i15, i32 0
118+
br label %common.ret
119+
}
120+
121+
; Function Attrs: convergent nocallback nofree nosync nounwind willreturn memory(none)
122+
declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg) #0
123+
124+
attributes #0 = { convergent nocallback nofree nosync nounwind willreturn memory(none) }

0 commit comments

Comments
 (0)