425 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -amdgpu-mfma-vgpr-form < %s | FileCheck %s3 4; After reassigning the MFMA to use AGPRs, we've alleviated enough5; register pressure to try eliminating the spill of %spill with the freed6; up VGPR.7define void @eliminate_spill_after_mfma_rewrite(i32 %x, i32 %y, <4 x i32> %arg, ptr addrspace(1) inreg %ptr) #0 {8; CHECK-LABEL: eliminate_spill_after_mfma_rewrite:9; CHECK: ; %bb.0:10; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)11; CHECK-NEXT: v_accvgpr_write_b32 a3, v512; CHECK-NEXT: v_accvgpr_write_b32 a2, v413; CHECK-NEXT: v_accvgpr_write_b32 a1, v314; CHECK-NEXT: v_accvgpr_write_b32 a0, v215; CHECK-NEXT: buffer_store_dword v40, off, s[0:3], s32 offset:188 ; 4-byte Folded Spill16; CHECK-NEXT: buffer_store_dword v41, off, s[0:3], s32 offset:184 ; 4-byte Folded Spill17; CHECK-NEXT: buffer_store_dword v42, off, s[0:3], s32 offset:180 ; 4-byte Folded Spill18; CHECK-NEXT: buffer_store_dword v43, off, s[0:3], s32 offset:176 ; 4-byte Folded Spill19; CHECK-NEXT: buffer_store_dword v44, off, s[0:3], s32 offset:172 ; 4-byte Folded Spill20; CHECK-NEXT: buffer_store_dword v45, off, s[0:3], s32 offset:168 ; 4-byte Folded Spill21; CHECK-NEXT: buffer_store_dword v46, off, s[0:3], s32 offset:164 ; 4-byte Folded Spill22; CHECK-NEXT: buffer_store_dword v47, off, s[0:3], s32 offset:160 ; 4-byte Folded Spill23; CHECK-NEXT: buffer_store_dword v56, off, s[0:3], s32 offset:156 ; 4-byte Folded Spill24; CHECK-NEXT: buffer_store_dword v57, off, s[0:3], s32 offset:152 ; 4-byte Folded Spill25; CHECK-NEXT: buffer_store_dword v58, off, s[0:3], s32 offset:148 ; 4-byte Folded Spill26; CHECK-NEXT: buffer_store_dword v59, off, s[0:3], s32 offset:144 ; 4-byte Folded Spill27; CHECK-NEXT: buffer_store_dword v60, off, s[0:3], s32 offset:140 ; 4-byte Folded Spill28; CHECK-NEXT: buffer_store_dword v61, off, s[0:3], s32 offset:136 ; 4-byte Folded Spill29; CHECK-NEXT: buffer_store_dword v62, off, s[0:3], s32 offset:132 ; 4-byte Folded Spill30; CHECK-NEXT: buffer_store_dword v63, off, s[0:3], s32 offset:128 ; 4-byte Folded Spill31; CHECK-NEXT: buffer_store_dword a32, off, s[0:3], s32 offset:124 ; 4-byte Folded Spill32; CHECK-NEXT: buffer_store_dword a33, off, s[0:3], s32 offset:120 ; 4-byte Folded Spill33; CHECK-NEXT: buffer_store_dword a34, off, s[0:3], s32 offset:116 ; 4-byte Folded Spill34; CHECK-NEXT: buffer_store_dword a35, off, s[0:3], s32 offset:112 ; 4-byte Folded Spill35; CHECK-NEXT: buffer_store_dword a36, off, s[0:3], s32 offset:108 ; 4-byte Folded Spill36; CHECK-NEXT: buffer_store_dword a37, off, s[0:3], s32 offset:104 ; 4-byte Folded Spill37; CHECK-NEXT: buffer_store_dword a38, off, s[0:3], s32 offset:100 ; 4-byte Folded Spill38; CHECK-NEXT: buffer_store_dword a39, off, s[0:3], s32 offset:96 ; 4-byte Folded Spill39; CHECK-NEXT: buffer_store_dword a40, off, s[0:3], s32 offset:92 ; 4-byte Folded Spill40; CHECK-NEXT: buffer_store_dword a41, off, s[0:3], s32 offset:88 ; 4-byte Folded Spill41; CHECK-NEXT: buffer_store_dword a42, off, s[0:3], s32 offset:84 ; 4-byte Folded Spill42; CHECK-NEXT: buffer_store_dword a43, off, s[0:3], s32 offset:80 ; 4-byte Folded Spill43; CHECK-NEXT: buffer_store_dword a44, off, s[0:3], s32 offset:76 ; 4-byte Folded Spill44; CHECK-NEXT: buffer_store_dword a45, off, s[0:3], s32 offset:72 ; 4-byte Folded Spill45; CHECK-NEXT: buffer_store_dword a46, off, s[0:3], s32 offset:68 ; 4-byte Folded Spill46; CHECK-NEXT: buffer_store_dword a47, off, s[0:3], s32 offset:64 ; 4-byte Folded Spill47; CHECK-NEXT: buffer_store_dword a48, off, s[0:3], s32 offset:60 ; 4-byte Folded Spill48; CHECK-NEXT: buffer_store_dword a49, off, s[0:3], s32 offset:56 ; 4-byte Folded Spill49; CHECK-NEXT: buffer_store_dword a50, off, s[0:3], s32 offset:52 ; 4-byte Folded Spill50; CHECK-NEXT: buffer_store_dword a51, off, s[0:3], s32 offset:48 ; 4-byte Folded Spill51; CHECK-NEXT: buffer_store_dword a52, off, s[0:3], s32 offset:44 ; 4-byte Folded Spill52; CHECK-NEXT: buffer_store_dword a53, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill53; CHECK-NEXT: buffer_store_dword a54, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill54; CHECK-NEXT: buffer_store_dword a55, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill55; CHECK-NEXT: buffer_store_dword a56, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill56; CHECK-NEXT: buffer_store_dword a57, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill57; CHECK-NEXT: buffer_store_dword a58, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill58; CHECK-NEXT: buffer_store_dword a59, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill59; CHECK-NEXT: buffer_store_dword a60, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill60; CHECK-NEXT: buffer_store_dword a61, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill61; CHECK-NEXT: buffer_store_dword a62, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill62; CHECK-NEXT: buffer_store_dword a63, off, s[0:3], s32 ; 4-byte Folded Spill63; CHECK-NEXT: v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[0:3]64; CHECK-NEXT: ;;#ASMSTART65; CHECK-NEXT: ; def v[32:63], v[0:31]66; CHECK-NEXT: ;;#ASMEND67; CHECK-NEXT: v_accvgpr_write_b32 a63, v3168; CHECK-NEXT: v_accvgpr_write_b32 a62, v3069; CHECK-NEXT: v_accvgpr_write_b32 a61, v2970; CHECK-NEXT: v_accvgpr_write_b32 a60, v2871; CHECK-NEXT: v_accvgpr_write_b32 a59, v2772; CHECK-NEXT: v_accvgpr_write_b32 a58, v2673; CHECK-NEXT: v_accvgpr_write_b32 a57, v2574; CHECK-NEXT: v_accvgpr_write_b32 a56, v2475; CHECK-NEXT: v_accvgpr_write_b32 a55, v2376; CHECK-NEXT: v_accvgpr_write_b32 a54, v2277; CHECK-NEXT: v_accvgpr_write_b32 a53, v2178; CHECK-NEXT: v_accvgpr_write_b32 a52, v2079; CHECK-NEXT: v_accvgpr_write_b32 a51, v1980; CHECK-NEXT: v_accvgpr_write_b32 a50, v1881; CHECK-NEXT: v_accvgpr_write_b32 a49, v1782; CHECK-NEXT: v_accvgpr_write_b32 a48, v1683; CHECK-NEXT: v_accvgpr_write_b32 a47, v1584; CHECK-NEXT: v_accvgpr_write_b32 a46, v1485; CHECK-NEXT: v_accvgpr_write_b32 a45, v1386; CHECK-NEXT: v_accvgpr_write_b32 a44, v1287; CHECK-NEXT: v_accvgpr_write_b32 a43, v1188; CHECK-NEXT: v_accvgpr_write_b32 a42, v1089; CHECK-NEXT: v_accvgpr_write_b32 a41, v990; CHECK-NEXT: v_accvgpr_write_b32 a40, v891; CHECK-NEXT: v_accvgpr_write_b32 a39, v792; CHECK-NEXT: v_accvgpr_write_b32 a38, v693; CHECK-NEXT: v_accvgpr_write_b32 a37, v594; CHECK-NEXT: v_accvgpr_write_b32 a36, v495; CHECK-NEXT: v_accvgpr_write_b32 a35, v396; CHECK-NEXT: v_accvgpr_write_b32 a34, v297; CHECK-NEXT: v_accvgpr_write_b32 a33, v198; CHECK-NEXT: v_accvgpr_write_b32 a32, v099; CHECK-NEXT: v_accvgpr_read_b32 v0, a0100; CHECK-NEXT: v_accvgpr_read_b32 v1, a1101; CHECK-NEXT: v_accvgpr_read_b32 v2, a2102; CHECK-NEXT: v_accvgpr_read_b32 v3, a3103; CHECK-NEXT: ;;#ASMSTART104; CHECK-NEXT: ; def v[6:9]105; CHECK-NEXT: ;;#ASMEND106; CHECK-NEXT: v_mov_b32_e32 v0, 0107; CHECK-NEXT: ;;#ASMSTART108; CHECK-NEXT: ; def a[0:31]109; CHECK-NEXT: ;;#ASMEND110; CHECK-NEXT: ;;#ASMSTART111; CHECK-NEXT: ; def a[0:31]112; CHECK-NEXT: ;;#ASMEND113; CHECK-NEXT: global_store_dwordx4 v0, v[60:63], s[16:17] offset:112114; CHECK-NEXT: s_waitcnt vmcnt(0)115; CHECK-NEXT: global_store_dwordx4 v0, v[56:59], s[16:17] offset:96116; CHECK-NEXT: s_waitcnt vmcnt(0)117; CHECK-NEXT: global_store_dwordx4 v0, v[52:55], s[16:17] offset:80118; CHECK-NEXT: s_waitcnt vmcnt(0)119; CHECK-NEXT: global_store_dwordx4 v0, v[48:51], s[16:17] offset:64120; CHECK-NEXT: s_waitcnt vmcnt(0)121; CHECK-NEXT: global_store_dwordx4 v0, v[44:47], s[16:17] offset:48122; CHECK-NEXT: s_waitcnt vmcnt(0)123; CHECK-NEXT: global_store_dwordx4 v0, v[40:43], s[16:17] offset:32124; CHECK-NEXT: s_waitcnt vmcnt(0)125; CHECK-NEXT: global_store_dwordx4 v0, v[36:39], s[16:17] offset:16126; CHECK-NEXT: s_waitcnt vmcnt(0)127; CHECK-NEXT: global_store_dwordx4 v0, v[32:35], s[16:17]128; CHECK-NEXT: s_waitcnt vmcnt(0)129; CHECK-NEXT: global_store_dwordx4 v0, a[56:59], s[16:17] offset:96130; CHECK-NEXT: s_waitcnt vmcnt(0)131; CHECK-NEXT: global_store_dwordx4 v0, a[60:63], s[16:17] offset:112132; CHECK-NEXT: s_waitcnt vmcnt(0)133; CHECK-NEXT: global_store_dwordx4 v0, a[48:51], s[16:17] offset:64134; CHECK-NEXT: s_waitcnt vmcnt(0)135; CHECK-NEXT: global_store_dwordx4 v0, a[52:55], s[16:17] offset:80136; CHECK-NEXT: s_waitcnt vmcnt(0)137; CHECK-NEXT: global_store_dwordx4 v0, a[40:43], s[16:17] offset:32138; CHECK-NEXT: s_waitcnt vmcnt(0)139; CHECK-NEXT: global_store_dwordx4 v0, a[44:47], s[16:17] offset:48140; CHECK-NEXT: s_waitcnt vmcnt(0)141; CHECK-NEXT: global_store_dwordx4 v0, a[32:35], s[16:17]142; CHECK-NEXT: s_waitcnt vmcnt(0)143; CHECK-NEXT: global_store_dwordx4 v0, a[36:39], s[16:17] offset:16144; CHECK-NEXT: s_waitcnt vmcnt(0)145; CHECK-NEXT: global_store_dwordx4 v0, v[6:9], s[16:17]146; CHECK-NEXT: s_waitcnt vmcnt(0)147; CHECK-NEXT: buffer_load_dword a63, off, s[0:3], s32 ; 4-byte Folded Reload148; CHECK-NEXT: buffer_load_dword a62, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload149; CHECK-NEXT: buffer_load_dword a61, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload150; CHECK-NEXT: buffer_load_dword a60, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload151; CHECK-NEXT: buffer_load_dword a59, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload152; CHECK-NEXT: buffer_load_dword a58, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload153; CHECK-NEXT: buffer_load_dword a57, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload154; CHECK-NEXT: buffer_load_dword a56, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload155; CHECK-NEXT: buffer_load_dword a55, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload156; CHECK-NEXT: buffer_load_dword a54, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload157; CHECK-NEXT: buffer_load_dword a53, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload158; CHECK-NEXT: buffer_load_dword a52, off, s[0:3], s32 offset:44 ; 4-byte Folded Reload159; CHECK-NEXT: buffer_load_dword a51, off, s[0:3], s32 offset:48 ; 4-byte Folded Reload160; CHECK-NEXT: buffer_load_dword a50, off, s[0:3], s32 offset:52 ; 4-byte Folded Reload161; CHECK-NEXT: buffer_load_dword a49, off, s[0:3], s32 offset:56 ; 4-byte Folded Reload162; CHECK-NEXT: buffer_load_dword a48, off, s[0:3], s32 offset:60 ; 4-byte Folded Reload163; CHECK-NEXT: buffer_load_dword a47, off, s[0:3], s32 offset:64 ; 4-byte Folded Reload164; CHECK-NEXT: buffer_load_dword a46, off, s[0:3], s32 offset:68 ; 4-byte Folded Reload165; CHECK-NEXT: buffer_load_dword a45, off, s[0:3], s32 offset:72 ; 4-byte Folded Reload166; CHECK-NEXT: buffer_load_dword a44, off, s[0:3], s32 offset:76 ; 4-byte Folded Reload167; CHECK-NEXT: buffer_load_dword a43, off, s[0:3], s32 offset:80 ; 4-byte Folded Reload168; CHECK-NEXT: buffer_load_dword a42, off, s[0:3], s32 offset:84 ; 4-byte Folded Reload169; CHECK-NEXT: buffer_load_dword a41, off, s[0:3], s32 offset:88 ; 4-byte Folded Reload170; CHECK-NEXT: buffer_load_dword a40, off, s[0:3], s32 offset:92 ; 4-byte Folded Reload171; CHECK-NEXT: buffer_load_dword a39, off, s[0:3], s32 offset:96 ; 4-byte Folded Reload172; CHECK-NEXT: buffer_load_dword a38, off, s[0:3], s32 offset:100 ; 4-byte Folded Reload173; CHECK-NEXT: buffer_load_dword a37, off, s[0:3], s32 offset:104 ; 4-byte Folded Reload174; CHECK-NEXT: buffer_load_dword a36, off, s[0:3], s32 offset:108 ; 4-byte Folded Reload175; CHECK-NEXT: buffer_load_dword a35, off, s[0:3], s32 offset:112 ; 4-byte Folded Reload176; CHECK-NEXT: buffer_load_dword a34, off, s[0:3], s32 offset:116 ; 4-byte Folded Reload177; CHECK-NEXT: buffer_load_dword a33, off, s[0:3], s32 offset:120 ; 4-byte Folded Reload178; CHECK-NEXT: buffer_load_dword a32, off, s[0:3], s32 offset:124 ; 4-byte Folded Reload179; CHECK-NEXT: buffer_load_dword v63, off, s[0:3], s32 offset:128 ; 4-byte Folded Reload180; CHECK-NEXT: buffer_load_dword v62, off, s[0:3], s32 offset:132 ; 4-byte Folded Reload181; CHECK-NEXT: buffer_load_dword v61, off, s[0:3], s32 offset:136 ; 4-byte Folded Reload182; CHECK-NEXT: buffer_load_dword v60, off, s[0:3], s32 offset:140 ; 4-byte Folded Reload183; CHECK-NEXT: buffer_load_dword v59, off, s[0:3], s32 offset:144 ; 4-byte Folded Reload184; CHECK-NEXT: buffer_load_dword v58, off, s[0:3], s32 offset:148 ; 4-byte Folded Reload185; CHECK-NEXT: buffer_load_dword v57, off, s[0:3], s32 offset:152 ; 4-byte Folded Reload186; CHECK-NEXT: buffer_load_dword v56, off, s[0:3], s32 offset:156 ; 4-byte Folded Reload187; CHECK-NEXT: buffer_load_dword v47, off, s[0:3], s32 offset:160 ; 4-byte Folded Reload188; CHECK-NEXT: buffer_load_dword v46, off, s[0:3], s32 offset:164 ; 4-byte Folded Reload189; CHECK-NEXT: buffer_load_dword v45, off, s[0:3], s32 offset:168 ; 4-byte Folded Reload190; CHECK-NEXT: buffer_load_dword v44, off, s[0:3], s32 offset:172 ; 4-byte Folded Reload191; CHECK-NEXT: buffer_load_dword v43, off, s[0:3], s32 offset:176 ; 4-byte Folded Reload192; CHECK-NEXT: buffer_load_dword v42, off, s[0:3], s32 offset:180 ; 4-byte Folded Reload193; CHECK-NEXT: buffer_load_dword v41, off, s[0:3], s32 offset:184 ; 4-byte Folded Reload194; CHECK-NEXT: buffer_load_dword v40, off, s[0:3], s32 offset:188 ; 4-byte Folded Reload195; CHECK-NEXT: s_waitcnt vmcnt(0)196; CHECK-NEXT: s_setpc_b64 s[30:31]197 %mai = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 %x, i32 %y, <4 x i32> %arg, i32 0, i32 0, i32 0)198 %v = call { <32 x i32>, <32 x i32> } asm sideeffect "; def $0, $1", "=v,=v"()199 %v0 = extractvalue { <32 x i32>, <32 x i32> } %v, 0200 %v1 = extractvalue { <32 x i32>, <32 x i32> } %v, 1201 %spill = call <4 x i32> asm sideeffect "; def $0", "=v,v"(<4 x i32> %mai)202 %a0 = call <32 x i32> asm sideeffect "; def $0", "=a"()203 %a1 = call <32 x i32> asm sideeffect "; def $0", "=a"()204 store volatile <32 x i32> %v0, ptr addrspace(1) %ptr205 store volatile <32 x i32> %v1, ptr addrspace(1) %ptr206 store volatile <4 x i32> %spill, ptr addrspace(1) %ptr207 ret void208}209 210; Same, except we fold out 2 spills from %spill0 and %spill1211define void @eliminate_spill_after_mfma_rewrite_x2(i32 %x, i32 %y, <4 x i32> %arg, ptr addrspace(1) inreg %ptr) #0 {212; CHECK-LABEL: eliminate_spill_after_mfma_rewrite_x2:213; CHECK: ; %bb.0:214; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)215; CHECK-NEXT: v_accvgpr_write_b32 a3, v5216; CHECK-NEXT: v_accvgpr_write_b32 a2, v4217; CHECK-NEXT: v_accvgpr_write_b32 a1, v3218; CHECK-NEXT: v_accvgpr_write_b32 a0, v2219; CHECK-NEXT: buffer_store_dword v40, off, s[0:3], s32 offset:188 ; 4-byte Folded Spill220; CHECK-NEXT: buffer_store_dword v41, off, s[0:3], s32 offset:184 ; 4-byte Folded Spill221; CHECK-NEXT: buffer_store_dword v42, off, s[0:3], s32 offset:180 ; 4-byte Folded Spill222; CHECK-NEXT: buffer_store_dword v43, off, s[0:3], s32 offset:176 ; 4-byte Folded Spill223; CHECK-NEXT: buffer_store_dword v44, off, s[0:3], s32 offset:172 ; 4-byte Folded Spill224; CHECK-NEXT: buffer_store_dword v45, off, s[0:3], s32 offset:168 ; 4-byte Folded Spill225; CHECK-NEXT: buffer_store_dword v46, off, s[0:3], s32 offset:164 ; 4-byte Folded Spill226; CHECK-NEXT: buffer_store_dword v47, off, s[0:3], s32 offset:160 ; 4-byte Folded Spill227; CHECK-NEXT: buffer_store_dword v56, off, s[0:3], s32 offset:156 ; 4-byte Folded Spill228; CHECK-NEXT: buffer_store_dword v57, off, s[0:3], s32 offset:152 ; 4-byte Folded Spill229; CHECK-NEXT: buffer_store_dword v58, off, s[0:3], s32 offset:148 ; 4-byte Folded Spill230; CHECK-NEXT: buffer_store_dword v59, off, s[0:3], s32 offset:144 ; 4-byte Folded Spill231; CHECK-NEXT: buffer_store_dword v60, off, s[0:3], s32 offset:140 ; 4-byte Folded Spill232; CHECK-NEXT: buffer_store_dword v61, off, s[0:3], s32 offset:136 ; 4-byte Folded Spill233; CHECK-NEXT: buffer_store_dword v62, off, s[0:3], s32 offset:132 ; 4-byte Folded Spill234; CHECK-NEXT: buffer_store_dword v63, off, s[0:3], s32 offset:128 ; 4-byte Folded Spill235; CHECK-NEXT: buffer_store_dword a32, off, s[0:3], s32 offset:124 ; 4-byte Folded Spill236; CHECK-NEXT: buffer_store_dword a33, off, s[0:3], s32 offset:120 ; 4-byte Folded Spill237; CHECK-NEXT: buffer_store_dword a34, off, s[0:3], s32 offset:116 ; 4-byte Folded Spill238; CHECK-NEXT: buffer_store_dword a35, off, s[0:3], s32 offset:112 ; 4-byte Folded Spill239; CHECK-NEXT: buffer_store_dword a36, off, s[0:3], s32 offset:108 ; 4-byte Folded Spill240; CHECK-NEXT: buffer_store_dword a37, off, s[0:3], s32 offset:104 ; 4-byte Folded Spill241; CHECK-NEXT: buffer_store_dword a38, off, s[0:3], s32 offset:100 ; 4-byte Folded Spill242; CHECK-NEXT: buffer_store_dword a39, off, s[0:3], s32 offset:96 ; 4-byte Folded Spill243; CHECK-NEXT: buffer_store_dword a40, off, s[0:3], s32 offset:92 ; 4-byte Folded Spill244; CHECK-NEXT: buffer_store_dword a41, off, s[0:3], s32 offset:88 ; 4-byte Folded Spill245; CHECK-NEXT: buffer_store_dword a42, off, s[0:3], s32 offset:84 ; 4-byte Folded Spill246; CHECK-NEXT: buffer_store_dword a43, off, s[0:3], s32 offset:80 ; 4-byte Folded Spill247; CHECK-NEXT: buffer_store_dword a44, off, s[0:3], s32 offset:76 ; 4-byte Folded Spill248; CHECK-NEXT: buffer_store_dword a45, off, s[0:3], s32 offset:72 ; 4-byte Folded Spill249; CHECK-NEXT: buffer_store_dword a46, off, s[0:3], s32 offset:68 ; 4-byte Folded Spill250; CHECK-NEXT: buffer_store_dword a47, off, s[0:3], s32 offset:64 ; 4-byte Folded Spill251; CHECK-NEXT: buffer_store_dword a48, off, s[0:3], s32 offset:60 ; 4-byte Folded Spill252; CHECK-NEXT: buffer_store_dword a49, off, s[0:3], s32 offset:56 ; 4-byte Folded Spill253; CHECK-NEXT: buffer_store_dword a50, off, s[0:3], s32 offset:52 ; 4-byte Folded Spill254; CHECK-NEXT: buffer_store_dword a51, off, s[0:3], s32 offset:48 ; 4-byte Folded Spill255; CHECK-NEXT: buffer_store_dword a52, off, s[0:3], s32 offset:44 ; 4-byte Folded Spill256; CHECK-NEXT: buffer_store_dword a53, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill257; CHECK-NEXT: buffer_store_dword a54, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill258; CHECK-NEXT: buffer_store_dword a55, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill259; CHECK-NEXT: buffer_store_dword a56, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill260; CHECK-NEXT: buffer_store_dword a57, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill261; CHECK-NEXT: buffer_store_dword a58, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill262; CHECK-NEXT: buffer_store_dword a59, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill263; CHECK-NEXT: buffer_store_dword a60, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill264; CHECK-NEXT: buffer_store_dword a61, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill265; CHECK-NEXT: buffer_store_dword a62, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill266; CHECK-NEXT: buffer_store_dword a63, off, s[0:3], s32 ; 4-byte Folded Spill267; CHECK-NEXT: v_mfma_i32_4x4x4i8 a[0:3], v0, v1, a[0:3]268; CHECK-NEXT: ;;#ASMSTART269; CHECK-NEXT: ; def v[32:63], v[0:31]270; CHECK-NEXT: ;;#ASMEND271; CHECK-NEXT: v_accvgpr_write_b32 a63, v31272; CHECK-NEXT: v_accvgpr_write_b32 a62, v30273; CHECK-NEXT: v_accvgpr_write_b32 a61, v29274; CHECK-NEXT: v_accvgpr_write_b32 a60, v28275; CHECK-NEXT: v_accvgpr_write_b32 a59, v27276; CHECK-NEXT: v_accvgpr_write_b32 a58, v26277; CHECK-NEXT: v_accvgpr_write_b32 a57, v25278; CHECK-NEXT: v_accvgpr_write_b32 a56, v24279; CHECK-NEXT: v_accvgpr_write_b32 a55, v23280; CHECK-NEXT: v_accvgpr_write_b32 a54, v22281; CHECK-NEXT: v_accvgpr_write_b32 a53, v21282; CHECK-NEXT: v_accvgpr_write_b32 a52, v20283; CHECK-NEXT: v_accvgpr_write_b32 a51, v19284; CHECK-NEXT: v_accvgpr_write_b32 a50, v18285; CHECK-NEXT: v_accvgpr_write_b32 a49, v17286; CHECK-NEXT: v_accvgpr_write_b32 a48, v16287; CHECK-NEXT: v_accvgpr_write_b32 a47, v15288; CHECK-NEXT: v_accvgpr_write_b32 a46, v14289; CHECK-NEXT: v_accvgpr_write_b32 a45, v13290; CHECK-NEXT: v_accvgpr_write_b32 a44, v12291; CHECK-NEXT: v_accvgpr_write_b32 a43, v11292; CHECK-NEXT: v_accvgpr_write_b32 a42, v10293; CHECK-NEXT: v_accvgpr_write_b32 a41, v9294; CHECK-NEXT: v_accvgpr_write_b32 a40, v8295; CHECK-NEXT: v_accvgpr_write_b32 a39, v7296; CHECK-NEXT: v_accvgpr_write_b32 a38, v6297; CHECK-NEXT: v_accvgpr_write_b32 a37, v5298; CHECK-NEXT: v_accvgpr_write_b32 a36, v4299; CHECK-NEXT: v_accvgpr_write_b32 a35, v3300; CHECK-NEXT: v_accvgpr_write_b32 a34, v2301; CHECK-NEXT: v_accvgpr_write_b32 a33, v1302; CHECK-NEXT: v_accvgpr_write_b32 a32, v0303; CHECK-NEXT: v_accvgpr_read_b32 v7, a3304; CHECK-NEXT: v_mov_b32_e32 v0, 0305; CHECK-NEXT: v_accvgpr_read_b32 v6, a2306; CHECK-NEXT: v_accvgpr_read_b32 v5, a1307; CHECK-NEXT: v_accvgpr_read_b32 v4, a0308; CHECK-NEXT: ;;#ASMSTART309; CHECK-NEXT: ; def v[8:11]310; CHECK-NEXT: ;;#ASMEND311; CHECK-NEXT: ;;#ASMSTART312; CHECK-NEXT: ; def v[12:15]313; CHECK-NEXT: ;;#ASMEND314; CHECK-NEXT: ;;#ASMSTART315; CHECK-NEXT: ; def a[0:31]316; CHECK-NEXT: ;;#ASMEND317; CHECK-NEXT: ;;#ASMSTART318; CHECK-NEXT: ; def a[0:31]319; CHECK-NEXT: ;;#ASMEND320; CHECK-NEXT: global_store_dwordx4 v0, v[60:63], s[16:17] offset:112321; CHECK-NEXT: s_waitcnt vmcnt(0)322; CHECK-NEXT: global_store_dwordx4 v0, v[56:59], s[16:17] offset:96323; CHECK-NEXT: s_waitcnt vmcnt(0)324; CHECK-NEXT: global_store_dwordx4 v0, v[52:55], s[16:17] offset:80325; CHECK-NEXT: s_waitcnt vmcnt(0)326; CHECK-NEXT: global_store_dwordx4 v0, v[48:51], s[16:17] offset:64327; CHECK-NEXT: s_waitcnt vmcnt(0)328; CHECK-NEXT: global_store_dwordx4 v0, v[44:47], s[16:17] offset:48329; CHECK-NEXT: s_waitcnt vmcnt(0)330; CHECK-NEXT: global_store_dwordx4 v0, v[40:43], s[16:17] offset:32331; CHECK-NEXT: s_waitcnt vmcnt(0)332; CHECK-NEXT: global_store_dwordx4 v0, v[36:39], s[16:17] offset:16333; CHECK-NEXT: s_waitcnt vmcnt(0)334; CHECK-NEXT: global_store_dwordx4 v0, v[32:35], s[16:17]335; CHECK-NEXT: s_waitcnt vmcnt(0)336; CHECK-NEXT: global_store_dwordx4 v0, a[56:59], s[16:17] offset:96337; CHECK-NEXT: s_waitcnt vmcnt(0)338; CHECK-NEXT: global_store_dwordx4 v0, a[60:63], s[16:17] offset:112339; CHECK-NEXT: s_waitcnt vmcnt(0)340; CHECK-NEXT: global_store_dwordx4 v0, a[48:51], s[16:17] offset:64341; CHECK-NEXT: s_waitcnt vmcnt(0)342; CHECK-NEXT: global_store_dwordx4 v0, a[52:55], s[16:17] offset:80343; CHECK-NEXT: s_waitcnt vmcnt(0)344; CHECK-NEXT: global_store_dwordx4 v0, a[40:43], s[16:17] offset:32345; CHECK-NEXT: s_waitcnt vmcnt(0)346; CHECK-NEXT: global_store_dwordx4 v0, a[44:47], s[16:17] offset:48347; CHECK-NEXT: s_waitcnt vmcnt(0)348; CHECK-NEXT: global_store_dwordx4 v0, a[32:35], s[16:17]349; CHECK-NEXT: s_waitcnt vmcnt(0)350; CHECK-NEXT: global_store_dwordx4 v0, a[36:39], s[16:17] offset:16351; CHECK-NEXT: s_waitcnt vmcnt(0)352; CHECK-NEXT: global_store_dwordx4 v0, v[8:11], s[16:17]353; CHECK-NEXT: s_waitcnt vmcnt(0)354; CHECK-NEXT: global_store_dwordx4 v0, v[12:15], s[16:17]355; CHECK-NEXT: s_waitcnt vmcnt(0)356; CHECK-NEXT: buffer_load_dword a63, off, s[0:3], s32 ; 4-byte Folded Reload357; CHECK-NEXT: buffer_load_dword a62, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload358; CHECK-NEXT: buffer_load_dword a61, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload359; CHECK-NEXT: buffer_load_dword a60, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload360; CHECK-NEXT: buffer_load_dword a59, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload361; CHECK-NEXT: buffer_load_dword a58, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload362; CHECK-NEXT: buffer_load_dword a57, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload363; CHECK-NEXT: buffer_load_dword a56, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload364; CHECK-NEXT: buffer_load_dword a55, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload365; CHECK-NEXT: buffer_load_dword a54, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload366; CHECK-NEXT: buffer_load_dword a53, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload367; CHECK-NEXT: buffer_load_dword a52, off, s[0:3], s32 offset:44 ; 4-byte Folded Reload368; CHECK-NEXT: buffer_load_dword a51, off, s[0:3], s32 offset:48 ; 4-byte Folded Reload369; CHECK-NEXT: buffer_load_dword a50, off, s[0:3], s32 offset:52 ; 4-byte Folded Reload370; CHECK-NEXT: buffer_load_dword a49, off, s[0:3], s32 offset:56 ; 4-byte Folded Reload371; CHECK-NEXT: buffer_load_dword a48, off, s[0:3], s32 offset:60 ; 4-byte Folded Reload372; CHECK-NEXT: buffer_load_dword a47, off, s[0:3], s32 offset:64 ; 4-byte Folded Reload373; CHECK-NEXT: buffer_load_dword a46, off, s[0:3], s32 offset:68 ; 4-byte Folded Reload374; CHECK-NEXT: buffer_load_dword a45, off, s[0:3], s32 offset:72 ; 4-byte Folded Reload375; CHECK-NEXT: buffer_load_dword a44, off, s[0:3], s32 offset:76 ; 4-byte Folded Reload376; CHECK-NEXT: buffer_load_dword a43, off, s[0:3], s32 offset:80 ; 4-byte Folded Reload377; CHECK-NEXT: buffer_load_dword a42, off, s[0:3], s32 offset:84 ; 4-byte Folded Reload378; CHECK-NEXT: buffer_load_dword a41, off, s[0:3], s32 offset:88 ; 4-byte Folded Reload379; CHECK-NEXT: buffer_load_dword a40, off, s[0:3], s32 offset:92 ; 4-byte Folded Reload380; CHECK-NEXT: buffer_load_dword a39, off, s[0:3], s32 offset:96 ; 4-byte Folded Reload381; CHECK-NEXT: buffer_load_dword a38, off, s[0:3], s32 offset:100 ; 4-byte Folded Reload382; CHECK-NEXT: buffer_load_dword a37, off, s[0:3], s32 offset:104 ; 4-byte Folded Reload383; CHECK-NEXT: buffer_load_dword a36, off, s[0:3], s32 offset:108 ; 4-byte Folded Reload384; CHECK-NEXT: buffer_load_dword a35, off, s[0:3], s32 offset:112 ; 4-byte Folded Reload385; CHECK-NEXT: buffer_load_dword a34, off, s[0:3], s32 offset:116 ; 4-byte Folded Reload386; CHECK-NEXT: buffer_load_dword a33, off, s[0:3], s32 offset:120 ; 4-byte Folded Reload387; CHECK-NEXT: buffer_load_dword a32, off, s[0:3], s32 offset:124 ; 4-byte Folded Reload388; CHECK-NEXT: buffer_load_dword v63, off, s[0:3], s32 offset:128 ; 4-byte Folded Reload389; CHECK-NEXT: buffer_load_dword v62, off, s[0:3], s32 offset:132 ; 4-byte Folded Reload390; CHECK-NEXT: buffer_load_dword v61, off, s[0:3], s32 offset:136 ; 4-byte Folded Reload391; CHECK-NEXT: buffer_load_dword v60, off, s[0:3], s32 offset:140 ; 4-byte Folded Reload392; CHECK-NEXT: buffer_load_dword v59, off, s[0:3], s32 offset:144 ; 4-byte Folded Reload393; CHECK-NEXT: buffer_load_dword v58, off, s[0:3], s32 offset:148 ; 4-byte Folded Reload394; CHECK-NEXT: buffer_load_dword v57, off, s[0:3], s32 offset:152 ; 4-byte Folded Reload395; CHECK-NEXT: buffer_load_dword v56, off, s[0:3], s32 offset:156 ; 4-byte Folded Reload396; CHECK-NEXT: buffer_load_dword v47, off, s[0:3], s32 offset:160 ; 4-byte Folded Reload397; CHECK-NEXT: buffer_load_dword v46, off, s[0:3], s32 offset:164 ; 4-byte Folded Reload398; CHECK-NEXT: buffer_load_dword v45, off, s[0:3], s32 offset:168 ; 4-byte Folded Reload399; CHECK-NEXT: buffer_load_dword v44, off, s[0:3], s32 offset:172 ; 4-byte Folded Reload400; CHECK-NEXT: buffer_load_dword v43, off, s[0:3], s32 offset:176 ; 4-byte Folded Reload401; CHECK-NEXT: buffer_load_dword v42, off, s[0:3], s32 offset:180 ; 4-byte Folded Reload402; CHECK-NEXT: buffer_load_dword v41, off, s[0:3], s32 offset:184 ; 4-byte Folded Reload403; CHECK-NEXT: buffer_load_dword v40, off, s[0:3], s32 offset:188 ; 4-byte Folded Reload404; CHECK-NEXT: s_waitcnt vmcnt(0)405; CHECK-NEXT: s_setpc_b64 s[30:31]406 %mai = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 %x, i32 %y, <4 x i32> %arg, i32 0, i32 0, i32 0)407 %v = call { <32 x i32>, <32 x i32> } asm sideeffect "; def $0, $1", "=v,=v"()408 %v0 = extractvalue { <32 x i32>, <32 x i32> } %v, 0409 %v1 = extractvalue { <32 x i32>, <32 x i32> } %v, 1410 %spill0 = call <4 x i32> asm sideeffect "; def $0", "=v,v"(<4 x i32> %mai)411 %spill1 = call <4 x i32> asm sideeffect "; def $0", "=v,v"(<4 x i32> %mai)412 %a0 = call <32 x i32> asm sideeffect "; def $0", "=a"()413 %a1 = call <32 x i32> asm sideeffect "; def $0", "=a"()414 store volatile <32 x i32> %v0, ptr addrspace(1) %ptr415 store volatile <32 x i32> %v1, ptr addrspace(1) %ptr416 store volatile <4 x i32> %spill0, ptr addrspace(1) %ptr417 store volatile <4 x i32> %spill1, ptr addrspace(1) %ptr418 ret void419}420 421declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32 immarg, i32 immarg, i32 immarg) #1422 423attributes #0 = { nounwind "amdgpu-waves-per-eu"="10,10" }424attributes #1 = { convergent nocallback nofree nosync nounwind willreturn memory(none) }425