1025 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY908 %s3; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY90A %s4; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -early-live-intervals < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY90A %s5; RUN: llc -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY942 %s6; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY90A-GISEL %s7; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -sgpr-regalloc=fast -vgpr-regalloc=fast < %s | FileCheck -enable-var-scope --check-prefixes=FAST90A %s8 9; This is better with 90a10 11; Check that Dst and SrcC of MFMA instructions reading more than 4 registers as SrcC12; is either completely disjoint or exactly the same, but does not alias.13 14declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)15declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32, i32, i32)16declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)17 18define amdgpu_kernel void @test_mfma_f32_32x32x1f32(ptr addrspace(1) %arg) #0 {19; GREEDY908-LABEL: test_mfma_f32_32x32x1f32:20; GREEDY908: ; %bb.0: ; %bb21; GREEDY908-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x2422; GREEDY908-NEXT: s_waitcnt lgkmcnt(0)23; GREEDY908-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x024; GREEDY908-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x4025; GREEDY908-NEXT: s_waitcnt lgkmcnt(0)26; GREEDY908-NEXT: v_mov_b32_e32 v0, s1627; GREEDY908-NEXT: v_mov_b32_e32 v1, s1728; GREEDY908-NEXT: v_mov_b32_e32 v2, s1829; GREEDY908-NEXT: v_accvgpr_write_b32 a0, v030; GREEDY908-NEXT: v_accvgpr_write_b32 a1, v131; GREEDY908-NEXT: v_mov_b32_e32 v0, s2232; GREEDY908-NEXT: v_mov_b32_e32 v1, s2333; GREEDY908-NEXT: v_accvgpr_write_b32 a2, v234; GREEDY908-NEXT: v_accvgpr_write_b32 a6, v035; GREEDY908-NEXT: v_accvgpr_write_b32 a7, v136; GREEDY908-NEXT: v_mov_b32_e32 v0, s2437; GREEDY908-NEXT: v_mov_b32_e32 v1, s2538; GREEDY908-NEXT: v_mov_b32_e32 v2, s2639; GREEDY908-NEXT: v_accvgpr_write_b32 a8, v040; GREEDY908-NEXT: v_accvgpr_write_b32 a9, v141; GREEDY908-NEXT: v_accvgpr_write_b32 a10, v242; GREEDY908-NEXT: v_mov_b32_e32 v0, s2743; GREEDY908-NEXT: v_mov_b32_e32 v1, s2844; GREEDY908-NEXT: v_mov_b32_e32 v2, s2945; GREEDY908-NEXT: v_accvgpr_write_b32 a11, v046; GREEDY908-NEXT: v_accvgpr_write_b32 a12, v147; GREEDY908-NEXT: v_accvgpr_write_b32 a13, v248; GREEDY908-NEXT: v_mov_b32_e32 v0, s3049; GREEDY908-NEXT: v_mov_b32_e32 v1, s3150; GREEDY908-NEXT: v_mov_b32_e32 v2, s051; GREEDY908-NEXT: v_accvgpr_write_b32 a14, v052; GREEDY908-NEXT: v_accvgpr_write_b32 a15, v153; GREEDY908-NEXT: v_accvgpr_write_b32 a16, v254; GREEDY908-NEXT: v_mov_b32_e32 v0, s155; GREEDY908-NEXT: v_mov_b32_e32 v1, s256; GREEDY908-NEXT: v_mov_b32_e32 v2, s357; GREEDY908-NEXT: v_accvgpr_write_b32 a17, v058; GREEDY908-NEXT: v_accvgpr_write_b32 a18, v159; GREEDY908-NEXT: v_accvgpr_write_b32 a19, v260; GREEDY908-NEXT: v_mov_b32_e32 v0, s461; GREEDY908-NEXT: v_mov_b32_e32 v1, s562; GREEDY908-NEXT: v_mov_b32_e32 v2, s663; GREEDY908-NEXT: v_accvgpr_write_b32 a20, v064; GREEDY908-NEXT: v_accvgpr_write_b32 a21, v165; GREEDY908-NEXT: v_accvgpr_write_b32 a22, v266; GREEDY908-NEXT: v_mov_b32_e32 v0, s767; GREEDY908-NEXT: v_mov_b32_e32 v1, s868; GREEDY908-NEXT: v_mov_b32_e32 v2, s969; GREEDY908-NEXT: v_mov_b32_e32 v3, s1970; GREEDY908-NEXT: v_accvgpr_write_b32 a23, v071; GREEDY908-NEXT: v_accvgpr_write_b32 a24, v172; GREEDY908-NEXT: v_accvgpr_write_b32 a25, v273; GREEDY908-NEXT: v_mov_b32_e32 v0, s1074; GREEDY908-NEXT: v_mov_b32_e32 v1, s1175; GREEDY908-NEXT: v_mov_b32_e32 v2, s1276; GREEDY908-NEXT: v_mov_b32_e32 v4, s2077; GREEDY908-NEXT: v_mov_b32_e32 v5, s2178; GREEDY908-NEXT: v_accvgpr_write_b32 a3, v379; GREEDY908-NEXT: v_accvgpr_write_b32 a26, v080; GREEDY908-NEXT: v_accvgpr_write_b32 a27, v181; GREEDY908-NEXT: v_accvgpr_write_b32 a28, v282; GREEDY908-NEXT: v_mov_b32_e32 v0, s1383; GREEDY908-NEXT: v_mov_b32_e32 v1, s1484; GREEDY908-NEXT: v_mov_b32_e32 v2, s1585; GREEDY908-NEXT: v_mov_b32_e32 v3, 1.086; GREEDY908-NEXT: v_accvgpr_write_b32 a4, v487; GREEDY908-NEXT: v_accvgpr_write_b32 a5, v588; GREEDY908-NEXT: v_accvgpr_write_b32 a29, v089; GREEDY908-NEXT: v_accvgpr_write_b32 a30, v190; GREEDY908-NEXT: v_accvgpr_write_b32 a31, v291; GREEDY908-NEXT: v_mov_b32_e32 v0, 2.092; GREEDY908-NEXT: v_mov_b32_e32 v4, 093; GREEDY908-NEXT: s_nop 094; GREEDY908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v3, v0, a[0:31]95; GREEDY908-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v3, v0, a[0:31]96; GREEDY908-NEXT: s_nop 1597; GREEDY908-NEXT: s_nop 198; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a3299; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a33100; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a34101; GREEDY908-NEXT: v_accvgpr_write_b32 a2, v2102; GREEDY908-NEXT: v_accvgpr_write_b32 a3, v6103; GREEDY908-NEXT: v_accvgpr_write_b32 a4, v1104; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a35105; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a36106; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a37107; GREEDY908-NEXT: v_accvgpr_write_b32 a5, v2108; GREEDY908-NEXT: v_accvgpr_write_b32 a6, v6109; GREEDY908-NEXT: v_accvgpr_write_b32 a7, v1110; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a38111; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a39112; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a40113; GREEDY908-NEXT: v_accvgpr_write_b32 a8, v2114; GREEDY908-NEXT: v_accvgpr_write_b32 a9, v6115; GREEDY908-NEXT: v_accvgpr_write_b32 a10, v1116; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a41117; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a42118; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a43119; GREEDY908-NEXT: v_accvgpr_write_b32 a11, v2120; GREEDY908-NEXT: v_accvgpr_write_b32 a12, v6121; GREEDY908-NEXT: v_accvgpr_write_b32 a13, v1122; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a44123; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a45124; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a46125; GREEDY908-NEXT: v_accvgpr_write_b32 a14, v2126; GREEDY908-NEXT: v_accvgpr_write_b32 a15, v6127; GREEDY908-NEXT: v_accvgpr_write_b32 a16, v1128; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a47129; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a48130; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a49131; GREEDY908-NEXT: v_accvgpr_write_b32 a17, v2132; GREEDY908-NEXT: v_accvgpr_write_b32 a18, v6133; GREEDY908-NEXT: v_accvgpr_write_b32 a19, v1134; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a50135; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a51136; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a52137; GREEDY908-NEXT: v_accvgpr_write_b32 a20, v2138; GREEDY908-NEXT: v_accvgpr_write_b32 a21, v6139; GREEDY908-NEXT: v_accvgpr_write_b32 a22, v1140; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a53141; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a54142; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a55143; GREEDY908-NEXT: v_accvgpr_write_b32 a23, v2144; GREEDY908-NEXT: v_accvgpr_write_b32 a24, v6145; GREEDY908-NEXT: v_accvgpr_write_b32 a25, v1146; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a56147; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a57148; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a58149; GREEDY908-NEXT: v_accvgpr_write_b32 a26, v2150; GREEDY908-NEXT: v_accvgpr_write_b32 a27, v6151; GREEDY908-NEXT: v_accvgpr_write_b32 a28, v1152; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a59153; GREEDY908-NEXT: v_accvgpr_read_b32 v6, a60154; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a61155; GREEDY908-NEXT: v_accvgpr_write_b32 a29, v2156; GREEDY908-NEXT: v_accvgpr_write_b32 a30, v6157; GREEDY908-NEXT: v_accvgpr_write_b32 a31, v1158; GREEDY908-NEXT: s_nop 0159; GREEDY908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v3, v0, a[0:31]160; GREEDY908-NEXT: s_nop 15161; GREEDY908-NEXT: s_nop 1162; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a27163; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a26164; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a25165; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a24166; GREEDY908-NEXT: s_nop 1167; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:96168; GREEDY908-NEXT: s_nop 0169; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a31170; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a30171; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a29172; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a28173; GREEDY908-NEXT: s_nop 1174; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:112175; GREEDY908-NEXT: s_nop 0176; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a19177; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a18178; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a17179; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a16180; GREEDY908-NEXT: s_nop 1181; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:64182; GREEDY908-NEXT: s_nop 0183; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a23184; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a22185; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a21186; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a20187; GREEDY908-NEXT: s_nop 1188; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:80189; GREEDY908-NEXT: s_nop 0190; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a11191; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a10192; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a9193; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a8194; GREEDY908-NEXT: s_nop 1195; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:32196; GREEDY908-NEXT: s_nop 0197; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a15198; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a14199; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a13200; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a12201; GREEDY908-NEXT: s_nop 1202; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:48203; GREEDY908-NEXT: s_nop 0204; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a3205; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a2206; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a1207; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a0208; GREEDY908-NEXT: s_nop 1209; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35]210; GREEDY908-NEXT: s_nop 0211; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a7212; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a6213; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a5214; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a4215; GREEDY908-NEXT: s_nop 1216; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:16217; GREEDY908-NEXT: s_endpgm218;219; GREEDY90A-LABEL: test_mfma_f32_32x32x1f32:220; GREEDY90A: ; %bb.0: ; %bb221; GREEDY90A-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24222; GREEDY90A-NEXT: v_mov_b32_e32 v0, 1.0223; GREEDY90A-NEXT: v_mov_b32_e32 v1, 2.0224; GREEDY90A-NEXT: v_mov_b32_e32 v2, 0225; GREEDY90A-NEXT: s_waitcnt lgkmcnt(0)226; GREEDY90A-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0227; GREEDY90A-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40228; GREEDY90A-NEXT: s_waitcnt lgkmcnt(0)229; GREEDY90A-NEXT: v_accvgpr_write_b32 a0, s16230; GREEDY90A-NEXT: v_accvgpr_write_b32 a1, s17231; GREEDY90A-NEXT: v_accvgpr_write_b32 a2, s18232; GREEDY90A-NEXT: v_accvgpr_write_b32 a3, s19233; GREEDY90A-NEXT: v_accvgpr_write_b32 a4, s20234; GREEDY90A-NEXT: v_accvgpr_write_b32 a5, s21235; GREEDY90A-NEXT: v_accvgpr_write_b32 a6, s22236; GREEDY90A-NEXT: v_accvgpr_write_b32 a7, s23237; GREEDY90A-NEXT: v_accvgpr_write_b32 a8, s24238; GREEDY90A-NEXT: v_accvgpr_write_b32 a9, s25239; GREEDY90A-NEXT: v_accvgpr_write_b32 a10, s26240; GREEDY90A-NEXT: v_accvgpr_write_b32 a11, s27241; GREEDY90A-NEXT: v_accvgpr_write_b32 a12, s28242; GREEDY90A-NEXT: v_accvgpr_write_b32 a13, s29243; GREEDY90A-NEXT: v_accvgpr_write_b32 a14, s30244; GREEDY90A-NEXT: v_accvgpr_write_b32 a15, s31245; GREEDY90A-NEXT: v_accvgpr_write_b32 a16, s0246; GREEDY90A-NEXT: v_accvgpr_write_b32 a17, s1247; GREEDY90A-NEXT: v_accvgpr_write_b32 a18, s2248; GREEDY90A-NEXT: v_accvgpr_write_b32 a19, s3249; GREEDY90A-NEXT: v_accvgpr_write_b32 a20, s4250; GREEDY90A-NEXT: v_accvgpr_write_b32 a21, s5251; GREEDY90A-NEXT: v_accvgpr_write_b32 a22, s6252; GREEDY90A-NEXT: v_accvgpr_write_b32 a23, s7253; GREEDY90A-NEXT: v_accvgpr_write_b32 a24, s8254; GREEDY90A-NEXT: v_accvgpr_write_b32 a25, s9255; GREEDY90A-NEXT: v_accvgpr_write_b32 a26, s10256; GREEDY90A-NEXT: v_accvgpr_write_b32 a27, s11257; GREEDY90A-NEXT: v_accvgpr_write_b32 a28, s12258; GREEDY90A-NEXT: v_accvgpr_write_b32 a29, s13259; GREEDY90A-NEXT: v_accvgpr_write_b32 a30, s14260; GREEDY90A-NEXT: v_accvgpr_write_b32 a31, s15261; GREEDY90A-NEXT: s_nop 1262; GREEDY90A-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]263; GREEDY90A-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v0, v1, a[0:31]264; GREEDY90A-NEXT: s_nop 15265; GREEDY90A-NEXT: s_nop 2266; GREEDY90A-NEXT: v_accvgpr_mov_b32 a2, a32267; GREEDY90A-NEXT: v_accvgpr_mov_b32 a3, a33268; GREEDY90A-NEXT: v_accvgpr_mov_b32 a4, a34269; GREEDY90A-NEXT: v_accvgpr_mov_b32 a5, a35270; GREEDY90A-NEXT: v_accvgpr_mov_b32 a6, a36271; GREEDY90A-NEXT: v_accvgpr_mov_b32 a7, a37272; GREEDY90A-NEXT: v_accvgpr_mov_b32 a8, a38273; GREEDY90A-NEXT: v_accvgpr_mov_b32 a9, a39274; GREEDY90A-NEXT: v_accvgpr_mov_b32 a10, a40275; GREEDY90A-NEXT: v_accvgpr_mov_b32 a11, a41276; GREEDY90A-NEXT: v_accvgpr_mov_b32 a12, a42277; GREEDY90A-NEXT: v_accvgpr_mov_b32 a13, a43278; GREEDY90A-NEXT: v_accvgpr_mov_b32 a14, a44279; GREEDY90A-NEXT: v_accvgpr_mov_b32 a15, a45280; GREEDY90A-NEXT: v_accvgpr_mov_b32 a16, a46281; GREEDY90A-NEXT: v_accvgpr_mov_b32 a17, a47282; GREEDY90A-NEXT: v_accvgpr_mov_b32 a18, a48283; GREEDY90A-NEXT: v_accvgpr_mov_b32 a19, a49284; GREEDY90A-NEXT: v_accvgpr_mov_b32 a20, a50285; GREEDY90A-NEXT: v_accvgpr_mov_b32 a21, a51286; GREEDY90A-NEXT: v_accvgpr_mov_b32 a22, a52287; GREEDY90A-NEXT: v_accvgpr_mov_b32 a23, a53288; GREEDY90A-NEXT: v_accvgpr_mov_b32 a24, a54289; GREEDY90A-NEXT: v_accvgpr_mov_b32 a25, a55290; GREEDY90A-NEXT: v_accvgpr_mov_b32 a26, a56291; GREEDY90A-NEXT: v_accvgpr_mov_b32 a27, a57292; GREEDY90A-NEXT: v_accvgpr_mov_b32 a28, a58293; GREEDY90A-NEXT: v_accvgpr_mov_b32 a29, a59294; GREEDY90A-NEXT: v_accvgpr_mov_b32 a30, a60295; GREEDY90A-NEXT: v_accvgpr_mov_b32 a31, a61296; GREEDY90A-NEXT: s_nop 1297; GREEDY90A-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]298; GREEDY90A-NEXT: s_nop 15299; GREEDY90A-NEXT: s_nop 2300; GREEDY90A-NEXT: global_store_dwordx4 v2, a[24:27], s[34:35] offset:96301; GREEDY90A-NEXT: global_store_dwordx4 v2, a[28:31], s[34:35] offset:112302; GREEDY90A-NEXT: global_store_dwordx4 v2, a[16:19], s[34:35] offset:64303; GREEDY90A-NEXT: global_store_dwordx4 v2, a[20:23], s[34:35] offset:80304; GREEDY90A-NEXT: global_store_dwordx4 v2, a[8:11], s[34:35] offset:32305; GREEDY90A-NEXT: global_store_dwordx4 v2, a[12:15], s[34:35] offset:48306; GREEDY90A-NEXT: global_store_dwordx4 v2, a[0:3], s[34:35]307; GREEDY90A-NEXT: global_store_dwordx4 v2, a[4:7], s[34:35] offset:16308; GREEDY90A-NEXT: s_endpgm309;310; GREEDY942-LABEL: test_mfma_f32_32x32x1f32:311; GREEDY942: ; %bb.0: ; %bb312; GREEDY942-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24313; GREEDY942-NEXT: v_mov_b32_e32 v0, 1.0314; GREEDY942-NEXT: v_mov_b32_e32 v1, 2.0315; GREEDY942-NEXT: v_mov_b32_e32 v2, 0316; GREEDY942-NEXT: s_waitcnt lgkmcnt(0)317; GREEDY942-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0318; GREEDY942-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40319; GREEDY942-NEXT: s_waitcnt lgkmcnt(0)320; GREEDY942-NEXT: v_accvgpr_write_b32 a0, s16321; GREEDY942-NEXT: v_accvgpr_write_b32 a1, s17322; GREEDY942-NEXT: v_accvgpr_write_b32 a2, s18323; GREEDY942-NEXT: v_accvgpr_write_b32 a3, s19324; GREEDY942-NEXT: v_accvgpr_write_b32 a4, s20325; GREEDY942-NEXT: v_accvgpr_write_b32 a5, s21326; GREEDY942-NEXT: v_accvgpr_write_b32 a6, s22327; GREEDY942-NEXT: v_accvgpr_write_b32 a7, s23328; GREEDY942-NEXT: v_accvgpr_write_b32 a8, s24329; GREEDY942-NEXT: v_accvgpr_write_b32 a9, s25330; GREEDY942-NEXT: v_accvgpr_write_b32 a10, s26331; GREEDY942-NEXT: v_accvgpr_write_b32 a11, s27332; GREEDY942-NEXT: v_accvgpr_write_b32 a12, s28333; GREEDY942-NEXT: v_accvgpr_write_b32 a13, s29334; GREEDY942-NEXT: v_accvgpr_write_b32 a14, s30335; GREEDY942-NEXT: v_accvgpr_write_b32 a15, s31336; GREEDY942-NEXT: v_accvgpr_write_b32 a16, s0337; GREEDY942-NEXT: v_accvgpr_write_b32 a17, s1338; GREEDY942-NEXT: v_accvgpr_write_b32 a18, s2339; GREEDY942-NEXT: v_accvgpr_write_b32 a19, s3340; GREEDY942-NEXT: v_accvgpr_write_b32 a20, s4341; GREEDY942-NEXT: v_accvgpr_write_b32 a21, s5342; GREEDY942-NEXT: v_accvgpr_write_b32 a22, s6343; GREEDY942-NEXT: v_accvgpr_write_b32 a23, s7344; GREEDY942-NEXT: v_accvgpr_write_b32 a24, s8345; GREEDY942-NEXT: v_accvgpr_write_b32 a25, s9346; GREEDY942-NEXT: v_accvgpr_write_b32 a26, s10347; GREEDY942-NEXT: v_accvgpr_write_b32 a27, s11348; GREEDY942-NEXT: v_accvgpr_write_b32 a28, s12349; GREEDY942-NEXT: v_accvgpr_write_b32 a29, s13350; GREEDY942-NEXT: v_accvgpr_write_b32 a30, s14351; GREEDY942-NEXT: v_accvgpr_write_b32 a31, s15352; GREEDY942-NEXT: s_nop 1353; GREEDY942-NEXT: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[0:31]354; GREEDY942-NEXT: v_mfma_f32_32x32x1_2b_f32 a[32:63], v0, v1, a[0:31]355; GREEDY942-NEXT: s_nop 15356; GREEDY942-NEXT: s_nop 1357; GREEDY942-NEXT: v_accvgpr_mov_b32 a2, a32358; GREEDY942-NEXT: v_accvgpr_mov_b32 a3, a33359; GREEDY942-NEXT: v_accvgpr_mov_b32 a4, a34360; GREEDY942-NEXT: v_accvgpr_mov_b32 a5, a35361; GREEDY942-NEXT: v_accvgpr_mov_b32 a6, a36362; GREEDY942-NEXT: v_accvgpr_mov_b32 a7, a37363; GREEDY942-NEXT: v_accvgpr_mov_b32 a8, a38364; GREEDY942-NEXT: v_accvgpr_mov_b32 a9, a39365; GREEDY942-NEXT: v_accvgpr_mov_b32 a10, a40366; GREEDY942-NEXT: v_accvgpr_mov_b32 a11, a41367; GREEDY942-NEXT: v_accvgpr_mov_b32 a12, a42368; GREEDY942-NEXT: v_accvgpr_mov_b32 a13, a43369; GREEDY942-NEXT: v_accvgpr_mov_b32 a14, a44370; GREEDY942-NEXT: v_accvgpr_mov_b32 a15, a45371; GREEDY942-NEXT: v_accvgpr_mov_b32 a16, a46372; GREEDY942-NEXT: v_accvgpr_mov_b32 a17, a47373; GREEDY942-NEXT: v_accvgpr_mov_b32 a18, a48374; GREEDY942-NEXT: v_accvgpr_mov_b32 a19, a49375; GREEDY942-NEXT: v_accvgpr_mov_b32 a20, a50376; GREEDY942-NEXT: v_accvgpr_mov_b32 a21, a51377; GREEDY942-NEXT: v_accvgpr_mov_b32 a22, a52378; GREEDY942-NEXT: v_accvgpr_mov_b32 a23, a53379; GREEDY942-NEXT: v_accvgpr_mov_b32 a24, a54380; GREEDY942-NEXT: v_accvgpr_mov_b32 a25, a55381; GREEDY942-NEXT: v_accvgpr_mov_b32 a26, a56382; GREEDY942-NEXT: v_accvgpr_mov_b32 a27, a57383; GREEDY942-NEXT: v_accvgpr_mov_b32 a28, a58384; GREEDY942-NEXT: v_accvgpr_mov_b32 a29, a59385; GREEDY942-NEXT: v_accvgpr_mov_b32 a30, a60386; GREEDY942-NEXT: v_accvgpr_mov_b32 a31, a61387; GREEDY942-NEXT: s_nop 1388; GREEDY942-NEXT: v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[0:31]389; GREEDY942-NEXT: s_nop 15390; GREEDY942-NEXT: s_nop 1391; GREEDY942-NEXT: global_store_dwordx4 v2, a[24:27], s[34:35] offset:96392; GREEDY942-NEXT: global_store_dwordx4 v2, a[28:31], s[34:35] offset:112393; GREEDY942-NEXT: global_store_dwordx4 v2, a[16:19], s[34:35] offset:64394; GREEDY942-NEXT: global_store_dwordx4 v2, a[20:23], s[34:35] offset:80395; GREEDY942-NEXT: global_store_dwordx4 v2, a[8:11], s[34:35] offset:32396; GREEDY942-NEXT: global_store_dwordx4 v2, a[12:15], s[34:35] offset:48397; GREEDY942-NEXT: global_store_dwordx4 v2, a[0:3], s[34:35]398; GREEDY942-NEXT: global_store_dwordx4 v2, a[4:7], s[34:35] offset:16399; GREEDY942-NEXT: s_endpgm400;401; GREEDY90A-GISEL-LABEL: test_mfma_f32_32x32x1f32:402; GREEDY90A-GISEL: ; %bb.0: ; %bb403; GREEDY90A-GISEL-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24404; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v0, 1.0405; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v1, 2.0406; GREEDY90A-GISEL-NEXT: s_waitcnt lgkmcnt(0)407; GREEDY90A-GISEL-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x0408; GREEDY90A-GISEL-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x40409; GREEDY90A-GISEL-NEXT: s_waitcnt lgkmcnt(0)410; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a0, s0411; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a16, s16412; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a1, s1413; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a2, s2414; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a3, s3415; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a4, s4416; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a5, s5417; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a6, s6418; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a7, s7419; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a8, s8420; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a9, s9421; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a10, s10422; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a11, s11423; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a12, s12424; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a13, s13425; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a14, s14426; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a15, s15427; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a17, s17428; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a18, s18429; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a19, s19430; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a20, s20431; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a21, s21432; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a22, s22433; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a23, s23434; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a24, s24435; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a25, s25436; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a26, s26437; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a27, s27438; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a28, s28439; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a29, s29440; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a30, s30441; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a31, s31442; GREEDY90A-GISEL-NEXT: s_nop 1443; GREEDY90A-GISEL-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]444; GREEDY90A-GISEL-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v0, v1, a[0:31]445; GREEDY90A-GISEL-NEXT: s_nop 15446; GREEDY90A-GISEL-NEXT: s_nop 2447; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a2, a32448; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a3, a33449; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a4, a34450; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a5, a35451; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a6, a36452; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a7, a37453; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a8, a38454; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a9, a39455; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a10, a40456; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a11, a41457; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a12, a42458; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a13, a43459; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a14, a44460; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a15, a45461; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a16, a46462; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a17, a47463; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a18, a48464; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a19, a49465; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a20, a50466; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a21, a51467; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a22, a52468; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a23, a53469; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a24, a54470; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a25, a55471; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a26, a56472; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a27, a57473; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a28, a58474; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a29, a59475; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a30, a60476; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a31, a61477; GREEDY90A-GISEL-NEXT: s_nop 1478; GREEDY90A-GISEL-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]479; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v0, 0480; GREEDY90A-GISEL-NEXT: s_nop 15481; GREEDY90A-GISEL-NEXT: s_nop 1482; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[0:3], s[34:35]483; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[4:7], s[34:35] offset:16484; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[8:11], s[34:35] offset:32485; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[12:15], s[34:35] offset:48486; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[16:19], s[34:35] offset:64487; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[20:23], s[34:35] offset:80488; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[24:27], s[34:35] offset:96489; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[28:31], s[34:35] offset:112490; GREEDY90A-GISEL-NEXT: s_endpgm491;492; FAST90A-LABEL: test_mfma_f32_32x32x1f32:493; FAST90A: ; %bb.0: ; %bb494; FAST90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24495; FAST90A-NEXT: v_mov_b32_e32 v1, 1.0496; FAST90A-NEXT: v_mov_b32_e32 v2, 2.0497; FAST90A-NEXT: v_mov_b32_e32 v0, 0498; FAST90A-NEXT: s_waitcnt lgkmcnt(0)499; FAST90A-NEXT: s_load_dwordx16 s[36:51], s[0:1], 0x0500; FAST90A-NEXT: s_load_dwordx16 s[4:19], s[0:1], 0x40501; FAST90A-NEXT: s_waitcnt lgkmcnt(0)502; FAST90A-NEXT: v_accvgpr_write_b32 a0, s36503; FAST90A-NEXT: v_accvgpr_write_b32 a1, s37504; FAST90A-NEXT: v_accvgpr_write_b32 a2, s38505; FAST90A-NEXT: v_accvgpr_write_b32 a3, s39506; FAST90A-NEXT: v_accvgpr_write_b32 a4, s40507; FAST90A-NEXT: v_accvgpr_write_b32 a5, s41508; FAST90A-NEXT: v_accvgpr_write_b32 a6, s42509; FAST90A-NEXT: v_accvgpr_write_b32 a7, s43510; FAST90A-NEXT: v_accvgpr_write_b32 a8, s44511; FAST90A-NEXT: v_accvgpr_write_b32 a9, s45512; FAST90A-NEXT: v_accvgpr_write_b32 a10, s46513; FAST90A-NEXT: v_accvgpr_write_b32 a11, s47514; FAST90A-NEXT: v_accvgpr_write_b32 a12, s48515; FAST90A-NEXT: v_accvgpr_write_b32 a13, s49516; FAST90A-NEXT: v_accvgpr_write_b32 a14, s50517; FAST90A-NEXT: v_accvgpr_write_b32 a15, s51518; FAST90A-NEXT: v_accvgpr_write_b32 a16, s4519; FAST90A-NEXT: v_accvgpr_write_b32 a17, s5520; FAST90A-NEXT: v_accvgpr_write_b32 a18, s6521; FAST90A-NEXT: v_accvgpr_write_b32 a19, s7522; FAST90A-NEXT: v_accvgpr_write_b32 a20, s8523; FAST90A-NEXT: v_accvgpr_write_b32 a21, s9524; FAST90A-NEXT: v_accvgpr_write_b32 a22, s10525; FAST90A-NEXT: v_accvgpr_write_b32 a23, s11526; FAST90A-NEXT: v_accvgpr_write_b32 a24, s12527; FAST90A-NEXT: v_accvgpr_write_b32 a25, s13528; FAST90A-NEXT: v_accvgpr_write_b32 a26, s14529; FAST90A-NEXT: v_accvgpr_write_b32 a27, s15530; FAST90A-NEXT: v_accvgpr_write_b32 a28, s16531; FAST90A-NEXT: v_accvgpr_write_b32 a29, s17532; FAST90A-NEXT: v_accvgpr_write_b32 a30, s18533; FAST90A-NEXT: v_accvgpr_write_b32 a31, s19534; FAST90A-NEXT: s_nop 1535; FAST90A-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]536; FAST90A-NEXT: v_mfma_f32_32x32x1f32 a[32:63], v1, v2, a[0:31]537; FAST90A-NEXT: s_nop 15538; FAST90A-NEXT: s_nop 2539; FAST90A-NEXT: v_accvgpr_mov_b32 a2, a32540; FAST90A-NEXT: v_accvgpr_mov_b32 a3, a33541; FAST90A-NEXT: v_accvgpr_mov_b32 a4, a34542; FAST90A-NEXT: v_accvgpr_mov_b32 a5, a35543; FAST90A-NEXT: v_accvgpr_mov_b32 a6, a36544; FAST90A-NEXT: v_accvgpr_mov_b32 a7, a37545; FAST90A-NEXT: v_accvgpr_mov_b32 a8, a38546; FAST90A-NEXT: v_accvgpr_mov_b32 a9, a39547; FAST90A-NEXT: v_accvgpr_mov_b32 a10, a40548; FAST90A-NEXT: v_accvgpr_mov_b32 a11, a41549; FAST90A-NEXT: v_accvgpr_mov_b32 a12, a42550; FAST90A-NEXT: v_accvgpr_mov_b32 a13, a43551; FAST90A-NEXT: v_accvgpr_mov_b32 a14, a44552; FAST90A-NEXT: v_accvgpr_mov_b32 a15, a45553; FAST90A-NEXT: v_accvgpr_mov_b32 a16, a46554; FAST90A-NEXT: v_accvgpr_mov_b32 a17, a47555; FAST90A-NEXT: v_accvgpr_mov_b32 a18, a48556; FAST90A-NEXT: v_accvgpr_mov_b32 a19, a49557; FAST90A-NEXT: v_accvgpr_mov_b32 a20, a50558; FAST90A-NEXT: v_accvgpr_mov_b32 a21, a51559; FAST90A-NEXT: v_accvgpr_mov_b32 a22, a52560; FAST90A-NEXT: v_accvgpr_mov_b32 a23, a53561; FAST90A-NEXT: v_accvgpr_mov_b32 a24, a54562; FAST90A-NEXT: v_accvgpr_mov_b32 a25, a55563; FAST90A-NEXT: v_accvgpr_mov_b32 a26, a56564; FAST90A-NEXT: v_accvgpr_mov_b32 a27, a57565; FAST90A-NEXT: v_accvgpr_mov_b32 a28, a58566; FAST90A-NEXT: v_accvgpr_mov_b32 a29, a59567; FAST90A-NEXT: v_accvgpr_mov_b32 a30, a60568; FAST90A-NEXT: v_accvgpr_mov_b32 a31, a61569; FAST90A-NEXT: s_nop 1570; FAST90A-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v2, a[0:31]571; FAST90A-NEXT: s_nop 15572; FAST90A-NEXT: s_nop 2573; FAST90A-NEXT: global_store_dwordx4 v0, a[24:27], s[0:1] offset:96574; FAST90A-NEXT: global_store_dwordx4 v0, a[28:31], s[0:1] offset:112575; FAST90A-NEXT: global_store_dwordx4 v0, a[16:19], s[0:1] offset:64576; FAST90A-NEXT: global_store_dwordx4 v0, a[20:23], s[0:1] offset:80577; FAST90A-NEXT: global_store_dwordx4 v0, a[8:11], s[0:1] offset:32578; FAST90A-NEXT: global_store_dwordx4 v0, a[12:15], s[0:1] offset:48579; FAST90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]580; FAST90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16581; FAST90A-NEXT: s_endpgm582bb:583 %in.1 = load <32 x float>, ptr addrspace(1) %arg584 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 0, i32 0, i32 0)585 %mai.2 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %mai.1, i32 0, i32 0, i32 0)586 %tmp.1 = shufflevector <32 x float> %mai.2, <32 x float> %mai.1, <32 x i32> <i32 32, i32 33, i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15, i32 16, i32 17, i32 18, i32 19, i32 20, i32 21, i32 22, i32 23, i32 24, i32 25, i32 26, i32 27, i32 28, i32 29>587 %mai.3 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %tmp.1, i32 0, i32 0, i32 0)588 store <32 x float> %mai.3, ptr addrspace(1) %arg589 ret void590}591 592define amdgpu_kernel void @test_mfma_f32_16x16x1f32(ptr addrspace(1) %arg) #0 {593; GREEDY908-LABEL: test_mfma_f32_16x16x1f32:594; GREEDY908: ; %bb.0: ; %bb595; GREEDY908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24596; GREEDY908-NEXT: v_mov_b32_e32 v0, 1.0597; GREEDY908-NEXT: v_mov_b32_e32 v4, 0598; GREEDY908-NEXT: s_waitcnt lgkmcnt(0)599; GREEDY908-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0600; GREEDY908-NEXT: s_waitcnt lgkmcnt(0)601; GREEDY908-NEXT: v_mov_b32_e32 v5, s15602; GREEDY908-NEXT: v_mov_b32_e32 v2, s14603; GREEDY908-NEXT: v_mov_b32_e32 v1, s13604; GREEDY908-NEXT: v_accvgpr_write_b32 a33, v5605; GREEDY908-NEXT: v_mov_b32_e32 v5, s12606; GREEDY908-NEXT: v_accvgpr_write_b32 a32, v2607; GREEDY908-NEXT: v_accvgpr_write_b32 a31, v1608; GREEDY908-NEXT: v_accvgpr_write_b32 a30, v5609; GREEDY908-NEXT: v_mov_b32_e32 v2, s11610; GREEDY908-NEXT: v_mov_b32_e32 v1, s10611; GREEDY908-NEXT: v_mov_b32_e32 v5, s9612; GREEDY908-NEXT: v_accvgpr_write_b32 a29, v2613; GREEDY908-NEXT: v_accvgpr_write_b32 a28, v1614; GREEDY908-NEXT: v_accvgpr_write_b32 a27, v5615; GREEDY908-NEXT: v_mov_b32_e32 v2, s8616; GREEDY908-NEXT: v_mov_b32_e32 v1, s7617; GREEDY908-NEXT: v_mov_b32_e32 v5, s6618; GREEDY908-NEXT: v_accvgpr_write_b32 a26, v2619; GREEDY908-NEXT: v_accvgpr_write_b32 a25, v1620; GREEDY908-NEXT: v_accvgpr_write_b32 a24, v5621; GREEDY908-NEXT: v_mov_b32_e32 v2, s5622; GREEDY908-NEXT: v_mov_b32_e32 v1, s4623; GREEDY908-NEXT: v_mov_b32_e32 v5, s3624; GREEDY908-NEXT: v_accvgpr_write_b32 a23, v2625; GREEDY908-NEXT: v_accvgpr_write_b32 a22, v1626; GREEDY908-NEXT: v_accvgpr_write_b32 a21, v5627; GREEDY908-NEXT: v_mov_b32_e32 v2, s2628; GREEDY908-NEXT: v_mov_b32_e32 v1, s1629; GREEDY908-NEXT: v_mov_b32_e32 v5, s0630; GREEDY908-NEXT: v_accvgpr_write_b32 a20, v2631; GREEDY908-NEXT: v_accvgpr_write_b32 a19, v1632; GREEDY908-NEXT: v_accvgpr_write_b32 a18, v5633; GREEDY908-NEXT: v_mov_b32_e32 v1, 2.0634; GREEDY908-NEXT: s_nop 1635; GREEDY908-NEXT: v_mfma_f32_16x16x1f32 a[18:33], v0, v1, a[18:33]636; GREEDY908-NEXT: v_mfma_f32_16x16x1f32 a[2:17], v0, v1, a[18:33]637; GREEDY908-NEXT: s_nop 8638; GREEDY908-NEXT: v_accvgpr_read_b32 v5, a18639; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a19640; GREEDY908-NEXT: s_nop 0641; GREEDY908-NEXT: v_accvgpr_write_b32 a0, v5642; GREEDY908-NEXT: v_accvgpr_write_b32 a1, v2643; GREEDY908-NEXT: s_nop 0644; GREEDY908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[0:15]645; GREEDY908-NEXT: s_nop 9646; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a15647; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a14648; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a13649; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a12650; GREEDY908-NEXT: s_nop 1651; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[16:17] offset:48652; GREEDY908-NEXT: s_nop 0653; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a11654; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a10655; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a9656; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a8657; GREEDY908-NEXT: s_nop 1658; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[16:17] offset:32659; GREEDY908-NEXT: s_nop 0660; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a7661; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a6662; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a5663; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a4664; GREEDY908-NEXT: s_nop 1665; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[16:17] offset:16666; GREEDY908-NEXT: s_nop 0667; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a3668; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a2669; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a1670; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a0671; GREEDY908-NEXT: s_nop 1672; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[16:17]673; GREEDY908-NEXT: s_endpgm674;675; GREEDY90A-LABEL: test_mfma_f32_16x16x1f32:676; GREEDY90A: ; %bb.0: ; %bb677; GREEDY90A-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24678; GREEDY90A-NEXT: v_mov_b32_e32 v0, 1.0679; GREEDY90A-NEXT: v_mov_b32_e32 v1, 2.0680; GREEDY90A-NEXT: s_waitcnt lgkmcnt(0)681; GREEDY90A-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0682; GREEDY90A-NEXT: s_waitcnt lgkmcnt(0)683; GREEDY90A-NEXT: v_accvgpr_write_b32 a0, s0684; GREEDY90A-NEXT: v_accvgpr_write_b32 a1, s1685; GREEDY90A-NEXT: v_accvgpr_write_b32 a2, s2686; GREEDY90A-NEXT: v_accvgpr_write_b32 a3, s3687; GREEDY90A-NEXT: v_accvgpr_write_b32 a4, s4688; GREEDY90A-NEXT: v_accvgpr_write_b32 a5, s5689; GREEDY90A-NEXT: v_accvgpr_write_b32 a6, s6690; GREEDY90A-NEXT: v_accvgpr_write_b32 a7, s7691; GREEDY90A-NEXT: v_accvgpr_write_b32 a8, s8692; GREEDY90A-NEXT: v_accvgpr_write_b32 a9, s9693; GREEDY90A-NEXT: v_accvgpr_write_b32 a10, s10694; GREEDY90A-NEXT: v_accvgpr_write_b32 a11, s11695; GREEDY90A-NEXT: v_accvgpr_write_b32 a12, s12696; GREEDY90A-NEXT: v_accvgpr_write_b32 a13, s13697; GREEDY90A-NEXT: v_accvgpr_write_b32 a14, s14698; GREEDY90A-NEXT: v_accvgpr_write_b32 a15, s15699; GREEDY90A-NEXT: s_nop 1700; GREEDY90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[0:15]701; GREEDY90A-NEXT: v_mfma_f32_16x16x1f32 a[16:31], v0, v1, a[0:15]702; GREEDY90A-NEXT: s_nop 10703; GREEDY90A-NEXT: v_accvgpr_mov_b32 a2, a16704; GREEDY90A-NEXT: v_accvgpr_mov_b32 a3, a17705; GREEDY90A-NEXT: v_accvgpr_mov_b32 a4, a18706; GREEDY90A-NEXT: v_accvgpr_mov_b32 a5, a19707; GREEDY90A-NEXT: v_accvgpr_mov_b32 a6, a20708; GREEDY90A-NEXT: v_accvgpr_mov_b32 a7, a21709; GREEDY90A-NEXT: v_accvgpr_mov_b32 a8, a22710; GREEDY90A-NEXT: v_accvgpr_mov_b32 a9, a23711; GREEDY90A-NEXT: v_accvgpr_mov_b32 a10, a24712; GREEDY90A-NEXT: v_accvgpr_mov_b32 a11, a25713; GREEDY90A-NEXT: v_accvgpr_mov_b32 a12, a26714; GREEDY90A-NEXT: v_accvgpr_mov_b32 a13, a27715; GREEDY90A-NEXT: v_accvgpr_mov_b32 a14, a28716; GREEDY90A-NEXT: v_accvgpr_mov_b32 a15, a29717; GREEDY90A-NEXT: s_nop 1718; GREEDY90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[0:15]719; GREEDY90A-NEXT: v_mov_b32_e32 v0, 0720; GREEDY90A-NEXT: s_nop 9721; GREEDY90A-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48722; GREEDY90A-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32723; GREEDY90A-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16724; GREEDY90A-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]725; GREEDY90A-NEXT: s_endpgm726;727; GREEDY942-LABEL: test_mfma_f32_16x16x1f32:728; GREEDY942: ; %bb.0: ; %bb729; GREEDY942-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24730; GREEDY942-NEXT: v_mov_b32_e32 v0, 1.0731; GREEDY942-NEXT: v_mov_b32_e32 v1, 2.0732; GREEDY942-NEXT: s_waitcnt lgkmcnt(0)733; GREEDY942-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0734; GREEDY942-NEXT: s_waitcnt lgkmcnt(0)735; GREEDY942-NEXT: v_accvgpr_write_b32 a0, s0736; GREEDY942-NEXT: v_accvgpr_write_b32 a1, s1737; GREEDY942-NEXT: v_accvgpr_write_b32 a2, s2738; GREEDY942-NEXT: v_accvgpr_write_b32 a3, s3739; GREEDY942-NEXT: v_accvgpr_write_b32 a4, s4740; GREEDY942-NEXT: v_accvgpr_write_b32 a5, s5741; GREEDY942-NEXT: v_accvgpr_write_b32 a6, s6742; GREEDY942-NEXT: v_accvgpr_write_b32 a7, s7743; GREEDY942-NEXT: v_accvgpr_write_b32 a8, s8744; GREEDY942-NEXT: v_accvgpr_write_b32 a9, s9745; GREEDY942-NEXT: v_accvgpr_write_b32 a10, s10746; GREEDY942-NEXT: v_accvgpr_write_b32 a11, s11747; GREEDY942-NEXT: v_accvgpr_write_b32 a12, s12748; GREEDY942-NEXT: v_accvgpr_write_b32 a13, s13749; GREEDY942-NEXT: v_accvgpr_write_b32 a14, s14750; GREEDY942-NEXT: v_accvgpr_write_b32 a15, s15751; GREEDY942-NEXT: s_nop 1752; GREEDY942-NEXT: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[0:15]753; GREEDY942-NEXT: v_mfma_f32_16x16x1_4b_f32 a[16:31], v0, v1, a[0:15]754; GREEDY942-NEXT: s_nop 9755; GREEDY942-NEXT: v_accvgpr_mov_b32 a2, a16756; GREEDY942-NEXT: v_accvgpr_mov_b32 a3, a17757; GREEDY942-NEXT: v_accvgpr_mov_b32 a4, a18758; GREEDY942-NEXT: v_accvgpr_mov_b32 a5, a19759; GREEDY942-NEXT: v_accvgpr_mov_b32 a6, a20760; GREEDY942-NEXT: v_accvgpr_mov_b32 a7, a21761; GREEDY942-NEXT: v_accvgpr_mov_b32 a8, a22762; GREEDY942-NEXT: v_accvgpr_mov_b32 a9, a23763; GREEDY942-NEXT: v_accvgpr_mov_b32 a10, a24764; GREEDY942-NEXT: v_accvgpr_mov_b32 a11, a25765; GREEDY942-NEXT: v_accvgpr_mov_b32 a12, a26766; GREEDY942-NEXT: v_accvgpr_mov_b32 a13, a27767; GREEDY942-NEXT: v_accvgpr_mov_b32 a14, a28768; GREEDY942-NEXT: v_accvgpr_mov_b32 a15, a29769; GREEDY942-NEXT: s_nop 1770; GREEDY942-NEXT: v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[0:15]771; GREEDY942-NEXT: v_mov_b32_e32 v0, 0772; GREEDY942-NEXT: s_nop 8773; GREEDY942-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48774; GREEDY942-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32775; GREEDY942-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16776; GREEDY942-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]777; GREEDY942-NEXT: s_endpgm778;779; GREEDY90A-GISEL-LABEL: test_mfma_f32_16x16x1f32:780; GREEDY90A-GISEL: ; %bb.0: ; %bb781; GREEDY90A-GISEL-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24782; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v0, 1.0783; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v1, 2.0784; GREEDY90A-GISEL-NEXT: s_waitcnt lgkmcnt(0)785; GREEDY90A-GISEL-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0786; GREEDY90A-GISEL-NEXT: s_waitcnt lgkmcnt(0)787; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a0, s0788; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a1, s1789; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a2, s2790; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a3, s3791; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a4, s4792; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a5, s5793; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a6, s6794; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a7, s7795; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a8, s8796; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a9, s9797; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a10, s10798; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a11, s11799; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a12, s12800; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a13, s13801; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a14, s14802; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a15, s15803; GREEDY90A-GISEL-NEXT: s_nop 1804; GREEDY90A-GISEL-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[0:15]805; GREEDY90A-GISEL-NEXT: v_mfma_f32_16x16x1f32 a[16:31], v0, v1, a[0:15]806; GREEDY90A-GISEL-NEXT: s_nop 10807; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a2, a16808; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a3, a17809; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a4, a18810; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a5, a19811; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a6, a20812; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a7, a21813; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a8, a22814; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a9, a23815; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a10, a24816; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a11, a25817; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a12, a26818; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a13, a27819; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a14, a28820; GREEDY90A-GISEL-NEXT: v_accvgpr_mov_b32 a15, a29821; GREEDY90A-GISEL-NEXT: s_nop 1822; GREEDY90A-GISEL-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[0:15]823; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v0, 0824; GREEDY90A-GISEL-NEXT: s_nop 9825; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]826; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16827; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32828; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48829; GREEDY90A-GISEL-NEXT: s_endpgm830;831; FAST90A-LABEL: test_mfma_f32_16x16x1f32:832; FAST90A: ; %bb.0: ; %bb833; FAST90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24834; FAST90A-NEXT: v_mov_b32_e32 v0, 1.0835; FAST90A-NEXT: v_mov_b32_e32 v1, 2.0836; FAST90A-NEXT: s_waitcnt lgkmcnt(0)837; FAST90A-NEXT: s_load_dwordx16 s[4:19], s[0:1], 0x0838; FAST90A-NEXT: s_waitcnt lgkmcnt(0)839; FAST90A-NEXT: v_accvgpr_write_b32 a0, s4840; FAST90A-NEXT: v_accvgpr_write_b32 a1, s5841; FAST90A-NEXT: v_accvgpr_write_b32 a2, s6842; FAST90A-NEXT: v_accvgpr_write_b32 a3, s7843; FAST90A-NEXT: v_accvgpr_write_b32 a4, s8844; FAST90A-NEXT: v_accvgpr_write_b32 a5, s9845; FAST90A-NEXT: v_accvgpr_write_b32 a6, s10846; FAST90A-NEXT: v_accvgpr_write_b32 a7, s11847; FAST90A-NEXT: v_accvgpr_write_b32 a8, s12848; FAST90A-NEXT: v_accvgpr_write_b32 a9, s13849; FAST90A-NEXT: v_accvgpr_write_b32 a10, s14850; FAST90A-NEXT: v_accvgpr_write_b32 a11, s15851; FAST90A-NEXT: v_accvgpr_write_b32 a12, s16852; FAST90A-NEXT: v_accvgpr_write_b32 a13, s17853; FAST90A-NEXT: v_accvgpr_write_b32 a14, s18854; FAST90A-NEXT: v_accvgpr_write_b32 a15, s19855; FAST90A-NEXT: s_nop 1856; FAST90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[0:15]857; FAST90A-NEXT: v_mfma_f32_16x16x1f32 a[16:31], v0, v1, a[0:15]858; FAST90A-NEXT: s_nop 10859; FAST90A-NEXT: v_accvgpr_mov_b32 a2, a16860; FAST90A-NEXT: v_accvgpr_mov_b32 a3, a17861; FAST90A-NEXT: v_accvgpr_mov_b32 a4, a18862; FAST90A-NEXT: v_accvgpr_mov_b32 a5, a19863; FAST90A-NEXT: v_accvgpr_mov_b32 a6, a20864; FAST90A-NEXT: v_accvgpr_mov_b32 a7, a21865; FAST90A-NEXT: v_accvgpr_mov_b32 a8, a22866; FAST90A-NEXT: v_accvgpr_mov_b32 a9, a23867; FAST90A-NEXT: v_accvgpr_mov_b32 a10, a24868; FAST90A-NEXT: v_accvgpr_mov_b32 a11, a25869; FAST90A-NEXT: v_accvgpr_mov_b32 a12, a26870; FAST90A-NEXT: v_accvgpr_mov_b32 a13, a27871; FAST90A-NEXT: v_accvgpr_mov_b32 a14, a28872; FAST90A-NEXT: v_accvgpr_mov_b32 a15, a29873; FAST90A-NEXT: s_nop 1874; FAST90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v0, v1, a[0:15]875; FAST90A-NEXT: v_mov_b32_e32 v0, 0876; FAST90A-NEXT: s_nop 9877; FAST90A-NEXT: global_store_dwordx4 v0, a[12:15], s[0:1] offset:48878; FAST90A-NEXT: global_store_dwordx4 v0, a[8:11], s[0:1] offset:32879; FAST90A-NEXT: global_store_dwordx4 v0, a[4:7], s[0:1] offset:16880; FAST90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]881; FAST90A-NEXT: s_endpgm882bb:883 %in.1 = load <16 x float>, ptr addrspace(1) %arg884 %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> %in.1, i32 0, i32 0, i32 0)885 %mai.2 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> %mai.1, i32 0, i32 0, i32 0)886 %tmp.1 = shufflevector <16 x float> %mai.2, <16 x float> %mai.1, <16 x i32> <i32 16, i32 17, i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13>887 %mai.3 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> %tmp.1, i32 0, i32 0, i32 0)888 store <16 x float> %mai.3, ptr addrspace(1) %arg889 ret void890}891 892; This instruction allows the overlap since it only read 4 registers.893 894define amdgpu_kernel void @test_mfma_f32_4x4x1f32(ptr addrspace(1) %arg) #0 {895; GREEDY908-LABEL: test_mfma_f32_4x4x1f32:896; GREEDY908: ; %bb.0: ; %bb897; GREEDY908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24898; GREEDY908-NEXT: v_mov_b32_e32 v0, 1.0899; GREEDY908-NEXT: v_mov_b32_e32 v1, 2.0900; GREEDY908-NEXT: v_mov_b32_e32 v4, 0901; GREEDY908-NEXT: s_waitcnt lgkmcnt(0)902; GREEDY908-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0903; GREEDY908-NEXT: s_waitcnt lgkmcnt(0)904; GREEDY908-NEXT: v_mov_b32_e32 v5, s0905; GREEDY908-NEXT: v_mov_b32_e32 v2, s1906; GREEDY908-NEXT: v_mov_b32_e32 v3, s2907; GREEDY908-NEXT: v_accvgpr_write_b32 a0, v5908; GREEDY908-NEXT: v_mov_b32_e32 v5, s3909; GREEDY908-NEXT: v_accvgpr_write_b32 a1, v2910; GREEDY908-NEXT: v_accvgpr_write_b32 a2, v3911; GREEDY908-NEXT: v_accvgpr_write_b32 a3, v5912; GREEDY908-NEXT: s_nop 0913; GREEDY908-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]914; GREEDY908-NEXT: v_mfma_f32_4x4x1f32 a[2:5], v0, v1, a[0:3]915; GREEDY908-NEXT: s_nop 1916; GREEDY908-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]917; GREEDY908-NEXT: s_nop 3918; GREEDY908-NEXT: v_accvgpr_read_b32 v0, a0919; GREEDY908-NEXT: v_accvgpr_read_b32 v1, a1920; GREEDY908-NEXT: v_accvgpr_read_b32 v2, a2921; GREEDY908-NEXT: v_accvgpr_read_b32 v3, a3922; GREEDY908-NEXT: s_nop 1923; GREEDY908-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]924; GREEDY908-NEXT: s_endpgm925;926; GREEDY90A-LABEL: test_mfma_f32_4x4x1f32:927; GREEDY90A: ; %bb.0: ; %bb928; GREEDY90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24929; GREEDY90A-NEXT: v_mov_b32_e32 v0, 1.0930; GREEDY90A-NEXT: v_mov_b32_e32 v1, 2.0931; GREEDY90A-NEXT: v_mov_b32_e32 v2, 0932; GREEDY90A-NEXT: s_waitcnt lgkmcnt(0)933; GREEDY90A-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0934; GREEDY90A-NEXT: s_waitcnt lgkmcnt(0)935; GREEDY90A-NEXT: v_accvgpr_write_b32 a0, s0936; GREEDY90A-NEXT: v_accvgpr_write_b32 a1, s1937; GREEDY90A-NEXT: v_accvgpr_write_b32 a2, s2938; GREEDY90A-NEXT: v_accvgpr_write_b32 a3, s3939; GREEDY90A-NEXT: s_nop 1940; GREEDY90A-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]941; GREEDY90A-NEXT: v_mfma_f32_4x4x1f32 a[2:5], v0, v1, a[0:3]942; GREEDY90A-NEXT: s_nop 1943; GREEDY90A-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]944; GREEDY90A-NEXT: s_nop 4945; GREEDY90A-NEXT: global_store_dwordx4 v2, a[0:3], s[6:7]946; GREEDY90A-NEXT: s_endpgm947;948; GREEDY942-LABEL: test_mfma_f32_4x4x1f32:949; GREEDY942: ; %bb.0: ; %bb950; GREEDY942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24951; GREEDY942-NEXT: v_mov_b32_e32 v0, 1.0952; GREEDY942-NEXT: v_mov_b32_e32 v1, 2.0953; GREEDY942-NEXT: v_mov_b32_e32 v2, 0954; GREEDY942-NEXT: s_waitcnt lgkmcnt(0)955; GREEDY942-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0956; GREEDY942-NEXT: s_waitcnt lgkmcnt(0)957; GREEDY942-NEXT: v_accvgpr_write_b32 a0, s0958; GREEDY942-NEXT: v_accvgpr_write_b32 a1, s1959; GREEDY942-NEXT: v_accvgpr_write_b32 a2, s2960; GREEDY942-NEXT: v_accvgpr_write_b32 a3, s3961; GREEDY942-NEXT: s_nop 1962; GREEDY942-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[0:3]963; GREEDY942-NEXT: s_nop 1964; GREEDY942-NEXT: v_mfma_f32_4x4x1_16b_f32 a[2:5], v0, v1, a[0:3]965; GREEDY942-NEXT: s_nop 1966; GREEDY942-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[0:3]967; GREEDY942-NEXT: s_nop 3968; GREEDY942-NEXT: global_store_dwordx4 v2, a[0:3], s[6:7]969; GREEDY942-NEXT: s_endpgm970;971; GREEDY90A-GISEL-LABEL: test_mfma_f32_4x4x1f32:972; GREEDY90A-GISEL: ; %bb.0: ; %bb973; GREEDY90A-GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24974; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v0, 1.0975; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v1, 2.0976; GREEDY90A-GISEL-NEXT: s_waitcnt lgkmcnt(0)977; GREEDY90A-GISEL-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0978; GREEDY90A-GISEL-NEXT: s_waitcnt lgkmcnt(0)979; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a0, s0980; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a1, s1981; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a2, s2982; GREEDY90A-GISEL-NEXT: v_accvgpr_write_b32 a3, s3983; GREEDY90A-GISEL-NEXT: s_nop 1984; GREEDY90A-GISEL-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]985; GREEDY90A-GISEL-NEXT: v_mfma_f32_4x4x1f32 a[2:5], v0, v1, a[0:3]986; GREEDY90A-GISEL-NEXT: s_nop 1987; GREEDY90A-GISEL-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]988; GREEDY90A-GISEL-NEXT: v_mov_b32_e32 v0, 0989; GREEDY90A-GISEL-NEXT: s_nop 3990; GREEDY90A-GISEL-NEXT: global_store_dwordx4 v0, a[0:3], s[6:7]991; GREEDY90A-GISEL-NEXT: s_endpgm992;993; FAST90A-LABEL: test_mfma_f32_4x4x1f32:994; FAST90A: ; %bb.0: ; %bb995; FAST90A-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24996; FAST90A-NEXT: v_mov_b32_e32 v1, 1.0997; FAST90A-NEXT: v_mov_b32_e32 v2, 2.0998; FAST90A-NEXT: v_mov_b32_e32 v0, 0999; FAST90A-NEXT: s_waitcnt lgkmcnt(0)1000; FAST90A-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x01001; FAST90A-NEXT: s_waitcnt lgkmcnt(0)1002; FAST90A-NEXT: v_accvgpr_write_b32 a0, s41003; FAST90A-NEXT: v_accvgpr_write_b32 a1, s51004; FAST90A-NEXT: v_accvgpr_write_b32 a2, s61005; FAST90A-NEXT: v_accvgpr_write_b32 a3, s71006; FAST90A-NEXT: s_nop 11007; FAST90A-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v1, v2, a[0:3]1008; FAST90A-NEXT: v_mfma_f32_4x4x1f32 a[2:5], v1, v2, a[0:3]1009; FAST90A-NEXT: s_nop 11010; FAST90A-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v1, v2, a[0:3]1011; FAST90A-NEXT: s_nop 41012; FAST90A-NEXT: global_store_dwordx4 v0, a[0:3], s[0:1]1013; FAST90A-NEXT: s_endpgm1014bb:1015 %in.1 = load <4 x float>, ptr addrspace(1) %arg1016 %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in.1, i32 0, i32 0, i32 0)1017 %mai.2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %mai.1, i32 0, i32 0, i32 0)1018 %tmp.1 = shufflevector <4 x float> %mai.1, <4 x float> %mai.2, <4 x i32> <i32 0, i32 1, i32 4, i32 5>1019 %mai.3 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %tmp.1, i32 0, i32 0, i32 0)1020 store <4 x float> %mai.3, ptr addrspace(1) %arg1021 ret void1022}1023 1024attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }1025