| blob | bf5843ea8047d5fe6c420df77fae370f748a4a46 |
1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
2 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 < %s | FileCheck -check-prefix=GFX908 %s
3 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefix=GFX90A %s
5 ; This testcase would fail on GFX908 due to not having a free VGPR available to
6 ; copy between AGPRs.
7 define void @no_free_vgprs_at_agpr_to_agpr_copy(float %v0, float %v1) #0 {
8 ; GFX908-LABEL: no_free_vgprs_at_agpr_to_agpr_copy:
9 ; GFX908: ; %bb.0:
10 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
11 ; GFX908-NEXT: v_mov_b32_e32 v32, v1
12 ; GFX908-NEXT: v_mov_b32_e32 v33, v0
13 ; GFX908-NEXT: ;;#ASMSTART
14 ; GFX908-NEXT: ; def v[0:31] a[0:15]
15 ; GFX908-NEXT: ;;#ASMEND
16 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a15
17 ; GFX908-NEXT: s_nop 1
18 ; GFX908-NEXT: v_accvgpr_write_b32 a31, v39
19 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a14
20 ; GFX908-NEXT: s_nop 1
21 ; GFX908-NEXT: v_accvgpr_write_b32 a30, v39
22 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a13
23 ; GFX908-NEXT: s_nop 1
24 ; GFX908-NEXT: v_accvgpr_write_b32 a29, v39
25 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a12
26 ; GFX908-NEXT: s_nop 1
27 ; GFX908-NEXT: v_accvgpr_write_b32 a28, v39
28 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a11
29 ; GFX908-NEXT: s_nop 1
30 ; GFX908-NEXT: v_accvgpr_write_b32 a27, v39
31 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a10
32 ; GFX908-NEXT: s_nop 1
33 ; GFX908-NEXT: v_accvgpr_write_b32 a26, v39
34 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a9
35 ; GFX908-NEXT: s_nop 1
36 ; GFX908-NEXT: v_accvgpr_write_b32 a25, v39
37 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a8
38 ; GFX908-NEXT: s_nop 1
39 ; GFX908-NEXT: v_accvgpr_write_b32 a24, v39
40 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a7
41 ; GFX908-NEXT: s_nop 1
42 ; GFX908-NEXT: v_accvgpr_write_b32 a23, v39
43 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a6
44 ; GFX908-NEXT: s_nop 1
45 ; GFX908-NEXT: v_accvgpr_write_b32 a22, v39
46 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a5
47 ; GFX908-NEXT: s_nop 1
48 ; GFX908-NEXT: v_accvgpr_write_b32 a21, v39
49 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a4
50 ; GFX908-NEXT: s_nop 1
51 ; GFX908-NEXT: v_accvgpr_write_b32 a20, v39
52 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a3
53 ; GFX908-NEXT: s_nop 1
54 ; GFX908-NEXT: v_accvgpr_write_b32 a19, v39
55 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a2
56 ; GFX908-NEXT: s_nop 1
57 ; GFX908-NEXT: v_accvgpr_write_b32 a18, v39
58 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a1
59 ; GFX908-NEXT: s_nop 1
60 ; GFX908-NEXT: v_accvgpr_write_b32 a17, v39
61 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a0
62 ; GFX908-NEXT: s_nop 1
63 ; GFX908-NEXT: v_accvgpr_write_b32 a16, v39
64 ; GFX908-NEXT: s_nop 0
65 ; GFX908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]
66 ; GFX908-NEXT: s_nop 7
67 ; GFX908-NEXT: s_nop 1
68 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a0 ; Reload Reuse
69 ; GFX908-NEXT: v_accvgpr_read_b32 v38, a11 ; Reload Reuse
70 ; GFX908-NEXT: v_accvgpr_read_b32 v37, a12 ; Reload Reuse
71 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 ; 4-byte Folded Spill
72 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a1 ; Reload Reuse
73 ; GFX908-NEXT: v_accvgpr_read_b32 v36, a13 ; Reload Reuse
74 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a14 ; Reload Reuse
75 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill
76 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a2 ; Reload Reuse
77 ; GFX908-NEXT: v_accvgpr_read_b32 v34, a15 ; Reload Reuse
78 ; GFX908-NEXT: s_nop 0
79 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill
80 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a3 ; Reload Reuse
81 ; GFX908-NEXT: s_nop 1
82 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill
83 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a4 ; Reload Reuse
84 ; GFX908-NEXT: s_nop 1
85 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill
86 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a5 ; Reload Reuse
87 ; GFX908-NEXT: s_nop 1
88 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill
89 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a6 ; Reload Reuse
90 ; GFX908-NEXT: s_nop 1
91 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill
92 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a7 ; Reload Reuse
93 ; GFX908-NEXT: s_nop 1
94 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill
95 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a8 ; Reload Reuse
96 ; GFX908-NEXT: s_nop 1
97 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill
98 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a9 ; Reload Reuse
99 ; GFX908-NEXT: s_nop 1
100 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill
101 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a10 ; Reload Reuse
102 ; GFX908-NEXT: s_nop 1
103 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill
104 ; GFX908-NEXT: ;;#ASMSTART
105 ; GFX908-NEXT: ; copy
106 ; GFX908-NEXT: ;;#ASMEND
107 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a1
108 ; GFX908-NEXT: s_nop 1
109 ; GFX908-NEXT: v_accvgpr_write_b32 a16, v39
110 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 ; 4-byte Folded Reload
111 ; GFX908-NEXT: s_waitcnt vmcnt(0)
112 ; GFX908-NEXT: v_accvgpr_write_b32 a0, v39 ; Reload Reuse
113 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload
114 ; GFX908-NEXT: v_accvgpr_write_b32 a11, v38 ; Reload Reuse
115 ; GFX908-NEXT: v_accvgpr_write_b32 a12, v37 ; Reload Reuse
116 ; GFX908-NEXT: v_accvgpr_write_b32 a13, v36 ; Reload Reuse
117 ; GFX908-NEXT: v_accvgpr_write_b32 a14, v35 ; Reload Reuse
118 ; GFX908-NEXT: s_waitcnt vmcnt(0)
119 ; GFX908-NEXT: v_accvgpr_write_b32 a1, v39 ; Reload Reuse
120 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload
121 ; GFX908-NEXT: s_waitcnt vmcnt(0)
122 ; GFX908-NEXT: v_accvgpr_write_b32 a2, v39 ; Reload Reuse
123 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload
124 ; GFX908-NEXT: s_waitcnt vmcnt(0)
125 ; GFX908-NEXT: v_accvgpr_write_b32 a3, v39 ; Reload Reuse
126 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload
127 ; GFX908-NEXT: s_waitcnt vmcnt(0)
128 ; GFX908-NEXT: v_accvgpr_write_b32 a4, v39 ; Reload Reuse
129 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload
130 ; GFX908-NEXT: s_waitcnt vmcnt(0)
131 ; GFX908-NEXT: v_accvgpr_write_b32 a5, v39 ; Reload Reuse
132 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload
133 ; GFX908-NEXT: s_waitcnt vmcnt(0)
134 ; GFX908-NEXT: v_accvgpr_write_b32 a6, v39 ; Reload Reuse
135 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload
136 ; GFX908-NEXT: s_waitcnt vmcnt(0)
137 ; GFX908-NEXT: v_accvgpr_write_b32 a7, v39 ; Reload Reuse
138 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload
139 ; GFX908-NEXT: s_waitcnt vmcnt(0)
140 ; GFX908-NEXT: v_accvgpr_write_b32 a8, v39 ; Reload Reuse
141 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload
142 ; GFX908-NEXT: s_waitcnt vmcnt(0)
143 ; GFX908-NEXT: v_accvgpr_write_b32 a9, v39 ; Reload Reuse
144 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload
145 ; GFX908-NEXT: s_waitcnt vmcnt(0)
146 ; GFX908-NEXT: v_accvgpr_write_b32 a10, v39 ; Reload Reuse
147 ; GFX908-NEXT: v_accvgpr_write_b32 a15, v34 ; Reload Reuse
148 ; GFX908-NEXT: ;;#ASMSTART
149 ; GFX908-NEXT: ; copy
150 ; GFX908-NEXT: ;;#ASMEND
151 ; GFX908-NEXT: v_accvgpr_read_b32 v32, a2
152 ; GFX908-NEXT: s_nop 1
153 ; GFX908-NEXT: v_accvgpr_write_b32 a3, v32
154 ; GFX908-NEXT: ;;#ASMSTART
155 ; GFX908-NEXT: ; use a3 v[0:31]
156 ; GFX908-NEXT: ;;#ASMEND
157 ; GFX908-NEXT: s_setpc_b64 s[30:31]
158 ;
159 ; GFX90A-LABEL: no_free_vgprs_at_agpr_to_agpr_copy:
160 ; GFX90A: ; %bb.0:
161 ; GFX90A-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
162 ; GFX90A-NEXT: v_mov_b32_e32 v33, v0
163 ; GFX90A-NEXT: v_mov_b32_e32 v32, v1
164 ; GFX90A-NEXT: ;;#ASMSTART
165 ; GFX90A-NEXT: ; def v[0:31] a[0:15]
166 ; GFX90A-NEXT: ;;#ASMEND
167 ; GFX90A-NEXT: v_accvgpr_mov_b32 a31, a15
168 ; GFX90A-NEXT: v_accvgpr_mov_b32 a30, a14
169 ; GFX90A-NEXT: v_accvgpr_mov_b32 a29, a13
170 ; GFX90A-NEXT: v_accvgpr_mov_b32 a28, a12
171 ; GFX90A-NEXT: v_accvgpr_mov_b32 a27, a11
172 ; GFX90A-NEXT: v_accvgpr_mov_b32 a26, a10
173 ; GFX90A-NEXT: v_accvgpr_mov_b32 a25, a9
174 ; GFX90A-NEXT: v_accvgpr_mov_b32 a24, a8
175 ; GFX90A-NEXT: v_accvgpr_mov_b32 a23, a7
176 ; GFX90A-NEXT: v_accvgpr_mov_b32 a22, a6
177 ; GFX90A-NEXT: v_accvgpr_mov_b32 a21, a5
178 ; GFX90A-NEXT: v_accvgpr_mov_b32 a20, a4
179 ; GFX90A-NEXT: v_accvgpr_mov_b32 a19, a3
180 ; GFX90A-NEXT: v_accvgpr_mov_b32 a18, a2
181 ; GFX90A-NEXT: v_accvgpr_mov_b32 a17, a1
182 ; GFX90A-NEXT: v_accvgpr_mov_b32 a16, a0
183 ; GFX90A-NEXT: s_nop 1
184 ; GFX90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]
185 ; GFX90A-NEXT: s_nop 7
186 ; GFX90A-NEXT: s_nop 2
187 ; GFX90A-NEXT: buffer_store_dword a0, off, s[0:3], s32 ; 4-byte Folded Spill
188 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
189 ; GFX90A-NEXT: buffer_store_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill
190 ; GFX90A-NEXT: buffer_store_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill
191 ; GFX90A-NEXT: buffer_store_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill
192 ; GFX90A-NEXT: buffer_store_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill
193 ; GFX90A-NEXT: buffer_store_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill
194 ; GFX90A-NEXT: buffer_store_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill
195 ; GFX90A-NEXT: buffer_store_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill
196 ; GFX90A-NEXT: buffer_store_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill
197 ; GFX90A-NEXT: buffer_store_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill
198 ; GFX90A-NEXT: v_accvgpr_read_b32 v39, a10 ; Reload Reuse
199 ; GFX90A-NEXT: v_accvgpr_read_b32 v38, a11 ; Reload Reuse
200 ; GFX90A-NEXT: v_accvgpr_read_b32 v37, a12 ; Reload Reuse
201 ; GFX90A-NEXT: v_accvgpr_read_b32 v36, a13 ; Reload Reuse
202 ; GFX90A-NEXT: v_accvgpr_read_b32 v35, a14 ; Reload Reuse
203 ; GFX90A-NEXT: v_accvgpr_read_b32 v34, a15 ; Reload Reuse
204 ; GFX90A-NEXT: ;;#ASMSTART
205 ; GFX90A-NEXT: ; copy
206 ; GFX90A-NEXT: ;;#ASMEND
207 ; GFX90A-NEXT: v_accvgpr_mov_b32 a16, a1
208 ; GFX90A-NEXT: buffer_load_dword a0, off, s[0:3], s32 ; 4-byte Folded Reload
209 ; GFX90A-NEXT: buffer_load_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload
210 ; GFX90A-NEXT: buffer_load_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload
211 ; GFX90A-NEXT: buffer_load_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload
212 ; GFX90A-NEXT: buffer_load_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload
213 ; GFX90A-NEXT: buffer_load_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload
214 ; GFX90A-NEXT: buffer_load_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload
215 ; GFX90A-NEXT: buffer_load_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload
216 ; GFX90A-NEXT: buffer_load_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload
217 ; GFX90A-NEXT: buffer_load_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload
218 ; GFX90A-NEXT: s_waitcnt vmcnt(9)
219 ; GFX90A-NEXT: v_accvgpr_write_b32 a10, v39 ; Reload Reuse
220 ; GFX90A-NEXT: v_accvgpr_write_b32 a11, v38 ; Reload Reuse
221 ; GFX90A-NEXT: v_accvgpr_write_b32 a12, v37 ; Reload Reuse
222 ; GFX90A-NEXT: v_accvgpr_write_b32 a13, v36 ; Reload Reuse
223 ; GFX90A-NEXT: v_accvgpr_write_b32 a14, v35 ; Reload Reuse
224 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
225 ; GFX90A-NEXT: v_accvgpr_write_b32 a15, v34 ; Reload Reuse
226 ; GFX90A-NEXT: ;;#ASMSTART
227 ; GFX90A-NEXT: ; copy
228 ; GFX90A-NEXT: ;;#ASMEND
229 ; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a2
230 ; GFX90A-NEXT: ;;#ASMSTART
231 ; GFX90A-NEXT: ; use a3 v[0:31]
232 ; GFX90A-NEXT: ;;#ASMEND
233 ; GFX90A-NEXT: s_setpc_b64 s[30:31]
234 %asm = call { <32 x i32>, <16 x float> } asm sideeffect "; def $0 $1","=${v[0:31]},=${a[0:15]}"()
235 %vgpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 0
236 %agpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 1
237 %mfma = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %v0, float %v1, <16 x float> %agpr0, i32 0, i32 0, i32 0)
238 %agpr1 = call i32 asm sideeffect "; copy ", "={a1},a,~{a[0:15]}"(<16 x float> %agpr0)
239 %agpr2 = call i32 asm sideeffect "; copy ", "={a2},a,{a[0:15]}"(i32 %agpr1, <16 x float> %mfma)
240 call void asm sideeffect "; use $0 $1","{a3},{v[0:31]}"(i32 %agpr2, <32 x i32> %vgpr0)
241 ret void
242 }
244 ; Check that we do make use of v32 if there are no AGPRs present in the function
245 define amdgpu_kernel void @no_agpr_no_reserve(ptr addrspace(1) %arg) #0 {
246 ; GFX908-LABEL: no_agpr_no_reserve:
247 ; GFX908: ; %bb.0:
248 ; GFX908-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
249 ; GFX908-NEXT: v_lshlrev_b32_e32 v0, 7, v0
250 ; GFX908-NEXT: s_waitcnt lgkmcnt(0)
251 ; GFX908-NEXT: global_load_dwordx4 v[1:4], v0, s[0:1] offset:16
252 ; GFX908-NEXT: global_load_dwordx4 v[5:8], v0, s[0:1]
253 ; GFX908-NEXT: global_load_dwordx4 v[9:12], v0, s[0:1] offset:48
254 ; GFX908-NEXT: global_load_dwordx4 v[13:16], v0, s[0:1] offset:32
255 ; GFX908-NEXT: global_load_dwordx4 v[17:20], v0, s[0:1] offset:80
256 ; GFX908-NEXT: global_load_dwordx4 v[21:24], v0, s[0:1] offset:64
257 ; GFX908-NEXT: global_load_dwordx4 v[25:28], v0, s[0:1] offset:112
258 ; GFX908-NEXT: global_load_dwordx4 v[29:32], v0, s[0:1] offset:96
259 ; GFX908-NEXT: s_waitcnt vmcnt(7)
260 ; GFX908-NEXT: v_add_u32_e32 v4, v4, v4
261 ; GFX908-NEXT: v_add_u32_e32 v3, v3, v3
262 ; GFX908-NEXT: v_add_u32_e32 v2, v2, v2
263 ; GFX908-NEXT: v_add_u32_e32 v1, v1, v1
264 ; GFX908-NEXT: s_waitcnt vmcnt(6)
265 ; GFX908-NEXT: v_add_u32_e32 v8, v8, v8
266 ; GFX908-NEXT: v_add_u32_e32 v7, v7, v7
267 ; GFX908-NEXT: v_add_u32_e32 v6, v6, v6
268 ; GFX908-NEXT: s_waitcnt vmcnt(0)
269 ; GFX908-NEXT: v_add_u32_e32 v32, v32, v32
270 ; GFX908-NEXT: v_add_u32_e32 v31, v31, v31
271 ; GFX908-NEXT: v_add_u32_e32 v30, v30, v30
272 ; GFX908-NEXT: v_add_u32_e32 v29, v29, v29
273 ; GFX908-NEXT: v_add_u32_e32 v5, v5, v5
274 ; GFX908-NEXT: v_add_u32_e32 v12, v12, v12
275 ; GFX908-NEXT: v_add_u32_e32 v11, v11, v11
276 ; GFX908-NEXT: v_add_u32_e32 v10, v10, v10
277 ; GFX908-NEXT: v_add_u32_e32 v9, v9, v9
278 ; GFX908-NEXT: v_add_u32_e32 v16, v16, v16
279 ; GFX908-NEXT: v_add_u32_e32 v15, v15, v15
280 ; GFX908-NEXT: v_add_u32_e32 v14, v14, v14
281 ; GFX908-NEXT: v_add_u32_e32 v13, v13, v13
282 ; GFX908-NEXT: v_add_u32_e32 v20, v20, v20
283 ; GFX908-NEXT: v_add_u32_e32 v19, v19, v19
284 ; GFX908-NEXT: v_add_u32_e32 v18, v18, v18
285 ; GFX908-NEXT: v_add_u32_e32 v17, v17, v17
286 ; GFX908-NEXT: v_add_u32_e32 v24, v24, v24
287 ; GFX908-NEXT: v_add_u32_e32 v23, v23, v23
288 ; GFX908-NEXT: v_add_u32_e32 v22, v22, v22
289 ; GFX908-NEXT: v_add_u32_e32 v21, v21, v21
290 ; GFX908-NEXT: v_add_u32_e32 v28, v28, v28
291 ; GFX908-NEXT: v_add_u32_e32 v27, v27, v27
292 ; GFX908-NEXT: v_add_u32_e32 v26, v26, v26
293 ; GFX908-NEXT: v_add_u32_e32 v25, v25, v25
294 ; GFX908-NEXT: global_store_dwordx4 v0, v[29:32], s[0:1] offset:96
295 ; GFX908-NEXT: global_store_dwordx4 v0, v[25:28], s[0:1] offset:112
296 ; GFX908-NEXT: global_store_dwordx4 v0, v[21:24], s[0:1] offset:64
297 ; GFX908-NEXT: global_store_dwordx4 v0, v[17:20], s[0:1] offset:80
298 ; GFX908-NEXT: global_store_dwordx4 v0, v[13:16], s[0:1] offset:32
299 ; GFX908-NEXT: global_store_dwordx4 v0, v[9:12], s[0:1] offset:48
300 ; GFX908-NEXT: global_store_dwordx4 v0, v[5:8], s[0:1]
301 ; GFX908-NEXT: global_store_dwordx4 v0, v[1:4], s[0:1] offset:16
302 ; GFX908-NEXT: s_endpgm
303 ;
304 ; GFX90A-LABEL: no_agpr_no_reserve:
305 ; GFX90A: ; %bb.0:
306 ; GFX90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
307 ; GFX90A-NEXT: v_lshlrev_b32_e32 v32, 7, v0
308 ; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
309 ; GFX90A-NEXT: global_load_dwordx4 v[0:3], v32, s[0:1] offset:16
310 ; GFX90A-NEXT: global_load_dwordx4 v[4:7], v32, s[0:1]
311 ; GFX90A-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:48
312 ; GFX90A-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:32
313 ; GFX90A-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:80
314 ; GFX90A-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:64
315 ; GFX90A-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:112
316 ; GFX90A-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:96
317 ; GFX90A-NEXT: s_waitcnt vmcnt(7)
318 ; GFX90A-NEXT: v_add_u32_e32 v3, v3, v3
319 ; GFX90A-NEXT: v_add_u32_e32 v2, v2, v2
320 ; GFX90A-NEXT: v_add_u32_e32 v1, v1, v1
321 ; GFX90A-NEXT: v_add_u32_e32 v0, v0, v0
322 ; GFX90A-NEXT: s_waitcnt vmcnt(6)
323 ; GFX90A-NEXT: v_add_u32_e32 v7, v7, v7
324 ; GFX90A-NEXT: v_add_u32_e32 v6, v6, v6
325 ; GFX90A-NEXT: v_add_u32_e32 v5, v5, v5
326 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
327 ; GFX90A-NEXT: v_add_u32_e32 v31, v31, v31
328 ; GFX90A-NEXT: v_add_u32_e32 v30, v30, v30
329 ; GFX90A-NEXT: v_add_u32_e32 v29, v29, v29
330 ; GFX90A-NEXT: v_add_u32_e32 v28, v28, v28
331 ; GFX90A-NEXT: v_add_u32_e32 v4, v4, v4
332 ; GFX90A-NEXT: v_add_u32_e32 v11, v11, v11
333 ; GFX90A-NEXT: v_add_u32_e32 v10, v10, v10
334 ; GFX90A-NEXT: v_add_u32_e32 v9, v9, v9
335 ; GFX90A-NEXT: v_add_u32_e32 v8, v8, v8
336 ; GFX90A-NEXT: v_add_u32_e32 v15, v15, v15
337 ; GFX90A-NEXT: v_add_u32_e32 v14, v14, v14
338 ; GFX90A-NEXT: v_add_u32_e32 v13, v13, v13
339 ; GFX90A-NEXT: v_add_u32_e32 v12, v12, v12
340 ; GFX90A-NEXT: v_add_u32_e32 v19, v19, v19
341 ; GFX90A-NEXT: v_add_u32_e32 v18, v18, v18
342 ; GFX90A-NEXT: v_add_u32_e32 v17, v17, v17
343 ; GFX90A-NEXT: v_add_u32_e32 v16, v16, v16
344 ; GFX90A-NEXT: v_add_u32_e32 v23, v23, v23
345 ; GFX90A-NEXT: v_add_u32_e32 v22, v22, v22
346 ; GFX90A-NEXT: v_add_u32_e32 v21, v21, v21
347 ; GFX90A-NEXT: v_add_u32_e32 v20, v20, v20
348 ; GFX90A-NEXT: v_add_u32_e32 v27, v27, v27
349 ; GFX90A-NEXT: v_add_u32_e32 v26, v26, v26
350 ; GFX90A-NEXT: v_add_u32_e32 v25, v25, v25
351 ; GFX90A-NEXT: v_add_u32_e32 v24, v24, v24
352 ; GFX90A-NEXT: global_store_dwordx4 v32, v[28:31], s[0:1] offset:96
353 ; GFX90A-NEXT: global_store_dwordx4 v32, v[24:27], s[0:1] offset:112
354 ; GFX90A-NEXT: global_store_dwordx4 v32, v[20:23], s[0:1] offset:64
355 ; GFX90A-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:80
356 ; GFX90A-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:32
357 ; GFX90A-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:48
358 ; GFX90A-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1]
359 ; GFX90A-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:16
360 ; GFX90A-NEXT: s_endpgm
361 %id = call i32 @llvm.amdgcn.workitem.id.x()
362 %gep = getelementptr inbounds <32 x i32>, ptr addrspace(1) %arg, i32 %id
363 %load = load <32 x i32>, ptr addrspace(1) %gep
364 %add = add <32 x i32> %load, %load
365 store <32 x i32> %add, ptr addrspace(1) %gep
366 ret void
367 }
369 ; FIXME: This case is broken. The asm value passed in v32 is live
370 ; through the range where the reserved def for the copy is introduced,
371 ; clobbering the user value.
372 define void @v32_asm_def_use(float %v0, float %v1) #0 {
373 ; GFX908-LABEL: v32_asm_def_use:
374 ; GFX908: ; %bb.0:
375 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
376 ; GFX908-NEXT: v_mov_b32_e32 v33, v1
377 ; GFX908-NEXT: v_mov_b32_e32 v34, v0
378 ; GFX908-NEXT: ;;#ASMSTART
379 ; GFX908-NEXT: ; def v[0:31] a[0:15]
380 ; GFX908-NEXT: ;;#ASMEND
381 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a15
382 ; GFX908-NEXT: ;;#ASMSTART
383 ; GFX908-NEXT: ; def v32
384 ; GFX908-NEXT: ;;#ASMEND
385 ; GFX908-NEXT: s_nop 1
386 ; GFX908-NEXT: v_accvgpr_write_b32 a31, v35
387 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a14
388 ; GFX908-NEXT: s_nop 1
389 ; GFX908-NEXT: v_accvgpr_write_b32 a30, v35
390 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a13
391 ; GFX908-NEXT: s_nop 1
392 ; GFX908-NEXT: v_accvgpr_write_b32 a29, v35
393 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a12
394 ; GFX908-NEXT: s_nop 1
395 ; GFX908-NEXT: v_accvgpr_write_b32 a28, v35
396 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a11
397 ; GFX908-NEXT: s_nop 1
398 ; GFX908-NEXT: v_accvgpr_write_b32 a27, v35
399 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a10
400 ; GFX908-NEXT: s_nop 1
401 ; GFX908-NEXT: v_accvgpr_write_b32 a26, v35
402 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a9
403 ; GFX908-NEXT: s_nop 1
404 ; GFX908-NEXT: v_accvgpr_write_b32 a25, v35
405 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a8
406 ; GFX908-NEXT: s_nop 1
407 ; GFX908-NEXT: v_accvgpr_write_b32 a24, v35
408 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a7
409 ; GFX908-NEXT: s_nop 1
410 ; GFX908-NEXT: v_accvgpr_write_b32 a23, v35
411 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a6
412 ; GFX908-NEXT: s_nop 1
413 ; GFX908-NEXT: v_accvgpr_write_b32 a22, v35
414 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a5
415 ; GFX908-NEXT: s_nop 1
416 ; GFX908-NEXT: v_accvgpr_write_b32 a21, v35
417 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a4
418 ; GFX908-NEXT: s_nop 1
419 ; GFX908-NEXT: v_accvgpr_write_b32 a20, v35
420 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a3
421 ; GFX908-NEXT: s_nop 1
422 ; GFX908-NEXT: v_accvgpr_write_b32 a19, v35
423 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a2
424 ; GFX908-NEXT: s_nop 1
425 ; GFX908-NEXT: v_accvgpr_write_b32 a18, v35
426 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a1
427 ; GFX908-NEXT: s_nop 1
428 ; GFX908-NEXT: v_accvgpr_write_b32 a17, v35
429 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a0
430 ; GFX908-NEXT: s_nop 1
431 ; GFX908-NEXT: v_accvgpr_write_b32 a16, v35
432 ; GFX908-NEXT: ;;#ASMSTART
433 ; GFX908-NEXT: ; copy
434 ; GFX908-NEXT: ;;#ASMEND
435 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a1
436 ; GFX908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v34, v33, a[16:31]
437 ; GFX908-NEXT: s_nop 0
438 ; GFX908-NEXT: v_accvgpr_write_b32 a32, v35
439 ; GFX908-NEXT: ;;#ASMSTART
440 ; GFX908-NEXT: ; copy
441 ; GFX908-NEXT: ;;#ASMEND
442 ; GFX908-NEXT: s_nop 7
443 ; GFX908-NEXT: v_accvgpr_read_b32 v33, a2
444 ; GFX908-NEXT: s_nop 1
445 ; GFX908-NEXT: v_accvgpr_write_b32 a3, v33
446 ; GFX908-NEXT: ;;#ASMSTART
447 ; GFX908-NEXT: ; use a3 v[0:31]
448 ; GFX908-NEXT: ;;#ASMEND
449 ; GFX908-NEXT: ;;#ASMSTART
450 ; GFX908-NEXT: ; use v32
451 ; GFX908-NEXT: ;;#ASMEND
452 ; GFX908-NEXT: s_setpc_b64 s[30:31]
453 ;
454 ; GFX90A-LABEL: v32_asm_def_use:
455 ; GFX90A: ; %bb.0:
456 ; GFX90A-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
457 ; GFX90A-NEXT: v_mov_b32_e32 v34, v0
458 ; GFX90A-NEXT: v_mov_b32_e32 v33, v1
459 ; GFX90A-NEXT: ;;#ASMSTART
460 ; GFX90A-NEXT: ; def v[0:31] a[0:15]
461 ; GFX90A-NEXT: ;;#ASMEND
462 ; GFX90A-NEXT: v_accvgpr_mov_b32 a31, a15
463 ; GFX90A-NEXT: v_accvgpr_mov_b32 a30, a14
464 ; GFX90A-NEXT: v_accvgpr_mov_b32 a29, a13
465 ; GFX90A-NEXT: v_accvgpr_mov_b32 a28, a12
466 ; GFX90A-NEXT: v_accvgpr_mov_b32 a27, a11
467 ; GFX90A-NEXT: v_accvgpr_mov_b32 a26, a10
468 ; GFX90A-NEXT: v_accvgpr_mov_b32 a25, a9
469 ; GFX90A-NEXT: v_accvgpr_mov_b32 a24, a8
470 ; GFX90A-NEXT: v_accvgpr_mov_b32 a23, a7
471 ; GFX90A-NEXT: v_accvgpr_mov_b32 a22, a6
472 ; GFX90A-NEXT: v_accvgpr_mov_b32 a21, a5
473 ; GFX90A-NEXT: v_accvgpr_mov_b32 a20, a4
474 ; GFX90A-NEXT: v_accvgpr_mov_b32 a19, a3
475 ; GFX90A-NEXT: v_accvgpr_mov_b32 a18, a2
476 ; GFX90A-NEXT: v_accvgpr_mov_b32 a17, a1
477 ; GFX90A-NEXT: v_accvgpr_mov_b32 a16, a0
478 ; GFX90A-NEXT: ;;#ASMSTART
479 ; GFX90A-NEXT: ; def v32
480 ; GFX90A-NEXT: ;;#ASMEND
481 ; GFX90A-NEXT: ;;#ASMSTART
482 ; GFX90A-NEXT: ; copy
483 ; GFX90A-NEXT: ;;#ASMEND
484 ; GFX90A-NEXT: v_accvgpr_read_b32 v35, a32 ; Reload Reuse
485 ; GFX90A-NEXT: v_accvgpr_mov_b32 a32, a1
486 ; GFX90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v34, v33, a[16:31]
487 ; GFX90A-NEXT: ;;#ASMSTART
488 ; GFX90A-NEXT: ; copy
489 ; GFX90A-NEXT: ;;#ASMEND
490 ; GFX90A-NEXT: v_accvgpr_write_b32 a32, v35 ; Reload Reuse
491 ; GFX90A-NEXT: s_nop 7
492 ; GFX90A-NEXT: s_nop 1
493 ; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a2
494 ; GFX90A-NEXT: ;;#ASMSTART
495 ; GFX90A-NEXT: ; use a3 v[0:31]
496 ; GFX90A-NEXT: ;;#ASMEND
497 ; GFX90A-NEXT: ;;#ASMSTART
498 ; GFX90A-NEXT: ; use v32
499 ; GFX90A-NEXT: ;;#ASMEND
500 ; GFX90A-NEXT: s_setpc_b64 s[30:31]
501 %asm = call { <32 x i32>, <16 x float> } asm sideeffect "; def $0 $1","=${v[0:31]},=${a[0:15]}"()
502 %vgpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 0
503 %agpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 1
504 %mfma = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %v0, float %v1, <16 x float> %agpr0, i32 0, i32 0, i32 0)
505 %v32 = call i32 asm sideeffect "; def $0","=${v32}"()
506 %agpr1 = call i32 asm sideeffect "; copy ", "={a1},a,~{a[0:15]}"(<16 x float> %agpr0)
507 %agpr2 = call i32 asm sideeffect "; copy ", "={a2},a,{a[0:15]}"(i32 %agpr1, <16 x float> %mfma)
508 call void asm sideeffect "; use $0 $1","{a3},{v[0:31]}"(i32 %agpr2, <32 x i32> %vgpr0)
509 call void asm sideeffect "; use $0","${v32}"(i32 %v32)
510 ret void
511 }
513 define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg2, i64 %arg3, <2 x half> %arg4, <2 x half> %arg5) #3 {
514 ; GFX908-LABEL: introduced_copy_to_sgpr:
515 ; GFX908: ; %bb.0: ; %bb
516 ; GFX908-NEXT: global_load_ushort v16, v[0:1], off glc
517 ; GFX908-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0
518 ; GFX908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x10
519 ; GFX908-NEXT: s_load_dword s9, s[4:5], 0x18
520 ; GFX908-NEXT: s_mov_b32 s8, 0
521 ; GFX908-NEXT: s_mov_b32 s5, s8
522 ; GFX908-NEXT: s_waitcnt lgkmcnt(0)
523 ; GFX908-NEXT: v_cvt_f32_u32_e32 v0, s3
524 ; GFX908-NEXT: s_sub_i32 s4, 0, s3
525 ; GFX908-NEXT: v_cvt_f32_f16_e32 v17, s9
526 ; GFX908-NEXT: v_mov_b32_e32 v19, 0
527 ; GFX908-NEXT: v_rcp_iflag_f32_e32 v2, v0
528 ; GFX908-NEXT: v_mov_b32_e32 v0, 0
529 ; GFX908-NEXT: v_mov_b32_e32 v1, 0
530 ; GFX908-NEXT: v_mul_f32_e32 v2, 0x4f7ffffe, v2
531 ; GFX908-NEXT: v_cvt_u32_f32_e32 v2, v2
532 ; GFX908-NEXT: v_readfirstlane_b32 s10, v2
533 ; GFX908-NEXT: s_mul_i32 s4, s4, s10
534 ; GFX908-NEXT: s_mul_hi_u32 s4, s10, s4
535 ; GFX908-NEXT: s_add_i32 s10, s10, s4
536 ; GFX908-NEXT: s_mul_hi_u32 s4, s2, s10
537 ; GFX908-NEXT: s_mul_i32 s10, s4, s3
538 ; GFX908-NEXT: s_sub_i32 s2, s2, s10
539 ; GFX908-NEXT: s_add_i32 s11, s4, 1
540 ; GFX908-NEXT: s_sub_i32 s10, s2, s3
541 ; GFX908-NEXT: s_cmp_ge_u32 s2, s3
542 ; GFX908-NEXT: s_cselect_b32 s4, s11, s4
543 ; GFX908-NEXT: s_cselect_b32 s2, s10, s2
544 ; GFX908-NEXT: s_add_i32 s10, s4, 1
545 ; GFX908-NEXT: s_cmp_ge_u32 s2, s3
546 ; GFX908-NEXT: s_cselect_b32 s4, s10, s4
547 ; GFX908-NEXT: s_lshr_b32 s9, s9, 16
548 ; GFX908-NEXT: s_lshl_b64 s[12:13], s[4:5], 5
549 ; GFX908-NEXT: v_cvt_f32_f16_e32 v18, s9
550 ; GFX908-NEXT: s_lshl_b64 s[2:3], s[0:1], 5
551 ; GFX908-NEXT: s_lshl_b64 s[10:11], s[6:7], 5
552 ; GFX908-NEXT: s_or_b32 s10, s10, 28
553 ; GFX908-NEXT: s_waitcnt vmcnt(0)
554 ; GFX908-NEXT: v_readfirstlane_b32 s5, v16
555 ; GFX908-NEXT: s_and_b32 s5, 0xffff, s5
556 ; GFX908-NEXT: s_mul_i32 s1, s1, s5
557 ; GFX908-NEXT: s_mul_hi_u32 s9, s0, s5
558 ; GFX908-NEXT: s_mul_i32 s0, s0, s5
559 ; GFX908-NEXT: s_add_i32 s1, s9, s1
560 ; GFX908-NEXT: s_lshl_b64 s[14:15], s[0:1], 5
561 ; GFX908-NEXT: s_branch .LBB3_2
562 ; GFX908-NEXT: .LBB3_1: ; %Flow20
563 ; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
564 ; GFX908-NEXT: s_andn2_b64 vcc, exec, s[0:1]
565 ; GFX908-NEXT: s_cbranch_vccz .LBB3_12
566 ; GFX908-NEXT: .LBB3_2: ; %bb9
567 ; GFX908-NEXT: ; =>This Loop Header: Depth=1
568 ; GFX908-NEXT: ; Child Loop BB3_5 Depth 2
569 ; GFX908-NEXT: s_mov_b64 s[16:17], -1
570 ; GFX908-NEXT: s_cbranch_scc0 .LBB3_10
571 ; GFX908-NEXT: ; %bb.3: ; %bb14
572 ; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
573 ; GFX908-NEXT: global_load_dwordx2 v[2:3], v[0:1], off
574 ; GFX908-NEXT: v_cmp_gt_i64_e64 s[0:1], s[6:7], -1
575 ; GFX908-NEXT: s_mov_b32 s9, s8
576 ; GFX908-NEXT: v_cndmask_b32_e64 v6, 0, 1, s[0:1]
577 ; GFX908-NEXT: v_mov_b32_e32 v4, s8
578 ; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v6
579 ; GFX908-NEXT: v_mov_b32_e32 v8, s8
580 ; GFX908-NEXT: v_mov_b32_e32 v6, s8
581 ; GFX908-NEXT: v_mov_b32_e32 v5, s9
582 ; GFX908-NEXT: v_mov_b32_e32 v9, s9
583 ; GFX908-NEXT: v_mov_b32_e32 v7, s9
584 ; GFX908-NEXT: v_cmp_lt_i64_e64 s[16:17], s[6:7], 0
585 ; GFX908-NEXT: v_mov_b32_e32 v11, v5
586 ; GFX908-NEXT: s_mov_b64 s[18:19], s[10:11]
587 ; GFX908-NEXT: v_mov_b32_e32 v10, v4
588 ; GFX908-NEXT: s_waitcnt vmcnt(0)
589 ; GFX908-NEXT: v_readfirstlane_b32 s5, v2
590 ; GFX908-NEXT: v_readfirstlane_b32 s9, v3
591 ; GFX908-NEXT: s_add_u32 s5, s5, 1
592 ; GFX908-NEXT: s_addc_u32 s9, s9, 0
593 ; GFX908-NEXT: s_mul_hi_u32 s21, s2, s5
594 ; GFX908-NEXT: s_mul_i32 s22, s3, s5
595 ; GFX908-NEXT: s_mul_i32 s20, s2, s5
596 ; GFX908-NEXT: s_mul_i32 s5, s2, s9
597 ; GFX908-NEXT: s_add_i32 s5, s21, s5
598 ; GFX908-NEXT: s_add_i32 s5, s5, s22
599 ; GFX908-NEXT: s_branch .LBB3_5
600 ; GFX908-NEXT: .LBB3_4: ; %bb58
601 ; GFX908-NEXT: ; in Loop: Header=BB3_5 Depth=2
602 ; 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
603 ; GFX908-NEXT: v_addc_co_u32_e32 v3, vcc, 0, v3, vcc
604 ; GFX908-NEXT: s_add_u32 s18, s18, s14
605 ; GFX908-NEXT: v_cmp_lt_i64_e64 s[24:25], -1, v[2:3]
606 ; GFX908-NEXT: s_addc_u32 s19, s19, s15
607 ; GFX908-NEXT: s_mov_b64 s[22:23], 0
608 ; GFX908-NEXT: s_andn2_b64 vcc, exec, s[24:25]
609 ; GFX908-NEXT: s_cbranch_vccz .LBB3_9
610 ; GFX908-NEXT: .LBB3_5: ; %bb16
611 ; GFX908-NEXT: ; Parent Loop BB3_2 Depth=1
612 ; GFX908-NEXT: ; => This Inner Loop Header: Depth=2
613 ; GFX908-NEXT: s_add_u32 s22, s18, s20
614 ; GFX908-NEXT: s_addc_u32 s23, s19, s5
615 ; GFX908-NEXT: global_load_dword v21, v19, s[22:23] offset:-12 glc
616 ; GFX908-NEXT: s_waitcnt vmcnt(0)
617 ; GFX908-NEXT: global_load_dword v20, v19, s[22:23] offset:-8 glc
618 ; GFX908-NEXT: s_waitcnt vmcnt(0)
619 ; GFX908-NEXT: global_load_dword v12, v19, s[22:23] offset:-4 glc
620 ; GFX908-NEXT: s_waitcnt vmcnt(0)
621 ; GFX908-NEXT: global_load_dword v12, v19, s[22:23] glc
622 ; GFX908-NEXT: s_waitcnt vmcnt(0)
623 ; GFX908-NEXT: ds_read_b64 v[12:13], v19
624 ; GFX908-NEXT: ds_read_b64 v[14:15], v0
625 ; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1]
626 ; GFX908-NEXT: s_waitcnt lgkmcnt(0)
627 ; GFX908-NEXT: s_cbranch_vccnz .LBB3_7
628 ; GFX908-NEXT: ; %bb.6: ; %bb51
629 ; GFX908-NEXT: ; in Loop: Header=BB3_5 Depth=2
630 ; GFX908-NEXT: v_cvt_f32_f16_sdwa v22, v21 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
631 ; GFX908-NEXT: v_cvt_f32_f16_e32 v21, v21
632 ; GFX908-NEXT: v_cvt_f32_f16_sdwa v23, v20 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
633 ; GFX908-NEXT: v_cvt_f32_f16_e32 v20, v20
634 ; GFX908-NEXT: v_add_f32_e32 v24, v17, v12
635 ; GFX908-NEXT: v_add_f32_e32 v25, v18, v13
636 ; GFX908-NEXT: v_add_f32_e32 v26, 0, v12
637 ; GFX908-NEXT: v_add_f32_e32 v27, 0, v13
638 ; GFX908-NEXT: v_add_f32_e32 v15, v22, v15
639 ; GFX908-NEXT: v_add_f32_e32 v14, v21, v14
640 ; GFX908-NEXT: v_add_f32_e32 v13, v23, v13
641 ; GFX908-NEXT: v_add_f32_e32 v12, v20, v12
642 ; GFX908-NEXT: v_add_f32_e32 v5, v5, v25
643 ; GFX908-NEXT: v_add_f32_e32 v4, v4, v24
644 ; GFX908-NEXT: v_add_f32_e32 v9, v9, v27
645 ; GFX908-NEXT: v_add_f32_e32 v8, v8, v26
646 ; GFX908-NEXT: v_add_f32_e32 v6, v6, v14
647 ; GFX908-NEXT: v_add_f32_e32 v7, v7, v15
648 ; GFX908-NEXT: v_add_f32_e32 v10, v10, v12
649 ; GFX908-NEXT: v_add_f32_e32 v11, v11, v13
650 ; GFX908-NEXT: s_mov_b64 s[22:23], -1
651 ; GFX908-NEXT: s_branch .LBB3_4
652 ; GFX908-NEXT: .LBB3_7: ; in Loop: Header=BB3_5 Depth=2
653 ; GFX908-NEXT: s_mov_b64 s[22:23], s[16:17]
654 ; GFX908-NEXT: s_andn2_b64 vcc, exec, s[22:23]
655 ; GFX908-NEXT: s_cbranch_vccz .LBB3_4
656 ; GFX908-NEXT: ; %bb.8: ; in Loop: Header=BB3_2 Depth=1
657 ; GFX908-NEXT: ; implicit-def: $vgpr10_vgpr11
658 ; GFX908-NEXT: ; implicit-def: $vgpr6_vgpr7
659 ; GFX908-NEXT: ; implicit-def: $vgpr8_vgpr9
660 ; GFX908-NEXT: ; implicit-def: $vgpr4_vgpr5
661 ; GFX908-NEXT: ; implicit-def: $vgpr2_vgpr3
662 ; GFX908-NEXT: ; implicit-def: $sgpr18_sgpr19
663 ; GFX908-NEXT: .LBB3_9: ; %loop.exit.guard
664 ; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
665 ; GFX908-NEXT: s_xor_b64 s[16:17], s[22:23], -1
666 ; GFX908-NEXT: .LBB3_10: ; %Flow19
667 ; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
668 ; GFX908-NEXT: s_mov_b64 s[0:1], -1
669 ; GFX908-NEXT: s_and_b64 vcc, exec, s[16:17]
670 ; GFX908-NEXT: s_cbranch_vccz .LBB3_1
671 ; GFX908-NEXT: ; %bb.11: ; %bb12
672 ; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1
673 ; GFX908-NEXT: s_add_u32 s6, s6, s4
674 ; GFX908-NEXT: s_addc_u32 s7, s7, 0
675 ; GFX908-NEXT: s_add_u32 s10, s10, s12
676 ; GFX908-NEXT: s_addc_u32 s11, s11, s13
677 ; GFX908-NEXT: s_mov_b64 s[0:1], 0
678 ; GFX908-NEXT: s_branch .LBB3_1
679 ; GFX908-NEXT: .LBB3_12: ; %DummyReturnBlock
680 ; GFX908-NEXT: s_endpgm
681 ;
682 ; GFX90A-LABEL: introduced_copy_to_sgpr:
683 ; GFX90A: ; %bb.0: ; %bb
684 ; GFX90A-NEXT: global_load_ushort v18, v[0:1], off glc
685 ; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0
686 ; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x10
687 ; GFX90A-NEXT: s_load_dword s9, s[4:5], 0x18
688 ; GFX90A-NEXT: s_mov_b32 s8, 0
689 ; GFX90A-NEXT: s_mov_b32 s5, s8
690 ; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
691 ; GFX90A-NEXT: v_cvt_f32_u32_e32 v0, s3
692 ; GFX90A-NEXT: s_sub_i32 s4, 0, s3
693 ; GFX90A-NEXT: v_mov_b32_e32 v19, 0
694 ; GFX90A-NEXT: v_rcp_iflag_f32_e32 v2, v0
695 ; GFX90A-NEXT: v_pk_mov_b32 v[0:1], 0, 0
696 ; GFX90A-NEXT: v_mul_f32_e32 v2, 0x4f7ffffe, v2
697 ; GFX90A-NEXT: v_cvt_u32_f32_e32 v3, v2
698 ; GFX90A-NEXT: v_cvt_f32_f16_e32 v2, s9
699 ; GFX90A-NEXT: v_readfirstlane_b32 s10, v3
700 ; GFX90A-NEXT: s_mul_i32 s4, s4, s10
701 ; GFX90A-NEXT: s_mul_hi_u32 s4, s10, s4
702 ; GFX90A-NEXT: s_add_i32 s10, s10, s4
703 ; GFX90A-NEXT: s_mul_hi_u32 s4, s2, s10
704 ; GFX90A-NEXT: s_mul_i32 s10, s4, s3
705 ; GFX90A-NEXT: s_sub_i32 s2, s2, s10
706 ; GFX90A-NEXT: s_add_i32 s11, s4, 1
707 ; GFX90A-NEXT: s_sub_i32 s10, s2, s3
708 ; GFX90A-NEXT: s_cmp_ge_u32 s2, s3
709 ; GFX90A-NEXT: s_cselect_b32 s4, s11, s4
710 ; GFX90A-NEXT: s_cselect_b32 s2, s10, s2
711 ; GFX90A-NEXT: s_add_i32 s10, s4, 1
712 ; GFX90A-NEXT: s_cmp_ge_u32 s2, s3
713 ; GFX90A-NEXT: s_cselect_b32 s4, s10, s4
714 ; GFX90A-NEXT: s_lshr_b32 s9, s9, 16
715 ; GFX90A-NEXT: s_lshl_b64 s[12:13], s[4:5], 5
716 ; GFX90A-NEXT: v_cvt_f32_f16_e32 v3, s9
717 ; GFX90A-NEXT: s_lshl_b64 s[2:3], s[0:1], 5
718 ; GFX90A-NEXT: s_lshl_b64 s[10:11], s[6:7], 5
719 ; GFX90A-NEXT: s_or_b32 s10, s10, 28
720 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
721 ; GFX90A-NEXT: v_readfirstlane_b32 s5, v18
722 ; GFX90A-NEXT: s_and_b32 s5, 0xffff, s5
723 ; GFX90A-NEXT: s_mul_i32 s1, s1, s5
724 ; GFX90A-NEXT: s_mul_hi_u32 s9, s0, s5
725 ; GFX90A-NEXT: s_mul_i32 s0, s0, s5
726 ; GFX90A-NEXT: s_add_i32 s1, s9, s1
727 ; GFX90A-NEXT: s_lshl_b64 s[14:15], s[0:1], 5
728 ; GFX90A-NEXT: s_branch .LBB3_2
729 ; GFX90A-NEXT: .LBB3_1: ; %Flow20
730 ; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
731 ; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[0:1]
732 ; GFX90A-NEXT: s_cbranch_vccz .LBB3_12
733 ; GFX90A-NEXT: .LBB3_2: ; %bb9
734 ; GFX90A-NEXT: ; =>This Loop Header: Depth=1
735 ; GFX90A-NEXT: ; Child Loop BB3_5 Depth 2
736 ; GFX90A-NEXT: s_mov_b64 s[16:17], -1
737 ; GFX90A-NEXT: s_cbranch_scc0 .LBB3_10
738 ; GFX90A-NEXT: ; %bb.3: ; %bb14
739 ; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
740 ; GFX90A-NEXT: global_load_dwordx2 v[4:5], v[0:1], off
741 ; GFX90A-NEXT: v_cmp_gt_i64_e64 s[0:1], s[6:7], -1
742 ; GFX90A-NEXT: s_mov_b32 s9, s8
743 ; GFX90A-NEXT: v_cndmask_b32_e64 v8, 0, 1, s[0:1]
744 ; GFX90A-NEXT: v_pk_mov_b32 v[6:7], s[8:9], s[8:9] op_sel:[0,1]
745 ; GFX90A-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v8
746 ; GFX90A-NEXT: v_pk_mov_b32 v[10:11], s[8:9], s[8:9] op_sel:[0,1]
747 ; GFX90A-NEXT: v_pk_mov_b32 v[8:9], s[8:9], s[8:9] op_sel:[0,1]
748 ; GFX90A-NEXT: v_cmp_lt_i64_e64 s[16:17], s[6:7], 0
749 ; GFX90A-NEXT: s_mov_b64 s[18:19], s[10:11]
750 ; GFX90A-NEXT: v_pk_mov_b32 v[12:13], v[6:7], v[6:7] op_sel:[0,1]
751 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
752 ; GFX90A-NEXT: v_readfirstlane_b32 s5, v4
753 ; GFX90A-NEXT: v_readfirstlane_b32 s9, v5
754 ; GFX90A-NEXT: s_add_u32 s5, s5, 1
755 ; GFX90A-NEXT: s_addc_u32 s9, s9, 0
756 ; GFX90A-NEXT: s_mul_hi_u32 s21, s2, s5
757 ; GFX90A-NEXT: s_mul_i32 s22, s3, s5
758 ; GFX90A-NEXT: s_mul_i32 s20, s2, s5
759 ; GFX90A-NEXT: s_mul_i32 s5, s2, s9
760 ; GFX90A-NEXT: s_add_i32 s5, s21, s5
761 ; GFX90A-NEXT: s_add_i32 s5, s5, s22
762 ; GFX90A-NEXT: s_branch .LBB3_5
763 ; GFX90A-NEXT: .LBB3_4: ; %bb58
764 ; GFX90A-NEXT: ; in Loop: Header=BB3_5 Depth=2
765 ; 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
766 ; GFX90A-NEXT: v_addc_co_u32_e32 v5, vcc, 0, v5, vcc
767 ; GFX90A-NEXT: s_add_u32 s18, s18, s14
768 ; GFX90A-NEXT: s_addc_u32 s19, s19, s15
769 ; GFX90A-NEXT: v_cmp_lt_i64_e64 s[24:25], -1, v[4:5]
770 ; GFX90A-NEXT: s_mov_b64 s[22:23], 0
771 ; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[24:25]
772 ; GFX90A-NEXT: s_cbranch_vccz .LBB3_9
773 ; GFX90A-NEXT: .LBB3_5: ; %bb16
774 ; GFX90A-NEXT: ; Parent Loop BB3_2 Depth=1
775 ; GFX90A-NEXT: ; => This Inner Loop Header: Depth=2
776 ; GFX90A-NEXT: s_add_u32 s22, s18, s20
777 ; GFX90A-NEXT: s_addc_u32 s23, s19, s5
778 ; GFX90A-NEXT: global_load_dword v21, v19, s[22:23] offset:-12 glc
779 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
780 ; GFX90A-NEXT: global_load_dword v20, v19, s[22:23] offset:-8 glc
781 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
782 ; GFX90A-NEXT: global_load_dword v14, v19, s[22:23] offset:-4 glc
783 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
784 ; GFX90A-NEXT: global_load_dword v14, v19, s[22:23] glc
785 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
786 ; GFX90A-NEXT: ds_read_b64 v[14:15], v19
787 ; GFX90A-NEXT: ds_read_b64 v[16:17], v0
788 ; GFX90A-NEXT: s_and_b64 vcc, exec, s[0:1]
789 ; GFX90A-NEXT: ; kill: killed $sgpr22 killed $sgpr23
790 ; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
791 ; GFX90A-NEXT: s_cbranch_vccnz .LBB3_7
792 ; GFX90A-NEXT: ; %bb.6: ; %bb51
793 ; GFX90A-NEXT: ; in Loop: Header=BB3_5 Depth=2
794 ; GFX90A-NEXT: v_cvt_f32_f16_sdwa v23, v21 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
795 ; GFX90A-NEXT: v_cvt_f32_f16_e32 v22, v21
796 ; GFX90A-NEXT: v_cvt_f32_f16_sdwa v21, v20 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
797 ; GFX90A-NEXT: v_cvt_f32_f16_e32 v20, v20
798 ; GFX90A-NEXT: v_pk_add_f32 v[24:25], v[2:3], v[14:15]
799 ; GFX90A-NEXT: v_pk_add_f32 v[26:27], v[14:15], 0 op_sel_hi:[1,0]
800 ; GFX90A-NEXT: v_pk_add_f32 v[16:17], v[22:23], v[16:17]
801 ; GFX90A-NEXT: v_pk_add_f32 v[14:15], v[20:21], v[14:15]
802 ; GFX90A-NEXT: v_pk_add_f32 v[6:7], v[6:7], v[24:25]
803 ; GFX90A-NEXT: v_pk_add_f32 v[10:11], v[10:11], v[26:27]
804 ; GFX90A-NEXT: v_pk_add_f32 v[8:9], v[8:9], v[16:17]
805 ; GFX90A-NEXT: v_pk_add_f32 v[12:13], v[12:13], v[14:15]
806 ; GFX90A-NEXT: s_mov_b64 s[22:23], -1
807 ; GFX90A-NEXT: s_branch .LBB3_4
808 ; GFX90A-NEXT: .LBB3_7: ; in Loop: Header=BB3_5 Depth=2
809 ; GFX90A-NEXT: s_mov_b64 s[22:23], s[16:17]
810 ; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[22:23]
811 ; GFX90A-NEXT: s_cbranch_vccz .LBB3_4
812 ; GFX90A-NEXT: ; %bb.8: ; in Loop: Header=BB3_2 Depth=1
813 ; GFX90A-NEXT: ; implicit-def: $vgpr12_vgpr13
814 ; GFX90A-NEXT: ; implicit-def: $vgpr8_vgpr9
815 ; GFX90A-NEXT: ; implicit-def: $vgpr10_vgpr11
816 ; GFX90A-NEXT: ; implicit-def: $vgpr6_vgpr7
817 ; GFX90A-NEXT: ; implicit-def: $vgpr4_vgpr5
818 ; GFX90A-NEXT: ; implicit-def: $sgpr18_sgpr19
819 ; GFX90A-NEXT: .LBB3_9: ; %loop.exit.guard
820 ; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
821 ; GFX90A-NEXT: s_xor_b64 s[16:17], s[22:23], -1
822 ; GFX90A-NEXT: .LBB3_10: ; %Flow19
823 ; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
824 ; GFX90A-NEXT: s_mov_b64 s[0:1], -1
825 ; GFX90A-NEXT: s_and_b64 vcc, exec, s[16:17]
826 ; GFX90A-NEXT: s_cbranch_vccz .LBB3_1
827 ; GFX90A-NEXT: ; %bb.11: ; %bb12
828 ; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1
829 ; GFX90A-NEXT: s_add_u32 s6, s6, s4
830 ; GFX90A-NEXT: s_addc_u32 s7, s7, 0
831 ; GFX90A-NEXT: s_add_u32 s10, s10, s12
832 ; GFX90A-NEXT: s_addc_u32 s11, s11, s13
833 ; GFX90A-NEXT: s_mov_b64 s[0:1], 0
834 ; GFX90A-NEXT: s_branch .LBB3_1
835 ; GFX90A-NEXT: .LBB3_12: ; %DummyReturnBlock
836 ; GFX90A-NEXT: s_endpgm
837 bb:
838 %i = load volatile i16, ptr addrspace(4) undef, align 2
839 %i6 = zext i16 %i to i64
840 %i7 = udiv i32 %arg1, %arg2
841 %i8 = zext i32 %i7 to i64
842 br label %bb9
844 bb9: ; preds = %bb12, %bb
845 %i10 = phi i64 [ %arg3, %bb ], [ %i13, %bb12 ]
846 br i1 undef, label %bb14, label %bb12
848 bb12: ; preds = %bb58, %bb9
849 %i13 = add nuw nsw i64 %i10, %i8
850 br label %bb9
852 bb14: ; preds = %bb9
853 %i11 = icmp slt i64 %i10, 0
854 %i15 = load i64, ptr addrspace(1) null, align 8
855 br label %bb16
857 bb16: ; preds = %bb58, %bb14
858 %i17 = phi i64 [ %i65, %bb58 ], [ %i15, %bb14 ]
859 %i18 = phi <2 x float> [ %i59, %bb58 ], [ zeroinitializer, %bb14 ]
860 %i19 = phi <2 x float> [ %i60, %bb58 ], [ zeroinitializer, %bb14 ]
861 %i20 = phi <2 x float> [ %i61, %bb58 ], [ zeroinitializer, %bb14 ]
862 %i21 = phi <2 x float> [ %i62, %bb58 ], [ zeroinitializer, %bb14 ]
863 %i22 = add nsw i64 %i17, 1
864 %i23 = mul nsw i64 %i22, %arg
865 %i24 = add nsw i64 %i23, %i10
866 %i25 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 8
867 %i27 = load volatile <2 x half>, ptr addrspace(1) %i25, align 16
868 %i28 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 10
869 %i30 = load volatile <2 x half>, ptr addrspace(1) %i28, align 4
870 %i31 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 12
871 %i33 = load volatile <2 x half>, ptr addrspace(1) %i31, align 8
872 %i34 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 14
873 %i36 = load volatile <2 x half>, ptr addrspace(1) %i34, align 4
874 %i43 = load volatile <2 x float>, ptr addrspace(3) null, align 8
875 %i46 = load volatile <2 x float>, ptr addrspace(3) undef, align 32
876 fence syncscope("workgroup") acquire
877 br i1 %i11, label %bb58, label %bb51
879 bb51: ; preds = %bb16
880 %i37 = fpext <2 x half> %arg4 to <2 x float>
881 %i39 = fpext <2 x half> %i27 to <2 x float>
882 %i40 = fpext <2 x half> %i30 to <2 x float>
883 %i41 = fpext <2 x half> %i33 to <2 x float>
884 %i42 = fpext <2 x half> %i36 to <2 x float>
885 %i44 = fadd contract <2 x float> %i37, %i43
886 %i45 = fadd contract <2 x float> %i43, zeroinitializer
887 %i47 = fadd contract <2 x float> %i39, %i46
888 %i48 = fadd contract <2 x float> %i40, %i43
889 %i49 = fadd contract <2 x float> %i41, zeroinitializer
890 %i50 = fadd contract <2 x float> %i42, zeroinitializer
891 %i52 = fadd contract <2 x float> %i18, %i44
892 %i53 = fadd contract <2 x float> %i19, %i45
893 %i54 = fadd contract <2 x float> %i20, %i47
894 %i55 = fadd contract <2 x float> %i21, %i48
895 %i56 = fadd contract <2 x float> %i49, zeroinitializer
896 %i57 = fadd contract <2 x float> %i50, zeroinitializer
897 br label %bb58
899 bb58: ; preds = %bb51, %bb16
900 %i59 = phi <2 x float> [ %i18, %bb16 ], [ %i52, %bb51 ]
901 %i60 = phi <2 x float> [ %i19, %bb16 ], [ %i53, %bb51 ]
902 %i61 = phi <2 x float> [ %i20, %bb16 ], [ %i54, %bb51 ]
903 %i62 = phi <2 x float> [ %i21, %bb16 ], [ %i55, %bb51 ]
904 %i63 = phi <2 x float> [ zeroinitializer, %bb16 ], [ %i56, %bb51 ]
905 %i64 = phi <2 x float> [ zeroinitializer, %bb16 ], [ %i57, %bb51 ]
906 %i65 = add nsw i64 %i17, %i6
907 %i66 = icmp slt i64 %i65, 0
908 br i1 %i66, label %bb16, label %bb12
909 }
911 ; This testcase would fail on GFX908 due to not having a free VGPR available to
912 ; copy SGPR to AGPR.
913 define void @no_free_vgprs_at_sgpr_to_agpr_copy(float %v0, float %v1) #0 {
914 ; GFX908-LABEL: no_free_vgprs_at_sgpr_to_agpr_copy:
915 ; GFX908: ; %bb.0:
916 ; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
917 ; GFX908-NEXT: v_mov_b32_e32 v32, v1
918 ; GFX908-NEXT: v_mov_b32_e32 v33, v0
919 ; GFX908-NEXT: ;;#ASMSTART
920 ; GFX908-NEXT: ; def v[0:31] s[0:15]
921 ; GFX908-NEXT: ;;#ASMEND
922 ; GFX908-NEXT: v_mov_b32_e32 v39, s15
923 ; GFX908-NEXT: s_nop 1
924 ; GFX908-NEXT: v_accvgpr_write_b32 a31, v39
925 ; GFX908-NEXT: v_mov_b32_e32 v39, s14
926 ; GFX908-NEXT: s_nop 1
927 ; GFX908-NEXT: v_accvgpr_write_b32 a30, v39
928 ; GFX908-NEXT: v_mov_b32_e32 v39, s13
929 ; GFX908-NEXT: s_nop 1
930 ; GFX908-NEXT: v_accvgpr_write_b32 a29, v39
931 ; GFX908-NEXT: v_mov_b32_e32 v39, s12
932 ; GFX908-NEXT: s_nop 1
933 ; GFX908-NEXT: v_accvgpr_write_b32 a28, v39
934 ; GFX908-NEXT: v_mov_b32_e32 v39, s11
935 ; GFX908-NEXT: s_nop 1
936 ; GFX908-NEXT: v_accvgpr_write_b32 a27, v39
937 ; GFX908-NEXT: v_mov_b32_e32 v39, s10
938 ; GFX908-NEXT: s_nop 1
939 ; GFX908-NEXT: v_accvgpr_write_b32 a26, v39
940 ; GFX908-NEXT: v_mov_b32_e32 v39, s9
941 ; GFX908-NEXT: s_nop 1
942 ; GFX908-NEXT: v_accvgpr_write_b32 a25, v39
943 ; GFX908-NEXT: v_mov_b32_e32 v39, s8
944 ; GFX908-NEXT: s_nop 1
945 ; GFX908-NEXT: v_accvgpr_write_b32 a24, v39
946 ; GFX908-NEXT: v_mov_b32_e32 v39, s7
947 ; GFX908-NEXT: s_nop 1
948 ; GFX908-NEXT: v_accvgpr_write_b32 a23, v39
949 ; GFX908-NEXT: v_mov_b32_e32 v39, s6
950 ; GFX908-NEXT: s_nop 1
951 ; GFX908-NEXT: v_accvgpr_write_b32 a22, v39
952 ; GFX908-NEXT: v_mov_b32_e32 v39, s5
953 ; GFX908-NEXT: s_nop 1
954 ; GFX908-NEXT: v_accvgpr_write_b32 a21, v39
955 ; GFX908-NEXT: v_mov_b32_e32 v39, s4
956 ; GFX908-NEXT: s_nop 1
957 ; GFX908-NEXT: v_accvgpr_write_b32 a20, v39
958 ; GFX908-NEXT: v_mov_b32_e32 v39, s3
959 ; GFX908-NEXT: s_nop 1
960 ; GFX908-NEXT: v_accvgpr_write_b32 a19, v39
961 ; GFX908-NEXT: v_mov_b32_e32 v39, s2
962 ; GFX908-NEXT: s_nop 1
963 ; GFX908-NEXT: v_accvgpr_write_b32 a18, v39
964 ; GFX908-NEXT: v_mov_b32_e32 v39, s1
965 ; GFX908-NEXT: s_nop 1
966 ; GFX908-NEXT: v_accvgpr_write_b32 a17, v39
967 ; GFX908-NEXT: v_mov_b32_e32 v39, s0
968 ; GFX908-NEXT: s_nop 1
969 ; GFX908-NEXT: v_accvgpr_write_b32 a16, v39
970 ; GFX908-NEXT: s_nop 0
971 ; GFX908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]
972 ; GFX908-NEXT: s_nop 7
973 ; GFX908-NEXT: s_nop 1
974 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a0 ; Reload Reuse
975 ; GFX908-NEXT: v_accvgpr_read_b32 v38, a11 ; Reload Reuse
976 ; GFX908-NEXT: v_accvgpr_read_b32 v37, a12 ; Reload Reuse
977 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 ; 4-byte Folded Spill
978 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a1 ; Reload Reuse
979 ; GFX908-NEXT: v_accvgpr_read_b32 v36, a13 ; Reload Reuse
980 ; GFX908-NEXT: v_accvgpr_read_b32 v35, a14 ; Reload Reuse
981 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill
982 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a2 ; Reload Reuse
983 ; GFX908-NEXT: v_accvgpr_read_b32 v34, a15 ; Reload Reuse
984 ; GFX908-NEXT: s_nop 0
985 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill
986 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a3 ; Reload Reuse
987 ; GFX908-NEXT: s_nop 1
988 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill
989 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a4 ; Reload Reuse
990 ; GFX908-NEXT: s_nop 1
991 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill
992 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a5 ; Reload Reuse
993 ; GFX908-NEXT: s_nop 1
994 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill
995 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a6 ; Reload Reuse
996 ; GFX908-NEXT: s_nop 1
997 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill
998 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a7 ; Reload Reuse
999 ; GFX908-NEXT: s_nop 1
1000 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill
1001 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a8 ; Reload Reuse
1002 ; GFX908-NEXT: s_nop 1
1003 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill
1004 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a9 ; Reload Reuse
1005 ; GFX908-NEXT: s_nop 1
1006 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill
1007 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a10 ; Reload Reuse
1008 ; GFX908-NEXT: s_nop 1
1009 ; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill
1010 ; GFX908-NEXT: ;;#ASMSTART
1011 ; GFX908-NEXT: ; copy
1012 ; GFX908-NEXT: ;;#ASMEND
1013 ; GFX908-NEXT: v_accvgpr_read_b32 v39, a1
1014 ; GFX908-NEXT: s_nop 1
1015 ; GFX908-NEXT: v_accvgpr_write_b32 a32, v39
1016 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 ; 4-byte Folded Reload
1017 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1018 ; GFX908-NEXT: v_accvgpr_write_b32 a0, v39 ; Reload Reuse
1019 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload
1020 ; GFX908-NEXT: v_accvgpr_write_b32 a11, v38 ; Reload Reuse
1021 ; GFX908-NEXT: v_accvgpr_write_b32 a12, v37 ; Reload Reuse
1022 ; GFX908-NEXT: v_accvgpr_write_b32 a13, v36 ; Reload Reuse
1023 ; GFX908-NEXT: v_accvgpr_write_b32 a14, v35 ; Reload Reuse
1024 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1025 ; GFX908-NEXT: v_accvgpr_write_b32 a1, v39 ; Reload Reuse
1026 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload
1027 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1028 ; GFX908-NEXT: v_accvgpr_write_b32 a2, v39 ; Reload Reuse
1029 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload
1030 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1031 ; GFX908-NEXT: v_accvgpr_write_b32 a3, v39 ; Reload Reuse
1032 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload
1033 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1034 ; GFX908-NEXT: v_accvgpr_write_b32 a4, v39 ; Reload Reuse
1035 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload
1036 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1037 ; GFX908-NEXT: v_accvgpr_write_b32 a5, v39 ; Reload Reuse
1038 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload
1039 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1040 ; GFX908-NEXT: v_accvgpr_write_b32 a6, v39 ; Reload Reuse
1041 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload
1042 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1043 ; GFX908-NEXT: v_accvgpr_write_b32 a7, v39 ; Reload Reuse
1044 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload
1045 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1046 ; GFX908-NEXT: v_accvgpr_write_b32 a8, v39 ; Reload Reuse
1047 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload
1048 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1049 ; GFX908-NEXT: v_accvgpr_write_b32 a9, v39 ; Reload Reuse
1050 ; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload
1051 ; GFX908-NEXT: s_waitcnt vmcnt(0)
1052 ; GFX908-NEXT: v_accvgpr_write_b32 a10, v39 ; Reload Reuse
1053 ; GFX908-NEXT: v_accvgpr_write_b32 a15, v34 ; Reload Reuse
1054 ; GFX908-NEXT: ;;#ASMSTART
1055 ; GFX908-NEXT: ; copy
1056 ; GFX908-NEXT: ;;#ASMEND
1057 ; GFX908-NEXT: v_accvgpr_read_b32 v32, a2
1058 ; GFX908-NEXT: s_nop 1
1059 ; GFX908-NEXT: v_accvgpr_write_b32 a3, v32
1060 ; GFX908-NEXT: ;;#ASMSTART
1061 ; GFX908-NEXT: ; use a3 v[0:31]
1062 ; GFX908-NEXT: ;;#ASMEND
1063 ; GFX908-NEXT: s_setpc_b64 s[30:31]
1064 ;
1065 ; GFX90A-LABEL: no_free_vgprs_at_sgpr_to_agpr_copy:
1066 ; GFX90A: ; %bb.0:
1067 ; GFX90A-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1068 ; GFX90A-NEXT: v_mov_b32_e32 v33, v0
1069 ; GFX90A-NEXT: v_mov_b32_e32 v32, v1
1070 ; GFX90A-NEXT: ;;#ASMSTART
1071 ; GFX90A-NEXT: ; def v[0:31] s[0:15]
1072 ; GFX90A-NEXT: ;;#ASMEND
1073 ; GFX90A-NEXT: v_accvgpr_write_b32 a31, s15
1074 ; GFX90A-NEXT: v_accvgpr_write_b32 a30, s14
1075 ; GFX90A-NEXT: v_accvgpr_write_b32 a29, s13
1076 ; GFX90A-NEXT: v_accvgpr_write_b32 a28, s12
1077 ; GFX90A-NEXT: v_accvgpr_write_b32 a27, s11
1078 ; GFX90A-NEXT: v_accvgpr_write_b32 a26, s10
1079 ; GFX90A-NEXT: v_accvgpr_write_b32 a25, s9
1080 ; GFX90A-NEXT: v_accvgpr_write_b32 a24, s8
1081 ; GFX90A-NEXT: v_accvgpr_write_b32 a23, s7
1082 ; GFX90A-NEXT: v_accvgpr_write_b32 a22, s6
1083 ; GFX90A-NEXT: v_accvgpr_write_b32 a21, s5
1084 ; GFX90A-NEXT: v_accvgpr_write_b32 a20, s4
1085 ; GFX90A-NEXT: v_accvgpr_write_b32 a19, s3
1086 ; GFX90A-NEXT: v_accvgpr_write_b32 a18, s2
1087 ; GFX90A-NEXT: v_accvgpr_write_b32 a17, s1
1088 ; GFX90A-NEXT: v_accvgpr_write_b32 a16, s0
1089 ; GFX90A-NEXT: v_accvgpr_read_b32 v34, a32 ; Reload Reuse
1090 ; GFX90A-NEXT: s_nop 0
1091 ; GFX90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]
1092 ; GFX90A-NEXT: s_nop 7
1093 ; GFX90A-NEXT: s_nop 2
1094 ; GFX90A-NEXT: buffer_store_dword a0, off, s[0:3], s32 ; 4-byte Folded Spill
1095 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
1096 ; GFX90A-NEXT: buffer_store_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill
1097 ; GFX90A-NEXT: buffer_store_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill
1098 ; GFX90A-NEXT: buffer_store_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill
1099 ; GFX90A-NEXT: buffer_store_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill
1100 ; GFX90A-NEXT: buffer_store_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill
1101 ; GFX90A-NEXT: buffer_store_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill
1102 ; GFX90A-NEXT: buffer_store_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill
1103 ; GFX90A-NEXT: buffer_store_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill
1104 ; GFX90A-NEXT: buffer_store_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill
1105 ; GFX90A-NEXT: buffer_store_dword a10, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill
1106 ; GFX90A-NEXT: v_accvgpr_read_b32 v39, a11 ; Reload Reuse
1107 ; GFX90A-NEXT: v_accvgpr_read_b32 v38, a12 ; Reload Reuse
1108 ; GFX90A-NEXT: v_accvgpr_read_b32 v37, a13 ; Reload Reuse
1109 ; GFX90A-NEXT: v_accvgpr_read_b32 v36, a14 ; Reload Reuse
1110 ; GFX90A-NEXT: v_accvgpr_read_b32 v35, a15 ; Reload Reuse
1111 ; GFX90A-NEXT: ;;#ASMSTART
1112 ; GFX90A-NEXT: ; copy
1113 ; GFX90A-NEXT: ;;#ASMEND
1114 ; GFX90A-NEXT: v_accvgpr_mov_b32 a32, a1
1115 ; GFX90A-NEXT: buffer_load_dword a0, off, s[0:3], s32 ; 4-byte Folded Reload
1116 ; GFX90A-NEXT: buffer_load_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload
1117 ; GFX90A-NEXT: buffer_load_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload
1118 ; GFX90A-NEXT: buffer_load_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload
1119 ; GFX90A-NEXT: buffer_load_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload
1120 ; GFX90A-NEXT: buffer_load_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload
1121 ; GFX90A-NEXT: buffer_load_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload
1122 ; GFX90A-NEXT: buffer_load_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload
1123 ; GFX90A-NEXT: buffer_load_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload
1124 ; GFX90A-NEXT: buffer_load_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload
1125 ; GFX90A-NEXT: buffer_load_dword a10, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload
1126 ; GFX90A-NEXT: s_waitcnt vmcnt(10)
1127 ; GFX90A-NEXT: v_accvgpr_write_b32 a11, v39 ; Reload Reuse
1128 ; GFX90A-NEXT: v_accvgpr_write_b32 a12, v38 ; Reload Reuse
1129 ; GFX90A-NEXT: v_accvgpr_write_b32 a13, v37 ; Reload Reuse
1130 ; GFX90A-NEXT: v_accvgpr_write_b32 a14, v36 ; Reload Reuse
1131 ; GFX90A-NEXT: s_waitcnt vmcnt(0)
1132 ; GFX90A-NEXT: v_accvgpr_write_b32 a15, v35 ; Reload Reuse
1133 ; GFX90A-NEXT: ;;#ASMSTART
1134 ; GFX90A-NEXT: ; copy
1135 ; GFX90A-NEXT: ;;#ASMEND
1136 ; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a2
1137 ; GFX90A-NEXT: ;;#ASMSTART
1138 ; GFX90A-NEXT: ; use a3 v[0:31]
1139 ; GFX90A-NEXT: ;;#ASMEND
1140 ; GFX90A-NEXT: v_accvgpr_write_b32 a32, v34 ; Reload Reuse
1141 ; GFX90A-NEXT: s_setpc_b64 s[30:31]
1142 %asm = call { <32 x i32>, <16 x float> } asm sideeffect "; def $0 $1","=${v[0:31]},=${s[0:15]}"()
1143 %vgpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 0
1144 %agpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 1
1145 %mfma = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %v0, float %v1, <16 x float> %agpr0, i32 0, i32 0, i32 0)
1146 %agpr1 = call i32 asm sideeffect "; copy ", "={a1},a,~{a[0:15]}"(<16 x float> %agpr0)
1147 %agpr2 = call i32 asm sideeffect "; copy ", "={a2},a,{a[0:15]}"(i32 %agpr1, <16 x float> %mfma)
1148 call void asm sideeffect "; use $0 $1","{a3},{v[0:31]}"(i32 %agpr2, <32 x i32> %vgpr0)
1149 ret void
1150 }
1152 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32 immarg, i32 immarg, i32 immarg) #1
1153 declare i32 @llvm.amdgcn.workitem.id.x() #2
1155 attributes #0 = { "amdgpu-waves-per-eu"="6,6" }
1156 attributes #1 = { convergent nounwind readnone willreturn }
1157 attributes #2 = { nounwind readnone willreturn }
1158 attributes #3 = { "amdgpu-waves-per-eu"="7,7" }