llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.ll

   1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
   2 ; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs -misched-cluster=0  < %s | FileCheck -check-prefix=GCN %s
   3 ; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs -misched-cluster=0 -amdgpu-igrouplp-exact-solver-max-branches=250000 < %s | FileCheck -check-prefix=EXACTCUTOFF %s
   4
   5 define 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
  21 entry:
  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
  29 define 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[0:1], 0x24
  33 ; GCN-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
  34 ; GCN-NEXT:    ; kill: killed $sgpr0_sgpr1
  35 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
  36 ; GCN-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
  37 ; GCN-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:16
  38 ; GCN-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:32
  39 ; GCN-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:48
  40 ; GCN-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:64
  41 ; GCN-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:80
  42 ; GCN-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:96
  43 ; GCN-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:112
  44 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(8) SyncID(0)
  45 ; GCN-NEXT:    s_waitcnt vmcnt(7)
  46 ; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
  47 ; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
  48 ; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
  49 ; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
  50 ; GCN-NEXT:    s_waitcnt vmcnt(6)
  51 ; GCN-NEXT:    v_mul_lo_u32 v7, v7, v7
  52 ; GCN-NEXT:    v_mul_lo_u32 v6, v6, v6
  53 ; GCN-NEXT:    v_mul_lo_u32 v5, v5, v5
  54 ; GCN-NEXT:    v_mul_lo_u32 v4, v4, v4
  55 ; GCN-NEXT:    s_waitcnt vmcnt(5)
  56 ; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
  57 ; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
  58 ; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
  59 ; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
  60 ; GCN-NEXT:    s_waitcnt vmcnt(4)
  61 ; GCN-NEXT:    v_mul_lo_u32 v15, v15, v15
  62 ; GCN-NEXT:    v_mul_lo_u32 v14, v14, v14
  63 ; GCN-NEXT:    v_mul_lo_u32 v13, v13, v13
  64 ; GCN-NEXT:    v_mul_lo_u32 v12, v12, v12
  65 ; GCN-NEXT:    s_waitcnt vmcnt(3)
  66 ; GCN-NEXT:    v_mul_lo_u32 v19, v19, v19
  67 ; GCN-NEXT:    v_mul_lo_u32 v18, v18, v18
  68 ; GCN-NEXT:    v_mul_lo_u32 v17, v17, v17
  69 ; GCN-NEXT:    v_mul_lo_u32 v16, v16, v16
  70 ; GCN-NEXT:    s_waitcnt vmcnt(2)
  71 ; GCN-NEXT:    v_mul_lo_u32 v23, v23, v23
  72 ; GCN-NEXT:    v_mul_lo_u32 v22, v22, v22
  73 ; GCN-NEXT:    v_mul_lo_u32 v21, v21, v21
  74 ; GCN-NEXT:    v_mul_lo_u32 v20, v20, v20
  75 ; GCN-NEXT:    s_waitcnt vmcnt(1)
  76 ; GCN-NEXT:    v_mul_lo_u32 v27, v27, v27
  77 ; GCN-NEXT:    v_mul_lo_u32 v26, v26, v26
  78 ; GCN-NEXT:    v_mul_lo_u32 v25, v25, v25
  79 ; GCN-NEXT:    v_mul_lo_u32 v24, v24, v24
  80 ; GCN-NEXT:    s_waitcnt vmcnt(0)
  81 ; GCN-NEXT:    v_mul_lo_u32 v31, v31, v31
  82 ; GCN-NEXT:    v_mul_lo_u32 v30, v30, v30
  83 ; GCN-NEXT:    v_mul_lo_u32 v29, v29, v29
  84 ; GCN-NEXT:    v_mul_lo_u32 v28, v28, v28
  85 ; GCN-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:112
  86 ; GCN-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:96
  87 ; GCN-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:80
  88 ; GCN-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:64
  89 ; GCN-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:48
  90 ; GCN-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:32
  91 ; GCN-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:16
  92 ; GCN-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
  93 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(30) SyncID(0)
  94 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
  95 ; GCN-NEXT:    s_endpgm
  96 ;
  97 ; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_READ_VALU_WRITE:
  98 ; EXACTCUTOFF:       ; %bb.0:
  99 ; EXACTCUTOFF-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
 100 ; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
 101 ; EXACTCUTOFF-NEXT:    ; kill: killed $sgpr0_sgpr1
 102 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
 103 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
 104 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:16
 105 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:32
 106 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:48
 107 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:64
 108 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:80
 109 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:96
 110 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:112
 111 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(8) SyncID(0)
 112 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(7)
 113 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
 114 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
 115 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
 116 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
 117 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(6)
 118 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v7, v7, v7
 119 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v6, v6, v6
 120 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v5, v5, v5
 121 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v4, v4, v4
 122 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(5)
 123 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
 124 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
 125 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
 126 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
 127 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(4)
 128 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v15, v15, v15
 129 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v14, v14, v14
 130 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v13, v13, v13
 131 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v12, v12, v12
 132 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(3)
 133 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v19, v19, v19
 134 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v18, v18, v18
 135 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v17, v17, v17
 136 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v16, v16, v16
 137 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(2)
 138 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v23, v23, v23
 139 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v22, v22, v22
 140 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v21, v21, v21
 141 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v20, v20, v20
 142 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
 143 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v27, v27, v27
 144 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v26, v26, v26
 145 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v25, v25, v25
 146 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v24, v24, v24
 147 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 148 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v31, v31, v31
 149 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v30, v30, v30
 150 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v29, v29, v29
 151 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v28, v28, v28
 152 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:112
 153 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:96
 154 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:80
 155 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:64
 156 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:48
 157 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:32
 158 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:16
 159 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
 160 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(30) SyncID(0)
 161 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
 162 ; EXACTCUTOFF-NEXT:    s_endpgm
 163   %tid = call i32 @llvm.amdgcn.workitem.id.x() #2
 164   %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid
 165   %load = load <32 x i32>, ptr addrspace(1) %gep1
 166   %mul = mul <32 x i32> %load, %load
 167   %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid
 168   store <32 x i32> %mul, ptr addrspace(1) %gep2
 169   ; 8 VMEM read
 170   call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 8, i32 0)
 171   ; 30 VALU
 172   call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 30, i32 0)
 173   ; 8 VMEM write
 174   call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 8, i32 0)
 175   ret void
 176 }
 177
 178 define amdgpu_kernel void @test_sched_group_barrier_pipeline_alternating_READ_VALU(ptr addrspace(1) noalias %in, ptr addrspace(1) noalias %out) #0 {
 179 ; GCN-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU:
 180 ; GCN:       ; %bb.0:
 181 ; GCN-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
 182 ; GCN-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
 183 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 184 ; GCN-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:16
 185 ; GCN-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:96
 186 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 187 ; GCN-NEXT:    s_waitcnt vmcnt(1)
 188 ; GCN-NEXT:    v_mul_lo_u32 v29, v29, v29
 189 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 190 ; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
 191 ; GCN-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
 192 ; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
 193 ; GCN-NEXT:    v_mul_lo_u32 v28, v28, v28
 194 ; GCN-NEXT:    v_mul_lo_u32 v31, v31, v31
 195 ; GCN-NEXT:    v_mul_lo_u32 v30, v30, v30
 196 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 197 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 198 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 199 ; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
 200 ; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
 201 ; GCN-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:112
 202 ; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
 203 ; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
 204 ; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
 205 ; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
 206 ; GCN-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:48
 207 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 208 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 209 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 210 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 211 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 212 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 213 ; GCN-NEXT:    s_waitcnt vmcnt(1)
 214 ; GCN-NEXT:    v_mul_lo_u32 v7, v7, v7
 215 ; GCN-NEXT:    v_mul_lo_u32 v6, v6, v6
 216 ; GCN-NEXT:    v_mul_lo_u32 v5, v5, v5
 217 ; GCN-NEXT:    v_mul_lo_u32 v4, v4, v4
 218 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 219 ; GCN-NEXT:    v_mul_lo_u32 v13, v13, v13
 220 ; GCN-NEXT:    v_mul_lo_u32 v15, v15, v15
 221 ; GCN-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:80
 222 ; GCN-NEXT:    v_mul_lo_u32 v14, v14, v14
 223 ; GCN-NEXT:    v_mul_lo_u32 v12, v12, v12
 224 ; GCN-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:64
 225 ; GCN-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:32
 226 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 227 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 228 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 229 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 230 ; GCN-NEXT:    s_waitcnt vmcnt(2)
 231 ; GCN-NEXT:    v_mul_lo_u32 v19, v19, v19
 232 ; GCN-NEXT:    v_mul_lo_u32 v18, v18, v18
 233 ; GCN-NEXT:    v_mul_lo_u32 v17, v17, v17
 234 ; GCN-NEXT:    s_waitcnt vmcnt(1)
 235 ; GCN-NEXT:    v_mul_lo_u32 v23, v23, v23
 236 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 237 ; GCN-NEXT:    v_mul_lo_u32 v27, v27, v27
 238 ; GCN-NEXT:    v_mul_lo_u32 v26, v26, v26
 239 ; GCN-NEXT:    v_mul_lo_u32 v25, v25, v25
 240 ; GCN-NEXT:    v_mul_lo_u32 v24, v24, v24
 241 ; GCN-NEXT:    v_mul_lo_u32 v22, v22, v22
 242 ; GCN-NEXT:    v_mul_lo_u32 v21, v21, v21
 243 ; GCN-NEXT:    v_mul_lo_u32 v20, v20, v20
 244 ; GCN-NEXT:    v_mul_lo_u32 v16, v16, v16
 245 ; GCN-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:112
 246 ; GCN-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:96
 247 ; GCN-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:80
 248 ; GCN-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:64
 249 ; GCN-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:48
 250 ; GCN-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:32
 251 ; GCN-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:16
 252 ; GCN-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
 253 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 254 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 255 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 256 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
 257 ; GCN-NEXT:    s_endpgm
 258 ;
 259 ; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU:
 260 ; EXACTCUTOFF:       ; %bb.0:
 261 ; EXACTCUTOFF-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
 262 ; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v32, 7, v0
 263 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
 264 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[28:31], v32, s[0:1] offset:16
 265 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v32, s[0:1] offset:96
 266 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 267 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
 268 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v29, v29, v29
 269 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 270 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
 271 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v32, s[0:1]
 272 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
 273 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v28, v28, v28
 274 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v31, v31, v31
 275 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v30, v30, v30
 276 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 277 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 278 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 279 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
 280 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
 281 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[4:7], v32, s[0:1] offset:112
 282 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
 283 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
 284 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
 285 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
 286 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[12:15], v32, s[0:1] offset:48
 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:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 290 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 291 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 292 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 293 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
 294 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v7, v7, v7
 295 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v6, v6, v6
 296 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v5, v5, v5
 297 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v4, v4, v4
 298 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 299 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v13, v13, v13
 300 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v15, v15, v15
 301 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[16:19], v32, s[0:1] offset:80
 302 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v14, v14, v14
 303 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v12, v12, v12
 304 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[20:23], v32, s[0:1] offset:64
 305 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[24:27], v32, s[0:1] offset:32
 306 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 307 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 308 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 309 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 310 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(2)
 311 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v19, v19, v19
 312 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v18, v18, v18
 313 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v17, v17, v17
 314 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
 315 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v23, v23, v23
 316 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 317 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v27, v27, v27
 318 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v26, v26, v26
 319 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v25, v25, v25
 320 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v24, v24, v24
 321 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v22, v22, v22
 322 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v21, v21, v21
 323 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v20, v20, v20
 324 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v16, v16, v16
 325 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[4:7], s[2:3] offset:112
 326 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[8:11], s[2:3] offset:96
 327 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[16:19], s[2:3] offset:80
 328 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[20:23], s[2:3] offset:64
 329 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[12:15], s[2:3] offset:48
 330 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[24:27], s[2:3] offset:32
 331 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[28:31], s[2:3] offset:16
 332 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v32, v[0:3], s[2:3]
 333 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 334 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 335 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 336 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(8) SyncID(0)
 337 ; EXACTCUTOFF-NEXT:    s_endpgm
 338   %tid = call i32 @llvm.amdgcn.workitem.id.x() #2
 339   %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid
 340   %load = load <32 x i32>, ptr addrspace(1) %gep1
 341   %mul = mul <32 x i32> %load, %load
 342   %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid
 343   store <32 x i32> %mul, ptr addrspace(1) %gep2
 344   ; 1 VMEM read
 345   call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
 346   ; 2 VALU
 347   call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
 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   ; 8 VMEM write
 377   call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 8, i32 0)
 378   ret void
 379 }
 380
 381 define amdgpu_kernel void @test_sched_group_barrier_pipeline_alternating_READ_VALU_WRITE(ptr addrspace(1) noalias %in, ptr addrspace(1) noalias %out) #0 {
 382 ; GCN-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU_WRITE:
 383 ; GCN:       ; %bb.0:
 384 ; GCN-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
 385 ; GCN-NEXT:    v_lshlrev_b32_e32 v16, 7, v0
 386 ; GCN-NEXT:    ; kill: killed $sgpr0_sgpr1
 387 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 388 ; GCN-NEXT:    global_load_dwordx4 v[12:15], v16, s[0:1] offset:32
 389 ; GCN-NEXT:    global_load_dwordx4 v[4:7], v16, s[0:1] offset:48
 390 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 391 ; GCN-NEXT:    s_waitcnt vmcnt(1)
 392 ; GCN-NEXT:    v_mul_lo_u32 v13, v13, v13
 393 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 394 ; GCN-NEXT:    v_mul_lo_u32 v7, v7, v7
 395 ; GCN-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1]
 396 ; GCN-NEXT:    v_mul_lo_u32 v6, v6, v6
 397 ; GCN-NEXT:    v_mul_lo_u32 v12, v12, v12
 398 ; GCN-NEXT:    v_mul_lo_u32 v15, v15, v15
 399 ; GCN-NEXT:    v_mul_lo_u32 v14, v14, v14
 400 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 401 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 402 ; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
 403 ; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
 404 ; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
 405 ; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
 406 ; GCN-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3]
 407 ; GCN-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:112
 408 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 409 ; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
 410 ; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
 411 ; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
 412 ; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
 413 ; GCN-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:112
 414 ; GCN-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:96
 415 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 416 ; GCN-NEXT:    v_mul_lo_u32 v3, v3, v3
 417 ; GCN-NEXT:    v_mul_lo_u32 v2, v2, v2
 418 ; GCN-NEXT:    v_mul_lo_u32 v1, v1, v1
 419 ; GCN-NEXT:    v_mul_lo_u32 v0, v0, v0
 420 ; GCN-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:96
 421 ; GCN-NEXT:    v_mul_lo_u32 v5, v5, v5
 422 ; GCN-NEXT:    v_mul_lo_u32 v4, v4, v4
 423 ; GCN-NEXT:    global_store_dwordx4 v16, v[4:7], s[2:3] offset:48
 424 ; GCN-NEXT:    global_load_dwordx4 v[4:7], v16, s[0:1] offset:64
 425 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 426 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 427 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 428 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 429 ; GCN-NEXT:    v_mul_lo_u32 v7, v7, v7
 430 ; GCN-NEXT:    v_mul_lo_u32 v6, v6, v6
 431 ; GCN-NEXT:    v_mul_lo_u32 v5, v5, v5
 432 ; GCN-NEXT:    v_mul_lo_u32 v4, v4, v4
 433 ; GCN-NEXT:    global_store_dwordx4 v16, v[4:7], s[2:3] offset:64
 434 ; GCN-NEXT:    global_store_dwordx4 v16, v[12:15], s[2:3] offset:32
 435 ; GCN-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:16
 436 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 437 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 438 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 439 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 440 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 441 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 442 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 443 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 444 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 445 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 446 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 447 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 448 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 449 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 450 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 451 ; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
 452 ; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
 453 ; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
 454 ; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
 455 ; GCN-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:16
 456 ; GCN-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:80
 457 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 458 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 459 ; GCN-NEXT:    v_mul_lo_u32 v11, v11, v11
 460 ; GCN-NEXT:    v_mul_lo_u32 v10, v10, v10
 461 ; GCN-NEXT:    v_mul_lo_u32 v9, v9, v9
 462 ; GCN-NEXT:    v_mul_lo_u32 v8, v8, v8
 463 ; GCN-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:80
 464 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 465 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 466 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 467 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 468 ; GCN-NEXT:    s_endpgm
 469 ;
 470 ; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_alternating_READ_VALU_WRITE:
 471 ; EXACTCUTOFF:       ; %bb.0:
 472 ; EXACTCUTOFF-NEXT:    s_load_dwordx4 s[0:3], s[0:1], 0x24
 473 ; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v16, 7, v0
 474 ; EXACTCUTOFF-NEXT:    ; kill: killed $sgpr0_sgpr1
 475 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
 476 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[12:15], v16, s[0:1] offset:32
 477 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[4:7], v16, s[0:1] offset:48
 478 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 479 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(1)
 480 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v13, v13, v13
 481 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 482 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v7, v7, v7
 483 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1]
 484 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v6, v6, v6
 485 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v12, v12, v12
 486 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v15, v15, v15
 487 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v14, v14, v14
 488 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 489 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 490 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
 491 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
 492 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
 493 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
 494 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3]
 495 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:112
 496 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 497 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
 498 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
 499 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
 500 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
 501 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:112
 502 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[0:3], v16, s[0:1] offset:96
 503 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 504 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v3, v3, v3
 505 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v2, v2, v2
 506 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v1, v1, v1
 507 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v0, v0, v0
 508 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[0:3], s[2:3] offset:96
 509 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v5, v5, v5
 510 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v4, v4, v4
 511 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[4:7], s[2:3] offset:48
 512 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[4:7], v16, s[0:1] offset:64
 513 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 514 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 515 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 516 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 517 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v7, v7, v7
 518 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v6, v6, v6
 519 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v5, v5, v5
 520 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v4, v4, v4
 521 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[4:7], s[2:3] offset:64
 522 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[12:15], s[2:3] offset:32
 523 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:16
 524 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 525 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 526 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 527 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 528 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 529 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 530 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 531 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 532 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 533 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 534 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000020) size(1) SyncID(0)
 535 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 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:    s_waitcnt vmcnt(0)
 539 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
 540 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
 541 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
 542 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
 543 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:16
 544 ; EXACTCUTOFF-NEXT:    global_load_dwordx4 v[8:11], v16, s[0:1] offset:80
 545 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000002) size(2) SyncID(0)
 546 ; EXACTCUTOFF-NEXT:    s_waitcnt vmcnt(0)
 547 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v11, v11, v11
 548 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v10, v10, v10
 549 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v9, v9, v9
 550 ; EXACTCUTOFF-NEXT:    v_mul_lo_u32 v8, v8, v8
 551 ; EXACTCUTOFF-NEXT:    global_store_dwordx4 v16, v[8:11], s[2:3] offset:80
 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:    ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)
 556 ; EXACTCUTOFF-NEXT:    s_endpgm
 557   %tid = call i32 @llvm.amdgcn.workitem.id.x() #2
 558   %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid
 559   %load = load <32 x i32>, ptr addrspace(1) %gep1
 560   %mul = mul <32 x i32> %load, %load
 561   %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid
 562   store <32 x i32> %mul, ptr addrspace(1) %gep2
 563   ; 1 VMEM read
 564   call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)
 565   ; 2 VALU
 566   call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)
 567   ; 1 VMEM write
 568   call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)
 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   ret void
 612 }
 613
 614 define amdgpu_kernel void @test_sched_group_barrier_pipeline_MFMA_cluster(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {
 615 ; GCN-LABEL: test_sched_group_barrier_pipeline_MFMA_cluster:
 616 ; GCN:       ; %bb.0: ; %entry
 617 ; GCN-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0x24
 618 ; GCN-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
 619 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 620 ; GCN-NEXT:    v_add_u32_e32 v1, s0, v0
 621 ; GCN-NEXT:    ds_read_b128 a[28:31], v1 offset:112
 622 ; GCN-NEXT:    ds_read_b128 a[24:27], v1 offset:96
 623 ; GCN-NEXT:    ds_read_b128 a[20:23], v1 offset:80
 624 ; GCN-NEXT:    ds_read_b128 a[16:19], v1 offset:64
 625 ; GCN-NEXT:    ds_read_b128 a[0:3], v1
 626 ; GCN-NEXT:    ds_read_b128 a[4:7], v1 offset:16
 627 ; GCN-NEXT:    ds_read_b128 a[8:11], v1 offset:32
 628 ; GCN-NEXT:    ds_read_b128 a[12:15], v1 offset:48
 629 ; GCN-NEXT:    ds_read_b128 a[60:63], v1 offset:8304
 630 ; GCN-NEXT:    ds_read_b128 a[56:59], v1 offset:8288
 631 ; GCN-NEXT:    ds_read_b128 a[52:55], v1 offset:8272
 632 ; GCN-NEXT:    ds_read_b128 a[48:51], v1 offset:8256
 633 ; GCN-NEXT:    ds_read_b128 a[44:47], v1 offset:8240
 634 ; GCN-NEXT:    ds_read_b128 a[40:43], v1 offset:8224
 635 ; GCN-NEXT:    ds_read_b128 a[36:39], v1 offset:8208
 636 ; GCN-NEXT:    ds_read_b128 a[32:35], v1 offset:8192
 637 ; GCN-NEXT:    v_add_u32_e32 v2, 0x6000, v1
 638 ; GCN-NEXT:    ds_read_b128 a[92:95], v1 offset:24688
 639 ; GCN-NEXT:    ds_read_b128 a[88:91], v1 offset:24672
 640 ; GCN-NEXT:    ds_read_b128 a[84:87], v1 offset:24656
 641 ; GCN-NEXT:    ds_read_b128 a[80:83], v1 offset:24640
 642 ; GCN-NEXT:    ds_read_b128 a[76:79], v1 offset:24624
 643 ; GCN-NEXT:    ds_read_b128 a[72:75], v1 offset:24608
 644 ; GCN-NEXT:    ds_read_b128 a[68:71], v1 offset:24592
 645 ; GCN-NEXT:    ds_read_b128 a[64:67], v1 offset:24576
 646 ; GCN-NEXT:    ds_read_b128 a[124:127], v1 offset:49264
 647 ; GCN-NEXT:    ds_read_b128 a[120:123], v1 offset:49248
 648 ; GCN-NEXT:    ds_read_b128 a[116:119], v1 offset:49232
 649 ; GCN-NEXT:    ds_read_b128 a[112:115], v1 offset:49216
 650 ; GCN-NEXT:    ds_read_b128 a[108:111], v1 offset:49200
 651 ; GCN-NEXT:    ds_read_b128 a[104:107], v1 offset:49184
 652 ; GCN-NEXT:    ds_read_b128 a[100:103], v1 offset:49168
 653 ; GCN-NEXT:    ds_read_b128 a[96:99], v1 offset:49152
 654 ; GCN-NEXT:    v_mov_b32_e32 v1, 1.0
 655 ; GCN-NEXT:    ds_read_b128 a[156:159], v2 offset:57456
 656 ; GCN-NEXT:    ds_read_b128 a[152:155], v2 offset:57440
 657 ; GCN-NEXT:    ds_read_b128 a[148:151], v2 offset:57424
 658 ; GCN-NEXT:    ds_read_b128 a[144:147], v2 offset:57408
 659 ; GCN-NEXT:    ds_read_b128 a[128:131], v2 offset:57344
 660 ; GCN-NEXT:    ds_read_b128 a[132:135], v2 offset:57360
 661 ; GCN-NEXT:    ds_read_b128 a[136:139], v2 offset:57376
 662 ; GCN-NEXT:    ds_read_b128 a[140:143], v2 offset:57392
 663 ; GCN-NEXT:    v_mov_b32_e32 v2, 2.0
 664 ; GCN-NEXT:    v_add_u32_e32 v0, s1, v0
 665 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(40) SyncID(0)
 666 ; GCN-NEXT:    s_waitcnt lgkmcnt(14)
 667 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]
 668 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v1, v2, a[32:63]
 669 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[64:95], v1, v2, a[64:95]
 670 ; GCN-NEXT:    s_waitcnt lgkmcnt(8)
 671 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[96:127], v1, v2, a[96:127]
 672 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 673 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[128:159], v1, v2, a[128:159]
 674 ; GCN-NEXT:    s_nop 7
 675 ; GCN-NEXT:    s_nop 4
 676 ; GCN-NEXT:    ds_write_b128 v0, a[28:31] offset:112
 677 ; GCN-NEXT:    ds_write_b128 v0, a[24:27] offset:96
 678 ; GCN-NEXT:    ds_write_b128 v0, a[20:23] offset:80
 679 ; GCN-NEXT:    ds_write_b128 v0, a[16:19] offset:64
 680 ; GCN-NEXT:    ds_write_b128 v0, a[12:15] offset:48
 681 ; GCN-NEXT:    ds_write_b128 v0, a[8:11] offset:32
 682 ; GCN-NEXT:    ds_write_b128 v0, a[4:7] offset:16
 683 ; GCN-NEXT:    ds_write_b128 v0, a[0:3]
 684 ; GCN-NEXT:    v_mov_b32_e32 v0, s1
 685 ; GCN-NEXT:    ds_write_b128 v0, a[56:59] offset:8288
 686 ; GCN-NEXT:    ds_write_b128 v0, a[60:63] offset:8304
 687 ; GCN-NEXT:    ds_write_b128 v0, a[48:51] offset:8256
 688 ; GCN-NEXT:    ds_write_b128 v0, a[52:55] offset:8272
 689 ; GCN-NEXT:    ds_write_b128 v0, a[40:43] offset:8224
 690 ; GCN-NEXT:    ds_write_b128 v0, a[44:47] offset:8240
 691 ; GCN-NEXT:    ds_write_b128 v0, a[32:35] offset:8192
 692 ; GCN-NEXT:    ds_write_b128 v0, a[36:39] offset:8208
 693 ; GCN-NEXT:    ds_write_b128 v0, a[88:91] offset:16480
 694 ; GCN-NEXT:    ds_write_b128 v0, a[92:95] offset:16496
 695 ; GCN-NEXT:    ds_write_b128 v0, a[80:83] offset:16448
 696 ; GCN-NEXT:    ds_write_b128 v0, a[84:87] offset:16464
 697 ; GCN-NEXT:    ds_write_b128 v0, a[72:75] offset:16416
 698 ; GCN-NEXT:    ds_write_b128 v0, a[76:79] offset:16432
 699 ; GCN-NEXT:    ds_write_b128 v0, a[64:67] offset:16384
 700 ; GCN-NEXT:    ds_write_b128 v0, a[68:71] offset:16400
 701 ; GCN-NEXT:    ds_write_b128 v0, a[120:123] offset:24672
 702 ; GCN-NEXT:    ds_write_b128 v0, a[124:127] offset:24688
 703 ; GCN-NEXT:    ds_write_b128 v0, a[112:115] offset:24640
 704 ; GCN-NEXT:    ds_write_b128 v0, a[116:119] offset:24656
 705 ; GCN-NEXT:    ds_write_b128 v0, a[104:107] offset:24608
 706 ; GCN-NEXT:    ds_write_b128 v0, a[108:111] offset:24624
 707 ; GCN-NEXT:    ds_write_b128 v0, a[96:99] offset:24576
 708 ; GCN-NEXT:    ds_write_b128 v0, a[100:103] offset:24592
 709 ; GCN-NEXT:    ds_write_b128 v0, a[152:155] offset:32864
 710 ; GCN-NEXT:    ds_write_b128 v0, a[156:159] offset:32880
 711 ; GCN-NEXT:    ds_write_b128 v0, a[144:147] offset:32832
 712 ; GCN-NEXT:    ds_write_b128 v0, a[148:151] offset:32848
 713 ; GCN-NEXT:    ds_write_b128 v0, a[136:139] offset:32800
 714 ; GCN-NEXT:    ds_write_b128 v0, a[140:143] offset:32816
 715 ; GCN-NEXT:    ds_write_b128 v0, a[128:131] offset:32768
 716 ; GCN-NEXT:    ds_write_b128 v0, a[132:135] offset:32784
 717 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(5) SyncID(0)
 718 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(40) SyncID(0)
 719 ; GCN-NEXT:    s_endpgm
 720 ;
 721 ; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_MFMA_cluster:
 722 ; EXACTCUTOFF:       ; %bb.0: ; %entry
 723 ; EXACTCUTOFF-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0x24
 724 ; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v0, 7, v0
 725 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
 726 ; EXACTCUTOFF-NEXT:    v_add_u32_e32 v1, s0, v0
 727 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v1 offset:112
 728 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v1 offset:96
 729 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v1 offset:80
 730 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v1 offset:64
 731 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v1
 732 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v1 offset:16
 733 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v1 offset:32
 734 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v1 offset:48
 735 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[60:63], v1 offset:8304
 736 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[56:59], v1 offset:8288
 737 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[52:55], v1 offset:8272
 738 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[48:51], v1 offset:8256
 739 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[44:47], v1 offset:8240
 740 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[40:43], v1 offset:8224
 741 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[36:39], v1 offset:8208
 742 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[32:35], v1 offset:8192
 743 ; EXACTCUTOFF-NEXT:    v_add_u32_e32 v2, 0x6000, v1
 744 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[92:95], v1 offset:24688
 745 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[88:91], v1 offset:24672
 746 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[84:87], v1 offset:24656
 747 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[80:83], v1 offset:24640
 748 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[76:79], v1 offset:24624
 749 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[72:75], v1 offset:24608
 750 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[68:71], v1 offset:24592
 751 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[64:67], v1 offset:24576
 752 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[124:127], v1 offset:49264
 753 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[120:123], v1 offset:49248
 754 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[116:119], v1 offset:49232
 755 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[112:115], v1 offset:49216
 756 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[108:111], v1 offset:49200
 757 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[104:107], v1 offset:49184
 758 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[100:103], v1 offset:49168
 759 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[96:99], v1 offset:49152
 760 ; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v1, 1.0
 761 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[156:159], v2 offset:57456
 762 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[152:155], v2 offset:57440
 763 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[148:151], v2 offset:57424
 764 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[144:147], v2 offset:57408
 765 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[128:131], v2 offset:57344
 766 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[132:135], v2 offset:57360
 767 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[136:139], v2 offset:57376
 768 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[140:143], v2 offset:57392
 769 ; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v2, 2.0
 770 ; EXACTCUTOFF-NEXT:    v_add_u32_e32 v0, s1, v0
 771 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(40) SyncID(0)
 772 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(14)
 773 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]
 774 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v1, v2, a[32:63]
 775 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[64:95], v1, v2, a[64:95]
 776 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(8)
 777 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[96:127], v1, v2, a[96:127]
 778 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
 779 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[128:159], v1, v2, a[128:159]
 780 ; EXACTCUTOFF-NEXT:    s_nop 7
 781 ; EXACTCUTOFF-NEXT:    s_nop 4
 782 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[28:31] offset:112
 783 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[24:27] offset:96
 784 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[20:23] offset:80
 785 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[16:19] offset:64
 786 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[12:15] offset:48
 787 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[8:11] offset:32
 788 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[4:7] offset:16
 789 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[0:3]
 790 ; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v0, s1
 791 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[56:59] offset:8288
 792 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[60:63] offset:8304
 793 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[48:51] offset:8256
 794 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[52:55] offset:8272
 795 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[40:43] offset:8224
 796 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[44:47] offset:8240
 797 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[32:35] offset:8192
 798 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[36:39] offset:8208
 799 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[88:91] offset:16480
 800 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[92:95] offset:16496
 801 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[80:83] offset:16448
 802 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[84:87] offset:16464
 803 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[72:75] offset:16416
 804 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[76:79] offset:16432
 805 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[64:67] offset:16384
 806 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[68:71] offset:16400
 807 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[120:123] offset:24672
 808 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[124:127] offset:24688
 809 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[112:115] offset:24640
 810 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[116:119] offset:24656
 811 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[104:107] offset:24608
 812 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[108:111] offset:24624
 813 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[96:99] offset:24576
 814 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[100:103] offset:24592
 815 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[152:155] offset:32864
 816 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[156:159] offset:32880
 817 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[144:147] offset:32832
 818 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[148:151] offset:32848
 819 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[136:139] offset:32800
 820 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[140:143] offset:32816
 821 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[128:131] offset:32768
 822 ; EXACTCUTOFF-NEXT:    ds_write_b128 v0, a[132:135] offset:32784
 823 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(5) SyncID(0)
 824 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(40) SyncID(0)
 825 ; EXACTCUTOFF-NEXT:    s_endpgm
 826 entry:
 827   %idx = call i32 @llvm.amdgcn.workitem.id.x()
 828   %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx
 829   %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr
 830   %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64
 831   %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr
 832   %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128
 833   %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr
 834   %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192
 835   %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr
 836   %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256
 837   %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr
 838   %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)
 839   %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)
 840   %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)
 841   %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)
 842   %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)
 843   %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx
 844   store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr
 845   %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64
 846   store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr
 847   %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128
 848   store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr
 849   %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192
 850   store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr
 851   %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256
 852   store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr
 853   ; 40 DS read
 854   call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 40, i32 0)
 855   ; 5 MFMA
 856   call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 5, i32 0)
 857   ; 40 DS write
 858   call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 40, i32 0)
 859   ret void
 860 }
 861
 862 define amdgpu_kernel void @test_sched_group_barrier_pipeline_MFMA_interleave(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {
 863 ; GCN-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:
 864 ; GCN:       ; %bb.0: ; %entry
 865 ; GCN-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0x24
 866 ; GCN-NEXT:    v_lshlrev_b32_e32 v1, 7, v0
 867 ; GCN-NEXT:    v_mov_b32_e32 v2, 1.0
 868 ; GCN-NEXT:    v_mov_b32_e32 v3, 2.0
 869 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 870 ; GCN-NEXT:    v_add_u32_e32 v0, s0, v1
 871 ; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:112
 872 ; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:96
 873 ; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:80
 874 ; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:64
 875 ; GCN-NEXT:    ds_read_b128 a[0:3], v0
 876 ; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:16
 877 ; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:32
 878 ; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:48
 879 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 880 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
 881 ; GCN-NEXT:    v_add_u32_e32 v1, s1, v1
 882 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
 883 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
 884 ; GCN-NEXT:    s_nop 7
 885 ; GCN-NEXT:    s_nop 7
 886 ; GCN-NEXT:    s_nop 1
 887 ; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:112
 888 ; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:96
 889 ; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:80
 890 ; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:64
 891 ; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:48
 892 ; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:32
 893 ; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:16
 894 ; GCN-NEXT:    ds_write_b128 v1, a[0:3]
 895 ; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:8304
 896 ; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:8288
 897 ; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:8272
 898 ; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:8256
 899 ; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:8240
 900 ; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:8224
 901 ; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:8208
 902 ; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:8192
 903 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 904 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
 905 ; GCN-NEXT:    v_mov_b32_e32 v1, s1
 906 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
 907 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
 908 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
 909 ; GCN-NEXT:    s_nop 7
 910 ; GCN-NEXT:    s_nop 7
 911 ; GCN-NEXT:    s_nop 1
 912 ; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:8288
 913 ; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:8304
 914 ; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:8256
 915 ; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:8272
 916 ; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:8224
 917 ; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:8240
 918 ; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:8192
 919 ; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:8208
 920 ; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:24688
 921 ; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:24672
 922 ; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:24656
 923 ; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:24640
 924 ; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:24624
 925 ; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:24608
 926 ; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:24592
 927 ; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:24576
 928 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 929 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
 930 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
 931 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
 932 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
 933 ; GCN-NEXT:    s_nop 7
 934 ; GCN-NEXT:    s_nop 7
 935 ; GCN-NEXT:    s_nop 2
 936 ; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:16480
 937 ; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:16496
 938 ; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:16448
 939 ; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:16464
 940 ; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:16416
 941 ; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:16432
 942 ; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:16384
 943 ; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:16400
 944 ; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:49264
 945 ; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:49248
 946 ; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:49232
 947 ; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:49216
 948 ; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:49200
 949 ; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:49184
 950 ; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:49168
 951 ; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:49152
 952 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 953 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
 954 ; GCN-NEXT:    v_add_u32_e32 v0, 0x6000, v0
 955 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
 956 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
 957 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
 958 ; GCN-NEXT:    s_nop 7
 959 ; GCN-NEXT:    s_nop 7
 960 ; GCN-NEXT:    s_nop 1
 961 ; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:24672
 962 ; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:24688
 963 ; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:24640
 964 ; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:24656
 965 ; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:24608
 966 ; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:24624
 967 ; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:24576
 968 ; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:24592
 969 ; GCN-NEXT:    ds_read_b128 a[28:31], v0 offset:57456
 970 ; GCN-NEXT:    ds_read_b128 a[24:27], v0 offset:57440
 971 ; GCN-NEXT:    ds_read_b128 a[20:23], v0 offset:57424
 972 ; GCN-NEXT:    ds_read_b128 a[16:19], v0 offset:57408
 973 ; GCN-NEXT:    ds_read_b128 a[0:3], v0 offset:57344
 974 ; GCN-NEXT:    ds_read_b128 a[4:7], v0 offset:57360
 975 ; GCN-NEXT:    ds_read_b128 a[8:11], v0 offset:57376
 976 ; GCN-NEXT:    ds_read_b128 a[12:15], v0 offset:57392
 977 ; GCN-NEXT:    s_waitcnt lgkmcnt(0)
 978 ; GCN-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
 979 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
 980 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
 981 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
 982 ; GCN-NEXT:    s_nop 7
 983 ; GCN-NEXT:    s_nop 7
 984 ; GCN-NEXT:    s_nop 2
 985 ; GCN-NEXT:    ds_write_b128 v1, a[24:27] offset:32864
 986 ; GCN-NEXT:    ds_write_b128 v1, a[28:31] offset:32880
 987 ; GCN-NEXT:    ds_write_b128 v1, a[16:19] offset:32832
 988 ; GCN-NEXT:    ds_write_b128 v1, a[20:23] offset:32848
 989 ; GCN-NEXT:    ds_write_b128 v1, a[8:11] offset:32800
 990 ; GCN-NEXT:    ds_write_b128 v1, a[12:15] offset:32816
 991 ; GCN-NEXT:    ds_write_b128 v1, a[0:3] offset:32768
 992 ; GCN-NEXT:    ds_write_b128 v1, a[4:7] offset:32784
 993 ; GCN-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
 994 ; GCN-NEXT:    s_endpgm
 995 ;
 996 ; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:
 997 ; EXACTCUTOFF:       ; %bb.0: ; %entry
 998 ; EXACTCUTOFF-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0x24
 999 ; EXACTCUTOFF-NEXT:    v_lshlrev_b32_e32 v1, 7, v0
1000 ; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v2, 1.0
1001 ; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v3, 2.0
1002 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1003 ; EXACTCUTOFF-NEXT:    v_add_u32_e32 v0, s0, v1
1004 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:112
1005 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:96
1006 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:80
1007 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:64
1008 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0
1009 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:16
1010 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:32
1011 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:48
1012 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1013 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1014 ; EXACTCUTOFF-NEXT:    v_add_u32_e32 v1, s1, v1
1015 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1016 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1017 ; EXACTCUTOFF-NEXT:    s_nop 7
1018 ; EXACTCUTOFF-NEXT:    s_nop 7
1019 ; EXACTCUTOFF-NEXT:    s_nop 1
1020 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:112
1021 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:96
1022 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:80
1023 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:64
1024 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:48
1025 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:32
1026 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:16
1027 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3]
1028 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:8304
1029 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:8288
1030 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:8272
1031 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:8256
1032 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:8240
1033 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:8224
1034 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:8208
1035 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:8192
1036 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1037 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1038 ; EXACTCUTOFF-NEXT:    v_mov_b32_e32 v1, s1
1039 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1040 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1041 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1042 ; EXACTCUTOFF-NEXT:    s_nop 7
1043 ; EXACTCUTOFF-NEXT:    s_nop 7
1044 ; EXACTCUTOFF-NEXT:    s_nop 1
1045 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:8288
1046 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:8304
1047 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:8256
1048 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:8272
1049 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:8224
1050 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:8240
1051 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:8192
1052 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:8208
1053 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:24688
1054 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:24672
1055 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:24656
1056 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:24640
1057 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:24624
1058 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:24608
1059 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:24592
1060 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:24576
1061 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1062 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1063 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1064 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1065 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1066 ; EXACTCUTOFF-NEXT:    s_nop 7
1067 ; EXACTCUTOFF-NEXT:    s_nop 7
1068 ; EXACTCUTOFF-NEXT:    s_nop 2
1069 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:16480
1070 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:16496
1071 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:16448
1072 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:16464
1073 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:16416
1074 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:16432
1075 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:16384
1076 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:16400
1077 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:49264
1078 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:49248
1079 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:49232
1080 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:49216
1081 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:49200
1082 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:49184
1083 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:49168
1084 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:49152
1085 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1086 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1087 ; EXACTCUTOFF-NEXT:    v_add_u32_e32 v0, 0x6000, v0
1088 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1089 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1090 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1091 ; EXACTCUTOFF-NEXT:    s_nop 7
1092 ; EXACTCUTOFF-NEXT:    s_nop 7
1093 ; EXACTCUTOFF-NEXT:    s_nop 1
1094 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:24672
1095 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:24688
1096 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:24640
1097 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:24656
1098 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:24608
1099 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:24624
1100 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:24576
1101 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:24592
1102 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[28:31], v0 offset:57456
1103 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[24:27], v0 offset:57440
1104 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[20:23], v0 offset:57424
1105 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[16:19], v0 offset:57408
1106 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[0:3], v0 offset:57344
1107 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[4:7], v0 offset:57360
1108 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[8:11], v0 offset:57376
1109 ; EXACTCUTOFF-NEXT:    ds_read_b128 a[12:15], v0 offset:57392
1110 ; EXACTCUTOFF-NEXT:    s_waitcnt lgkmcnt(0)
1111 ; EXACTCUTOFF-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]
1112 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1113 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)
1114 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)
1115 ; EXACTCUTOFF-NEXT:    s_nop 7
1116 ; EXACTCUTOFF-NEXT:    s_nop 7
1117 ; EXACTCUTOFF-NEXT:    s_nop 2
1118 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[24:27] offset:32864
1119 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[28:31] offset:32880
1120 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[16:19] offset:32832
1121 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[20:23] offset:32848
1122 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[8:11] offset:32800
1123 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[12:15] offset:32816
1124 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[0:3] offset:32768
1125 ; EXACTCUTOFF-NEXT:    ds_write_b128 v1, a[4:7] offset:32784
1126 ; EXACTCUTOFF-NEXT:    ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)
1127 ; EXACTCUTOFF-NEXT:    s_endpgm
1128 entry:
1129   %idx = call i32 @llvm.amdgcn.workitem.id.x()
1130   %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx
1131   %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr
1132   %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64
1133   %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr
1134   %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128
1135   %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr
1136   %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192
1137   %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr
1138   %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256
1139   %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr
1140   %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)
1141   %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)
1142   %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)
1143   %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)
1144   %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)
1145   %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx
1146   store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr
1147   %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64
1148   store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr
1149   %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128
1150   store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr
1151   %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192
1152   store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr
1153   %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256
1154   store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr
1155   ; 8 DS read
1156   call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1157   ; 1 MFMA
1158   call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1159   ; 8 DS write
1160   call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1161   ; 8 DS read
1162   call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1163   ; 1 MFMA
1164   call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1165   ; 8 DS write
1166   call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1167   ; 8 DS read
1168   call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1169   ; 1 MFMA
1170   call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1171   ; 8 DS write
1172   call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1173   ; 8 DS read
1174   call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1175   ; 1 MFMA
1176   call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1177   ; 8 DS write
1178   call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1179   ; 8 DS read
1180   call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)
1181   ; 1 MFMA
1182   call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
1183   ; 8 DS write
1184   call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)
1185   ret void
1186 }
1187
1188
1189
1190 declare i32 @llvm.amdgcn.workitem.id.x() #2
1191 declare void @llvm.amdgcn.sched.group.barrier(i32, i32, i32) #1
1192 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32) #1
1193
1194 attributes #0 = { nounwind "amdgpu-flat-work-group-size"="1,256" }
1195 attributes #1 = { nounwind }
1196 attributes #2 = { nounwind readnone speculatable }