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