xref: /llvm-project/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.ll (revision 6548b6354d1d990e1c98736f5e7c3de876bedc8e) 
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
2; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs -misched-cluster=0  < %s | FileCheck -check-prefix=GCN %s
3; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs -misched-cluster=0 -amdgpu-igrouplp-exact-solver-max-branches=250000 < %s | FileCheck -check-prefix=EXACTCUTOFF %s
4
5define amdgpu_kernel void @test_sched_group_barrier() #0 {
6; GCN-LABEL: test_sched_group_barrier:
7; GCN:       ; %bb.0: ; %entry
8; GCN-NEXT:    ; sched_group_barrier mask(0x00000000) size(1) SyncID(2)
9; GCN-NEXT:    ; sched_group_barrier mask(0x00000001) size(2) SyncID(4)
10; GCN-NEXT:    ; sched_group_barrier mask(0x00000004) size(8) SyncID(16)
11; GCN-NEXT:    ; sched_group_barrier mask(0x0000000F) size(10000) SyncID(-1)
12; GCN-NEXT:    s_endpgm
13;
14; EXACTCUTOFF-LABEL: test_sched_group_barrier:
15; EXACTCUTOFF:       ; %bb.0: ; %entry
16; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000000) size(1) SyncID(2)
17; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000001) size(2) SyncID(4)
18; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000004) size(8) SyncID(16)
19; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x0000000F) size(10000) SyncID(-1)
20; EXACTCUTOFF-NEXT:    s_endpgm
21entry:
22  call void @llvm.amdgcn.sched.group.barrier(i32 0, i32 1, i32 2) #1
23  call void @llvm.amdgcn.sched.group.barrier(i32 1, i32 2, i32 4) #1
24  call void @llvm.amdgcn.sched.group.barrier(i32 4, i32 8, i32 16) #1
25  call void @llvm.amdgcn.sched.group.barrier(i32 15, i32 10000, i32 -1) #1
26  ret void
27}
28
29define amdgpu_kernel void @test_sched_group_barrier_pipeline_READ_VALU_WRITE(ptr addrspace(1) noalias %in, ptr addrspace(1) noalias %out) #0 {
30; GCN-LABEL: test_sched_group_barrier_pipeline_READ_VALU_WRITE:
31; GCN:       ; %bb.0:
32; GCN-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
33; GCN-NEXT:    v_and_b32_e32 v0, 0x3ff, v0
34; GCN-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
35; GCN-NEXT:    ; kill: killed $sgpr0_sgpr1
36; GCN-NEXT:    s_waitcnt lgkmcnt(0)
37; GCN-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
38; GCN-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:16
39; GCN-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:32
40; GCN-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:48
41; GCN-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:64
42; GCN-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:80
43; GCN-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:96
44; GCN-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:112
45; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(8) SyncID(0)
46; GCN-NEXT:    s_waitcnt vmcnt(7)
47; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
48; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
49; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
50; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
51; GCN-NEXT:    s_waitcnt vmcnt(6)
52; GCN-NEXT:    v_mul_lo_u32 v7, v7, v7
53; GCN-NEXT:    v_mul_lo_u32 v6, v6, v6
54; GCN-NEXT:    v_mul_lo_u32 v5, v5, v5
55; GCN-NEXT:    v_mul_lo_u32 v4, v4, v4
56; GCN-NEXT:    s_waitcnt vmcnt(5)
57; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
58; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
59; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
60; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
61; GCN-NEXT:    s_waitcnt vmcnt(4)
62; GCN-NEXT:    v_mul_lo_u32 v15, v15, v15
63; GCN-NEXT:    v_mul_lo_u32 v14, v14, v14
64; GCN-NEXT:    v_mul_lo_u32 v13, v13, v13
65; GCN-NEXT:    v_mul_lo_u32 v12, v12, v12
66; GCN-NEXT:    s_waitcnt vmcnt(3)
67; GCN-NEXT:    v_mul_lo_u32 v19, v19, v19
68; GCN-NEXT:    v_mul_lo_u32 v18, v18, v18
69; GCN-NEXT:    v_mul_lo_u32 v17, v17, v17
70; GCN-NEXT:    v_mul_lo_u32 v16, v16, v16
71; GCN-NEXT:    s_waitcnt vmcnt(2)
72; GCN-NEXT:    v_mul_lo_u32 v23, v23, v23
73; GCN-NEXT:    v_mul_lo_u32 v22, v22, v22
74; GCN-NEXT:    v_mul_lo_u32 v21, v21, v21
75; GCN-NEXT:    v_mul_lo_u32 v20, v20, v20
76; GCN-NEXT:    s_waitcnt vmcnt(1)
77; GCN-NEXT:    v_mul_lo_u32 v27, v27, v27
78; GCN-NEXT:    v_mul_lo_u32 v26, v26, v26
79; GCN-NEXT:    v_mul_lo_u32 v25, v25, v25
80; GCN-NEXT:    v_mul_lo_u32 v24, v24, v24
81; GCN-NEXT:    s_waitcnt vmcnt(0)
82; GCN-NEXT:    v_mul_lo_u32 v31, v31, v31
83; GCN-NEXT:    v_mul_lo_u32 v30, v30, v30
84; GCN-NEXT:    v_mul_lo_u32 v29, v29, v29
85; GCN-NEXT:    v_mul_lo_u32 v28, v28, v28
86; GCN-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:112
87; GCN-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:96
88; GCN-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:80
89; GCN-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:64
90; GCN-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:48
91; GCN-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:32
92; GCN-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:16
93; GCN-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
94; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(30) SyncID(0)
95; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
96; GCN-NEXT:    s_endpgm
97;
98; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_READ_VALU_WRITE:
99; EXACTCUTOFF:       ; %bb.0:
100; EXACTCUTOFF-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
101; EXACTCUTOFF-NEXT:    v_and_b32_e32 v0, 0x3ff, v0
102; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
103; EXACTCUTOFF-NEXT:    ; kill: killed $sgpr0_sgpr1
104; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
105; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
106; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:16
107; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:32
108; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:48
109; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:64
110; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:80
111; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:96
112; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:112
113; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(8) SyncID(0)
114; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(7)
115; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
116; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
117; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
118; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
119; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(6)
120; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v7, v7, v7
121; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v6, v6, v6
122; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v5, v5, v5
123; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v4, v4, v4
124; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(5)
125; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
126; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
127; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
128; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
129; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(4)
130; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v15, v15, v15
131; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v14, v14, v14
132; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v13, v13, v13
133; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v12, v12, v12
134; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(3)
135; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v19, v19, v19
136; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v18, v18, v18
137; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v17, v17, v17
138; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v16, v16, v16
139; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(2)
140; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v23, v23, v23
141; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v22, v22, v22
142; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v21, v21, v21
143; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v20, v20, v20
144; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
145; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v27, v27, v27
146; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v26, v26, v26
147; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v25, v25, v25
148; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v24, v24, v24
149; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
150; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v31, v31, v31
151; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v30, v30, v30
152; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v29, v29, v29
153; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v28, v28, v28
154; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:112
155; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:96
156; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:80
157; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:64
158; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:48
159; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:32
160; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:16
161; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
162; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(30) SyncID(0)
163; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
164; EXACTCUTOFF-NEXT:    s_endpgm
165  %tid = call i32 @llvm.amdgcn.workitem.id.x() #2
166  %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid
167  %load = load <32 x i32>, ptr addrspace(1) %gep1
168  %mul = mul <32 x i32> %load, %load
169  %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid
170  store <32 x i32> %mul, ptr addrspace(1) %gep2
171  ; 8 VMEM read
172  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 8, i32 0)
173  ; 30 VALU
174  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 30, i32 0)
175  ; 8 VMEM write
176  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 8, i32 0)
177  ret void
178}
179
180define amdgpu_kernel void @test_sched_group_barrier_pipeline_alternating_READ_VALU(ptr addrspace(1) noalias %in, ptr addrspace(1) noalias %out) #0 {
181; GCN-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU:
182; GCN:       ; %bb.0:
183; GCN-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
184; GCN-NEXT:    v_and_b32_e32 v0, 0x3ff, v0
185; GCN-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
186; GCN-NEXT:    s_waitcnt lgkmcnt(0)
187; GCN-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:16
188; GCN-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
189; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
190; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
191; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
192; GCN-NEXT:    s_waitcnt vmcnt(1)
193; GCN-NEXT:    v_mul_lo_u32 v29, v29, v29
194; GCN-NEXT:    s_waitcnt vmcnt(0)
195; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
196; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
197; GCN-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:112
198; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
199; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
200; GCN-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:96
201; GCN-NEXT:    v_mul_lo_u32 v28, v28, v28
202; GCN-NEXT:    v_mul_lo_u32 v31, v31, v31
203; GCN-NEXT:    v_mul_lo_u32 v30, v30, v30
204; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
205; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
206; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
207; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
208; GCN-NEXT:    s_waitcnt vmcnt(1)
209; GCN-NEXT:    v_mul_lo_u32 v7, v7, v7
210; GCN-NEXT:    v_mul_lo_u32 v6, v6, v6
211; GCN-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:80
212; GCN-NEXT:    v_mul_lo_u32 v5, v5, v5
213; GCN-NEXT:    v_mul_lo_u32 v4, v4, v4
214; GCN-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:48
215; GCN-NEXT:    s_waitcnt vmcnt(2)
216; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
217; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
218; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
219; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
220; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
221; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
222; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
223; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
224; GCN-NEXT:    s_waitcnt vmcnt(1)
225; GCN-NEXT:    v_mul_lo_u32 v15, v15, v15
226; GCN-NEXT:    v_mul_lo_u32 v14, v14, v14
227; GCN-NEXT:    v_mul_lo_u32 v13, v13, v13
228; GCN-NEXT:    s_waitcnt vmcnt(0)
229; GCN-NEXT:    v_mul_lo_u32 v19, v19, v19
230; GCN-NEXT:    v_mul_lo_u32 v18, v18, v18
231; GCN-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:64
232; GCN-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:32
233; GCN-NEXT:    v_mul_lo_u32 v17, v17, v17
234; GCN-NEXT:    v_mul_lo_u32 v16, v16, v16
235; GCN-NEXT:    v_mul_lo_u32 v12, v12, v12
236; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
237; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
238; GCN-NEXT:    s_waitcnt vmcnt(1)
239; GCN-NEXT:    v_mul_lo_u32 v23, v23, v23
240; GCN-NEXT:    s_waitcnt vmcnt(0)
241; GCN-NEXT:    v_mul_lo_u32 v27, v27, v27
242; GCN-NEXT:    v_mul_lo_u32 v26, v26, v26
243; GCN-NEXT:    v_mul_lo_u32 v25, v25, v25
244; GCN-NEXT:    v_mul_lo_u32 v24, v24, v24
245; GCN-NEXT:    v_mul_lo_u32 v22, v22, v22
246; GCN-NEXT:    v_mul_lo_u32 v21, v21, v21
247; GCN-NEXT:    v_mul_lo_u32 v20, v20, v20
248; GCN-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:112
249; GCN-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:96
250; GCN-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:80
251; GCN-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:64
252; GCN-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:48
253; GCN-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:32
254; GCN-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:16
255; GCN-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
256; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
257; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
258; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
259; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
260; GCN-NEXT:    s_endpgm
261;
262; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU:
263; EXACTCUTOFF:       ; %bb.0:
264; EXACTCUTOFF-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
265; EXACTCUTOFF-NEXT:    v_and_b32_e32 v0, 0x3ff, v0
266; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
267; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
268; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:16
269; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
270; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
271; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
272; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
273; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
274; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v29, v29, v29
275; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
276; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
277; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
278; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:112
279; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
280; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
281; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:96
282; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v28, v28, v28
283; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v31, v31, v31
284; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v30, v30, v30
285; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
286; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
287; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
288; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
289; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
290; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v7, v7, v7
291; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v6, v6, v6
292; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:80
293; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v5, v5, v5
294; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v4, v4, v4
295; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:48
296; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(2)
297; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
298; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
299; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
300; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
301; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
302; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
303; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
304; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
305; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
306; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v15, v15, v15
307; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v14, v14, v14
308; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v13, v13, v13
309; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
310; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v19, v19, v19
311; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v18, v18, v18
312; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:64
313; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:32
314; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v17, v17, v17
315; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v16, v16, v16
316; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v12, v12, v12
317; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
318; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
319; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
320; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v23, v23, v23
321; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
322; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v27, v27, v27
323; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v26, v26, v26
324; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v25, v25, v25
325; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v24, v24, v24
326; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v22, v22, v22
327; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v21, v21, v21
328; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v20, v20, v20
329; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:112
330; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:96
331; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:80
332; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:64
333; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:48
334; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:32
335; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:16
336; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
337; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
338; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
339; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
340; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
341; EXACTCUTOFF-NEXT:    s_endpgm
342  %tid = call i32 @llvm.amdgcn.workitem.id.x() #2
343  %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid
344  %load = load <32 x i32>, ptr addrspace(1) %gep1
345  %mul = mul <32 x i32> %load, %load
346  %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid
347  store <32 x i32> %mul, ptr addrspace(1) %gep2
348  ; 1 VMEM read
349  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
350  ; 2 VALU
351  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
352  ; 1 VMEM read
353  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
354  ; 2 VALU
355  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
356  ; 1 VMEM read
357  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
358  ; 2 VALU
359  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
360  ; 1 VMEM read
361  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
362  ; 2 VALU
363  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
364  ; 1 VMEM read
365  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
366  ; 2 VALU
367  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
368  ; 1 VMEM read
369  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
370  ; 2 VALU
371  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
372  ; 1 VMEM read
373  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
374  ; 2 VALU
375  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
376  ; 1 VMEM read
377  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
378  ; 2 VALU
379  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
380  ; 8 VMEM write
381  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 8, i32 0)
382  ret void
383}
384
385define amdgpu_kernel void @test_sched_group_barrier_pipeline_alternating_READ_VALU_WRITE(ptr addrspace(1) noalias %in, ptr addrspace(1) noalias %out) #0 {
386; GCN-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU_WRITE:
387; GCN:       ; %bb.0:
388; GCN-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
389; GCN-NEXT:    v_and_b32_e32 v0, 0x3ff, v0
390; GCN-NEXT:    v_lshlrev_b32_e32 v16, 7, v0
391; GCN-NEXT:    ; kill: killed $sgpr0_sgpr1
392; GCN-NEXT:    s_waitcnt lgkmcnt(0)
393; GCN-NEXT:    global_load_dwordx4 v[12:15], v16, s[0:1] offset:32
394; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
395; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
396; GCN-NEXT:    s_waitcnt vmcnt(0)
397; GCN-NEXT:    v_mul_lo_u32 v13, v13, v13
398; GCN-NEXT:    v_mul_lo_u32 v12, v12, v12
399; GCN-NEXT:    v_mul_lo_u32 v15, v15, v15
400; GCN-NEXT:    v_mul_lo_u32 v14, v14, v14
401; GCN-NEXT:    global_store_dwordx4 v16, v[12:15], s[2:3] offset:32
402; GCN-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1]
403; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
404; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
405; GCN-NEXT:    s_waitcnt vmcnt(0)
406; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
407; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
408; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
409; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
410; GCN-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3]
411; GCN-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:112
412; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
413; GCN-NEXT:    s_waitcnt vmcnt(0)
414; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
415; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
416; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
417; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
418; GCN-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:112
419; GCN-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:96
420; GCN-NEXT:    s_waitcnt vmcnt(0)
421; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
422; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
423; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
424; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
425; GCN-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:96
426; GCN-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:80
427; GCN-NEXT:    s_waitcnt vmcnt(0)
428; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
429; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
430; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
431; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
432; GCN-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:80
433; GCN-NEXT:    global_load_dwordx4 v[4:7], v16, s[0:1] offset:48
434; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
435; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
436; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
437; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
438; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
439; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
440; GCN-NEXT:    s_waitcnt vmcnt(0)
441; GCN-NEXT:    v_mul_lo_u32 v7, v7, v7
442; GCN-NEXT:    v_mul_lo_u32 v6, v6, v6
443; GCN-NEXT:    v_mul_lo_u32 v5, v5, v5
444; GCN-NEXT:    v_mul_lo_u32 v4, v4, v4
445; GCN-NEXT:    global_store_dwordx4 v16, v[4:7], s[2:3] offset:48
446; GCN-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:16
447; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
448; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
449; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
450; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
451; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
452; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
453; GCN-NEXT:    s_waitcnt vmcnt(0)
454; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
455; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
456; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
457; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
458; GCN-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:16
459; GCN-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:64
460; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
461; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
462; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
463; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
464; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
465; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
466; GCN-NEXT:    s_waitcnt vmcnt(0)
467; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
468; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
469; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
470; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
471; GCN-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:64
472; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
473; GCN-NEXT:    s_endpgm
474;
475; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU_WRITE:
476; EXACTCUTOFF:       ; %bb.0:
477; EXACTCUTOFF-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
478; EXACTCUTOFF-NEXT:    v_and_b32_e32 v0, 0x3ff, v0
479; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v16, 7, v0
480; EXACTCUTOFF-NEXT:    ; kill: killed $sgpr0_sgpr1
481; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
482; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[12:15], v16, s[0:1] offset:32
483; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
484; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
485; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
486; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v13, v13, v13
487; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v12, v12, v12
488; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v15, v15, v15
489; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v14, v14, v14
490; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[12:15], s[2:3] offset:32
491; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1]
492; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
493; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
494; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
495; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
496; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
497; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
498; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
499; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3]
500; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:112
501; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
502; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
503; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
504; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
505; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
506; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
507; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:112
508; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:96
509; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
510; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
511; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
512; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
513; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
514; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:96
515; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:80
516; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
517; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
518; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
519; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
520; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
521; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:80
522; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[4:7], v16, s[0:1] offset:48
523; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
524; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
525; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
526; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
527; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
528; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
529; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
530; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v7, v7, v7
531; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v6, v6, v6
532; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v5, v5, v5
533; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v4, v4, v4
534; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[4:7], s[2:3] offset:48
535; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:16
536; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
537; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
538; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
539; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
540; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
541; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
542; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
543; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
544; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
545; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
546; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
547; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:16
548; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:64
549; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
550; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
551; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
552; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
553; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
554; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
555; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
556; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
557; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
558; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
559; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
560; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:64
561; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
562; EXACTCUTOFF-NEXT:    s_endpgm
563  %tid = call i32 @llvm.amdgcn.workitem.id.x() #2
564  %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid
565  %load = load <32 x i32>, ptr addrspace(1) %gep1
566  %mul = mul <32 x i32> %load, %load
567  %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid
568  store <32 x i32> %mul, ptr addrspace(1) %gep2
569  ; 1 VMEM read
570  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
571  ; 2 VALU
572  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
573  ; 1 VMEM write
574  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
575  ; 1 VMEM read
576  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
577  ; 2 VALU
578  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
579  ; 1 VMEM write
580  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
581  ; 1 VMEM read
582  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
583  ; 2 VALU
584  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
585  ; 1 VMEM write
586  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
587  ; 1 VMEM read
588  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
589  ; 2 VALU
590  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
591  ; 1 VMEM write
592  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
593  ; 1 VMEM read
594  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
595  ; 2 VALU
596  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
597  ; 1 VMEM write
598  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
599  ; 1 VMEM read
600  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
601  ; 2 VALU
602  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
603  ; 1 VMEM write
604  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
605  ; 1 VMEM read
606  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
607  ; 2 VALU
608  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
609  ; 1 VMEM write
610  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
611  ; 1 VMEM read
612  call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
613  ; 2 VALU
614  call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
615  ; 1 VMEM write
616  call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
617  ret void
618}
619
620define amdgpu_kernel void @test_sched_group_barrier_pipeline_MFMA_cluster(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {
621; GCN-LABEL: test_sched_group_barrier_pipeline_MFMA_cluster:
622; GCN:       ; %bb.0: ; %entry
623; GCN-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x24
624; GCN-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
625; GCN-NEXT:    v_and_b32_e32 v0, 0x1ff80, v0
626; GCN-NEXT:    s_waitcnt lgkmcnt(0)
627; GCN-NEXT:    v_add_u32_e32 v1, s0, v0
628; GCN-NEXT:    ds_read_b128 a[28:31], v1 offset:112
629; GCN-NEXT:    ds_read_b128 a[24:27], v1 offset:96
630; GCN-NEXT:    ds_read_b128 a[20:23], v1 offset:80
631; GCN-NEXT:    ds_read_b128 a[16:19], v1 offset:64
632; GCN-NEXT:    ds_read_b128 a[0:3], v1
633; GCN-NEXT:    ds_read_b128 a[4:7], v1 offset:16
634; GCN-NEXT:    ds_read_b128 a[8:11], v1 offset:32
635; GCN-NEXT:    ds_read_b128 a[12:15], v1 offset:48
636; GCN-NEXT:    ds_read_b128 a[60:63], v1 offset:8304
637; GCN-NEXT:    ds_read_b128 a[56:59], v1 offset:8288
638; GCN-NEXT:    ds_read_b128 a[52:55], v1 offset:8272
639; GCN-NEXT:    ds_read_b128 a[48:51], v1 offset:8256
640; GCN-NEXT:    ds_read_b128 a[44:47], v1 offset:8240
641; GCN-NEXT:    ds_read_b128 a[40:43], v1 offset:8224
642; GCN-NEXT:    ds_read_b128 a[36:39], v1 offset:8208
643; GCN-NEXT:    ds_read_b128 a[32:35], v1 offset:8192
644; GCN-NEXT:    v_add_u32_e32 v2, 0x6000, v1
645; GCN-NEXT:    ds_read_b128 a[92:95], v1 offset:24688
646; GCN-NEXT:    ds_read_b128 a[88:91], v1 offset:24672
647; GCN-NEXT:    ds_read_b128 a[84:87], v1 offset:24656
648; GCN-NEXT:    ds_read_b128 a[80:83], v1 offset:24640
649; GCN-NEXT:    ds_read_b128 a[76:79], v1 offset:24624
650; GCN-NEXT:    ds_read_b128 a[72:75], v1 offset:24608
651; GCN-NEXT:    ds_read_b128 a[68:71], v1 offset:24592
652; GCN-NEXT:    ds_read_b128 a[64:67], v1 offset:24576
653; GCN-NEXT:    ds_read_b128 a[124:127], v1 offset:49264
654; GCN-NEXT:    ds_read_b128 a[120:123], v1 offset:49248
655; GCN-NEXT:    ds_read_b128 a[116:119], v1 offset:49232
656; GCN-NEXT:    ds_read_b128 a[112:115], v1 offset:49216
657; GCN-NEXT:    ds_read_b128 a[108:111], v1 offset:49200
658; GCN-NEXT:    ds_read_b128 a[104:107], v1 offset:49184
659; GCN-NEXT:    ds_read_b128 a[100:103], v1 offset:49168
660; GCN-NEXT:    ds_read_b128 a[96:99], v1 offset:49152
661; GCN-NEXT:    v_mov_b32_e32 v1, 1.0
662; GCN-NEXT:    ds_read_b128 a[156:159], v2 offset:57456
663; GCN-NEXT:    ds_read_b128 a[152:155], v2 offset:57440
664; GCN-NEXT:    ds_read_b128 a[148:151], v2 offset:57424
665; GCN-NEXT:    ds_read_b128 a[144:147], v2 offset:57408
666; GCN-NEXT:    ds_read_b128 a[128:131], v2 offset:57344
667; GCN-NEXT:    ds_read_b128 a[132:135], v2 offset:57360
668; GCN-NEXT:    ds_read_b128 a[136:139], v2 offset:57376
669; GCN-NEXT:    ds_read_b128 a[140:143], v2 offset:57392
670; GCN-NEXT:    v_mov_b32_e32 v2, 2.0
671; GCN-NEXT:    v_add_u32_e32 v0, s1, v0
672; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(40) SyncID(0)
673; GCN-NEXT:    s_waitcnt lgkmcnt(14)
674; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]
675; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v1, v2, a[32:63]
676; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[64:95], v1, v2, a[64:95]
677; GCN-NEXT:    s_waitcnt lgkmcnt(8)
678; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[96:127], v1, v2, a[96:127]
679; GCN-NEXT:    s_waitcnt lgkmcnt(0)
680; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[128:159], v1, v2, a[128:159]
681; GCN-NEXT:    s_nop 7
682; GCN-NEXT:    s_nop 4
683; GCN-NEXT:    ds_write_b128 v0, a[28:31] offset:112
684; GCN-NEXT:    ds_write_b128 v0, a[24:27] offset:96
685; GCN-NEXT:    ds_write_b128 v0, a[20:23] offset:80
686; GCN-NEXT:    ds_write_b128 v0, a[16:19] offset:64
687; GCN-NEXT:    ds_write_b128 v0, a[12:15] offset:48
688; GCN-NEXT:    ds_write_b128 v0, a[8:11] offset:32
689; GCN-NEXT:    ds_write_b128 v0, a[4:7] offset:16
690; GCN-NEXT:    ds_write_b128 v0, a[0:3]
691; GCN-NEXT:    v_mov_b32_e32 v0, s1
692; GCN-NEXT:    ds_write_b128 v0, a[56:59] offset:8288
693; GCN-NEXT:    ds_write_b128 v0, a[60:63] offset:8304
694; GCN-NEXT:    ds_write_b128 v0, a[48:51] offset:8256
695; GCN-NEXT:    ds_write_b128 v0, a[52:55] offset:8272
696; GCN-NEXT:    ds_write_b128 v0, a[40:43] offset:8224
697; GCN-NEXT:    ds_write_b128 v0, a[44:47] offset:8240
698; GCN-NEXT:    ds_write_b128 v0, a[32:35] offset:8192
699; GCN-NEXT:    ds_write_b128 v0, a[36:39] offset:8208
700; GCN-NEXT:    ds_write_b128 v0, a[88:91] offset:16480
701; GCN-NEXT:    ds_write_b128 v0, a[92:95] offset:16496
702; GCN-NEXT:    ds_write_b128 v0, a[80:83] offset:16448
703; GCN-NEXT:    ds_write_b128 v0, a[84:87] offset:16464
704; GCN-NEXT:    ds_write_b128 v0, a[72:75] offset:16416
705; GCN-NEXT:    ds_write_b128 v0, a[76:79] offset:16432
706; GCN-NEXT:    ds_write_b128 v0, a[64:67] offset:16384
707; GCN-NEXT:    ds_write_b128 v0, a[68:71] offset:16400
708; GCN-NEXT:    ds_write_b128 v0, a[120:123] offset:24672
709; GCN-NEXT:    ds_write_b128 v0, a[124:127] offset:24688
710; GCN-NEXT:    ds_write_b128 v0, a[112:115] offset:24640
711; GCN-NEXT:    ds_write_b128 v0, a[116:119] offset:24656
712; GCN-NEXT:    ds_write_b128 v0, a[104:107] offset:24608
713; GCN-NEXT:    ds_write_b128 v0, a[108:111] offset:24624
714; GCN-NEXT:    ds_write_b128 v0, a[96:99] offset:24576
715; GCN-NEXT:    ds_write_b128 v0, a[100:103] offset:24592
716; GCN-NEXT:    ds_write_b128 v0, a[152:155] offset:32864
717; GCN-NEXT:    ds_write_b128 v0, a[156:159] offset:32880
718; GCN-NEXT:    ds_write_b128 v0, a[144:147] offset:32832
719; GCN-NEXT:    ds_write_b128 v0, a[148:151] offset:32848
720; GCN-NEXT:    ds_write_b128 v0, a[136:139] offset:32800
721; GCN-NEXT:    ds_write_b128 v0, a[140:143] offset:32816
722; GCN-NEXT:    ds_write_b128 v0, a[128:131] offset:32768
723; GCN-NEXT:    ds_write_b128 v0, a[132:135] offset:32784
724; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(5) SyncID(0)
725; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(40) SyncID(0)
726; GCN-NEXT:    s_endpgm
727;
728; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_MFMA_cluster:
729; EXACTCUTOFF:       ; %bb.0: ; %entry
730; EXACTCUTOFF-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x24
731; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
732; EXACTCUTOFF-NEXT:    v_and_b32_e32 v0, 0x1ff80, v0
733; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
734; EXACTCUTOFF-NEXT:    v_add_u32_e32 v1, s0, v0
735; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v1 offset:112
736; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v1 offset:96
737; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v1 offset:80
738; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v1 offset:64
739; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v1
740; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v1 offset:16
741; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v1 offset:32
742; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v1 offset:48
743; EXACTCUTOFF-NEXT:    ds_read_b128 a[60:63], v1 offset:8304
744; EXACTCUTOFF-NEXT:    ds_read_b128 a[56:59], v1 offset:8288
745; EXACTCUTOFF-NEXT:    ds_read_b128 a[52:55], v1 offset:8272
746; EXACTCUTOFF-NEXT:    ds_read_b128 a[48:51], v1 offset:8256
747; EXACTCUTOFF-NEXT:    ds_read_b128 a[44:47], v1 offset:8240
748; EXACTCUTOFF-NEXT:    ds_read_b128 a[40:43], v1 offset:8224
749; EXACTCUTOFF-NEXT:    ds_read_b128 a[36:39], v1 offset:8208
750; EXACTCUTOFF-NEXT:    ds_read_b128 a[32:35], v1 offset:8192
751; EXACTCUTOFF-NEXT:    v_add_u32_e32 v2, 0x6000, v1
752; EXACTCUTOFF-NEXT:    ds_read_b128 a[92:95], v1 offset:24688
753; EXACTCUTOFF-NEXT:    ds_read_b128 a[88:91], v1 offset:24672
754; EXACTCUTOFF-NEXT:    ds_read_b128 a[84:87], v1 offset:24656
755; EXACTCUTOFF-NEXT:    ds_read_b128 a[80:83], v1 offset:24640
756; EXACTCUTOFF-NEXT:    ds_read_b128 a[76:79], v1 offset:24624
757; EXACTCUTOFF-NEXT:    ds_read_b128 a[72:75], v1 offset:24608
758; EXACTCUTOFF-NEXT:    ds_read_b128 a[68:71], v1 offset:24592
759; EXACTCUTOFF-NEXT:    ds_read_b128 a[64:67], v1 offset:24576
760; EXACTCUTOFF-NEXT:    ds_read_b128 a[124:127], v1 offset:49264
761; EXACTCUTOFF-NEXT:    ds_read_b128 a[120:123], v1 offset:49248
762; EXACTCUTOFF-NEXT:    ds_read_b128 a[116:119], v1 offset:49232
763; EXACTCUTOFF-NEXT:    ds_read_b128 a[112:115], v1 offset:49216
764; EXACTCUTOFF-NEXT:    ds_read_b128 a[108:111], v1 offset:49200
765; EXACTCUTOFF-NEXT:    ds_read_b128 a[104:107], v1 offset:49184
766; EXACTCUTOFF-NEXT:    ds_read_b128 a[100:103], v1 offset:49168
767; EXACTCUTOFF-NEXT:    ds_read_b128 a[96:99], v1 offset:49152
768; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v1, 1.0
769; EXACTCUTOFF-NEXT:    ds_read_b128 a[156:159], v2 offset:57456
770; EXACTCUTOFF-NEXT:    ds_read_b128 a[152:155], v2 offset:57440
771; EXACTCUTOFF-NEXT:    ds_read_b128 a[148:151], v2 offset:57424
772; EXACTCUTOFF-NEXT:    ds_read_b128 a[144:147], v2 offset:57408
773; EXACTCUTOFF-NEXT:    ds_read_b128 a[128:131], v2 offset:57344
774; EXACTCUTOFF-NEXT:    ds_read_b128 a[132:135], v2 offset:57360
775; EXACTCUTOFF-NEXT:    ds_read_b128 a[136:139], v2 offset:57376
776; EXACTCUTOFF-NEXT:    ds_read_b128 a[140:143], v2 offset:57392
777; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v2, 2.0
778; EXACTCUTOFF-NEXT:    v_add_u32_e32 v0, s1, v0
779; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(40) SyncID(0)
780; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(14)
781; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]
782; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v1, v2, a[32:63]
783; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[64:95], v1, v2, a[64:95]
784; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(8)
785; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[96:127], v1, v2, a[96:127]
786; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
787; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[128:159], v1, v2, a[128:159]
788; EXACTCUTOFF-NEXT:    s_nop 7
789; EXACTCUTOFF-NEXT:    s_nop 4
790; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[28:31] offset:112
791; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[24:27] offset:96
792; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[20:23] offset:80
793; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[16:19] offset:64
794; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[12:15] offset:48
795; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[8:11] offset:32
796; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[4:7] offset:16
797; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[0:3]
798; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v0, s1
799; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[56:59] offset:8288
800; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[60:63] offset:8304
801; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[48:51] offset:8256
802; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[52:55] offset:8272
803; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[40:43] offset:8224
804; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[44:47] offset:8240
805; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[32:35] offset:8192
806; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[36:39] offset:8208
807; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[88:91] offset:16480
808; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[92:95] offset:16496
809; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[80:83] offset:16448
810; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[84:87] offset:16464
811; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[72:75] offset:16416
812; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[76:79] offset:16432
813; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[64:67] offset:16384
814; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[68:71] offset:16400
815; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[120:123] offset:24672
816; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[124:127] offset:24688
817; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[112:115] offset:24640
818; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[116:119] offset:24656
819; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[104:107] offset:24608
820; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[108:111] offset:24624
821; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[96:99] offset:24576
822; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[100:103] offset:24592
823; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[152:155] offset:32864
824; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[156:159] offset:32880
825; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[144:147] offset:32832
826; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[148:151] offset:32848
827; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[136:139] offset:32800
828; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[140:143] offset:32816
829; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[128:131] offset:32768
830; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[132:135] offset:32784
831; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(5) SyncID(0)
832; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(40) SyncID(0)
833; EXACTCUTOFF-NEXT:    s_endpgm
834entry:
835  %idx = call i32 @llvm.amdgcn.workitem.id.x()
836  %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx
837  %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr
838  %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64
839  %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr
840  %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128
841  %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr
842  %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192
843  %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr
844  %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256
845  %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr
846  %mai.0 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.0, i32 0, i32 0, i32 0)
847  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.1, i32 0, i32 0, i32 0)
848  %mai.2 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.2, i32 0, i32 0, i32 0)
849  %mai.3 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.3, i32 0, i32 0, i32 0)
850  %mai.4 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.4, i32 0, i32 0, i32 0)
851  %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx
852  store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr
853  %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64
854  store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr
855  %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128
856  store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr
857  %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192
858  store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr
859  %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256
860  store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr
861  ; 40 DS read
862  call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 40, i32 0)
863  ; 5 MFMA
864  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 5, i32 0)
865  ; 40 DS write
866  call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 40, i32 0)
867  ret void
868}
869
870define amdgpu_kernel void @test_sched_group_barrier_pipeline_MFMA_interleave(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {
871; GCN-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:
872; GCN:       ; %bb.0: ; %entry
873; GCN-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x24
874; GCN-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
875; GCN-NEXT:    v_and_b32_e32 v1, 0x1ff80, v0
876; GCN-NEXT:    v_mov_b32_e32 v2, 1.0
877; GCN-NEXT:    v_mov_b32_e32 v3, 2.0
878; GCN-NEXT:    s_waitcnt lgkmcnt(0)
879; GCN-NEXT:    v_add_u32_e32 v0, s0, v1
880; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:112
881; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:96
882; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:80
883; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:64
884; GCN-NEXT:    ds_read_b128 a[0:3], v0
885; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:16
886; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:32
887; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:48
888; GCN-NEXT:    s_waitcnt lgkmcnt(0)
889; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
890; GCN-NEXT:    v_add_u32_e32 v1, s1, v1
891; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
892; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
893; GCN-NEXT:    s_nop 7
894; GCN-NEXT:    s_nop 7
895; GCN-NEXT:    s_nop 1
896; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:112
897; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:96
898; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:80
899; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:64
900; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:48
901; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:32
902; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:16
903; GCN-NEXT:    ds_write_b128 v1, a[0:3]
904; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:8304
905; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:8288
906; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:8272
907; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:8256
908; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:8240
909; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:8224
910; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:8208
911; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:8192
912; GCN-NEXT:    s_waitcnt lgkmcnt(0)
913; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
914; GCN-NEXT:    v_mov_b32_e32 v1, s1
915; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
916; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
917; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
918; GCN-NEXT:    s_nop 7
919; GCN-NEXT:    s_nop 7
920; GCN-NEXT:    s_nop 1
921; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:8288
922; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:8304
923; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:8256
924; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:8272
925; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:8224
926; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:8240
927; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:8192
928; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:8208
929; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:24688
930; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:24672
931; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:24656
932; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:24640
933; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:24624
934; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:24608
935; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:24592
936; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:24576
937; GCN-NEXT:    s_waitcnt lgkmcnt(0)
938; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
939; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
940; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
941; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
942; GCN-NEXT:    s_nop 7
943; GCN-NEXT:    s_nop 7
944; GCN-NEXT:    s_nop 2
945; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:16480
946; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:16496
947; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:16448
948; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:16464
949; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:16416
950; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:16432
951; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:16384
952; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:16400
953; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:49264
954; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:49248
955; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:49232
956; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:49216
957; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:49200
958; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:49184
959; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:49168
960; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:49152
961; GCN-NEXT:    s_waitcnt lgkmcnt(0)
962; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
963; GCN-NEXT:    v_add_u32_e32 v0, 0x6000, v0
964; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
965; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
966; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
967; GCN-NEXT:    s_nop 7
968; GCN-NEXT:    s_nop 7
969; GCN-NEXT:    s_nop 1
970; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:24672
971; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:24688
972; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:24640
973; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:24656
974; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:24608
975; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:24624
976; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:24576
977; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:24592
978; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:57456
979; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:57440
980; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:57424
981; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:57408
982; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:57344
983; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:57360
984; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:57376
985; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:57392
986; GCN-NEXT:    s_waitcnt lgkmcnt(0)
987; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
988; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
989; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
990; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
991; GCN-NEXT:    s_nop 7
992; GCN-NEXT:    s_nop 7
993; GCN-NEXT:    s_nop 2
994; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:32864
995; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:32880
996; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:32832
997; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:32848
998; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:32800
999; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:32816
1000; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:32768
1001; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:32784
1002; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1003; GCN-NEXT:    s_endpgm
1004;
1005; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:
1006; EXACTCUTOFF:       ; %bb.0: ; %entry
1007; EXACTCUTOFF-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x24
1008; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
1009; EXACTCUTOFF-NEXT:    v_and_b32_e32 v1, 0x1ff80, v0
1010; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v2, 1.0
1011; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v3, 2.0
1012; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1013; EXACTCUTOFF-NEXT:    v_add_u32_e32 v0, s0, v1
1014; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:112
1015; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:96
1016; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:80
1017; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:64
1018; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0
1019; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:16
1020; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:32
1021; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:48
1022; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1023; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1024; EXACTCUTOFF-NEXT:    v_add_u32_e32 v1, s1, v1
1025; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1026; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1027; EXACTCUTOFF-NEXT:    s_nop 7
1028; EXACTCUTOFF-NEXT:    s_nop 7
1029; EXACTCUTOFF-NEXT:    s_nop 1
1030; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:112
1031; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:96
1032; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:80
1033; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:64
1034; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:48
1035; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:32
1036; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:16
1037; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3]
1038; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:8304
1039; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:8288
1040; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:8272
1041; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:8256
1042; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:8240
1043; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:8224
1044; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:8208
1045; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:8192
1046; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1047; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1048; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v1, s1
1049; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1050; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1051; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1052; EXACTCUTOFF-NEXT:    s_nop 7
1053; EXACTCUTOFF-NEXT:    s_nop 7
1054; EXACTCUTOFF-NEXT:    s_nop 1
1055; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:8288
1056; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:8304
1057; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:8256
1058; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:8272
1059; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:8224
1060; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:8240
1061; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:8192
1062; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:8208
1063; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:24688
1064; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:24672
1065; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:24656
1066; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:24640
1067; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:24624
1068; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:24608
1069; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:24592
1070; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:24576
1071; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1072; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1073; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1074; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1075; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1076; EXACTCUTOFF-NEXT:    s_nop 7
1077; EXACTCUTOFF-NEXT:    s_nop 7
1078; EXACTCUTOFF-NEXT:    s_nop 2
1079; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:16480
1080; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:16496
1081; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:16448
1082; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:16464
1083; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:16416
1084; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:16432
1085; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:16384
1086; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:16400
1087; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:49264
1088; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:49248
1089; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:49232
1090; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:49216
1091; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:49200
1092; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:49184
1093; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:49168
1094; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:49152
1095; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1096; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1097; EXACTCUTOFF-NEXT:    v_add_u32_e32 v0, 0x6000, v0
1098; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1099; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1100; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1101; EXACTCUTOFF-NEXT:    s_nop 7
1102; EXACTCUTOFF-NEXT:    s_nop 7
1103; EXACTCUTOFF-NEXT:    s_nop 1
1104; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:24672
1105; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:24688
1106; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:24640
1107; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:24656
1108; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:24608
1109; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:24624
1110; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:24576
1111; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:24592
1112; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:57456
1113; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:57440
1114; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:57424
1115; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:57408
1116; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:57344
1117; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:57360
1118; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:57376
1119; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:57392
1120; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1121; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1122; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1123; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1124; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1125; EXACTCUTOFF-NEXT:    s_nop 7
1126; EXACTCUTOFF-NEXT:    s_nop 7
1127; EXACTCUTOFF-NEXT:    s_nop 2
1128; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:32864
1129; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:32880
1130; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:32832
1131; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:32848
1132; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:32800
1133; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:32816
1134; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:32768
1135; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:32784
1136; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1137; EXACTCUTOFF-NEXT:    s_endpgm
1138entry:
1139  %idx = call i32 @llvm.amdgcn.workitem.id.x()
1140  %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx
1141  %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr
1142  %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64
1143  %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr
1144  %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128
1145  %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr
1146  %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192
1147  %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr
1148  %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256
1149  %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr
1150  %mai.0 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.0, i32 0, i32 0, i32 0)
1151  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.1, i32 0, i32 0, i32 0)
1152  %mai.2 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.2, i32 0, i32 0, i32 0)
1153  %mai.3 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.3, i32 0, i32 0, i32 0)
1154  %mai.4 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %load.4, i32 0, i32 0, i32 0)
1155  %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx
1156  store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr
1157  %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64
1158  store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr
1159  %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128
1160  store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr
1161  %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192
1162  store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr
1163  %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256
1164  store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr
1165  ; 8 DS read
1166  call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1167  ; 1 MFMA
1168  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1169  ; 8 DS write
1170  call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1171  ; 8 DS read
1172  call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1173  ; 1 MFMA
1174  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1175  ; 8 DS write
1176  call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1177  ; 8 DS read
1178  call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1179  ; 1 MFMA
1180  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1181  ; 8 DS write
1182  call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1183  ; 8 DS read
1184  call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1185  ; 1 MFMA
1186  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1187  ; 8 DS write
1188  call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1189  ; 8 DS read
1190  call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1191  ; 1 MFMA
1192  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1193  ; 8 DS write
1194  call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1195  ret void
1196}
1197
1198define amdgpu_kernel void @test_sched_group_barrier_pipeline_interleave_EXP_MFMA(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out, <5 x float> %in1) #0 {
1199; GCN-LABEL: test_sched_group_barrier_pipeline_interleave_EXP_MFMA:
1200; GCN:       ; %bb.0: ; %entry
1201; GCN-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x44
1202; GCN-NEXT:    v_mov_b32_e32 v3, 0x3fb8aa3b
1203; GCN-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
1204; GCN-NEXT:    v_mov_b32_e32 v7, 0x32a5705f
1205; GCN-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
1206; GCN-NEXT:    s_waitcnt lgkmcnt(0)
1207; GCN-NEXT:    v_mul_f32_e32 v4, s0, v3
1208; GCN-NEXT:    v_rndne_f32_e32 v5, v4
1209; GCN-NEXT:    v_sub_f32_e32 v6, v4, v5
1210; GCN-NEXT:    v_fma_f32 v4, s0, v3, -v4
1211; GCN-NEXT:    v_fmac_f32_e32 v4, s0, v7
1212; GCN-NEXT:    v_add_f32_e32 v4, v6, v4
1213; GCN-NEXT:    v_exp_f32_e32 v4, v4
1214; GCN-NEXT:    v_cvt_i32_f32_e32 v5, v5
1215; GCN-NEXT:    v_and_b32_e32 v0, 0x1ff80, v0
1216; GCN-NEXT:    v_add_u32_e32 v1, s6, v0
1217; GCN-NEXT:    ds_read_b128 a[28:31], v1 offset:112
1218; GCN-NEXT:    ds_read_b128 a[24:27], v1 offset:96
1219; GCN-NEXT:    ds_read_b128 a[20:23], v1 offset:80
1220; GCN-NEXT:    ds_read_b128 a[16:19], v1 offset:64
1221; GCN-NEXT:    ds_read_b128 a[0:3], v1
1222; GCN-NEXT:    ds_read_b128 a[4:7], v1 offset:16
1223; GCN-NEXT:    ds_read_b128 a[8:11], v1 offset:32
1224; GCN-NEXT:    ds_read_b128 a[12:15], v1 offset:48
1225; GCN-NEXT:    v_mov_b32_e32 v9, 1.0
1226; GCN-NEXT:    v_ldexp_f32 v4, v4, v5
1227; GCN-NEXT:    v_mov_b32_e32 v5, 0xc2ce8ed0
1228; GCN-NEXT:    v_mul_f32_e32 v10, s1, v3
1229; GCN-NEXT:    v_cmp_nlt_f32_e32 vcc, s0, v5
1230; GCN-NEXT:    v_mov_b32_e32 v6, 0x42b17218
1231; GCN-NEXT:    v_rndne_f32_e32 v11, v10
1232; GCN-NEXT:    v_cndmask_b32_e32 v4, 0, v4, vcc
1233; GCN-NEXT:    v_cmp_ngt_f32_e32 vcc, s0, v6
1234; GCN-NEXT:    v_mov_b32_e32 v8, 0x7f800000
1235; GCN-NEXT:    v_sub_f32_e32 v12, v10, v11
1236; GCN-NEXT:    v_fma_f32 v10, s1, v3, -v10
1237; GCN-NEXT:    v_cndmask_b32_e32 v4, v8, v4, vcc
1238; GCN-NEXT:    v_fmac_f32_e32 v10, s1, v7
1239; GCN-NEXT:    ds_read_b128 a[60:63], v1 offset:8304
1240; GCN-NEXT:    s_waitcnt lgkmcnt(1)
1241; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v9, v4, a[0:31]
1242; GCN-NEXT:    v_add_f32_e32 v4, v12, v10
1243; GCN-NEXT:    v_exp_f32_e32 v4, v4
1244; GCN-NEXT:    v_cvt_i32_f32_e32 v10, v11
1245; GCN-NEXT:    ds_read_b128 a[56:59], v1 offset:8288
1246; GCN-NEXT:    ds_read_b128 a[52:55], v1 offset:8272
1247; GCN-NEXT:    ds_read_b128 a[48:51], v1 offset:8256
1248; GCN-NEXT:    ds_read_b128 a[44:47], v1 offset:8240
1249; GCN-NEXT:    ds_read_b128 a[40:43], v1 offset:8224
1250; GCN-NEXT:    ds_read_b128 a[36:39], v1 offset:8208
1251; GCN-NEXT:    ds_read_b128 a[32:35], v1 offset:8192
1252; GCN-NEXT:    v_ldexp_f32 v4, v4, v10
1253; GCN-NEXT:    v_cmp_nlt_f32_e32 vcc, s1, v5
1254; GCN-NEXT:    v_cndmask_b32_e32 v4, 0, v4, vcc
1255; GCN-NEXT:    v_cmp_ngt_f32_e32 vcc, s1, v6
1256; GCN-NEXT:    v_cndmask_b32_e32 v4, v8, v4, vcc
1257; GCN-NEXT:    v_mul_f32_e32 v10, s2, v3
1258; GCN-NEXT:    v_rndne_f32_e32 v11, v10
1259; GCN-NEXT:    s_waitcnt lgkmcnt(0)
1260; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v9, v4, a[32:63]
1261; GCN-NEXT:    v_fma_f32 v4, s2, v3, -v10
1262; GCN-NEXT:    v_sub_f32_e32 v12, v10, v11
1263; GCN-NEXT:    v_fmac_f32_e32 v4, s2, v7
1264; GCN-NEXT:    v_add_f32_e32 v4, v12, v4
1265; GCN-NEXT:    v_exp_f32_e32 v4, v4
1266; GCN-NEXT:    v_cvt_i32_f32_e32 v10, v11
1267; GCN-NEXT:    ds_read_b128 a[92:95], v1 offset:24688
1268; GCN-NEXT:    ds_read_b128 a[88:91], v1 offset:24672
1269; GCN-NEXT:    ds_read_b128 a[84:87], v1 offset:24656
1270; GCN-NEXT:    ds_read_b128 a[80:83], v1 offset:24640
1271; GCN-NEXT:    ds_read_b128 a[76:79], v1 offset:24624
1272; GCN-NEXT:    ds_read_b128 a[72:75], v1 offset:24608
1273; GCN-NEXT:    ds_read_b128 a[68:71], v1 offset:24592
1274; GCN-NEXT:    ds_read_b128 a[64:67], v1 offset:24576
1275; GCN-NEXT:    v_add_u32_e32 v2, 0x6000, v1
1276; GCN-NEXT:    ds_read_b128 a[124:127], v1 offset:49264
1277; GCN-NEXT:    ds_read_b128 a[120:123], v1 offset:49248
1278; GCN-NEXT:    ds_read_b128 a[116:119], v1 offset:49232
1279; GCN-NEXT:    ds_read_b128 a[112:115], v1 offset:49216
1280; GCN-NEXT:    ds_read_b128 a[108:111], v1 offset:49200
1281; GCN-NEXT:    ds_read_b128 a[104:107], v1 offset:49184
1282; GCN-NEXT:    ds_read_b128 a[100:103], v1 offset:49168
1283; GCN-NEXT:    ds_read_b128 a[96:99], v1 offset:49152
1284; GCN-NEXT:    v_ldexp_f32 v1, v4, v10
1285; GCN-NEXT:    v_cmp_nlt_f32_e32 vcc, s2, v5
1286; GCN-NEXT:    v_cndmask_b32_e32 v1, 0, v1, vcc
1287; GCN-NEXT:    v_cmp_ngt_f32_e32 vcc, s2, v6
1288; GCN-NEXT:    v_mul_f32_e32 v4, s3, v3
1289; GCN-NEXT:    v_cndmask_b32_e32 v1, v8, v1, vcc
1290; GCN-NEXT:    v_rndne_f32_e32 v10, v4
1291; GCN-NEXT:    s_load_dword s8, s[4:5], 0x54
1292; GCN-NEXT:    s_waitcnt lgkmcnt(0)
1293; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[64:95], v9, v1, a[64:95]
1294; GCN-NEXT:    v_sub_f32_e32 v1, v4, v10
1295; GCN-NEXT:    v_fma_f32 v4, s3, v3, -v4
1296; GCN-NEXT:    v_fmac_f32_e32 v4, s3, v7
1297; GCN-NEXT:    v_add_f32_e32 v1, v1, v4
1298; GCN-NEXT:    v_exp_f32_e32 v1, v1
1299; GCN-NEXT:    v_cvt_i32_f32_e32 v4, v10
1300; GCN-NEXT:    v_cmp_nlt_f32_e32 vcc, s3, v5
1301; GCN-NEXT:    ds_read_b128 a[156:159], v2 offset:57456
1302; GCN-NEXT:    ds_read_b128 a[152:155], v2 offset:57440
1303; GCN-NEXT:    v_ldexp_f32 v1, v1, v4
1304; GCN-NEXT:    v_cndmask_b32_e32 v1, 0, v1, vcc
1305; GCN-NEXT:    v_cmp_ngt_f32_e32 vcc, s3, v6
1306; GCN-NEXT:    v_cndmask_b32_e32 v1, v8, v1, vcc
1307; GCN-NEXT:    v_mul_f32_e32 v4, s8, v3
1308; GCN-NEXT:    v_fma_f32 v3, s8, v3, -v4
1309; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[96:127], v9, v1, a[96:127]
1310; GCN-NEXT:    v_rndne_f32_e32 v1, v4
1311; GCN-NEXT:    v_sub_f32_e32 v10, v4, v1
1312; GCN-NEXT:    v_fmac_f32_e32 v3, s8, v7
1313; GCN-NEXT:    v_add_f32_e32 v3, v10, v3
1314; GCN-NEXT:    v_exp_f32_e32 v3, v3
1315; GCN-NEXT:    v_cvt_i32_f32_e32 v1, v1
1316; GCN-NEXT:    ds_read_b128 a[148:151], v2 offset:57424
1317; GCN-NEXT:    ds_read_b128 a[144:147], v2 offset:57408
1318; GCN-NEXT:    ds_read_b128 a[128:131], v2 offset:57344
1319; GCN-NEXT:    ds_read_b128 a[132:135], v2 offset:57360
1320; GCN-NEXT:    ds_read_b128 a[136:139], v2 offset:57376
1321; GCN-NEXT:    ds_read_b128 a[140:143], v2 offset:57392
1322; GCN-NEXT:    v_ldexp_f32 v1, v3, v1
1323; GCN-NEXT:    v_cmp_nlt_f32_e32 vcc, s8, v5
1324; GCN-NEXT:    v_cndmask_b32_e32 v1, 0, v1, vcc
1325; GCN-NEXT:    v_cmp_ngt_f32_e32 vcc, s8, v6
1326; GCN-NEXT:    v_cndmask_b32_e32 v1, v8, v1, vcc
1327; GCN-NEXT:    v_add_u32_e32 v0, s7, v0
1328; GCN-NEXT:    ds_write_b128 v0, a[28:31] offset:112
1329; GCN-NEXT:    s_waitcnt lgkmcnt(1)
1330; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[128:159], v9, v1, a[128:159]
1331; GCN-NEXT:    ds_write_b128 v0, a[24:27] offset:96
1332; GCN-NEXT:    ds_write_b128 v0, a[20:23] offset:80
1333; GCN-NEXT:    ds_write_b128 v0, a[16:19] offset:64
1334; GCN-NEXT:    ds_write_b128 v0, a[12:15] offset:48
1335; GCN-NEXT:    ds_write_b128 v0, a[8:11] offset:32
1336; GCN-NEXT:    ds_write_b128 v0, a[4:7] offset:16
1337; GCN-NEXT:    ds_write_b128 v0, a[0:3]
1338; GCN-NEXT:    v_mov_b32_e32 v0, s7
1339; GCN-NEXT:    ; kill: killed $sgpr4_sgpr5
1340; GCN-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1341; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1342; GCN-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1343; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1344; GCN-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1345; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1346; GCN-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1347; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1348; GCN-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1349; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1350; GCN-NEXT:    ds_write_b128 v0, a[56:59] offset:8288
1351; GCN-NEXT:    ds_write_b128 v0, a[60:63] offset:8304
1352; GCN-NEXT:    ds_write_b128 v0, a[48:51] offset:8256
1353; GCN-NEXT:    ds_write_b128 v0, a[52:55] offset:8272
1354; GCN-NEXT:    ds_write_b128 v0, a[40:43] offset:8224
1355; GCN-NEXT:    ds_write_b128 v0, a[44:47] offset:8240
1356; GCN-NEXT:    ds_write_b128 v0, a[32:35] offset:8192
1357; GCN-NEXT:    ds_write_b128 v0, a[36:39] offset:8208
1358; GCN-NEXT:    ds_write_b128 v0, a[88:91] offset:16480
1359; GCN-NEXT:    ds_write_b128 v0, a[92:95] offset:16496
1360; GCN-NEXT:    ds_write_b128 v0, a[80:83] offset:16448
1361; GCN-NEXT:    ds_write_b128 v0, a[84:87] offset:16464
1362; GCN-NEXT:    ds_write_b128 v0, a[72:75] offset:16416
1363; GCN-NEXT:    ds_write_b128 v0, a[76:79] offset:16432
1364; GCN-NEXT:    ds_write_b128 v0, a[64:67] offset:16384
1365; GCN-NEXT:    ds_write_b128 v0, a[68:71] offset:16400
1366; GCN-NEXT:    ds_write_b128 v0, a[120:123] offset:24672
1367; GCN-NEXT:    ds_write_b128 v0, a[124:127] offset:24688
1368; GCN-NEXT:    ds_write_b128 v0, a[112:115] offset:24640
1369; GCN-NEXT:    ds_write_b128 v0, a[116:119] offset:24656
1370; GCN-NEXT:    ds_write_b128 v0, a[104:107] offset:24608
1371; GCN-NEXT:    ds_write_b128 v0, a[108:111] offset:24624
1372; GCN-NEXT:    ds_write_b128 v0, a[96:99] offset:24576
1373; GCN-NEXT:    ds_write_b128 v0, a[100:103] offset:24592
1374; GCN-NEXT:    ds_write_b128 v0, a[152:155] offset:32864
1375; GCN-NEXT:    ds_write_b128 v0, a[156:159] offset:32880
1376; GCN-NEXT:    ds_write_b128 v0, a[144:147] offset:32832
1377; GCN-NEXT:    ds_write_b128 v0, a[148:151] offset:32848
1378; GCN-NEXT:    ds_write_b128 v0, a[136:139] offset:32800
1379; GCN-NEXT:    ds_write_b128 v0, a[140:143] offset:32816
1380; GCN-NEXT:    ds_write_b128 v0, a[128:131] offset:32768
1381; GCN-NEXT:    ds_write_b128 v0, a[132:135] offset:32784
1382; GCN-NEXT:    s_endpgm
1383;
1384; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_interleave_EXP_MFMA:
1385; EXACTCUTOFF:       ; %bb.0: ; %entry
1386; EXACTCUTOFF-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x44
1387; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v3, 0x3fb8aa3b
1388; EXACTCUTOFF-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
1389; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v7, 0x32a5705f
1390; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
1391; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1392; EXACTCUTOFF-NEXT:    v_mul_f32_e32 v4, s0, v3
1393; EXACTCUTOFF-NEXT:    v_rndne_f32_e32 v5, v4
1394; EXACTCUTOFF-NEXT:    v_sub_f32_e32 v6, v4, v5
1395; EXACTCUTOFF-NEXT:    v_fma_f32 v4, s0, v3, -v4
1396; EXACTCUTOFF-NEXT:    v_fmac_f32_e32 v4, s0, v7
1397; EXACTCUTOFF-NEXT:    v_add_f32_e32 v4, v6, v4
1398; EXACTCUTOFF-NEXT:    v_exp_f32_e32 v4, v4
1399; EXACTCUTOFF-NEXT:    v_cvt_i32_f32_e32 v5, v5
1400; EXACTCUTOFF-NEXT:    v_and_b32_e32 v0, 0x1ff80, v0
1401; EXACTCUTOFF-NEXT:    v_add_u32_e32 v1, s6, v0
1402; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v1 offset:112
1403; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v1 offset:96
1404; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v1 offset:80
1405; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v1 offset:64
1406; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v1
1407; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v1 offset:16
1408; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v1 offset:32
1409; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v1 offset:48
1410; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v9, 1.0
1411; EXACTCUTOFF-NEXT:    v_ldexp_f32 v4, v4, v5
1412; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v5, 0xc2ce8ed0
1413; EXACTCUTOFF-NEXT:    v_mul_f32_e32 v10, s1, v3
1414; EXACTCUTOFF-NEXT:    v_cmp_nlt_f32_e32 vcc, s0, v5
1415; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v6, 0x42b17218
1416; EXACTCUTOFF-NEXT:    v_rndne_f32_e32 v11, v10
1417; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v4, 0, v4, vcc
1418; EXACTCUTOFF-NEXT:    v_cmp_ngt_f32_e32 vcc, s0, v6
1419; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v8, 0x7f800000
1420; EXACTCUTOFF-NEXT:    v_sub_f32_e32 v12, v10, v11
1421; EXACTCUTOFF-NEXT:    v_fma_f32 v10, s1, v3, -v10
1422; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v4, v8, v4, vcc
1423; EXACTCUTOFF-NEXT:    v_fmac_f32_e32 v10, s1, v7
1424; EXACTCUTOFF-NEXT:    ds_read_b128 a[60:63], v1 offset:8304
1425; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(1)
1426; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v9, v4, a[0:31]
1427; EXACTCUTOFF-NEXT:    v_add_f32_e32 v4, v12, v10
1428; EXACTCUTOFF-NEXT:    v_exp_f32_e32 v4, v4
1429; EXACTCUTOFF-NEXT:    v_cvt_i32_f32_e32 v10, v11
1430; EXACTCUTOFF-NEXT:    ds_read_b128 a[56:59], v1 offset:8288
1431; EXACTCUTOFF-NEXT:    ds_read_b128 a[52:55], v1 offset:8272
1432; EXACTCUTOFF-NEXT:    ds_read_b128 a[48:51], v1 offset:8256
1433; EXACTCUTOFF-NEXT:    ds_read_b128 a[44:47], v1 offset:8240
1434; EXACTCUTOFF-NEXT:    ds_read_b128 a[40:43], v1 offset:8224
1435; EXACTCUTOFF-NEXT:    ds_read_b128 a[36:39], v1 offset:8208
1436; EXACTCUTOFF-NEXT:    ds_read_b128 a[32:35], v1 offset:8192
1437; EXACTCUTOFF-NEXT:    v_ldexp_f32 v4, v4, v10
1438; EXACTCUTOFF-NEXT:    v_cmp_nlt_f32_e32 vcc, s1, v5
1439; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v4, 0, v4, vcc
1440; EXACTCUTOFF-NEXT:    v_cmp_ngt_f32_e32 vcc, s1, v6
1441; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v4, v8, v4, vcc
1442; EXACTCUTOFF-NEXT:    v_mul_f32_e32 v10, s2, v3
1443; EXACTCUTOFF-NEXT:    v_rndne_f32_e32 v11, v10
1444; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1445; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v9, v4, a[32:63]
1446; EXACTCUTOFF-NEXT:    v_fma_f32 v4, s2, v3, -v10
1447; EXACTCUTOFF-NEXT:    v_sub_f32_e32 v12, v10, v11
1448; EXACTCUTOFF-NEXT:    v_fmac_f32_e32 v4, s2, v7
1449; EXACTCUTOFF-NEXT:    v_add_f32_e32 v4, v12, v4
1450; EXACTCUTOFF-NEXT:    v_exp_f32_e32 v4, v4
1451; EXACTCUTOFF-NEXT:    v_cvt_i32_f32_e32 v10, v11
1452; EXACTCUTOFF-NEXT:    ds_read_b128 a[92:95], v1 offset:24688
1453; EXACTCUTOFF-NEXT:    ds_read_b128 a[88:91], v1 offset:24672
1454; EXACTCUTOFF-NEXT:    ds_read_b128 a[84:87], v1 offset:24656
1455; EXACTCUTOFF-NEXT:    ds_read_b128 a[80:83], v1 offset:24640
1456; EXACTCUTOFF-NEXT:    ds_read_b128 a[76:79], v1 offset:24624
1457; EXACTCUTOFF-NEXT:    ds_read_b128 a[72:75], v1 offset:24608
1458; EXACTCUTOFF-NEXT:    ds_read_b128 a[68:71], v1 offset:24592
1459; EXACTCUTOFF-NEXT:    ds_read_b128 a[64:67], v1 offset:24576
1460; EXACTCUTOFF-NEXT:    v_add_u32_e32 v2, 0x6000, v1
1461; EXACTCUTOFF-NEXT:    ds_read_b128 a[124:127], v1 offset:49264
1462; EXACTCUTOFF-NEXT:    ds_read_b128 a[120:123], v1 offset:49248
1463; EXACTCUTOFF-NEXT:    ds_read_b128 a[116:119], v1 offset:49232
1464; EXACTCUTOFF-NEXT:    ds_read_b128 a[112:115], v1 offset:49216
1465; EXACTCUTOFF-NEXT:    ds_read_b128 a[108:111], v1 offset:49200
1466; EXACTCUTOFF-NEXT:    ds_read_b128 a[104:107], v1 offset:49184
1467; EXACTCUTOFF-NEXT:    ds_read_b128 a[100:103], v1 offset:49168
1468; EXACTCUTOFF-NEXT:    ds_read_b128 a[96:99], v1 offset:49152
1469; EXACTCUTOFF-NEXT:    v_ldexp_f32 v1, v4, v10
1470; EXACTCUTOFF-NEXT:    v_cmp_nlt_f32_e32 vcc, s2, v5
1471; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v1, 0, v1, vcc
1472; EXACTCUTOFF-NEXT:    v_cmp_ngt_f32_e32 vcc, s2, v6
1473; EXACTCUTOFF-NEXT:    v_mul_f32_e32 v4, s3, v3
1474; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v1, v8, v1, vcc
1475; EXACTCUTOFF-NEXT:    v_rndne_f32_e32 v10, v4
1476; EXACTCUTOFF-NEXT:    s_load_dword s8, s[4:5], 0x54
1477; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1478; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[64:95], v9, v1, a[64:95]
1479; EXACTCUTOFF-NEXT:    v_sub_f32_e32 v1, v4, v10
1480; EXACTCUTOFF-NEXT:    v_fma_f32 v4, s3, v3, -v4
1481; EXACTCUTOFF-NEXT:    v_fmac_f32_e32 v4, s3, v7
1482; EXACTCUTOFF-NEXT:    v_add_f32_e32 v1, v1, v4
1483; EXACTCUTOFF-NEXT:    v_exp_f32_e32 v1, v1
1484; EXACTCUTOFF-NEXT:    v_cvt_i32_f32_e32 v4, v10
1485; EXACTCUTOFF-NEXT:    v_cmp_nlt_f32_e32 vcc, s3, v5
1486; EXACTCUTOFF-NEXT:    ds_read_b128 a[156:159], v2 offset:57456
1487; EXACTCUTOFF-NEXT:    ds_read_b128 a[152:155], v2 offset:57440
1488; EXACTCUTOFF-NEXT:    v_ldexp_f32 v1, v1, v4
1489; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v1, 0, v1, vcc
1490; EXACTCUTOFF-NEXT:    v_cmp_ngt_f32_e32 vcc, s3, v6
1491; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v1, v8, v1, vcc
1492; EXACTCUTOFF-NEXT:    v_mul_f32_e32 v4, s8, v3
1493; EXACTCUTOFF-NEXT:    v_fma_f32 v3, s8, v3, -v4
1494; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[96:127], v9, v1, a[96:127]
1495; EXACTCUTOFF-NEXT:    v_rndne_f32_e32 v1, v4
1496; EXACTCUTOFF-NEXT:    v_sub_f32_e32 v10, v4, v1
1497; EXACTCUTOFF-NEXT:    v_fmac_f32_e32 v3, s8, v7
1498; EXACTCUTOFF-NEXT:    v_add_f32_e32 v3, v10, v3
1499; EXACTCUTOFF-NEXT:    v_exp_f32_e32 v3, v3
1500; EXACTCUTOFF-NEXT:    v_cvt_i32_f32_e32 v1, v1
1501; EXACTCUTOFF-NEXT:    ds_read_b128 a[148:151], v2 offset:57424
1502; EXACTCUTOFF-NEXT:    ds_read_b128 a[144:147], v2 offset:57408
1503; EXACTCUTOFF-NEXT:    ds_read_b128 a[128:131], v2 offset:57344
1504; EXACTCUTOFF-NEXT:    ds_read_b128 a[132:135], v2 offset:57360
1505; EXACTCUTOFF-NEXT:    ds_read_b128 a[136:139], v2 offset:57376
1506; EXACTCUTOFF-NEXT:    ds_read_b128 a[140:143], v2 offset:57392
1507; EXACTCUTOFF-NEXT:    v_ldexp_f32 v1, v3, v1
1508; EXACTCUTOFF-NEXT:    v_cmp_nlt_f32_e32 vcc, s8, v5
1509; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v1, 0, v1, vcc
1510; EXACTCUTOFF-NEXT:    v_cmp_ngt_f32_e32 vcc, s8, v6
1511; EXACTCUTOFF-NEXT:    v_cndmask_b32_e32 v1, v8, v1, vcc
1512; EXACTCUTOFF-NEXT:    v_add_u32_e32 v0, s7, v0
1513; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[28:31] offset:112
1514; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(1)
1515; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[128:159], v9, v1, a[128:159]
1516; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[24:27] offset:96
1517; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[20:23] offset:80
1518; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[16:19] offset:64
1519; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[12:15] offset:48
1520; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[8:11] offset:32
1521; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[4:7] offset:16
1522; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[0:3]
1523; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v0, s7
1524; EXACTCUTOFF-NEXT:    ; kill: killed $sgpr4_sgpr5
1525; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1526; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1527; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1528; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1529; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1530; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1531; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1532; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1533; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)
1534; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1535; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[56:59] offset:8288
1536; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[60:63] offset:8304
1537; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[48:51] offset:8256
1538; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[52:55] offset:8272
1539; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[40:43] offset:8224
1540; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[44:47] offset:8240
1541; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[32:35] offset:8192
1542; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[36:39] offset:8208
1543; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[88:91] offset:16480
1544; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[92:95] offset:16496
1545; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[80:83] offset:16448
1546; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[84:87] offset:16464
1547; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[72:75] offset:16416
1548; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[76:79] offset:16432
1549; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[64:67] offset:16384
1550; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[68:71] offset:16400
1551; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[120:123] offset:24672
1552; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[124:127] offset:24688
1553; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[112:115] offset:24640
1554; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[116:119] offset:24656
1555; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[104:107] offset:24608
1556; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[108:111] offset:24624
1557; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[96:99] offset:24576
1558; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[100:103] offset:24592
1559; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[152:155] offset:32864
1560; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[156:159] offset:32880
1561; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[144:147] offset:32832
1562; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[148:151] offset:32848
1563; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[136:139] offset:32800
1564; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[140:143] offset:32816
1565; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[128:131] offset:32768
1566; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[132:135] offset:32784
1567; EXACTCUTOFF-NEXT:    s_endpgm
1568entry:
1569  %idx = call i32 @llvm.amdgcn.workitem.id.x()
1570  %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx
1571  %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr
1572  %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64
1573  %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr
1574  %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128
1575  %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr
1576  %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192
1577  %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr
1578  %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256
1579  %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr
1580  %el0 = extractelement <5 x float> %in1, i32 0
1581  %el1 = extractelement <5 x float> %in1, i32 1
1582  %el2 = extractelement <5 x float> %in1, i32 2
1583  %el3 = extractelement <5 x float> %in1, i32 3
1584  %el4 = extractelement <5 x float> %in1, i32 4
1585  %exp0 = tail call float @llvm.exp.f32(float %el0)
1586  %exp1 = tail call float @llvm.exp.f32(float %el1)
1587  %exp2 = tail call float @llvm.exp.f32(float %el2)
1588  %exp3 = tail call float @llvm.exp.f32(float %el3)
1589  %exp4 = tail call float @llvm.exp.f32(float %el4)
1590  %mai.0 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float %exp0, <32 x float> %load.0, i32 0, i32 0, i32 0)
1591  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float %exp1, <32 x float> %load.1, i32 0, i32 0, i32 0)
1592  %mai.2 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float %exp2, <32 x float> %load.2, i32 0, i32 0, i32 0)
1593  %mai.3 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float %exp3, <32 x float> %load.3, i32 0, i32 0, i32 0)
1594  %mai.4 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float %exp4, <32 x float> %load.4, i32 0, i32 0, i32 0)
1595  ; 1 TRANS
1596  call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)
1597  ; 1 MFMA
1598  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1599  ; 1 TRANS
1600  call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)
1601  ; 1 MFMA
1602  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1603  ; 1 TRANS
1604  call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)
1605  ; 1 MFMA
1606  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1607  ; 1 TRANS
1608  call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)
1609  ; 1 MFMA
1610  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1611  ; 1 TRANS
1612  call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)
1613  ; 1 MFMA
1614  call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1615  %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx
1616  store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr
1617  %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64
1618  store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr
1619  %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128
1620  store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr
1621  %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192
1622  store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr
1623  %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256
1624  store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr
1625  ret void
1626}
1627
1628declare i32 @llvm.amdgcn.workitem.id.x() #2
1629declare void @llvm.amdgcn.sched.group.barrier(i32, i32, i32) #1
1630declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32) #1
1631declare float @llvm.exp.f32(float) #2
1632
1633attributes #0 = { nounwind "amdgpu-flat-work-group-size"="1,256" }
1634attributes #1 = { nounwind }
1635attributes #2 = { nounwind readnone speculatable }
1636