904 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -misched-cluster=0 -misched=gcn-iterative-minreg < %s | FileCheck -check-prefix=GCN-MINREG %s3; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -misched-cluster=0 -misched=gcn-iterative-max-occupancy-experimental < %s | FileCheck -check-prefix=GCN-MAXOCC %s4; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -misched-cluster=0 -misched=gcn-iterative-ilp < %s | FileCheck -check-prefix=GCN-ILP %s5 6define amdgpu_kernel void @test_sched_group_barrier_pipeline_MFMA_interleave(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {7; GCN-MINREG-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:8; GCN-MINREG: ; %bb.0: ; %entry9; GCN-MINREG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x2410; GCN-MINREG-NEXT: v_lshlrev_b32_e32 v0, 7, v011; GCN-MINREG-NEXT: v_and_b32_e32 v0, 0x1ff80, v012; GCN-MINREG-NEXT: v_mov_b32_e32 v2, 1.013; GCN-MINREG-NEXT: v_mov_b32_e32 v1, 2.014; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)15; GCN-MINREG-NEXT: v_add_u32_e32 v4, s0, v016; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v4 offset:11217; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v4 offset:9618; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v4 offset:8019; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v4 offset:6420; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v421; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v4 offset:1622; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v4 offset:3223; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v4 offset:4824; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)25; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v1, a[0:31]26; GCN-MINREG-NEXT: v_add_u32_e32 v5, s1, v027; GCN-MINREG-NEXT: v_mov_b32_e32 v0, s128; GCN-MINREG-NEXT: v_add_u32_e32 v3, 0x6000, v429; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)30; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)31; GCN-MINREG-NEXT: s_nop 1532; GCN-MINREG-NEXT: ds_write_b128 v5, a[28:31] offset:11233; GCN-MINREG-NEXT: ds_write_b128 v5, a[24:27] offset:9634; GCN-MINREG-NEXT: ds_write_b128 v5, a[20:23] offset:8035; GCN-MINREG-NEXT: ds_write_b128 v5, a[16:19] offset:6436; GCN-MINREG-NEXT: ds_write_b128 v5, a[12:15] offset:4837; GCN-MINREG-NEXT: ds_write_b128 v5, a[8:11] offset:3238; GCN-MINREG-NEXT: ds_write_b128 v5, a[4:7] offset:1639; GCN-MINREG-NEXT: ds_write_b128 v5, a[0:3]40; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v4 offset:830441; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v4 offset:828842; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v4 offset:827243; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v4 offset:825644; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v4 offset:824045; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v4 offset:822446; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v4 offset:820847; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v4 offset:819248; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)49; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v1, a[0:31]50; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)51; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)52; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)53; GCN-MINREG-NEXT: s_nop 1554; GCN-MINREG-NEXT: s_nop 255; GCN-MINREG-NEXT: ds_write_b128 v0, a[24:27] offset:828856; GCN-MINREG-NEXT: ds_write_b128 v0, a[28:31] offset:830457; GCN-MINREG-NEXT: ds_write_b128 v0, a[16:19] offset:825658; GCN-MINREG-NEXT: ds_write_b128 v0, a[20:23] offset:827259; GCN-MINREG-NEXT: ds_write_b128 v0, a[8:11] offset:822460; GCN-MINREG-NEXT: ds_write_b128 v0, a[12:15] offset:824061; GCN-MINREG-NEXT: ds_write_b128 v0, a[0:3] offset:819262; GCN-MINREG-NEXT: ds_write_b128 v0, a[4:7] offset:820863; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v4 offset:2468864; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v4 offset:2467265; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v4 offset:2465666; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v4 offset:2464067; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v4 offset:2462468; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v4 offset:2460869; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v4 offset:2459270; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v4 offset:2457671; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)72; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v1, a[0:31]73; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)74; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)75; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)76; GCN-MINREG-NEXT: s_nop 1577; GCN-MINREG-NEXT: s_nop 278; GCN-MINREG-NEXT: ds_write_b128 v0, a[24:27] offset:1648079; GCN-MINREG-NEXT: ds_write_b128 v0, a[28:31] offset:1649680; GCN-MINREG-NEXT: ds_write_b128 v0, a[16:19] offset:1644881; GCN-MINREG-NEXT: ds_write_b128 v0, a[20:23] offset:1646482; GCN-MINREG-NEXT: ds_write_b128 v0, a[8:11] offset:1641683; GCN-MINREG-NEXT: ds_write_b128 v0, a[12:15] offset:1643284; GCN-MINREG-NEXT: ds_write_b128 v0, a[0:3] offset:1638485; GCN-MINREG-NEXT: ds_write_b128 v0, a[4:7] offset:1640086; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v4 offset:4926487; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v4 offset:4924888; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v4 offset:4923289; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v4 offset:4921690; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v4 offset:4920091; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v4 offset:4918492; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v4 offset:4916893; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v4 offset:4915294; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)95; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v1, a[0:31]96; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)97; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)98; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)99; GCN-MINREG-NEXT: s_nop 15100; GCN-MINREG-NEXT: s_nop 2101; GCN-MINREG-NEXT: ds_write_b128 v0, a[24:27] offset:24672102; GCN-MINREG-NEXT: ds_write_b128 v0, a[28:31] offset:24688103; GCN-MINREG-NEXT: ds_write_b128 v0, a[16:19] offset:24640104; GCN-MINREG-NEXT: ds_write_b128 v0, a[20:23] offset:24656105; GCN-MINREG-NEXT: ds_write_b128 v0, a[8:11] offset:24608106; GCN-MINREG-NEXT: ds_write_b128 v0, a[12:15] offset:24624107; GCN-MINREG-NEXT: ds_write_b128 v0, a[0:3] offset:24576108; GCN-MINREG-NEXT: ds_write_b128 v0, a[4:7] offset:24592109; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v3 offset:57456110; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v3 offset:57440111; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v3 offset:57424112; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v3 offset:57408113; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v3 offset:57344114; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v3 offset:57360115; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v3 offset:57376116; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v3 offset:57392117; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)118; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v1, a[0:31]119; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)120; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)121; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)122; GCN-MINREG-NEXT: s_nop 15123; GCN-MINREG-NEXT: s_nop 2124; GCN-MINREG-NEXT: ds_write_b128 v0, a[24:27] offset:32864125; GCN-MINREG-NEXT: ds_write_b128 v0, a[28:31] offset:32880126; GCN-MINREG-NEXT: ds_write_b128 v0, a[16:19] offset:32832127; GCN-MINREG-NEXT: ds_write_b128 v0, a[20:23] offset:32848128; GCN-MINREG-NEXT: ds_write_b128 v0, a[8:11] offset:32800129; GCN-MINREG-NEXT: ds_write_b128 v0, a[12:15] offset:32816130; GCN-MINREG-NEXT: ds_write_b128 v0, a[0:3] offset:32768131; GCN-MINREG-NEXT: ds_write_b128 v0, a[4:7] offset:32784132; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)133; GCN-MINREG-NEXT: s_endpgm134;135; GCN-MAXOCC-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:136; GCN-MAXOCC: ; %bb.0: ; %entry137; GCN-MAXOCC-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24138; GCN-MAXOCC-NEXT: v_lshlrev_b32_e32 v0, 7, v0139; GCN-MAXOCC-NEXT: v_and_b32_e32 v1, 0x1ff80, v0140; GCN-MAXOCC-NEXT: v_mov_b32_e32 v2, 1.0141; GCN-MAXOCC-NEXT: v_mov_b32_e32 v3, 2.0142; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)143; GCN-MAXOCC-NEXT: v_add_u32_e32 v0, s0, v1144; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:112145; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:96146; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:80147; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:64148; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0149; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:16150; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:32151; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:48152; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)153; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]154; GCN-MAXOCC-NEXT: v_add_u32_e32 v1, s1, v1155; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)156; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)157; GCN-MAXOCC-NEXT: s_nop 15158; GCN-MAXOCC-NEXT: s_nop 1159; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[28:31] offset:112160; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[24:27] offset:96161; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[20:23] offset:80162; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[16:19] offset:64163; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[12:15] offset:48164; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[8:11] offset:32165; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[4:7] offset:16166; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[0:3]167; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:8304168; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:8288169; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:8272170; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:8256171; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:8240172; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:8224173; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:8208174; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:8192175; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)176; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]177; GCN-MAXOCC-NEXT: v_mov_b32_e32 v1, s1178; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)179; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)180; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)181; GCN-MAXOCC-NEXT: s_nop 15182; GCN-MAXOCC-NEXT: s_nop 1183; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[24:27] offset:8288184; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[28:31] offset:8304185; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[16:19] offset:8256186; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[20:23] offset:8272187; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[8:11] offset:8224188; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[12:15] offset:8240189; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[0:3] offset:8192190; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[4:7] offset:8208191; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:24688192; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:24672193; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:24656194; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:24640195; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:24624196; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:24608197; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:24592198; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:24576199; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)200; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]201; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)202; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)203; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)204; GCN-MAXOCC-NEXT: s_nop 15205; GCN-MAXOCC-NEXT: s_nop 2206; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[24:27] offset:16480207; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[28:31] offset:16496208; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[16:19] offset:16448209; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[20:23] offset:16464210; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[8:11] offset:16416211; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[12:15] offset:16432212; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[0:3] offset:16384213; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[4:7] offset:16400214; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:49264215; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:49248216; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:49232217; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:49216218; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:49200219; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:49184220; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:49168221; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:49152222; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)223; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]224; GCN-MAXOCC-NEXT: v_add_u32_e32 v0, 0x6000, v0225; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)226; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)227; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)228; GCN-MAXOCC-NEXT: s_nop 15229; GCN-MAXOCC-NEXT: s_nop 1230; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[24:27] offset:24672231; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[28:31] offset:24688232; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[16:19] offset:24640233; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[20:23] offset:24656234; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[8:11] offset:24608235; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[12:15] offset:24624236; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[0:3] offset:24576237; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[4:7] offset:24592238; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:57456239; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:57440240; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:57424241; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:57408242; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:57344243; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:57360244; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:57376245; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:57392246; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)247; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]248; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)249; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)250; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)251; GCN-MAXOCC-NEXT: s_nop 15252; GCN-MAXOCC-NEXT: s_nop 2253; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[24:27] offset:32864254; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[28:31] offset:32880255; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[16:19] offset:32832256; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[20:23] offset:32848257; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[8:11] offset:32800258; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[12:15] offset:32816259; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[0:3] offset:32768260; GCN-MAXOCC-NEXT: ds_write_b128 v1, a[4:7] offset:32784261; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)262; GCN-MAXOCC-NEXT: s_endpgm263;264; GCN-ILP-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave:265; GCN-ILP: ; %bb.0: ; %entry266; GCN-ILP-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24267; GCN-ILP-NEXT: v_lshlrev_b32_e32 v0, 7, v0268; GCN-ILP-NEXT: v_and_b32_e32 v0, 0x1ff80, v0269; GCN-ILP-NEXT: v_mov_b32_e32 v1, 1.0270; GCN-ILP-NEXT: v_mov_b32_e32 v2, 2.0271; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)272; GCN-ILP-NEXT: v_add_u32_e32 v3, s0, v0273; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:48274; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:32275; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:16276; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3277; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:64278; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:80279; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:96280; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:112281; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)282; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]283; GCN-ILP-NEXT: v_add_u32_e32 v0, s1, v0284; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)285; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)286; GCN-ILP-NEXT: s_nop 15287; GCN-ILP-NEXT: s_nop 1288; GCN-ILP-NEXT: ds_write_b128 v0, a[28:31] offset:112289; GCN-ILP-NEXT: ds_write_b128 v0, a[24:27] offset:96290; GCN-ILP-NEXT: ds_write_b128 v0, a[20:23] offset:80291; GCN-ILP-NEXT: ds_write_b128 v0, a[16:19] offset:64292; GCN-ILP-NEXT: ds_write_b128 v0, a[12:15] offset:48293; GCN-ILP-NEXT: ds_write_b128 v0, a[8:11] offset:32294; GCN-ILP-NEXT: ds_write_b128 v0, a[4:7] offset:16295; GCN-ILP-NEXT: ds_write_b128 v0, a[0:3]296; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:8192297; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:8208298; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:8224299; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:8240300; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:8256301; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:8272302; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:8288303; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:8304304; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)305; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]306; GCN-ILP-NEXT: v_mov_b32_e32 v0, s1307; GCN-ILP-NEXT: s_nop 15308; GCN-ILP-NEXT: s_nop 1309; GCN-ILP-NEXT: ds_write_b128 v0, a[24:27] offset:8288310; GCN-ILP-NEXT: ds_write_b128 v0, a[28:31] offset:8304311; GCN-ILP-NEXT: ds_write_b128 v0, a[16:19] offset:8256312; GCN-ILP-NEXT: ds_write_b128 v0, a[20:23] offset:8272313; GCN-ILP-NEXT: ds_write_b128 v0, a[8:11] offset:8224314; GCN-ILP-NEXT: ds_write_b128 v0, a[12:15] offset:8240315; GCN-ILP-NEXT: ds_write_b128 v0, a[0:3] offset:8192316; GCN-ILP-NEXT: ds_write_b128 v0, a[4:7] offset:8208317; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:24576318; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:24592319; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:24608320; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:24624321; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:24640322; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:24656323; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:24672324; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:24688325; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)326; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]327; GCN-ILP-NEXT: s_nop 15328; GCN-ILP-NEXT: s_nop 2329; GCN-ILP-NEXT: ds_write_b128 v0, a[4:7] offset:16400330; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:49168331; GCN-ILP-NEXT: ds_write_b128 v0, a[0:3] offset:16384332; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:49152333; GCN-ILP-NEXT: ds_write_b128 v0, a[12:15] offset:16432334; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:49200335; GCN-ILP-NEXT: ds_write_b128 v0, a[8:11] offset:16416336; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:49184337; GCN-ILP-NEXT: ds_write_b128 v0, a[20:23] offset:16464338; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:49232339; GCN-ILP-NEXT: ds_write_b128 v0, a[16:19] offset:16448340; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:49216341; GCN-ILP-NEXT: ds_write_b128 v0, a[28:31] offset:16496342; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:49264343; GCN-ILP-NEXT: ds_write_b128 v0, a[24:27] offset:16480344; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:49248345; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)346; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]347; GCN-ILP-NEXT: v_add_u32_e32 v3, 0x6000, v3348; GCN-ILP-NEXT: s_nop 15349; GCN-ILP-NEXT: s_nop 1350; GCN-ILP-NEXT: ds_write_b128 v0, a[4:7] offset:24592351; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:57360352; GCN-ILP-NEXT: ds_write_b128 v0, a[0:3] offset:24576353; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:57344354; GCN-ILP-NEXT: ds_write_b128 v0, a[12:15] offset:24624355; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:57392356; GCN-ILP-NEXT: ds_write_b128 v0, a[8:11] offset:24608357; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:57376358; GCN-ILP-NEXT: ds_write_b128 v0, a[20:23] offset:24656359; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:57424360; GCN-ILP-NEXT: ds_write_b128 v0, a[16:19] offset:24640361; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:57408362; GCN-ILP-NEXT: ds_write_b128 v0, a[28:31] offset:24688363; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:57456364; GCN-ILP-NEXT: ds_write_b128 v0, a[24:27] offset:24672365; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:57440366; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)367; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]368; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)369; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)370; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)371; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)372; GCN-ILP-NEXT: s_nop 15373; GCN-ILP-NEXT: s_nop 2374; GCN-ILP-NEXT: ds_write_b128 v0, a[24:27] offset:32864375; GCN-ILP-NEXT: ds_write_b128 v0, a[28:31] offset:32880376; GCN-ILP-NEXT: ds_write_b128 v0, a[16:19] offset:32832377; GCN-ILP-NEXT: ds_write_b128 v0, a[20:23] offset:32848378; GCN-ILP-NEXT: ds_write_b128 v0, a[8:11] offset:32800379; GCN-ILP-NEXT: ds_write_b128 v0, a[12:15] offset:32816380; GCN-ILP-NEXT: ds_write_b128 v0, a[0:3] offset:32768381; GCN-ILP-NEXT: ds_write_b128 v0, a[4:7] offset:32784382; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)383; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)384; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)385; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)386; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)387; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)388; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)389; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)390; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)391; GCN-ILP-NEXT: s_endpgm392entry:393 %idx = call i32 @llvm.amdgcn.workitem.id.x()394 %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx395 %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr396 %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64397 %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr398 %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128399 %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr400 %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192401 %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr402 %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256403 %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr404 %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)405 %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)406 %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)407 %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)408 %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)409 %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx410 store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr411 %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64412 store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr413 %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128414 store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr415 %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192416 store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr417 %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256418 store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr419 ; 8 DS read420 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)421 ; 1 MFMA422 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)423 ; 8 DS write424 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)425 ; 8 DS read426 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)427 ; 1 MFMA428 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)429 ; 8 DS write430 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)431 ; 8 DS read432 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)433 ; 1 MFMA434 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)435 ; 8 DS write436 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)437 ; 8 DS read438 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)439 ; 1 MFMA440 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)441 ; 8 DS write442 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)443 ; 8 DS read444 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)445 ; 1 MFMA446 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)447 ; 8 DS write448 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)449 ret void450}451 452 453define amdgpu_kernel void @test_sched_group_barrier_pipeline_MFMA_interleave_split_region(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {454; GCN-MINREG-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave_split_region:455; GCN-MINREG: ; %bb.0: ; %entry456; GCN-MINREG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24457; GCN-MINREG-NEXT: v_lshlrev_b32_e32 v0, 7, v0458; GCN-MINREG-NEXT: v_and_b32_e32 v2, 0x1ff80, v0459; GCN-MINREG-NEXT: v_mov_b32_e32 v1, 1.0460; GCN-MINREG-NEXT: v_mov_b32_e32 v0, 2.0461; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)462; GCN-MINREG-NEXT: v_add_u32_e32 v3, s0, v2463; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v3 offset:112464; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v3 offset:96465; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v3 offset:80466; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v3 offset:64467; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v3468; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v3 offset:16469; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v3 offset:32470; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v3 offset:48471; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)472; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v0, a[0:31]473; GCN-MINREG-NEXT: v_add_u32_e32 v2, s1, v2474; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)475; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)476; GCN-MINREG-NEXT: s_nop 15477; GCN-MINREG-NEXT: s_nop 1478; GCN-MINREG-NEXT: ds_write_b128 v2, a[28:31] offset:112479; GCN-MINREG-NEXT: ds_write_b128 v2, a[24:27] offset:96480; GCN-MINREG-NEXT: ds_write_b128 v2, a[20:23] offset:80481; GCN-MINREG-NEXT: ds_write_b128 v2, a[16:19] offset:64482; GCN-MINREG-NEXT: ds_write_b128 v2, a[12:15] offset:48483; GCN-MINREG-NEXT: ds_write_b128 v2, a[8:11] offset:32484; GCN-MINREG-NEXT: ds_write_b128 v2, a[4:7] offset:16485; GCN-MINREG-NEXT: ds_write_b128 v2, a[0:3]486; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v3 offset:8304487; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v3 offset:8288488; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v3 offset:8272489; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v3 offset:8256490; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v3 offset:8240491; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v3 offset:8224492; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v3 offset:8208493; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v3 offset:8192494; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)495; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v0, a[0:31]496; GCN-MINREG-NEXT: v_mov_b32_e32 v2, s1497; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)498; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)499; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)500; GCN-MINREG-NEXT: s_nop 15501; GCN-MINREG-NEXT: s_nop 1502; GCN-MINREG-NEXT: ds_write_b128 v2, a[24:27] offset:8288503; GCN-MINREG-NEXT: ds_write_b128 v2, a[28:31] offset:8304504; GCN-MINREG-NEXT: ds_write_b128 v2, a[16:19] offset:8256505; GCN-MINREG-NEXT: ds_write_b128 v2, a[20:23] offset:8272506; GCN-MINREG-NEXT: ds_write_b128 v2, a[8:11] offset:8224507; GCN-MINREG-NEXT: ds_write_b128 v2, a[12:15] offset:8240508; GCN-MINREG-NEXT: ds_write_b128 v2, a[0:3] offset:8192509; GCN-MINREG-NEXT: ds_write_b128 v2, a[4:7] offset:8208510; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)511; GCN-MINREG-NEXT: ; sched_barrier mask(0x00000000)512; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v3 offset:24688513; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v3 offset:24672514; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v3 offset:24656515; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v3 offset:24640516; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v3 offset:24576517; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v3 offset:24592518; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v3 offset:24608519; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v3 offset:24624520; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)521; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v0, a[0:31]522; GCN-MINREG-NEXT: v_add_u32_e32 v4, 0x6000, v3523; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)524; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)525; GCN-MINREG-NEXT: s_nop 15526; GCN-MINREG-NEXT: s_nop 1527; GCN-MINREG-NEXT: ds_write_b128 v2, a[28:31] offset:16496528; GCN-MINREG-NEXT: ds_write_b128 v2, a[24:27] offset:16480529; GCN-MINREG-NEXT: ds_write_b128 v2, a[20:23] offset:16464530; GCN-MINREG-NEXT: ds_write_b128 v2, a[16:19] offset:16448531; GCN-MINREG-NEXT: ds_write_b128 v2, a[12:15] offset:16432532; GCN-MINREG-NEXT: ds_write_b128 v2, a[8:11] offset:16416533; GCN-MINREG-NEXT: ds_write_b128 v2, a[4:7] offset:16400534; GCN-MINREG-NEXT: ds_write_b128 v2, a[0:3] offset:16384535; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v3 offset:49264536; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v3 offset:49248537; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v3 offset:49232538; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v3 offset:49216539; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v3 offset:49200540; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v3 offset:49184541; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v3 offset:49168542; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v3 offset:49152543; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)544; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v0, a[0:31]545; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)546; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)547; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)548; GCN-MINREG-NEXT: s_nop 15549; GCN-MINREG-NEXT: s_nop 2550; GCN-MINREG-NEXT: ds_write_b128 v2, a[28:31] offset:24688551; GCN-MINREG-NEXT: ds_write_b128 v2, a[24:27] offset:24672552; GCN-MINREG-NEXT: ds_write_b128 v2, a[20:23] offset:24656553; GCN-MINREG-NEXT: ds_write_b128 v2, a[16:19] offset:24640554; GCN-MINREG-NEXT: ds_write_b128 v2, a[12:15] offset:24624555; GCN-MINREG-NEXT: ds_write_b128 v2, a[8:11] offset:24608556; GCN-MINREG-NEXT: ds_write_b128 v2, a[4:7] offset:24592557; GCN-MINREG-NEXT: ds_write_b128 v2, a[0:3] offset:24576558; GCN-MINREG-NEXT: ds_read_b128 a[28:31], v4 offset:57456559; GCN-MINREG-NEXT: ds_read_b128 a[24:27], v4 offset:57440560; GCN-MINREG-NEXT: ds_read_b128 a[20:23], v4 offset:57424561; GCN-MINREG-NEXT: ds_read_b128 a[16:19], v4 offset:57408562; GCN-MINREG-NEXT: ds_read_b128 a[0:3], v4 offset:57344563; GCN-MINREG-NEXT: ds_read_b128 a[4:7], v4 offset:57360564; GCN-MINREG-NEXT: ds_read_b128 a[8:11], v4 offset:57376565; GCN-MINREG-NEXT: ds_read_b128 a[12:15], v4 offset:57392566; GCN-MINREG-NEXT: s_waitcnt lgkmcnt(0)567; GCN-MINREG-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v0, a[0:31]568; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)569; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)570; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)571; GCN-MINREG-NEXT: s_nop 15572; GCN-MINREG-NEXT: s_nop 2573; GCN-MINREG-NEXT: ds_write_b128 v2, a[28:31] offset:32880574; GCN-MINREG-NEXT: ds_write_b128 v2, a[24:27] offset:32864575; GCN-MINREG-NEXT: ds_write_b128 v2, a[20:23] offset:32848576; GCN-MINREG-NEXT: ds_write_b128 v2, a[16:19] offset:32832577; GCN-MINREG-NEXT: ds_write_b128 v2, a[12:15] offset:32816578; GCN-MINREG-NEXT: ds_write_b128 v2, a[8:11] offset:32800579; GCN-MINREG-NEXT: ds_write_b128 v2, a[4:7] offset:32784580; GCN-MINREG-NEXT: ds_write_b128 v2, a[0:3] offset:32768581; GCN-MINREG-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)582; GCN-MINREG-NEXT: s_endpgm583;584; GCN-MAXOCC-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave_split_region:585; GCN-MAXOCC: ; %bb.0: ; %entry586; GCN-MAXOCC-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24587; GCN-MAXOCC-NEXT: v_lshlrev_b32_e32 v0, 7, v0588; GCN-MAXOCC-NEXT: v_and_b32_e32 v3, 0x1ff80, v0589; GCN-MAXOCC-NEXT: v_mov_b32_e32 v1, 1.0590; GCN-MAXOCC-NEXT: v_mov_b32_e32 v2, 2.0591; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)592; GCN-MAXOCC-NEXT: v_add_u32_e32 v0, s0, v3593; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:112594; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:96595; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:80596; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:64597; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0598; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:16599; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:32600; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:48601; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)602; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]603; GCN-MAXOCC-NEXT: v_add_u32_e32 v3, s1, v3604; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)605; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)606; GCN-MAXOCC-NEXT: s_nop 15607; GCN-MAXOCC-NEXT: s_nop 1608; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[28:31] offset:112609; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[24:27] offset:96610; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[20:23] offset:80611; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[16:19] offset:64612; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[12:15] offset:48613; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[8:11] offset:32614; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[4:7] offset:16615; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[0:3]616; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:8304617; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:8288618; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:8272619; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:8256620; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:8240621; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:8224622; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:8208623; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:8192624; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)625; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]626; GCN-MAXOCC-NEXT: v_mov_b32_e32 v3, s1627; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)628; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)629; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)630; GCN-MAXOCC-NEXT: s_nop 15631; GCN-MAXOCC-NEXT: s_nop 1632; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[24:27] offset:8288633; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[28:31] offset:8304634; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[16:19] offset:8256635; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[20:23] offset:8272636; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[8:11] offset:8224637; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[12:15] offset:8240638; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[0:3] offset:8192639; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[4:7] offset:8208640; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)641; GCN-MAXOCC-NEXT: ; sched_barrier mask(0x00000000)642; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:24688643; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:24672644; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:24656645; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:24640646; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:24576647; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:24592648; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:24608649; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:24624650; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)651; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]652; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)653; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)654; GCN-MAXOCC-NEXT: s_nop 15655; GCN-MAXOCC-NEXT: s_nop 2656; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[28:31] offset:16496657; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[24:27] offset:16480658; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[20:23] offset:16464659; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[16:19] offset:16448660; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[12:15] offset:16432661; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[8:11] offset:16416662; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[4:7] offset:16400663; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[0:3] offset:16384664; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:49264665; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:49248666; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:49232667; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:49216668; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:49200669; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:49184670; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:49168671; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:49152672; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)673; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]674; GCN-MAXOCC-NEXT: v_add_u32_e32 v0, 0x6000, v0675; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)676; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)677; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)678; GCN-MAXOCC-NEXT: s_nop 15679; GCN-MAXOCC-NEXT: s_nop 1680; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[28:31] offset:24688681; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[24:27] offset:24672682; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[20:23] offset:24656683; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[16:19] offset:24640684; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[12:15] offset:24624685; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[8:11] offset:24608686; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[4:7] offset:24592687; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[0:3] offset:24576688; GCN-MAXOCC-NEXT: ds_read_b128 a[28:31], v0 offset:57456689; GCN-MAXOCC-NEXT: ds_read_b128 a[24:27], v0 offset:57440690; GCN-MAXOCC-NEXT: ds_read_b128 a[20:23], v0 offset:57424691; GCN-MAXOCC-NEXT: ds_read_b128 a[16:19], v0 offset:57408692; GCN-MAXOCC-NEXT: ds_read_b128 a[0:3], v0 offset:57344693; GCN-MAXOCC-NEXT: ds_read_b128 a[4:7], v0 offset:57360694; GCN-MAXOCC-NEXT: ds_read_b128 a[8:11], v0 offset:57376695; GCN-MAXOCC-NEXT: ds_read_b128 a[12:15], v0 offset:57392696; GCN-MAXOCC-NEXT: s_waitcnt lgkmcnt(0)697; GCN-MAXOCC-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]698; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)699; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)700; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)701; GCN-MAXOCC-NEXT: s_nop 15702; GCN-MAXOCC-NEXT: s_nop 2703; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[28:31] offset:32880704; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[24:27] offset:32864705; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[20:23] offset:32848706; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[16:19] offset:32832707; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[12:15] offset:32816708; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[8:11] offset:32800709; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[4:7] offset:32784710; GCN-MAXOCC-NEXT: ds_write_b128 v3, a[0:3] offset:32768711; GCN-MAXOCC-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)712; GCN-MAXOCC-NEXT: s_endpgm713;714; GCN-ILP-LABEL: test_sched_group_barrier_pipeline_MFMA_interleave_split_region:715; GCN-ILP: ; %bb.0: ; %entry716; GCN-ILP-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24717; GCN-ILP-NEXT: v_lshlrev_b32_e32 v0, 7, v0718; GCN-ILP-NEXT: v_and_b32_e32 v2, 0x1ff80, v0719; GCN-ILP-NEXT: v_mov_b32_e32 v0, 1.0720; GCN-ILP-NEXT: v_mov_b32_e32 v1, 2.0721; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)722; GCN-ILP-NEXT: v_add_u32_e32 v3, s0, v2723; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:48724; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:32725; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:16726; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3727; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:64728; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:80729; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:96730; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:112731; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)732; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]733; GCN-ILP-NEXT: v_add_u32_e32 v2, s1, v2734; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)735; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)736; GCN-ILP-NEXT: s_nop 15737; GCN-ILP-NEXT: s_nop 1738; GCN-ILP-NEXT: ds_write_b128 v2, a[0:3]739; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:8192740; GCN-ILP-NEXT: ds_write_b128 v2, a[4:7] offset:16741; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:8208742; GCN-ILP-NEXT: ds_write_b128 v2, a[8:11] offset:32743; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:8224744; GCN-ILP-NEXT: ds_write_b128 v2, a[12:15] offset:48745; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:8240746; GCN-ILP-NEXT: ds_write_b128 v2, a[16:19] offset:64747; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:8256748; GCN-ILP-NEXT: ds_write_b128 v2, a[20:23] offset:80749; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:8272750; GCN-ILP-NEXT: ds_write_b128 v2, a[24:27] offset:96751; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:8288752; GCN-ILP-NEXT: ds_write_b128 v2, a[28:31] offset:112753; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:8304754; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)755; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]756; GCN-ILP-NEXT: v_mov_b32_e32 v2, s1757; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)758; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)759; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)760; GCN-ILP-NEXT: s_nop 15761; GCN-ILP-NEXT: s_nop 1762; GCN-ILP-NEXT: ds_write_b128 v2, a[24:27] offset:8288763; GCN-ILP-NEXT: ds_write_b128 v2, a[28:31] offset:8304764; GCN-ILP-NEXT: ds_write_b128 v2, a[16:19] offset:8256765; GCN-ILP-NEXT: ds_write_b128 v2, a[20:23] offset:8272766; GCN-ILP-NEXT: ds_write_b128 v2, a[8:11] offset:8224767; GCN-ILP-NEXT: ds_write_b128 v2, a[12:15] offset:8240768; GCN-ILP-NEXT: ds_write_b128 v2, a[0:3] offset:8192769; GCN-ILP-NEXT: ds_write_b128 v2, a[4:7] offset:8208770; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)771; GCN-ILP-NEXT: ; sched_barrier mask(0x00000000)772; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:24624773; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:24608774; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:24592775; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:24576776; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:24640777; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:24656778; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:24672779; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:24688780; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)781; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]782; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)783; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)784; GCN-ILP-NEXT: s_nop 15785; GCN-ILP-NEXT: s_nop 2786; GCN-ILP-NEXT: ds_write_b128 v2, a[28:31] offset:16496787; GCN-ILP-NEXT: ds_write_b128 v2, a[24:27] offset:16480788; GCN-ILP-NEXT: ds_write_b128 v2, a[20:23] offset:16464789; GCN-ILP-NEXT: ds_write_b128 v2, a[16:19] offset:16448790; GCN-ILP-NEXT: ds_write_b128 v2, a[12:15] offset:16432791; GCN-ILP-NEXT: ds_write_b128 v2, a[8:11] offset:16416792; GCN-ILP-NEXT: ds_write_b128 v2, a[4:7] offset:16400793; GCN-ILP-NEXT: ds_write_b128 v2, a[0:3] offset:16384794; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:49152795; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:49168796; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:49184797; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:49200798; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:49216799; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:49232800; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:49248801; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:49264802; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)803; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]804; GCN-ILP-NEXT: v_add_u32_e32 v3, 0x6000, v3805; GCN-ILP-NEXT: s_nop 15806; GCN-ILP-NEXT: s_nop 1807; GCN-ILP-NEXT: ds_write_b128 v2, a[0:3] offset:24576808; GCN-ILP-NEXT: ds_read_b128 a[0:3], v3 offset:57344809; GCN-ILP-NEXT: ds_write_b128 v2, a[4:7] offset:24592810; GCN-ILP-NEXT: ds_read_b128 a[4:7], v3 offset:57360811; GCN-ILP-NEXT: ds_write_b128 v2, a[8:11] offset:24608812; GCN-ILP-NEXT: ds_read_b128 a[8:11], v3 offset:57376813; GCN-ILP-NEXT: ds_write_b128 v2, a[12:15] offset:24624814; GCN-ILP-NEXT: ds_read_b128 a[12:15], v3 offset:57392815; GCN-ILP-NEXT: ds_write_b128 v2, a[16:19] offset:24640816; GCN-ILP-NEXT: ds_read_b128 a[16:19], v3 offset:57408817; GCN-ILP-NEXT: ds_write_b128 v2, a[20:23] offset:24656818; GCN-ILP-NEXT: ds_read_b128 a[20:23], v3 offset:57424819; GCN-ILP-NEXT: ds_write_b128 v2, a[24:27] offset:24672820; GCN-ILP-NEXT: ds_read_b128 a[24:27], v3 offset:57440821; GCN-ILP-NEXT: ds_write_b128 v2, a[28:31] offset:24688822; GCN-ILP-NEXT: ds_read_b128 a[28:31], v3 offset:57456823; GCN-ILP-NEXT: s_waitcnt lgkmcnt(0)824; GCN-ILP-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]825; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)826; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)827; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)828; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000100) size(8) SyncID(0)829; GCN-ILP-NEXT: s_nop 15830; GCN-ILP-NEXT: s_nop 2831; GCN-ILP-NEXT: ds_write_b128 v2, a[28:31] offset:32880832; GCN-ILP-NEXT: ds_write_b128 v2, a[24:27] offset:32864833; GCN-ILP-NEXT: ds_write_b128 v2, a[20:23] offset:32848834; GCN-ILP-NEXT: ds_write_b128 v2, a[16:19] offset:32832835; GCN-ILP-NEXT: ds_write_b128 v2, a[12:15] offset:32816836; GCN-ILP-NEXT: ds_write_b128 v2, a[8:11] offset:32800837; GCN-ILP-NEXT: ds_write_b128 v2, a[4:7] offset:32784838; GCN-ILP-NEXT: ds_write_b128 v2, a[0:3] offset:32768839; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)840; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000008) size(1) SyncID(0)841; GCN-ILP-NEXT: ; sched_group_barrier mask(0x00000200) size(8) SyncID(0)842; GCN-ILP-NEXT: s_endpgm843entry:844 %idx = call i32 @llvm.amdgcn.workitem.id.x()845 %load.0.addr = getelementptr <32 x float>, ptr addrspace(3) %in, i32 %idx846 %load.0 = load <32 x float>, ptr addrspace(3) %load.0.addr847 %load.1.addr = getelementptr <32 x float>, ptr addrspace(3) %load.0.addr, i32 64848 %load.1 = load <32 x float>, ptr addrspace(3) %load.1.addr849 %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)850 %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)851 %store.0.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 %idx852 store <32 x float> %mai.0, ptr addrspace(3) %store.0.addr853 %store.1.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 64854 store <32 x float> %mai.1, ptr addrspace(3) %store.1.addr855 ; 8 DS read856 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)857 ; 1 MFMA858 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)859 ; 8 DS write860 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)861 ; 8 DS read862 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)863 ; 1 MFMA864 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)865 ; 8 DS write866 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)867 ; parition the regions868 call void @llvm.amdgcn.sched.barrier(i32 0)869 %load.2.addr = getelementptr <32 x float>, ptr addrspace(3) %load.1.addr, i32 128870 %load.2 = load <32 x float>, ptr addrspace(3) %load.2.addr871 %load.3.addr = getelementptr <32 x float>, ptr addrspace(3) %load.2.addr, i32 192872 %load.3 = load <32 x float>, ptr addrspace(3) %load.3.addr873 %load.4.addr = getelementptr <32 x float>, ptr addrspace(3) %load.3.addr, i32 256874 %load.4 = load <32 x float>, ptr addrspace(3) %load.4.addr875 %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)876 %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)877 %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)878 %store.2.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 128879 store <32 x float> %mai.2, ptr addrspace(3) %store.2.addr880 %store.3.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 192881 store <32 x float> %mai.3, ptr addrspace(3) %store.3.addr882 %store.4.addr = getelementptr <32 x float>, ptr addrspace(3) %out, i32 256883 store <32 x float> %mai.4, ptr addrspace(3) %store.4.addr884 ; 8 DS read885 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)886 ; 1 MFMA887 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)888 ; 8 DS write889 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)890 ; 8 DS read891 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)892 ; 1 MFMA893 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)894 ; 8 DS write895 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)896 ; 8 DS read897 call void @llvm.amdgcn.sched.group.barrier(i32 256, i32 8, i32 0)898 ; 1 MFMA899 call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)900 ; 8 DS write901 call void @llvm.amdgcn.sched.group.barrier(i32 512, i32 8, i32 0)902 ret void903}904