1624 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py2; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -misched-cluster=0 < %s | FileCheck -check-prefix=GCN %s3; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -misched-cluster=0 -amdgpu-igrouplp-exact-solver-max-branches=250000 < %s | FileCheck -check-prefix=EXACTCUTOFF %s4 5define amdgpu_kernel void @test_sched_group_barrier() #0 {6; GCN-LABEL: test_sched_group_barrier:7; GCN: ; %bb.0: ; %entry8; 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_endpgm13;14; EXACTCUTOFF-LABEL: test_sched_group_barrier:15; EXACTCUTOFF: ; %bb.0: ; %entry16; 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_endpgm21entry:22 call void @llvm.amdgcn.sched.group.barrier(i32 0, i32 1, i32 2) #123 call void @llvm.amdgcn.sched.group.barrier(i32 1, i32 2, i32 4) #124 call void @llvm.amdgcn.sched.group.barrier(i32 4, i32 8, i32 16) #125 call void @llvm.amdgcn.sched.group.barrier(i32 15, i32 10000, i32 -1) #126 ret void27}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], 0x2433; GCN-NEXT: v_and_b32_e32 v0, 0x3ff, v034; GCN-NEXT: v_lshlrev_b32_e32 v32, 7, v035; GCN-NEXT: ; kill: killed $sgpr0_sgpr136; 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:1639; GCN-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:3240; GCN-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:4841; GCN-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:6442; GCN-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:8043; GCN-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:9644; GCN-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:11245; 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, v348; GCN-NEXT: v_mul_lo_u32 v2, v2, v249; GCN-NEXT: v_mul_lo_u32 v1, v1, v150; GCN-NEXT: v_mul_lo_u32 v0, v0, v051; GCN-NEXT: s_waitcnt vmcnt(6)52; GCN-NEXT: v_mul_lo_u32 v7, v7, v753; GCN-NEXT: v_mul_lo_u32 v6, v6, v654; GCN-NEXT: v_mul_lo_u32 v5, v5, v555; GCN-NEXT: v_mul_lo_u32 v4, v4, v456; GCN-NEXT: s_waitcnt vmcnt(5)57; GCN-NEXT: v_mul_lo_u32 v11, v11, v1158; GCN-NEXT: v_mul_lo_u32 v10, v10, v1059; GCN-NEXT: v_mul_lo_u32 v9, v9, v960; GCN-NEXT: v_mul_lo_u32 v8, v8, v861; GCN-NEXT: s_waitcnt vmcnt(4)62; GCN-NEXT: v_mul_lo_u32 v15, v15, v1563; GCN-NEXT: v_mul_lo_u32 v14, v14, v1464; GCN-NEXT: v_mul_lo_u32 v13, v13, v1365; GCN-NEXT: v_mul_lo_u32 v12, v12, v1266; GCN-NEXT: s_waitcnt vmcnt(3)67; GCN-NEXT: v_mul_lo_u32 v19, v19, v1968; GCN-NEXT: v_mul_lo_u32 v18, v18, v1869; GCN-NEXT: v_mul_lo_u32 v17, v17, v1770; GCN-NEXT: v_mul_lo_u32 v16, v16, v1671; GCN-NEXT: s_waitcnt vmcnt(2)72; GCN-NEXT: v_mul_lo_u32 v23, v23, v2373; GCN-NEXT: v_mul_lo_u32 v22, v22, v2274; GCN-NEXT: v_mul_lo_u32 v21, v21, v2175; GCN-NEXT: v_mul_lo_u32 v20, v20, v2076; GCN-NEXT: s_waitcnt vmcnt(1)77; GCN-NEXT: v_mul_lo_u32 v27, v27, v2778; GCN-NEXT: v_mul_lo_u32 v26, v26, v2679; GCN-NEXT: v_mul_lo_u32 v25, v25, v2580; GCN-NEXT: v_mul_lo_u32 v24, v24, v2481; GCN-NEXT: s_waitcnt vmcnt(0)82; GCN-NEXT: v_mul_lo_u32 v31, v31, v3183; GCN-NEXT: v_mul_lo_u32 v30, v30, v3084; GCN-NEXT: v_mul_lo_u32 v29, v29, v2985; GCN-NEXT: v_mul_lo_u32 v28, v28, v2886; GCN-NEXT: global_store_dwordx4 v32, v[28:31], s[2:3] offset:11287; GCN-NEXT: global_store_dwordx4 v32, v[24:27], s[2:3] offset:9688; GCN-NEXT: global_store_dwordx4 v32, v[20:23], s[2:3] offset:8089; GCN-NEXT: global_store_dwordx4 v32, v[16:19], s[2:3] offset:6490; GCN-NEXT: global_store_dwordx4 v32, v[12:15], s[2:3] offset:4891; GCN-NEXT: global_store_dwordx4 v32, v[8:11], s[2:3] offset:3292; GCN-NEXT: global_store_dwordx4 v32, v[4:7], s[2:3] offset:1693; 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_endpgm97;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], 0x24101; EXACTCUTOFF-NEXT: v_and_b32_e32 v0, 0x3ff, v0102; EXACTCUTOFF-NEXT: v_lshlrev_b32_e32 v32, 7, v0103; EXACTCUTOFF-NEXT: ; kill: killed $sgpr0_sgpr1104; 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:16107; EXACTCUTOFF-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:32108; EXACTCUTOFF-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:48109; EXACTCUTOFF-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:64110; EXACTCUTOFF-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:80111; EXACTCUTOFF-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:96112; EXACTCUTOFF-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:112113; 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, v3116; EXACTCUTOFF-NEXT: v_mul_lo_u32 v2, v2, v2117; EXACTCUTOFF-NEXT: v_mul_lo_u32 v1, v1, v1118; EXACTCUTOFF-NEXT: v_mul_lo_u32 v0, v0, v0119; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(6)120; EXACTCUTOFF-NEXT: v_mul_lo_u32 v7, v7, v7121; EXACTCUTOFF-NEXT: v_mul_lo_u32 v6, v6, v6122; EXACTCUTOFF-NEXT: v_mul_lo_u32 v5, v5, v5123; EXACTCUTOFF-NEXT: v_mul_lo_u32 v4, v4, v4124; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(5)125; EXACTCUTOFF-NEXT: v_mul_lo_u32 v11, v11, v11126; EXACTCUTOFF-NEXT: v_mul_lo_u32 v10, v10, v10127; EXACTCUTOFF-NEXT: v_mul_lo_u32 v9, v9, v9128; EXACTCUTOFF-NEXT: v_mul_lo_u32 v8, v8, v8129; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(4)130; EXACTCUTOFF-NEXT: v_mul_lo_u32 v15, v15, v15131; EXACTCUTOFF-NEXT: v_mul_lo_u32 v14, v14, v14132; EXACTCUTOFF-NEXT: v_mul_lo_u32 v13, v13, v13133; EXACTCUTOFF-NEXT: v_mul_lo_u32 v12, v12, v12134; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(3)135; EXACTCUTOFF-NEXT: v_mul_lo_u32 v19, v19, v19136; EXACTCUTOFF-NEXT: v_mul_lo_u32 v18, v18, v18137; EXACTCUTOFF-NEXT: v_mul_lo_u32 v17, v17, v17138; EXACTCUTOFF-NEXT: v_mul_lo_u32 v16, v16, v16139; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(2)140; EXACTCUTOFF-NEXT: v_mul_lo_u32 v23, v23, v23141; EXACTCUTOFF-NEXT: v_mul_lo_u32 v22, v22, v22142; EXACTCUTOFF-NEXT: v_mul_lo_u32 v21, v21, v21143; EXACTCUTOFF-NEXT: v_mul_lo_u32 v20, v20, v20144; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(1)145; EXACTCUTOFF-NEXT: v_mul_lo_u32 v27, v27, v27146; EXACTCUTOFF-NEXT: v_mul_lo_u32 v26, v26, v26147; EXACTCUTOFF-NEXT: v_mul_lo_u32 v25, v25, v25148; EXACTCUTOFF-NEXT: v_mul_lo_u32 v24, v24, v24149; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(0)150; EXACTCUTOFF-NEXT: v_mul_lo_u32 v31, v31, v31151; EXACTCUTOFF-NEXT: v_mul_lo_u32 v30, v30, v30152; EXACTCUTOFF-NEXT: v_mul_lo_u32 v29, v29, v29153; EXACTCUTOFF-NEXT: v_mul_lo_u32 v28, v28, v28154; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[28:31], s[2:3] offset:112155; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[24:27], s[2:3] offset:96156; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[20:23], s[2:3] offset:80157; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[16:19], s[2:3] offset:64158; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[12:15], s[2:3] offset:48159; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[8:11], s[2:3] offset:32160; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[4:7], s[2:3] offset:16161; 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_endpgm165 %tid = call i32 @llvm.amdgcn.workitem.id.x() #2166 %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid167 %load = load <32 x i32>, ptr addrspace(1) %gep1168 %mul = mul <32 x i32> %load, %load169 %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid170 store <32 x i32> %mul, ptr addrspace(1) %gep2171 ; 8 VMEM read172 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 8, i32 0)173 ; 30 VALU174 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 30, i32 0)175 ; 8 VMEM write176 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 8, i32 0)177 ret void178}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], 0x24184; GCN-NEXT: v_and_b32_e32 v0, 0x3ff, v0185; GCN-NEXT: v_lshlrev_b32_e32 v32, 7, v0186; GCN-NEXT: s_waitcnt lgkmcnt(0)187; GCN-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:16188; 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, v29194; GCN-NEXT: s_waitcnt vmcnt(0)195; GCN-NEXT: v_mul_lo_u32 v3, v3, v3196; GCN-NEXT: v_mul_lo_u32 v2, v2, v2197; GCN-NEXT: global_load_dwordx4 v[4:7], v32, s[0:1] offset:112198; GCN-NEXT: v_mul_lo_u32 v1, v1, v1199; GCN-NEXT: v_mul_lo_u32 v0, v0, v0200; GCN-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:96201; GCN-NEXT: v_mul_lo_u32 v28, v28, v28202; GCN-NEXT: v_mul_lo_u32 v31, v31, v31203; GCN-NEXT: v_mul_lo_u32 v30, v30, v30204; 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, v7210; GCN-NEXT: v_mul_lo_u32 v6, v6, v6211; GCN-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:80212; GCN-NEXT: v_mul_lo_u32 v5, v5, v5213; GCN-NEXT: v_mul_lo_u32 v4, v4, v4214; GCN-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:48215; GCN-NEXT: s_waitcnt vmcnt(2)216; GCN-NEXT: v_mul_lo_u32 v11, v11, v11217; GCN-NEXT: v_mul_lo_u32 v10, v10, v10218; GCN-NEXT: v_mul_lo_u32 v9, v9, v9219; GCN-NEXT: v_mul_lo_u32 v8, v8, v8220; 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, v15226; GCN-NEXT: v_mul_lo_u32 v14, v14, v14227; GCN-NEXT: v_mul_lo_u32 v13, v13, v13228; GCN-NEXT: s_waitcnt vmcnt(0)229; GCN-NEXT: v_mul_lo_u32 v19, v19, v19230; GCN-NEXT: v_mul_lo_u32 v18, v18, v18231; GCN-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:64232; GCN-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:32233; GCN-NEXT: v_mul_lo_u32 v17, v17, v17234; GCN-NEXT: v_mul_lo_u32 v16, v16, v16235; GCN-NEXT: v_mul_lo_u32 v12, v12, v12236; 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, v23240; GCN-NEXT: s_waitcnt vmcnt(0)241; GCN-NEXT: v_mul_lo_u32 v27, v27, v27242; GCN-NEXT: v_mul_lo_u32 v26, v26, v26243; GCN-NEXT: v_mul_lo_u32 v25, v25, v25244; GCN-NEXT: v_mul_lo_u32 v24, v24, v24245; GCN-NEXT: v_mul_lo_u32 v22, v22, v22246; GCN-NEXT: v_mul_lo_u32 v21, v21, v21247; GCN-NEXT: v_mul_lo_u32 v20, v20, v20248; GCN-NEXT: global_store_dwordx4 v32, v[4:7], s[2:3] offset:112249; GCN-NEXT: global_store_dwordx4 v32, v[8:11], s[2:3] offset:96250; GCN-NEXT: global_store_dwordx4 v32, v[12:15], s[2:3] offset:80251; GCN-NEXT: global_store_dwordx4 v32, v[20:23], s[2:3] offset:64252; GCN-NEXT: global_store_dwordx4 v32, v[16:19], s[2:3] offset:48253; GCN-NEXT: global_store_dwordx4 v32, v[24:27], s[2:3] offset:32254; GCN-NEXT: global_store_dwordx4 v32, v[28:31], s[2:3] offset:16255; 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_endpgm261;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], 0x24265; EXACTCUTOFF-NEXT: v_and_b32_e32 v0, 0x3ff, v0266; EXACTCUTOFF-NEXT: v_lshlrev_b32_e32 v32, 7, v0267; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)268; EXACTCUTOFF-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:16269; 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, v29275; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(0)276; EXACTCUTOFF-NEXT: v_mul_lo_u32 v3, v3, v3277; EXACTCUTOFF-NEXT: v_mul_lo_u32 v2, v2, v2278; EXACTCUTOFF-NEXT: global_load_dwordx4 v[4:7], v32, s[0:1] offset:112279; EXACTCUTOFF-NEXT: v_mul_lo_u32 v1, v1, v1280; EXACTCUTOFF-NEXT: v_mul_lo_u32 v0, v0, v0281; EXACTCUTOFF-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:96282; EXACTCUTOFF-NEXT: v_mul_lo_u32 v28, v28, v28283; EXACTCUTOFF-NEXT: v_mul_lo_u32 v31, v31, v31284; EXACTCUTOFF-NEXT: v_mul_lo_u32 v30, v30, v30285; 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, v7291; EXACTCUTOFF-NEXT: v_mul_lo_u32 v6, v6, v6292; EXACTCUTOFF-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:80293; EXACTCUTOFF-NEXT: v_mul_lo_u32 v5, v5, v5294; EXACTCUTOFF-NEXT: v_mul_lo_u32 v4, v4, v4295; EXACTCUTOFF-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:48296; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(2)297; EXACTCUTOFF-NEXT: v_mul_lo_u32 v11, v11, v11298; EXACTCUTOFF-NEXT: v_mul_lo_u32 v10, v10, v10299; EXACTCUTOFF-NEXT: v_mul_lo_u32 v9, v9, v9300; EXACTCUTOFF-NEXT: v_mul_lo_u32 v8, v8, v8301; 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, v15307; EXACTCUTOFF-NEXT: v_mul_lo_u32 v14, v14, v14308; EXACTCUTOFF-NEXT: v_mul_lo_u32 v13, v13, v13309; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(0)310; EXACTCUTOFF-NEXT: v_mul_lo_u32 v19, v19, v19311; EXACTCUTOFF-NEXT: v_mul_lo_u32 v18, v18, v18312; EXACTCUTOFF-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:64313; EXACTCUTOFF-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:32314; EXACTCUTOFF-NEXT: v_mul_lo_u32 v17, v17, v17315; EXACTCUTOFF-NEXT: v_mul_lo_u32 v16, v16, v16316; EXACTCUTOFF-NEXT: v_mul_lo_u32 v12, v12, v12317; 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, v23321; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(0)322; EXACTCUTOFF-NEXT: v_mul_lo_u32 v27, v27, v27323; EXACTCUTOFF-NEXT: v_mul_lo_u32 v26, v26, v26324; EXACTCUTOFF-NEXT: v_mul_lo_u32 v25, v25, v25325; EXACTCUTOFF-NEXT: v_mul_lo_u32 v24, v24, v24326; EXACTCUTOFF-NEXT: v_mul_lo_u32 v22, v22, v22327; EXACTCUTOFF-NEXT: v_mul_lo_u32 v21, v21, v21328; EXACTCUTOFF-NEXT: v_mul_lo_u32 v20, v20, v20329; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[4:7], s[2:3] offset:112330; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[8:11], s[2:3] offset:96331; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[12:15], s[2:3] offset:80332; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[20:23], s[2:3] offset:64333; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[16:19], s[2:3] offset:48334; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[24:27], s[2:3] offset:32335; EXACTCUTOFF-NEXT: global_store_dwordx4 v32, v[28:31], s[2:3] offset:16336; 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_endpgm342 %tid = call i32 @llvm.amdgcn.workitem.id.x() #2343 %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid344 %load = load <32 x i32>, ptr addrspace(1) %gep1345 %mul = mul <32 x i32> %load, %load346 %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid347 store <32 x i32> %mul, ptr addrspace(1) %gep2348 ; 1 VMEM read349 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)350 ; 2 VALU351 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)352 ; 1 VMEM read353 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)354 ; 2 VALU355 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)356 ; 1 VMEM read357 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)358 ; 2 VALU359 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)360 ; 1 VMEM read361 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)362 ; 2 VALU363 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)364 ; 1 VMEM read365 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)366 ; 2 VALU367 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)368 ; 1 VMEM read369 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)370 ; 2 VALU371 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)372 ; 1 VMEM read373 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)374 ; 2 VALU375 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)376 ; 1 VMEM read377 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)378 ; 2 VALU379 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)380 ; 8 VMEM write381 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 8, i32 0)382 ret void383}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], 0x24389; GCN-NEXT: v_and_b32_e32 v0, 0x3ff, v0390; GCN-NEXT: v_lshlrev_b32_e32 v16, 7, v0391; GCN-NEXT: ; kill: killed $sgpr0_sgpr1392; GCN-NEXT: s_waitcnt lgkmcnt(0)393; GCN-NEXT: global_load_dwordx4 v[12:15], v16, s[0:1] offset:32394; 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, v13398; GCN-NEXT: v_mul_lo_u32 v12, v12, v12399; GCN-NEXT: v_mul_lo_u32 v15, v15, v15400; GCN-NEXT: v_mul_lo_u32 v14, v14, v14401; GCN-NEXT: global_store_dwordx4 v16, v[12:15], s[2:3] offset:32402; 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, v3407; GCN-NEXT: v_mul_lo_u32 v2, v2, v2408; GCN-NEXT: v_mul_lo_u32 v1, v1, v1409; GCN-NEXT: v_mul_lo_u32 v0, v0, v0410; 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:112412; 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, v3415; GCN-NEXT: v_mul_lo_u32 v2, v2, v2416; GCN-NEXT: v_mul_lo_u32 v1, v1, v1417; GCN-NEXT: v_mul_lo_u32 v0, v0, v0418; GCN-NEXT: global_store_dwordx4 v16, v[0:3], s[2:3] offset:112419; GCN-NEXT: global_load_dwordx4 v[0:3], v16, s[0:1] offset:96420; GCN-NEXT: s_waitcnt vmcnt(0)421; GCN-NEXT: v_mul_lo_u32 v3, v3, v3422; GCN-NEXT: v_mul_lo_u32 v2, v2, v2423; GCN-NEXT: v_mul_lo_u32 v1, v1, v1424; GCN-NEXT: v_mul_lo_u32 v0, v0, v0425; GCN-NEXT: global_store_dwordx4 v16, v[0:3], s[2:3] offset:96426; GCN-NEXT: global_load_dwordx4 v[0:3], v16, s[0:1] offset:80427; GCN-NEXT: s_waitcnt vmcnt(0)428; GCN-NEXT: v_mul_lo_u32 v3, v3, v3429; GCN-NEXT: v_mul_lo_u32 v2, v2, v2430; GCN-NEXT: v_mul_lo_u32 v1, v1, v1431; GCN-NEXT: v_mul_lo_u32 v0, v0, v0432; GCN-NEXT: global_store_dwordx4 v16, v[0:3], s[2:3] offset:80433; GCN-NEXT: global_load_dwordx4 v[4:7], v16, s[0:1] offset:48434; 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, v7442; GCN-NEXT: v_mul_lo_u32 v6, v6, v6443; GCN-NEXT: v_mul_lo_u32 v5, v5, v5444; GCN-NEXT: v_mul_lo_u32 v4, v4, v4445; GCN-NEXT: global_store_dwordx4 v16, v[4:7], s[2:3] offset:48446; GCN-NEXT: global_load_dwordx4 v[8:11], v16, s[0:1] offset:16447; 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, v9455; GCN-NEXT: v_mul_lo_u32 v8, v8, v8456; GCN-NEXT: v_mul_lo_u32 v11, v11, v11457; GCN-NEXT: v_mul_lo_u32 v10, v10, v10458; GCN-NEXT: global_store_dwordx4 v16, v[8:11], s[2:3] offset:16459; GCN-NEXT: global_load_dwordx4 v[8:11], v16, s[0:1] offset:64460; 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, v11468; GCN-NEXT: v_mul_lo_u32 v10, v10, v10469; GCN-NEXT: v_mul_lo_u32 v9, v9, v9470; GCN-NEXT: v_mul_lo_u32 v8, v8, v8471; GCN-NEXT: global_store_dwordx4 v16, v[8:11], s[2:3] offset:64472; GCN-NEXT: ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)473; GCN-NEXT: s_endpgm474;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], 0x24478; EXACTCUTOFF-NEXT: v_and_b32_e32 v0, 0x3ff, v0479; EXACTCUTOFF-NEXT: v_lshlrev_b32_e32 v16, 7, v0480; EXACTCUTOFF-NEXT: ; kill: killed $sgpr0_sgpr1481; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)482; EXACTCUTOFF-NEXT: global_load_dwordx4 v[12:15], v16, s[0:1] offset:32483; 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, v13487; EXACTCUTOFF-NEXT: v_mul_lo_u32 v12, v12, v12488; EXACTCUTOFF-NEXT: v_mul_lo_u32 v15, v15, v15489; EXACTCUTOFF-NEXT: v_mul_lo_u32 v14, v14, v14490; EXACTCUTOFF-NEXT: global_store_dwordx4 v16, v[12:15], s[2:3] offset:32491; 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, v3496; EXACTCUTOFF-NEXT: v_mul_lo_u32 v2, v2, v2497; EXACTCUTOFF-NEXT: v_mul_lo_u32 v1, v1, v1498; EXACTCUTOFF-NEXT: v_mul_lo_u32 v0, v0, v0499; 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:112501; 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, v3504; EXACTCUTOFF-NEXT: v_mul_lo_u32 v2, v2, v2505; EXACTCUTOFF-NEXT: v_mul_lo_u32 v1, v1, v1506; EXACTCUTOFF-NEXT: v_mul_lo_u32 v0, v0, v0507; EXACTCUTOFF-NEXT: global_store_dwordx4 v16, v[0:3], s[2:3] offset:112508; EXACTCUTOFF-NEXT: global_load_dwordx4 v[0:3], v16, s[0:1] offset:96509; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(0)510; EXACTCUTOFF-NEXT: v_mul_lo_u32 v3, v3, v3511; EXACTCUTOFF-NEXT: v_mul_lo_u32 v2, v2, v2512; EXACTCUTOFF-NEXT: v_mul_lo_u32 v1, v1, v1513; EXACTCUTOFF-NEXT: v_mul_lo_u32 v0, v0, v0514; EXACTCUTOFF-NEXT: global_store_dwordx4 v16, v[0:3], s[2:3] offset:96515; EXACTCUTOFF-NEXT: global_load_dwordx4 v[0:3], v16, s[0:1] offset:80516; EXACTCUTOFF-NEXT: s_waitcnt vmcnt(0)517; EXACTCUTOFF-NEXT: v_mul_lo_u32 v3, v3, v3518; EXACTCUTOFF-NEXT: v_mul_lo_u32 v2, v2, v2519; EXACTCUTOFF-NEXT: v_mul_lo_u32 v1, v1, v1520; EXACTCUTOFF-NEXT: v_mul_lo_u32 v0, v0, v0521; EXACTCUTOFF-NEXT: global_store_dwordx4 v16, v[0:3], s[2:3] offset:80522; EXACTCUTOFF-NEXT: global_load_dwordx4 v[4:7], v16, s[0:1] offset:48523; 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, v7531; EXACTCUTOFF-NEXT: v_mul_lo_u32 v6, v6, v6532; EXACTCUTOFF-NEXT: v_mul_lo_u32 v5, v5, v5533; EXACTCUTOFF-NEXT: v_mul_lo_u32 v4, v4, v4534; EXACTCUTOFF-NEXT: global_store_dwordx4 v16, v[4:7], s[2:3] offset:48535; EXACTCUTOFF-NEXT: global_load_dwordx4 v[8:11], v16, s[0:1] offset:16536; 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, v9544; EXACTCUTOFF-NEXT: v_mul_lo_u32 v8, v8, v8545; EXACTCUTOFF-NEXT: v_mul_lo_u32 v11, v11, v11546; EXACTCUTOFF-NEXT: v_mul_lo_u32 v10, v10, v10547; EXACTCUTOFF-NEXT: global_store_dwordx4 v16, v[8:11], s[2:3] offset:16548; EXACTCUTOFF-NEXT: global_load_dwordx4 v[8:11], v16, s[0:1] offset:64549; 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, v11557; EXACTCUTOFF-NEXT: v_mul_lo_u32 v10, v10, v10558; EXACTCUTOFF-NEXT: v_mul_lo_u32 v9, v9, v9559; EXACTCUTOFF-NEXT: v_mul_lo_u32 v8, v8, v8560; EXACTCUTOFF-NEXT: global_store_dwordx4 v16, v[8:11], s[2:3] offset:64561; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000040) size(1) SyncID(0)562; EXACTCUTOFF-NEXT: s_endpgm563 %tid = call i32 @llvm.amdgcn.workitem.id.x() #2564 %gep1 = getelementptr <32 x i32>, ptr addrspace(1) %in, i32 %tid565 %load = load <32 x i32>, ptr addrspace(1) %gep1566 %mul = mul <32 x i32> %load, %load567 %gep2 = getelementptr <32 x i32>, ptr addrspace(1) %out, i32 %tid568 store <32 x i32> %mul, ptr addrspace(1) %gep2569 ; 1 VMEM read570 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)571 ; 2 VALU572 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)573 ; 1 VMEM write574 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)575 ; 1 VMEM read576 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)577 ; 2 VALU578 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)579 ; 1 VMEM write580 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)581 ; 1 VMEM read582 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)583 ; 2 VALU584 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)585 ; 1 VMEM write586 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)587 ; 1 VMEM read588 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)589 ; 2 VALU590 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)591 ; 1 VMEM write592 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)593 ; 1 VMEM read594 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)595 ; 2 VALU596 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)597 ; 1 VMEM write598 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)599 ; 1 VMEM read600 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)601 ; 2 VALU602 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)603 ; 1 VMEM write604 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)605 ; 1 VMEM read606 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)607 ; 2 VALU608 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)609 ; 1 VMEM write610 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)611 ; 1 VMEM read612 call void @llvm.amdgcn.sched.group.barrier(i32 32, i32 1, i32 0)613 ; 2 VALU614 call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 2, i32 0)615 ; 1 VMEM write616 call void @llvm.amdgcn.sched.group.barrier(i32 64, i32 1, i32 0)617 ret void618}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: ; %entry623; GCN-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24624; GCN-NEXT: v_lshlrev_b32_e32 v0, 7, v0625; GCN-NEXT: v_and_b32_e32 v0, 0x1ff80, v0626; GCN-NEXT: v_mov_b32_e32 v2, 1.0627; GCN-NEXT: v_mov_b32_e32 v1, 2.0628; GCN-NEXT: s_waitcnt lgkmcnt(0)629; GCN-NEXT: v_add_u32_e32 v3, s0, v0630; GCN-NEXT: ds_read_b128 a[60:63], v3 offset:112631; GCN-NEXT: ds_read_b128 a[56:59], v3 offset:96632; GCN-NEXT: ds_read_b128 a[52:55], v3 offset:80633; GCN-NEXT: ds_read_b128 a[48:51], v3 offset:64634; GCN-NEXT: ds_read_b128 a[32:35], v3635; GCN-NEXT: ds_read_b128 a[36:39], v3 offset:16636; GCN-NEXT: ds_read_b128 a[40:43], v3 offset:32637; GCN-NEXT: ds_read_b128 a[44:47], v3 offset:48638; GCN-NEXT: v_add_u32_e32 v4, 0x6000, v3639; GCN-NEXT: ds_read_b128 a[28:31], v3 offset:8304640; GCN-NEXT: ds_read_b128 a[24:27], v3 offset:8288641; GCN-NEXT: ds_read_b128 a[20:23], v3 offset:8272642; GCN-NEXT: ds_read_b128 a[16:19], v3 offset:8256643; GCN-NEXT: ds_read_b128 a[12:15], v3 offset:8240644; GCN-NEXT: ds_read_b128 a[8:11], v3 offset:8224645; GCN-NEXT: ds_read_b128 a[4:7], v3 offset:8208646; GCN-NEXT: ds_read_b128 a[0:3], v3 offset:8192647; GCN-NEXT: ds_read_b128 a[92:95], v3 offset:24688648; GCN-NEXT: ds_read_b128 a[88:91], v3 offset:24672649; GCN-NEXT: ds_read_b128 a[84:87], v3 offset:24656650; GCN-NEXT: ds_read_b128 a[80:83], v3 offset:24640651; GCN-NEXT: ds_read_b128 a[76:79], v3 offset:24624652; GCN-NEXT: ds_read_b128 a[72:75], v3 offset:24608653; GCN-NEXT: ds_read_b128 a[68:71], v3 offset:24592654; GCN-NEXT: ds_read_b128 a[64:67], v3 offset:24576655; GCN-NEXT: ds_read_b128 a[124:127], v3 offset:49264656; GCN-NEXT: ds_read_b128 a[120:123], v3 offset:49248657; GCN-NEXT: ds_read_b128 a[116:119], v3 offset:49232658; GCN-NEXT: ds_read_b128 a[112:115], v3 offset:49216659; GCN-NEXT: ds_read_b128 a[108:111], v3 offset:49200660; GCN-NEXT: ds_read_b128 a[104:107], v3 offset:49184661; GCN-NEXT: ds_read_b128 a[100:103], v3 offset:49168662; GCN-NEXT: ds_read_b128 a[96:99], v3 offset:49152663; GCN-NEXT: ds_read_b128 a[156:159], v4 offset:57456664; GCN-NEXT: ds_read_b128 a[152:155], v4 offset:57440665; GCN-NEXT: ds_read_b128 a[148:151], v4 offset:57424666; GCN-NEXT: ds_read_b128 a[144:147], v4 offset:57408667; GCN-NEXT: ds_read_b128 a[128:131], v4 offset:57344668; GCN-NEXT: ds_read_b128 a[132:135], v4 offset:57360669; GCN-NEXT: ds_read_b128 a[136:139], v4 offset:57376670; GCN-NEXT: ds_read_b128 a[140:143], v4 offset:57392671; GCN-NEXT: s_waitcnt lgkmcnt(14)672; GCN-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v2, v1, a[32:63]673; GCN-NEXT: v_add_u32_e32 v0, s1, v0674; GCN-NEXT: ; sched_group_barrier mask(0x00000100) size(40) SyncID(0)675; GCN-NEXT: s_waitcnt lgkmcnt(8)676; GCN-NEXT: v_mfma_f32_32x32x1f32 a[96:127], v2, v1, a[96:127]677; GCN-NEXT: v_mfma_f32_32x32x1f32 a[64:95], v2, v1, a[64:95]678; GCN-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v1, a[0:31]679; GCN-NEXT: s_waitcnt lgkmcnt(0)680; GCN-NEXT: v_mfma_f32_32x32x1f32 a[128:159], v2, v1, a[128:159]681; GCN-NEXT: s_nop 11682; GCN-NEXT: ds_write_b128 v0, a[60:63] offset:112683; GCN-NEXT: ds_write_b128 v0, a[56:59] offset:96684; GCN-NEXT: ds_write_b128 v0, a[52:55] offset:80685; GCN-NEXT: ds_write_b128 v0, a[48:51] offset:64686; GCN-NEXT: ds_write_b128 v0, a[44:47] offset:48687; GCN-NEXT: ds_write_b128 v0, a[40:43] offset:32688; GCN-NEXT: ds_write_b128 v0, a[36:39] offset:16689; GCN-NEXT: ds_write_b128 v0, a[32:35]690; GCN-NEXT: v_mov_b32_e32 v0, s1691; GCN-NEXT: ds_write_b128 v0, a[24:27] offset:8288692; GCN-NEXT: ds_write_b128 v0, a[28:31] offset:8304693; GCN-NEXT: ds_write_b128 v0, a[16:19] offset:8256694; GCN-NEXT: ds_write_b128 v0, a[20:23] offset:8272695; GCN-NEXT: ds_write_b128 v0, a[8:11] offset:8224696; GCN-NEXT: ds_write_b128 v0, a[12:15] offset:8240697; GCN-NEXT: ds_write_b128 v0, a[0:3] offset:8192698; GCN-NEXT: ds_write_b128 v0, a[4:7] offset:8208699; GCN-NEXT: ds_write_b128 v0, a[88:91] offset:16480700; GCN-NEXT: ds_write_b128 v0, a[92:95] offset:16496701; GCN-NEXT: ds_write_b128 v0, a[80:83] offset:16448702; GCN-NEXT: ds_write_b128 v0, a[84:87] offset:16464703; GCN-NEXT: ds_write_b128 v0, a[72:75] offset:16416704; GCN-NEXT: ds_write_b128 v0, a[76:79] offset:16432705; GCN-NEXT: ds_write_b128 v0, a[64:67] offset:16384706; GCN-NEXT: ds_write_b128 v0, a[68:71] offset:16400707; GCN-NEXT: ds_write_b128 v0, a[120:123] offset:24672708; GCN-NEXT: ds_write_b128 v0, a[124:127] offset:24688709; GCN-NEXT: ds_write_b128 v0, a[112:115] offset:24640710; GCN-NEXT: ds_write_b128 v0, a[116:119] offset:24656711; GCN-NEXT: ds_write_b128 v0, a[104:107] offset:24608712; GCN-NEXT: ds_write_b128 v0, a[108:111] offset:24624713; GCN-NEXT: ds_write_b128 v0, a[96:99] offset:24576714; GCN-NEXT: ds_write_b128 v0, a[100:103] offset:24592715; GCN-NEXT: ds_write_b128 v0, a[152:155] offset:32864716; GCN-NEXT: ds_write_b128 v0, a[156:159] offset:32880717; GCN-NEXT: ds_write_b128 v0, a[144:147] offset:32832718; GCN-NEXT: ds_write_b128 v0, a[148:151] offset:32848719; GCN-NEXT: ds_write_b128 v0, a[136:139] offset:32800720; GCN-NEXT: ds_write_b128 v0, a[140:143] offset:32816721; GCN-NEXT: ds_write_b128 v0, a[128:131] offset:32768722; GCN-NEXT: ds_write_b128 v0, a[132:135] offset:32784723; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(5) SyncID(0)724; GCN-NEXT: ; sched_group_barrier mask(0x00000200) size(40) SyncID(0)725; GCN-NEXT: s_endpgm726;727; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_MFMA_cluster:728; EXACTCUTOFF: ; %bb.0: ; %entry729; EXACTCUTOFF-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24730; EXACTCUTOFF-NEXT: v_lshlrev_b32_e32 v0, 7, v0731; EXACTCUTOFF-NEXT: v_and_b32_e32 v0, 0x1ff80, v0732; EXACTCUTOFF-NEXT: v_mov_b32_e32 v2, 1.0733; EXACTCUTOFF-NEXT: v_mov_b32_e32 v1, 2.0734; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)735; EXACTCUTOFF-NEXT: v_add_u32_e32 v3, s0, v0736; EXACTCUTOFF-NEXT: ds_read_b128 a[60:63], v3 offset:112737; EXACTCUTOFF-NEXT: ds_read_b128 a[56:59], v3 offset:96738; EXACTCUTOFF-NEXT: ds_read_b128 a[52:55], v3 offset:80739; EXACTCUTOFF-NEXT: ds_read_b128 a[48:51], v3 offset:64740; EXACTCUTOFF-NEXT: ds_read_b128 a[32:35], v3741; EXACTCUTOFF-NEXT: ds_read_b128 a[36:39], v3 offset:16742; EXACTCUTOFF-NEXT: ds_read_b128 a[40:43], v3 offset:32743; EXACTCUTOFF-NEXT: ds_read_b128 a[44:47], v3 offset:48744; EXACTCUTOFF-NEXT: v_add_u32_e32 v4, 0x6000, v3745; EXACTCUTOFF-NEXT: ds_read_b128 a[28:31], v3 offset:8304746; EXACTCUTOFF-NEXT: ds_read_b128 a[24:27], v3 offset:8288747; EXACTCUTOFF-NEXT: ds_read_b128 a[20:23], v3 offset:8272748; EXACTCUTOFF-NEXT: ds_read_b128 a[16:19], v3 offset:8256749; EXACTCUTOFF-NEXT: ds_read_b128 a[12:15], v3 offset:8240750; EXACTCUTOFF-NEXT: ds_read_b128 a[8:11], v3 offset:8224751; EXACTCUTOFF-NEXT: ds_read_b128 a[4:7], v3 offset:8208752; EXACTCUTOFF-NEXT: ds_read_b128 a[0:3], v3 offset:8192753; EXACTCUTOFF-NEXT: ds_read_b128 a[92:95], v3 offset:24688754; EXACTCUTOFF-NEXT: ds_read_b128 a[88:91], v3 offset:24672755; EXACTCUTOFF-NEXT: ds_read_b128 a[84:87], v3 offset:24656756; EXACTCUTOFF-NEXT: ds_read_b128 a[80:83], v3 offset:24640757; EXACTCUTOFF-NEXT: ds_read_b128 a[76:79], v3 offset:24624758; EXACTCUTOFF-NEXT: ds_read_b128 a[72:75], v3 offset:24608759; EXACTCUTOFF-NEXT: ds_read_b128 a[68:71], v3 offset:24592760; EXACTCUTOFF-NEXT: ds_read_b128 a[64:67], v3 offset:24576761; EXACTCUTOFF-NEXT: ds_read_b128 a[124:127], v3 offset:49264762; EXACTCUTOFF-NEXT: ds_read_b128 a[120:123], v3 offset:49248763; EXACTCUTOFF-NEXT: ds_read_b128 a[116:119], v3 offset:49232764; EXACTCUTOFF-NEXT: ds_read_b128 a[112:115], v3 offset:49216765; EXACTCUTOFF-NEXT: ds_read_b128 a[108:111], v3 offset:49200766; EXACTCUTOFF-NEXT: ds_read_b128 a[104:107], v3 offset:49184767; EXACTCUTOFF-NEXT: ds_read_b128 a[100:103], v3 offset:49168768; EXACTCUTOFF-NEXT: ds_read_b128 a[96:99], v3 offset:49152769; EXACTCUTOFF-NEXT: ds_read_b128 a[156:159], v4 offset:57456770; EXACTCUTOFF-NEXT: ds_read_b128 a[152:155], v4 offset:57440771; EXACTCUTOFF-NEXT: ds_read_b128 a[148:151], v4 offset:57424772; EXACTCUTOFF-NEXT: ds_read_b128 a[144:147], v4 offset:57408773; EXACTCUTOFF-NEXT: ds_read_b128 a[128:131], v4 offset:57344774; EXACTCUTOFF-NEXT: ds_read_b128 a[132:135], v4 offset:57360775; EXACTCUTOFF-NEXT: ds_read_b128 a[136:139], v4 offset:57376776; EXACTCUTOFF-NEXT: ds_read_b128 a[140:143], v4 offset:57392777; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(14)778; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v2, v1, a[32:63]779; EXACTCUTOFF-NEXT: v_add_u32_e32 v0, s1, v0780; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000100) size(40) SyncID(0)781; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(8)782; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[96:127], v2, v1, a[96:127]783; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[64:95], v2, v1, a[64:95]784; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v1, a[0:31]785; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)786; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[128:159], v2, v1, a[128:159]787; EXACTCUTOFF-NEXT: s_nop 11788; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[60:63] offset:112789; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[56:59] offset:96790; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[52:55] offset:80791; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[48:51] offset:64792; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[44:47] offset:48793; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[40:43] offset:32794; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[36:39] offset:16795; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[32:35]796; EXACTCUTOFF-NEXT: v_mov_b32_e32 v0, s1797; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[24:27] offset:8288798; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[28:31] offset:8304799; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[16:19] offset:8256800; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[20:23] offset:8272801; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[8:11] offset:8224802; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[12:15] offset:8240803; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[0:3] offset:8192804; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[4:7] offset:8208805; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[88:91] offset:16480806; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[92:95] offset:16496807; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[80:83] offset:16448808; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[84:87] offset:16464809; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[72:75] offset:16416810; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[76:79] offset:16432811; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[64:67] offset:16384812; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[68:71] offset:16400813; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[120:123] offset:24672814; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[124:127] offset:24688815; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[112:115] offset:24640816; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[116:119] offset:24656817; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[104:107] offset:24608818; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[108:111] offset:24624819; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[96:99] offset:24576820; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[100:103] offset:24592821; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[152:155] offset:32864822; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[156:159] offset:32880823; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[144:147] offset:32832824; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[148:151] offset:32848825; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[136:139] offset:32800826; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[140:143] offset:32816827; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[128:131] offset:32768828; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[132:135] offset:32784829; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(5) SyncID(0)830; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000200) size(40) SyncID(0)831; EXACTCUTOFF-NEXT: s_endpgm832entry:833 %idx = call i32 @llvm.amdgcn.workitem.id.x()834 %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx835 %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr836 %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64837 %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr838 %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128839 %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr840 %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192841 %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr842 %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256843 %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr844 %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)845 %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)846 %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)847 %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)848 %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)849 %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx850 store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr851 %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64852 store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr853 %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128854 store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr855 %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192856 store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr857 %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256858 store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr859 ; 40 DS read860 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 40, i32 0)861 ; 5 MFMA862 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 5, i32 0)863 ; 40 DS write864 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 40, i32 0)865 ret void866}867 868define amdgpu_kernel void @test_sched_group_barrier_pipeline_MFMA_interleave(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {869; GCN-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:870; GCN: ; %bb.0: ; %entry871; GCN-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24872; GCN-NEXT: v_lshlrev_b32_e32 v0, 7, v0873; GCN-NEXT: v_and_b32_e32 v0, 0x1ff80, v0874; GCN-NEXT: v_mov_b32_e32 v2, 1.0875; GCN-NEXT: v_mov_b32_e32 v3, 2.0876; GCN-NEXT: s_waitcnt lgkmcnt(0)877; GCN-NEXT: v_add_u32_e32 v1, s0, v0878; GCN-NEXT: ds_read_b128 a[28:31], v1 offset:112879; GCN-NEXT: ds_read_b128 a[24:27], v1 offset:96880; GCN-NEXT: ds_read_b128 a[20:23], v1 offset:80881; GCN-NEXT: ds_read_b128 a[16:19], v1 offset:64882; GCN-NEXT: ds_read_b128 a[0:3], v1883; GCN-NEXT: ds_read_b128 a[4:7], v1 offset:16884; GCN-NEXT: ds_read_b128 a[8:11], v1 offset:32885; GCN-NEXT: ds_read_b128 a[12:15], v1 offset:48886; GCN-NEXT: s_waitcnt lgkmcnt(0)887; GCN-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]888; GCN-NEXT: v_add_u32_e32 v0, s1, v0889; GCN-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)890; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)891; GCN-NEXT: s_nop 15892; GCN-NEXT: s_nop 1893; GCN-NEXT: ds_write_b128 v0, a[28:31] offset:112894; GCN-NEXT: ds_write_b128 v0, a[24:27] offset:96895; GCN-NEXT: ds_write_b128 v0, a[20:23] offset:80896; GCN-NEXT: ds_write_b128 v0, a[16:19] offset:64897; GCN-NEXT: ds_write_b128 v0, a[12:15] offset:48898; GCN-NEXT: ds_write_b128 v0, a[8:11] offset:32899; GCN-NEXT: ds_write_b128 v0, a[4:7] offset:16900; GCN-NEXT: ds_write_b128 v0, a[0:3]901; GCN-NEXT: ds_read_b128 a[28:31], v1 offset:8304902; GCN-NEXT: ds_read_b128 a[24:27], v1 offset:8288903; GCN-NEXT: ds_read_b128 a[20:23], v1 offset:8272904; GCN-NEXT: ds_read_b128 a[16:19], v1 offset:8256905; GCN-NEXT: ds_read_b128 a[12:15], v1 offset:8240906; GCN-NEXT: ds_read_b128 a[8:11], v1 offset:8224907; GCN-NEXT: ds_read_b128 a[4:7], v1 offset:8208908; GCN-NEXT: ds_read_b128 a[0:3], v1 offset:8192909; GCN-NEXT: s_waitcnt lgkmcnt(0)910; GCN-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]911; GCN-NEXT: v_mov_b32_e32 v0, s1912; GCN-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)913; GCN-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)914; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)915; GCN-NEXT: s_nop 15916; GCN-NEXT: s_nop 1917; GCN-NEXT: ds_write_b128 v0, a[24:27] offset:8288918; GCN-NEXT: ds_write_b128 v0, a[28:31] offset:8304919; GCN-NEXT: ds_write_b128 v0, a[16:19] offset:8256920; GCN-NEXT: ds_write_b128 v0, a[20:23] offset:8272921; GCN-NEXT: ds_write_b128 v0, a[8:11] offset:8224922; GCN-NEXT: ds_write_b128 v0, a[12:15] offset:8240923; GCN-NEXT: ds_write_b128 v0, a[0:3] offset:8192924; GCN-NEXT: ds_write_b128 v0, a[4:7] offset:8208925; GCN-NEXT: ds_read_b128 a[28:31], v1 offset:24688926; GCN-NEXT: ds_read_b128 a[24:27], v1 offset:24672927; GCN-NEXT: ds_read_b128 a[20:23], v1 offset:24656928; GCN-NEXT: ds_read_b128 a[16:19], v1 offset:24640929; GCN-NEXT: ds_read_b128 a[12:15], v1 offset:24624930; GCN-NEXT: ds_read_b128 a[8:11], v1 offset:24608931; GCN-NEXT: ds_read_b128 a[4:7], v1 offset:24592932; GCN-NEXT: ds_read_b128 a[0:3], v1 offset:24576933; GCN-NEXT: s_waitcnt lgkmcnt(0)934; GCN-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]935; GCN-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)936; GCN-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)937; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)938; GCN-NEXT: s_nop 15939; GCN-NEXT: s_nop 2940; GCN-NEXT: ds_write_b128 v0, a[24:27] offset:16480941; GCN-NEXT: ds_write_b128 v0, a[28:31] offset:16496942; GCN-NEXT: ds_write_b128 v0, a[16:19] offset:16448943; GCN-NEXT: ds_write_b128 v0, a[20:23] offset:16464944; GCN-NEXT: ds_write_b128 v0, a[8:11] offset:16416945; GCN-NEXT: ds_write_b128 v0, a[12:15] offset:16432946; GCN-NEXT: ds_write_b128 v0, a[0:3] offset:16384947; GCN-NEXT: ds_write_b128 v0, a[4:7] offset:16400948; GCN-NEXT: ds_read_b128 a[28:31], v1 offset:49264949; GCN-NEXT: ds_read_b128 a[24:27], v1 offset:49248950; GCN-NEXT: ds_read_b128 a[20:23], v1 offset:49232951; GCN-NEXT: ds_read_b128 a[16:19], v1 offset:49216952; GCN-NEXT: ds_read_b128 a[12:15], v1 offset:49200953; GCN-NEXT: ds_read_b128 a[8:11], v1 offset:49184954; GCN-NEXT: ds_read_b128 a[4:7], v1 offset:49168955; GCN-NEXT: ds_read_b128 a[0:3], v1 offset:49152956; GCN-NEXT: s_waitcnt lgkmcnt(0)957; GCN-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]958; GCN-NEXT: v_add_u32_e32 v1, 0x6000, v1959; GCN-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)960; GCN-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)961; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)962; GCN-NEXT: s_nop 15963; GCN-NEXT: s_nop 1964; GCN-NEXT: ds_write_b128 v0, a[24:27] offset:24672965; GCN-NEXT: ds_write_b128 v0, a[28:31] offset:24688966; GCN-NEXT: ds_write_b128 v0, a[16:19] offset:24640967; GCN-NEXT: ds_write_b128 v0, a[20:23] offset:24656968; GCN-NEXT: ds_write_b128 v0, a[8:11] offset:24608969; GCN-NEXT: ds_write_b128 v0, a[12:15] offset:24624970; GCN-NEXT: ds_write_b128 v0, a[0:3] offset:24576971; GCN-NEXT: ds_write_b128 v0, a[4:7] offset:24592972; GCN-NEXT: ds_read_b128 a[28:31], v1 offset:57456973; GCN-NEXT: ds_read_b128 a[24:27], v1 offset:57440974; GCN-NEXT: ds_read_b128 a[20:23], v1 offset:57424975; GCN-NEXT: ds_read_b128 a[16:19], v1 offset:57408976; GCN-NEXT: ds_read_b128 a[0:3], v1 offset:57344977; GCN-NEXT: ds_read_b128 a[4:7], v1 offset:57360978; GCN-NEXT: ds_read_b128 a[8:11], v1 offset:57376979; GCN-NEXT: ds_read_b128 a[12:15], v1 offset:57392980; GCN-NEXT: s_waitcnt lgkmcnt(0)981; GCN-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]982; GCN-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)983; GCN-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)984; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)985; GCN-NEXT: s_nop 15986; GCN-NEXT: s_nop 2987; GCN-NEXT: ds_write_b128 v0, a[24:27] offset:32864988; GCN-NEXT: ds_write_b128 v0, a[28:31] offset:32880989; GCN-NEXT: ds_write_b128 v0, a[16:19] offset:32832990; GCN-NEXT: ds_write_b128 v0, a[20:23] offset:32848991; GCN-NEXT: ds_write_b128 v0, a[8:11] offset:32800992; GCN-NEXT: ds_write_b128 v0, a[12:15] offset:32816993; GCN-NEXT: ds_write_b128 v0, a[0:3] offset:32768994; GCN-NEXT: ds_write_b128 v0, a[4:7] offset:32784995; GCN-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)996; GCN-NEXT: s_endpgm997;998; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:999; EXACTCUTOFF: ; %bb.0: ; %entry1000; EXACTCUTOFF-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x241001; EXACTCUTOFF-NEXT: v_lshlrev_b32_e32 v0, 7, v01002; EXACTCUTOFF-NEXT: v_and_b32_e32 v0, 0x1ff80, v01003; EXACTCUTOFF-NEXT: v_mov_b32_e32 v2, 1.01004; EXACTCUTOFF-NEXT: v_mov_b32_e32 v3, 2.01005; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1006; EXACTCUTOFF-NEXT: v_add_u32_e32 v1, s0, v01007; EXACTCUTOFF-NEXT: ds_read_b128 a[28:31], v1 offset:1121008; EXACTCUTOFF-NEXT: ds_read_b128 a[24:27], v1 offset:961009; EXACTCUTOFF-NEXT: ds_read_b128 a[20:23], v1 offset:801010; EXACTCUTOFF-NEXT: ds_read_b128 a[16:19], v1 offset:641011; EXACTCUTOFF-NEXT: ds_read_b128 a[0:3], v11012; EXACTCUTOFF-NEXT: ds_read_b128 a[4:7], v1 offset:161013; EXACTCUTOFF-NEXT: ds_read_b128 a[8:11], v1 offset:321014; EXACTCUTOFF-NEXT: ds_read_b128 a[12:15], v1 offset:481015; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1016; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]1017; EXACTCUTOFF-NEXT: v_add_u32_e32 v0, s1, v01018; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)1019; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1020; EXACTCUTOFF-NEXT: s_nop 151021; EXACTCUTOFF-NEXT: s_nop 11022; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[28:31] offset:1121023; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[24:27] offset:961024; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[20:23] offset:801025; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[16:19] offset:641026; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[12:15] offset:481027; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[8:11] offset:321028; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[4:7] offset:161029; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[0:3]1030; EXACTCUTOFF-NEXT: ds_read_b128 a[28:31], v1 offset:83041031; EXACTCUTOFF-NEXT: ds_read_b128 a[24:27], v1 offset:82881032; EXACTCUTOFF-NEXT: ds_read_b128 a[20:23], v1 offset:82721033; EXACTCUTOFF-NEXT: ds_read_b128 a[16:19], v1 offset:82561034; EXACTCUTOFF-NEXT: ds_read_b128 a[12:15], v1 offset:82401035; EXACTCUTOFF-NEXT: ds_read_b128 a[8:11], v1 offset:82241036; EXACTCUTOFF-NEXT: ds_read_b128 a[4:7], v1 offset:82081037; EXACTCUTOFF-NEXT: ds_read_b128 a[0:3], v1 offset:81921038; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1039; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]1040; EXACTCUTOFF-NEXT: v_mov_b32_e32 v0, s11041; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)1042; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)1043; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1044; EXACTCUTOFF-NEXT: s_nop 151045; EXACTCUTOFF-NEXT: s_nop 11046; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[24:27] offset:82881047; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[28:31] offset:83041048; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[16:19] offset:82561049; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[20:23] offset:82721050; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[8:11] offset:82241051; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[12:15] offset:82401052; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[0:3] offset:81921053; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[4:7] offset:82081054; EXACTCUTOFF-NEXT: ds_read_b128 a[28:31], v1 offset:246881055; EXACTCUTOFF-NEXT: ds_read_b128 a[24:27], v1 offset:246721056; EXACTCUTOFF-NEXT: ds_read_b128 a[20:23], v1 offset:246561057; EXACTCUTOFF-NEXT: ds_read_b128 a[16:19], v1 offset:246401058; EXACTCUTOFF-NEXT: ds_read_b128 a[12:15], v1 offset:246241059; EXACTCUTOFF-NEXT: ds_read_b128 a[8:11], v1 offset:246081060; EXACTCUTOFF-NEXT: ds_read_b128 a[4:7], v1 offset:245921061; EXACTCUTOFF-NEXT: ds_read_b128 a[0:3], v1 offset:245761062; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1063; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]1064; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)1065; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)1066; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1067; EXACTCUTOFF-NEXT: s_nop 151068; EXACTCUTOFF-NEXT: s_nop 21069; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[24:27] offset:164801070; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[28:31] offset:164961071; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[16:19] offset:164481072; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[20:23] offset:164641073; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[8:11] offset:164161074; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[12:15] offset:164321075; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[0:3] offset:163841076; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[4:7] offset:164001077; EXACTCUTOFF-NEXT: ds_read_b128 a[28:31], v1 offset:492641078; EXACTCUTOFF-NEXT: ds_read_b128 a[24:27], v1 offset:492481079; EXACTCUTOFF-NEXT: ds_read_b128 a[20:23], v1 offset:492321080; EXACTCUTOFF-NEXT: ds_read_b128 a[16:19], v1 offset:492161081; EXACTCUTOFF-NEXT: ds_read_b128 a[12:15], v1 offset:492001082; EXACTCUTOFF-NEXT: ds_read_b128 a[8:11], v1 offset:491841083; EXACTCUTOFF-NEXT: ds_read_b128 a[4:7], v1 offset:491681084; EXACTCUTOFF-NEXT: ds_read_b128 a[0:3], v1 offset:491521085; 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 v1, 0x6000, v11088; 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 151092; EXACTCUTOFF-NEXT: s_nop 11093; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[24:27] offset:246721094; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[28:31] offset:246881095; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[16:19] offset:246401096; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[20:23] offset:246561097; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[8:11] offset:246081098; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[12:15] offset:246241099; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[0:3] offset:245761100; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[4:7] offset:245921101; EXACTCUTOFF-NEXT: ds_read_b128 a[28:31], v1 offset:574561102; EXACTCUTOFF-NEXT: ds_read_b128 a[24:27], v1 offset:574401103; EXACTCUTOFF-NEXT: ds_read_b128 a[20:23], v1 offset:574241104; EXACTCUTOFF-NEXT: ds_read_b128 a[16:19], v1 offset:574081105; EXACTCUTOFF-NEXT: ds_read_b128 a[0:3], v1 offset:573441106; EXACTCUTOFF-NEXT: ds_read_b128 a[4:7], v1 offset:573601107; EXACTCUTOFF-NEXT: ds_read_b128 a[8:11], v1 offset:573761108; EXACTCUTOFF-NEXT: ds_read_b128 a[12:15], v1 offset:573921109; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1110; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]1111; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)1112; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)1113; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1114; EXACTCUTOFF-NEXT: s_nop 151115; EXACTCUTOFF-NEXT: s_nop 21116; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[24:27] offset:328641117; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[28:31] offset:328801118; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[16:19] offset:328321119; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[20:23] offset:328481120; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[8:11] offset:328001121; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[12:15] offset:328161122; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[0:3] offset:327681123; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[4:7] offset:327841124; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)1125; EXACTCUTOFF-NEXT: s_endpgm1126entry:1127 %idx = call i32 @llvm.amdgcn.workitem.id.x()1128 %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx1129 %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr1130 %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 641131 %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr1132 %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 1281133 %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr1134 %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 1921135 %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr1136 %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 2561137 %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr1138 %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)1139 %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)1140 %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)1141 %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)1142 %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)1143 %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx1144 store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr1145 %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 641146 store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr1147 %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 1281148 store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr1149 %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 1921150 store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr1151 %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 2561152 store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr1153 ; 8 DS read1154 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)1155 ; 1 MFMA1156 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1157 ; 8 DS write1158 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)1159 ; 8 DS read1160 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)1161 ; 1 MFMA1162 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1163 ; 8 DS write1164 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)1165 ; 8 DS read1166 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)1167 ; 1 MFMA1168 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1169 ; 8 DS write1170 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)1171 ; 8 DS read1172 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)1173 ; 1 MFMA1174 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1175 ; 8 DS write1176 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)1177 ; 8 DS read1178 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)1179 ; 1 MFMA1180 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1181 ; 8 DS write1182 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)1183 ret void1184}1185 1186define 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 {1187; GCN-LABEL: test_sched_group_barrier_pipeline_interleave_EXP_MFMA:1188; GCN: ; %bb.0: ; %entry1189; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x441190; GCN-NEXT: v_mov_b32_e32 v3, 0x3fb8aa3b1191; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x241192; GCN-NEXT: v_mov_b32_e32 v7, 0x32a5705f1193; GCN-NEXT: v_lshlrev_b32_e32 v0, 7, v01194; GCN-NEXT: s_waitcnt lgkmcnt(0)1195; GCN-NEXT: v_mul_f32_e32 v4, s0, v31196; GCN-NEXT: v_rndne_f32_e32 v5, v41197; GCN-NEXT: v_sub_f32_e32 v6, v4, v51198; GCN-NEXT: v_fma_f32 v4, s0, v3, -v41199; GCN-NEXT: v_fmac_f32_e32 v4, s0, v71200; GCN-NEXT: v_add_f32_e32 v4, v6, v41201; GCN-NEXT: v_exp_f32_e32 v4, v41202; GCN-NEXT: v_cvt_i32_f32_e32 v5, v51203; GCN-NEXT: v_and_b32_e32 v0, 0x1ff80, v01204; GCN-NEXT: v_add_u32_e32 v1, s6, v01205; GCN-NEXT: ds_read_b128 a[60:63], v1 offset:1121206; GCN-NEXT: ds_read_b128 a[56:59], v1 offset:961207; GCN-NEXT: ds_read_b128 a[52:55], v1 offset:801208; GCN-NEXT: ds_read_b128 a[48:51], v1 offset:641209; GCN-NEXT: ds_read_b128 a[32:35], v11210; GCN-NEXT: ds_read_b128 a[36:39], v1 offset:161211; GCN-NEXT: ds_read_b128 a[40:43], v1 offset:321212; GCN-NEXT: ds_read_b128 a[44:47], v1 offset:481213; GCN-NEXT: v_mov_b32_e32 v9, 1.01214; GCN-NEXT: v_ldexp_f32 v4, v4, v51215; GCN-NEXT: v_mov_b32_e32 v5, 0xc2ce8ed01216; GCN-NEXT: v_mul_f32_e32 v10, s1, v31217; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s0, v51218; GCN-NEXT: v_mov_b32_e32 v6, 0x42b172181219; GCN-NEXT: v_rndne_f32_e32 v11, v101220; GCN-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc1221; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s0, v61222; GCN-NEXT: v_mov_b32_e32 v8, 0x7f8000001223; GCN-NEXT: v_sub_f32_e32 v12, v10, v111224; GCN-NEXT: v_fma_f32 v10, s1, v3, -v101225; GCN-NEXT: v_cndmask_b32_e32 v4, v8, v4, vcc1226; GCN-NEXT: v_fmac_f32_e32 v10, s1, v71227; GCN-NEXT: ds_read_b128 a[28:31], v1 offset:83041228; GCN-NEXT: s_waitcnt lgkmcnt(1)1229; GCN-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v9, v4, a[32:63]1230; GCN-NEXT: v_add_f32_e32 v4, v12, v101231; GCN-NEXT: v_exp_f32_e32 v4, v41232; GCN-NEXT: v_cvt_i32_f32_e32 v10, v111233; GCN-NEXT: ds_read_b128 a[24:27], v1 offset:82881234; GCN-NEXT: ds_read_b128 a[20:23], v1 offset:82721235; GCN-NEXT: ds_read_b128 a[16:19], v1 offset:82561236; GCN-NEXT: ds_read_b128 a[12:15], v1 offset:82401237; GCN-NEXT: ds_read_b128 a[8:11], v1 offset:82241238; GCN-NEXT: ds_read_b128 a[4:7], v1 offset:82081239; GCN-NEXT: ds_read_b128 a[0:3], v1 offset:81921240; GCN-NEXT: v_ldexp_f32 v4, v4, v101241; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s1, v51242; GCN-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc1243; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s1, v61244; GCN-NEXT: v_cndmask_b32_e32 v4, v8, v4, vcc1245; GCN-NEXT: v_mul_f32_e32 v10, s2, v31246; GCN-NEXT: v_rndne_f32_e32 v11, v101247; GCN-NEXT: s_waitcnt lgkmcnt(0)1248; GCN-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v9, v4, a[0:31]1249; GCN-NEXT: v_fma_f32 v4, s2, v3, -v101250; GCN-NEXT: v_sub_f32_e32 v12, v10, v111251; GCN-NEXT: v_fmac_f32_e32 v4, s2, v71252; GCN-NEXT: v_add_f32_e32 v4, v12, v41253; GCN-NEXT: v_exp_f32_e32 v4, v41254; GCN-NEXT: v_cvt_i32_f32_e32 v10, v111255; GCN-NEXT: ds_read_b128 a[92:95], v1 offset:246881256; GCN-NEXT: ds_read_b128 a[88:91], v1 offset:246721257; GCN-NEXT: ds_read_b128 a[84:87], v1 offset:246561258; GCN-NEXT: ds_read_b128 a[80:83], v1 offset:246401259; GCN-NEXT: ds_read_b128 a[76:79], v1 offset:246241260; GCN-NEXT: ds_read_b128 a[72:75], v1 offset:246081261; GCN-NEXT: ds_read_b128 a[68:71], v1 offset:245921262; GCN-NEXT: ds_read_b128 a[64:67], v1 offset:245761263; GCN-NEXT: v_add_u32_e32 v2, 0x6000, v11264; GCN-NEXT: ds_read_b128 a[124:127], v1 offset:492641265; GCN-NEXT: ds_read_b128 a[120:123], v1 offset:492481266; GCN-NEXT: ds_read_b128 a[116:119], v1 offset:492321267; GCN-NEXT: ds_read_b128 a[112:115], v1 offset:492161268; GCN-NEXT: ds_read_b128 a[108:111], v1 offset:492001269; GCN-NEXT: ds_read_b128 a[104:107], v1 offset:491841270; GCN-NEXT: ds_read_b128 a[100:103], v1 offset:491681271; GCN-NEXT: ds_read_b128 a[96:99], v1 offset:491521272; GCN-NEXT: v_ldexp_f32 v1, v4, v101273; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s2, v51274; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc1275; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s2, v61276; GCN-NEXT: v_mul_f32_e32 v4, s3, v31277; GCN-NEXT: v_cndmask_b32_e32 v1, v8, v1, vcc1278; GCN-NEXT: v_rndne_f32_e32 v10, v41279; GCN-NEXT: s_load_dword s8, s[4:5], 0x541280; GCN-NEXT: s_waitcnt lgkmcnt(0)1281; GCN-NEXT: v_mfma_f32_32x32x1f32 a[64:95], v9, v1, a[64:95]1282; GCN-NEXT: v_sub_f32_e32 v1, v4, v101283; GCN-NEXT: v_fma_f32 v4, s3, v3, -v41284; GCN-NEXT: v_fmac_f32_e32 v4, s3, v71285; GCN-NEXT: v_add_f32_e32 v1, v1, v41286; GCN-NEXT: v_exp_f32_e32 v1, v11287; GCN-NEXT: v_cvt_i32_f32_e32 v4, v101288; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s3, v51289; GCN-NEXT: ds_read_b128 a[156:159], v2 offset:574561290; GCN-NEXT: ds_read_b128 a[152:155], v2 offset:574401291; GCN-NEXT: v_ldexp_f32 v1, v1, v41292; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc1293; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s3, v61294; GCN-NEXT: v_cndmask_b32_e32 v1, v8, v1, vcc1295; GCN-NEXT: v_mul_f32_e32 v4, s8, v31296; GCN-NEXT: v_fma_f32 v3, s8, v3, -v41297; GCN-NEXT: v_mfma_f32_32x32x1f32 a[96:127], v9, v1, a[96:127]1298; GCN-NEXT: v_rndne_f32_e32 v1, v41299; GCN-NEXT: v_sub_f32_e32 v10, v4, v11300; GCN-NEXT: v_fmac_f32_e32 v3, s8, v71301; GCN-NEXT: v_add_f32_e32 v3, v10, v31302; GCN-NEXT: v_exp_f32_e32 v3, v31303; GCN-NEXT: v_cvt_i32_f32_e32 v1, v11304; GCN-NEXT: ds_read_b128 a[148:151], v2 offset:574241305; GCN-NEXT: ds_read_b128 a[144:147], v2 offset:574081306; GCN-NEXT: ds_read_b128 a[128:131], v2 offset:573441307; GCN-NEXT: ds_read_b128 a[132:135], v2 offset:573601308; GCN-NEXT: ds_read_b128 a[136:139], v2 offset:573761309; GCN-NEXT: ds_read_b128 a[140:143], v2 offset:573921310; GCN-NEXT: v_ldexp_f32 v1, v3, v11311; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s8, v51312; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc1313; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s8, v61314; GCN-NEXT: v_cndmask_b32_e32 v1, v8, v1, vcc1315; GCN-NEXT: v_add_u32_e32 v0, s7, v01316; GCN-NEXT: ds_write_b128 v0, a[60:63] offset:1121317; GCN-NEXT: s_waitcnt lgkmcnt(1)1318; GCN-NEXT: v_mfma_f32_32x32x1f32 a[128:159], v9, v1, a[128:159]1319; GCN-NEXT: ds_write_b128 v0, a[56:59] offset:961320; GCN-NEXT: ds_write_b128 v0, a[52:55] offset:801321; GCN-NEXT: ds_write_b128 v0, a[48:51] offset:641322; GCN-NEXT: ds_write_b128 v0, a[44:47] offset:481323; GCN-NEXT: ds_write_b128 v0, a[40:43] offset:321324; GCN-NEXT: ds_write_b128 v0, a[36:39] offset:161325; GCN-NEXT: ds_write_b128 v0, a[32:35]1326; GCN-NEXT: v_mov_b32_e32 v0, s71327; GCN-NEXT: ; kill: killed $sgpr4_sgpr51328; GCN-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1329; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1330; GCN-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1331; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1332; GCN-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1333; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1334; GCN-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1335; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1336; GCN-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1337; GCN-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1338; GCN-NEXT: ds_write_b128 v0, a[24:27] offset:82881339; GCN-NEXT: ds_write_b128 v0, a[28:31] offset:83041340; GCN-NEXT: ds_write_b128 v0, a[16:19] offset:82561341; GCN-NEXT: ds_write_b128 v0, a[20:23] offset:82721342; GCN-NEXT: ds_write_b128 v0, a[8:11] offset:82241343; GCN-NEXT: ds_write_b128 v0, a[12:15] offset:82401344; GCN-NEXT: ds_write_b128 v0, a[0:3] offset:81921345; GCN-NEXT: ds_write_b128 v0, a[4:7] offset:82081346; GCN-NEXT: ds_write_b128 v0, a[88:91] offset:164801347; GCN-NEXT: ds_write_b128 v0, a[92:95] offset:164961348; GCN-NEXT: ds_write_b128 v0, a[80:83] offset:164481349; GCN-NEXT: ds_write_b128 v0, a[84:87] offset:164641350; GCN-NEXT: ds_write_b128 v0, a[72:75] offset:164161351; GCN-NEXT: ds_write_b128 v0, a[76:79] offset:164321352; GCN-NEXT: ds_write_b128 v0, a[64:67] offset:163841353; GCN-NEXT: ds_write_b128 v0, a[68:71] offset:164001354; GCN-NEXT: ds_write_b128 v0, a[120:123] offset:246721355; GCN-NEXT: ds_write_b128 v0, a[124:127] offset:246881356; GCN-NEXT: ds_write_b128 v0, a[112:115] offset:246401357; GCN-NEXT: ds_write_b128 v0, a[116:119] offset:246561358; GCN-NEXT: ds_write_b128 v0, a[104:107] offset:246081359; GCN-NEXT: ds_write_b128 v0, a[108:111] offset:246241360; GCN-NEXT: ds_write_b128 v0, a[96:99] offset:245761361; GCN-NEXT: ds_write_b128 v0, a[100:103] offset:245921362; GCN-NEXT: ds_write_b128 v0, a[152:155] offset:328641363; GCN-NEXT: ds_write_b128 v0, a[156:159] offset:328801364; GCN-NEXT: ds_write_b128 v0, a[144:147] offset:328321365; GCN-NEXT: ds_write_b128 v0, a[148:151] offset:328481366; GCN-NEXT: ds_write_b128 v0, a[136:139] offset:328001367; GCN-NEXT: ds_write_b128 v0, a[140:143] offset:328161368; GCN-NEXT: ds_write_b128 v0, a[128:131] offset:327681369; GCN-NEXT: ds_write_b128 v0, a[132:135] offset:327841370; GCN-NEXT: s_endpgm1371;1372; EXACTCUTOFF-LABEL: test_sched_group_barrier_pipeline_interleave_EXP_MFMA:1373; EXACTCUTOFF: ; %bb.0: ; %entry1374; EXACTCUTOFF-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x441375; EXACTCUTOFF-NEXT: v_mov_b32_e32 v3, 0x3fb8aa3b1376; EXACTCUTOFF-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x241377; EXACTCUTOFF-NEXT: v_mov_b32_e32 v7, 0x32a5705f1378; EXACTCUTOFF-NEXT: v_lshlrev_b32_e32 v0, 7, v01379; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1380; EXACTCUTOFF-NEXT: v_mul_f32_e32 v4, s0, v31381; EXACTCUTOFF-NEXT: v_rndne_f32_e32 v5, v41382; EXACTCUTOFF-NEXT: v_sub_f32_e32 v6, v4, v51383; EXACTCUTOFF-NEXT: v_fma_f32 v4, s0, v3, -v41384; EXACTCUTOFF-NEXT: v_fmac_f32_e32 v4, s0, v71385; EXACTCUTOFF-NEXT: v_add_f32_e32 v4, v6, v41386; EXACTCUTOFF-NEXT: v_exp_f32_e32 v4, v41387; EXACTCUTOFF-NEXT: v_cvt_i32_f32_e32 v5, v51388; EXACTCUTOFF-NEXT: v_and_b32_e32 v0, 0x1ff80, v01389; EXACTCUTOFF-NEXT: v_add_u32_e32 v1, s6, v01390; EXACTCUTOFF-NEXT: ds_read_b128 a[60:63], v1 offset:1121391; EXACTCUTOFF-NEXT: ds_read_b128 a[56:59], v1 offset:961392; EXACTCUTOFF-NEXT: ds_read_b128 a[52:55], v1 offset:801393; EXACTCUTOFF-NEXT: ds_read_b128 a[48:51], v1 offset:641394; EXACTCUTOFF-NEXT: ds_read_b128 a[32:35], v11395; EXACTCUTOFF-NEXT: ds_read_b128 a[36:39], v1 offset:161396; EXACTCUTOFF-NEXT: ds_read_b128 a[40:43], v1 offset:321397; EXACTCUTOFF-NEXT: ds_read_b128 a[44:47], v1 offset:481398; EXACTCUTOFF-NEXT: v_mov_b32_e32 v9, 1.01399; EXACTCUTOFF-NEXT: v_ldexp_f32 v4, v4, v51400; EXACTCUTOFF-NEXT: v_mov_b32_e32 v5, 0xc2ce8ed01401; EXACTCUTOFF-NEXT: v_mul_f32_e32 v10, s1, v31402; EXACTCUTOFF-NEXT: v_cmp_nlt_f32_e32 vcc, s0, v51403; EXACTCUTOFF-NEXT: v_mov_b32_e32 v6, 0x42b172181404; EXACTCUTOFF-NEXT: v_rndne_f32_e32 v11, v101405; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc1406; EXACTCUTOFF-NEXT: v_cmp_ngt_f32_e32 vcc, s0, v61407; EXACTCUTOFF-NEXT: v_mov_b32_e32 v8, 0x7f8000001408; EXACTCUTOFF-NEXT: v_sub_f32_e32 v12, v10, v111409; EXACTCUTOFF-NEXT: v_fma_f32 v10, s1, v3, -v101410; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v4, v8, v4, vcc1411; EXACTCUTOFF-NEXT: v_fmac_f32_e32 v10, s1, v71412; EXACTCUTOFF-NEXT: ds_read_b128 a[28:31], v1 offset:83041413; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(1)1414; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v9, v4, a[32:63]1415; EXACTCUTOFF-NEXT: v_add_f32_e32 v4, v12, v101416; EXACTCUTOFF-NEXT: v_exp_f32_e32 v4, v41417; EXACTCUTOFF-NEXT: v_cvt_i32_f32_e32 v10, v111418; EXACTCUTOFF-NEXT: ds_read_b128 a[24:27], v1 offset:82881419; EXACTCUTOFF-NEXT: ds_read_b128 a[20:23], v1 offset:82721420; EXACTCUTOFF-NEXT: ds_read_b128 a[16:19], v1 offset:82561421; EXACTCUTOFF-NEXT: ds_read_b128 a[12:15], v1 offset:82401422; EXACTCUTOFF-NEXT: ds_read_b128 a[8:11], v1 offset:82241423; EXACTCUTOFF-NEXT: ds_read_b128 a[4:7], v1 offset:82081424; EXACTCUTOFF-NEXT: ds_read_b128 a[0:3], v1 offset:81921425; EXACTCUTOFF-NEXT: v_ldexp_f32 v4, v4, v101426; EXACTCUTOFF-NEXT: v_cmp_nlt_f32_e32 vcc, s1, v51427; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v4, 0, v4, vcc1428; EXACTCUTOFF-NEXT: v_cmp_ngt_f32_e32 vcc, s1, v61429; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v4, v8, v4, vcc1430; EXACTCUTOFF-NEXT: v_mul_f32_e32 v10, s2, v31431; EXACTCUTOFF-NEXT: v_rndne_f32_e32 v11, v101432; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1433; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v9, v4, a[0:31]1434; EXACTCUTOFF-NEXT: v_fma_f32 v4, s2, v3, -v101435; EXACTCUTOFF-NEXT: v_sub_f32_e32 v12, v10, v111436; EXACTCUTOFF-NEXT: v_fmac_f32_e32 v4, s2, v71437; EXACTCUTOFF-NEXT: v_add_f32_e32 v4, v12, v41438; EXACTCUTOFF-NEXT: v_exp_f32_e32 v4, v41439; EXACTCUTOFF-NEXT: v_cvt_i32_f32_e32 v10, v111440; EXACTCUTOFF-NEXT: ds_read_b128 a[92:95], v1 offset:246881441; EXACTCUTOFF-NEXT: ds_read_b128 a[88:91], v1 offset:246721442; EXACTCUTOFF-NEXT: ds_read_b128 a[84:87], v1 offset:246561443; EXACTCUTOFF-NEXT: ds_read_b128 a[80:83], v1 offset:246401444; EXACTCUTOFF-NEXT: ds_read_b128 a[76:79], v1 offset:246241445; EXACTCUTOFF-NEXT: ds_read_b128 a[72:75], v1 offset:246081446; EXACTCUTOFF-NEXT: ds_read_b128 a[68:71], v1 offset:245921447; EXACTCUTOFF-NEXT: ds_read_b128 a[64:67], v1 offset:245761448; EXACTCUTOFF-NEXT: v_add_u32_e32 v2, 0x6000, v11449; EXACTCUTOFF-NEXT: ds_read_b128 a[124:127], v1 offset:492641450; EXACTCUTOFF-NEXT: ds_read_b128 a[120:123], v1 offset:492481451; EXACTCUTOFF-NEXT: ds_read_b128 a[116:119], v1 offset:492321452; EXACTCUTOFF-NEXT: ds_read_b128 a[112:115], v1 offset:492161453; EXACTCUTOFF-NEXT: ds_read_b128 a[108:111], v1 offset:492001454; EXACTCUTOFF-NEXT: ds_read_b128 a[104:107], v1 offset:491841455; EXACTCUTOFF-NEXT: ds_read_b128 a[100:103], v1 offset:491681456; EXACTCUTOFF-NEXT: ds_read_b128 a[96:99], v1 offset:491521457; EXACTCUTOFF-NEXT: v_ldexp_f32 v1, v4, v101458; EXACTCUTOFF-NEXT: v_cmp_nlt_f32_e32 vcc, s2, v51459; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc1460; EXACTCUTOFF-NEXT: v_cmp_ngt_f32_e32 vcc, s2, v61461; EXACTCUTOFF-NEXT: v_mul_f32_e32 v4, s3, v31462; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v1, v8, v1, vcc1463; EXACTCUTOFF-NEXT: v_rndne_f32_e32 v10, v41464; EXACTCUTOFF-NEXT: s_load_dword s8, s[4:5], 0x541465; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(0)1466; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[64:95], v9, v1, a[64:95]1467; EXACTCUTOFF-NEXT: v_sub_f32_e32 v1, v4, v101468; EXACTCUTOFF-NEXT: v_fma_f32 v4, s3, v3, -v41469; EXACTCUTOFF-NEXT: v_fmac_f32_e32 v4, s3, v71470; EXACTCUTOFF-NEXT: v_add_f32_e32 v1, v1, v41471; EXACTCUTOFF-NEXT: v_exp_f32_e32 v1, v11472; EXACTCUTOFF-NEXT: v_cvt_i32_f32_e32 v4, v101473; EXACTCUTOFF-NEXT: v_cmp_nlt_f32_e32 vcc, s3, v51474; EXACTCUTOFF-NEXT: ds_read_b128 a[156:159], v2 offset:574561475; EXACTCUTOFF-NEXT: ds_read_b128 a[152:155], v2 offset:574401476; EXACTCUTOFF-NEXT: v_ldexp_f32 v1, v1, v41477; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc1478; EXACTCUTOFF-NEXT: v_cmp_ngt_f32_e32 vcc, s3, v61479; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v1, v8, v1, vcc1480; EXACTCUTOFF-NEXT: v_mul_f32_e32 v4, s8, v31481; EXACTCUTOFF-NEXT: v_fma_f32 v3, s8, v3, -v41482; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[96:127], v9, v1, a[96:127]1483; EXACTCUTOFF-NEXT: v_rndne_f32_e32 v1, v41484; EXACTCUTOFF-NEXT: v_sub_f32_e32 v10, v4, v11485; EXACTCUTOFF-NEXT: v_fmac_f32_e32 v3, s8, v71486; EXACTCUTOFF-NEXT: v_add_f32_e32 v3, v10, v31487; EXACTCUTOFF-NEXT: v_exp_f32_e32 v3, v31488; EXACTCUTOFF-NEXT: v_cvt_i32_f32_e32 v1, v11489; EXACTCUTOFF-NEXT: ds_read_b128 a[148:151], v2 offset:574241490; EXACTCUTOFF-NEXT: ds_read_b128 a[144:147], v2 offset:574081491; EXACTCUTOFF-NEXT: ds_read_b128 a[128:131], v2 offset:573441492; EXACTCUTOFF-NEXT: ds_read_b128 a[132:135], v2 offset:573601493; EXACTCUTOFF-NEXT: ds_read_b128 a[136:139], v2 offset:573761494; EXACTCUTOFF-NEXT: ds_read_b128 a[140:143], v2 offset:573921495; EXACTCUTOFF-NEXT: v_ldexp_f32 v1, v3, v11496; EXACTCUTOFF-NEXT: v_cmp_nlt_f32_e32 vcc, s8, v51497; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc1498; EXACTCUTOFF-NEXT: v_cmp_ngt_f32_e32 vcc, s8, v61499; EXACTCUTOFF-NEXT: v_cndmask_b32_e32 v1, v8, v1, vcc1500; EXACTCUTOFF-NEXT: v_add_u32_e32 v0, s7, v01501; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[60:63] offset:1121502; EXACTCUTOFF-NEXT: s_waitcnt lgkmcnt(1)1503; EXACTCUTOFF-NEXT: v_mfma_f32_32x32x1f32 a[128:159], v9, v1, a[128:159]1504; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[56:59] offset:961505; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[52:55] offset:801506; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[48:51] offset:641507; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[44:47] offset:481508; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[40:43] offset:321509; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[36:39] offset:161510; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[32:35]1511; EXACTCUTOFF-NEXT: v_mov_b32_e32 v0, s71512; EXACTCUTOFF-NEXT: ; kill: killed $sgpr4_sgpr51513; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1514; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1515; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1516; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1517; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1518; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1519; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1520; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1521; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000400) size(1) SyncID(0)1522; EXACTCUTOFF-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)1523; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[24:27] offset:82881524; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[28:31] offset:83041525; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[16:19] offset:82561526; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[20:23] offset:82721527; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[8:11] offset:82241528; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[12:15] offset:82401529; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[0:3] offset:81921530; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[4:7] offset:82081531; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[88:91] offset:164801532; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[92:95] offset:164961533; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[80:83] offset:164481534; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[84:87] offset:164641535; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[72:75] offset:164161536; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[76:79] offset:164321537; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[64:67] offset:163841538; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[68:71] offset:164001539; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[120:123] offset:246721540; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[124:127] offset:246881541; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[112:115] offset:246401542; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[116:119] offset:246561543; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[104:107] offset:246081544; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[108:111] offset:246241545; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[96:99] offset:245761546; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[100:103] offset:245921547; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[152:155] offset:328641548; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[156:159] offset:328801549; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[144:147] offset:328321550; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[148:151] offset:328481551; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[136:139] offset:328001552; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[140:143] offset:328161553; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[128:131] offset:327681554; EXACTCUTOFF-NEXT: ds_write_b128 v0, a[132:135] offset:327841555; EXACTCUTOFF-NEXT: s_endpgm1556entry:1557 %idx = call i32 @llvm.amdgcn.workitem.id.x()1558 %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx1559 %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr1560 %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 641561 %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr1562 %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 1281563 %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr1564 %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 1921565 %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr1566 %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 2561567 %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr1568 %el0 = extractelement <5 x float> %in1, i32 01569 %el1 = extractelement <5 x float> %in1, i32 11570 %el2 = extractelement <5 x float> %in1, i32 21571 %el3 = extractelement <5 x float> %in1, i32 31572 %el4 = extractelement <5 x float> %in1, i32 41573 %exp0 = tail call float @llvm.exp.f32(float %el0)1574 %exp1 = tail call float @llvm.exp.f32(float %el1)1575 %exp2 = tail call float @llvm.exp.f32(float %el2)1576 %exp3 = tail call float @llvm.exp.f32(float %el3)1577 %exp4 = tail call float @llvm.exp.f32(float %el4)1578 %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)1579 %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)1580 %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)1581 %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)1582 %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)1583 ; 1 TRANS1584 call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)1585 ; 1 MFMA1586 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1587 ; 1 TRANS1588 call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)1589 ; 1 MFMA1590 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1591 ; 1 TRANS1592 call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)1593 ; 1 MFMA1594 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1595 ; 1 TRANS1596 call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)1597 ; 1 MFMA1598 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1599 ; 1 TRANS1600 call void @llvm.amdgcn.sched.group.barrier(i32 1024, i32 1, i32 0)1601 ; 1 MFMA1602 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)1603 %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx1604 store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr1605 %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 641606 store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr1607 %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 1281608 store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr1609 %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 1921610 store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr1611 %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 2561612 store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr1613 ret void1614}1615 1616declare i32 @llvm.amdgcn.workitem.id.x() #21617declare void @llvm.amdgcn.sched.group.barrier(i32, i32, i32) #11618declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32) #11619declare float @llvm.exp.f32(float) #21620 1621attributes #0 = { nounwind "amdgpu-flat-work-group-size"="1,256" }1622attributes #1 = { nounwind }1623attributes #2 = { nounwind readnone speculatable }1624