1135 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck --check-prefixes=GCN,GFX908 %s3; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck --check-prefixes=GCN,GFX90A %s4; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck --check-prefixes=GCN,GFX90A %s5; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GCN,GFX90A %s6; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GCN,GFX90A %s7 8declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)9 10define amdgpu_kernel void @test_mfma_f32_32x32x1f32_vgpr(ptr addrspace(1) %arg) #0 {11; GFX908-LABEL: test_mfma_f32_32x32x1f32_vgpr:12; GFX908: ; %bb.0: ; %bb13; GFX908-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x2414; GFX908-NEXT: v_mov_b32_e32 v4, 015; GFX908-NEXT: s_waitcnt lgkmcnt(0)16; GFX908-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x017; GFX908-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x4018; GFX908-NEXT: s_waitcnt lgkmcnt(0)19; GFX908-NEXT: v_mov_b32_e32 v0, s1620; GFX908-NEXT: v_mov_b32_e32 v1, s1721; GFX908-NEXT: v_mov_b32_e32 v2, s1822; GFX908-NEXT: v_accvgpr_write_b32 a0, v023; GFX908-NEXT: v_accvgpr_write_b32 a1, v124; GFX908-NEXT: v_accvgpr_write_b32 a2, v225; GFX908-NEXT: v_mov_b32_e32 v0, s2126; GFX908-NEXT: v_mov_b32_e32 v1, s2227; GFX908-NEXT: v_mov_b32_e32 v2, s2328; GFX908-NEXT: v_accvgpr_write_b32 a5, v029; GFX908-NEXT: v_accvgpr_write_b32 a6, v130; GFX908-NEXT: v_accvgpr_write_b32 a7, v231; GFX908-NEXT: v_mov_b32_e32 v0, s2432; GFX908-NEXT: v_mov_b32_e32 v1, s2533; GFX908-NEXT: v_mov_b32_e32 v2, s2634; GFX908-NEXT: v_accvgpr_write_b32 a8, v035; GFX908-NEXT: v_accvgpr_write_b32 a9, v136; GFX908-NEXT: v_accvgpr_write_b32 a10, v237; GFX908-NEXT: v_mov_b32_e32 v0, s2738; GFX908-NEXT: v_mov_b32_e32 v1, s2839; GFX908-NEXT: v_mov_b32_e32 v2, s2940; GFX908-NEXT: v_accvgpr_write_b32 a11, v041; GFX908-NEXT: v_accvgpr_write_b32 a12, v142; GFX908-NEXT: v_accvgpr_write_b32 a13, v243; GFX908-NEXT: v_mov_b32_e32 v0, s3044; GFX908-NEXT: v_mov_b32_e32 v1, s3145; GFX908-NEXT: v_mov_b32_e32 v2, s046; GFX908-NEXT: v_accvgpr_write_b32 a14, v047; GFX908-NEXT: v_accvgpr_write_b32 a15, v148; GFX908-NEXT: v_accvgpr_write_b32 a16, v249; GFX908-NEXT: v_mov_b32_e32 v0, s150; GFX908-NEXT: v_mov_b32_e32 v1, s251; GFX908-NEXT: v_mov_b32_e32 v2, s352; GFX908-NEXT: v_accvgpr_write_b32 a17, v053; GFX908-NEXT: v_accvgpr_write_b32 a18, v154; GFX908-NEXT: v_accvgpr_write_b32 a19, v255; GFX908-NEXT: v_mov_b32_e32 v0, s456; GFX908-NEXT: v_mov_b32_e32 v1, s557; GFX908-NEXT: v_mov_b32_e32 v2, s658; GFX908-NEXT: v_accvgpr_write_b32 a20, v059; GFX908-NEXT: v_accvgpr_write_b32 a21, v160; GFX908-NEXT: v_accvgpr_write_b32 a22, v261; GFX908-NEXT: v_mov_b32_e32 v0, s762; GFX908-NEXT: v_mov_b32_e32 v1, s863; GFX908-NEXT: v_mov_b32_e32 v2, s964; GFX908-NEXT: v_mov_b32_e32 v3, s1965; GFX908-NEXT: v_accvgpr_write_b32 a23, v066; GFX908-NEXT: v_accvgpr_write_b32 a24, v167; GFX908-NEXT: v_accvgpr_write_b32 a25, v268; GFX908-NEXT: v_mov_b32_e32 v0, s1069; GFX908-NEXT: v_mov_b32_e32 v1, s1170; GFX908-NEXT: v_mov_b32_e32 v2, s1271; GFX908-NEXT: v_mov_b32_e32 v5, s2072; GFX908-NEXT: v_accvgpr_write_b32 a3, v373; GFX908-NEXT: v_accvgpr_write_b32 a26, v074; GFX908-NEXT: v_accvgpr_write_b32 a27, v175; GFX908-NEXT: v_accvgpr_write_b32 a28, v276; GFX908-NEXT: v_mov_b32_e32 v0, s1377; GFX908-NEXT: v_mov_b32_e32 v1, s1478; GFX908-NEXT: v_mov_b32_e32 v2, s1579; GFX908-NEXT: v_mov_b32_e32 v3, 1.080; GFX908-NEXT: v_accvgpr_write_b32 a4, v581; GFX908-NEXT: v_accvgpr_write_b32 a29, v082; GFX908-NEXT: v_accvgpr_write_b32 a30, v183; GFX908-NEXT: v_accvgpr_write_b32 a31, v284; GFX908-NEXT: v_mov_b32_e32 v0, 2.085; GFX908-NEXT: s_nop 186; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v3, v0, a[0:31]87; GFX908-NEXT: s_nop 1588; GFX908-NEXT: s_nop 189; GFX908-NEXT: v_accvgpr_read_b32 v3, a2790; GFX908-NEXT: v_accvgpr_read_b32 v2, a2691; GFX908-NEXT: v_accvgpr_read_b32 v1, a2592; GFX908-NEXT: v_accvgpr_read_b32 v0, a2493; GFX908-NEXT: s_nop 194; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:9695; GFX908-NEXT: s_nop 096; GFX908-NEXT: v_accvgpr_read_b32 v3, a3197; GFX908-NEXT: v_accvgpr_read_b32 v2, a3098; GFX908-NEXT: v_accvgpr_read_b32 v1, a2999; GFX908-NEXT: v_accvgpr_read_b32 v0, a28100; GFX908-NEXT: s_nop 1101; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:112102; GFX908-NEXT: s_nop 0103; GFX908-NEXT: v_accvgpr_read_b32 v3, a19104; GFX908-NEXT: v_accvgpr_read_b32 v2, a18105; GFX908-NEXT: v_accvgpr_read_b32 v1, a17106; GFX908-NEXT: v_accvgpr_read_b32 v0, a16107; GFX908-NEXT: s_nop 1108; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:64109; GFX908-NEXT: s_nop 0110; GFX908-NEXT: v_accvgpr_read_b32 v3, a23111; GFX908-NEXT: v_accvgpr_read_b32 v2, a22112; GFX908-NEXT: v_accvgpr_read_b32 v1, a21113; GFX908-NEXT: v_accvgpr_read_b32 v0, a20114; GFX908-NEXT: s_nop 1115; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:80116; GFX908-NEXT: s_nop 0117; GFX908-NEXT: v_accvgpr_read_b32 v3, a11118; GFX908-NEXT: v_accvgpr_read_b32 v2, a10119; GFX908-NEXT: v_accvgpr_read_b32 v1, a9120; GFX908-NEXT: v_accvgpr_read_b32 v0, a8121; GFX908-NEXT: s_nop 1122; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:32123; GFX908-NEXT: s_nop 0124; GFX908-NEXT: v_accvgpr_read_b32 v3, a15125; GFX908-NEXT: v_accvgpr_read_b32 v2, a14126; GFX908-NEXT: v_accvgpr_read_b32 v1, a13127; GFX908-NEXT: v_accvgpr_read_b32 v0, a12128; GFX908-NEXT: s_nop 1129; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:48130; GFX908-NEXT: s_nop 0131; GFX908-NEXT: v_accvgpr_read_b32 v3, a3132; GFX908-NEXT: v_accvgpr_read_b32 v2, a2133; GFX908-NEXT: v_accvgpr_read_b32 v1, a1134; GFX908-NEXT: v_accvgpr_read_b32 v0, a0135; GFX908-NEXT: s_nop 1136; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35]137; GFX908-NEXT: s_nop 0138; GFX908-NEXT: v_accvgpr_read_b32 v3, a7139; GFX908-NEXT: v_accvgpr_read_b32 v2, a6140; GFX908-NEXT: v_accvgpr_read_b32 v1, a5141; GFX908-NEXT: v_accvgpr_read_b32 v0, a4142; GFX908-NEXT: s_nop 1143; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:16144; GFX908-NEXT: s_endpgm145bb:146 %in.1 = load <32 x float>, ptr addrspace(1) %arg147 %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)148 store <32 x float> %mai.1, ptr addrspace(1) %arg149 ret void150}151 152define amdgpu_kernel void @test_mfma_f32_32x32x1f32_agpr(ptr addrspace(1) %arg) #2 {153; GFX908-LABEL: test_mfma_f32_32x32x1f32_agpr:154; GFX908: ; %bb.0: ; %bb155; GFX908-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24156; GFX908-NEXT: v_mov_b32_e32 v4, 0157; GFX908-NEXT: s_waitcnt lgkmcnt(0)158; GFX908-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0159; GFX908-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40160; GFX908-NEXT: s_waitcnt lgkmcnt(0)161; GFX908-NEXT: v_mov_b32_e32 v0, s16162; GFX908-NEXT: v_mov_b32_e32 v1, s17163; GFX908-NEXT: v_mov_b32_e32 v2, s18164; GFX908-NEXT: v_accvgpr_write_b32 a0, v0165; GFX908-NEXT: v_accvgpr_write_b32 a1, v1166; GFX908-NEXT: v_accvgpr_write_b32 a2, v2167; GFX908-NEXT: v_mov_b32_e32 v0, s21168; GFX908-NEXT: v_mov_b32_e32 v1, s22169; GFX908-NEXT: v_mov_b32_e32 v2, s23170; GFX908-NEXT: v_accvgpr_write_b32 a5, v0171; GFX908-NEXT: v_accvgpr_write_b32 a6, v1172; GFX908-NEXT: v_accvgpr_write_b32 a7, v2173; GFX908-NEXT: v_mov_b32_e32 v0, s24174; GFX908-NEXT: v_mov_b32_e32 v1, s25175; GFX908-NEXT: v_mov_b32_e32 v2, s26176; GFX908-NEXT: v_accvgpr_write_b32 a8, v0177; GFX908-NEXT: v_accvgpr_write_b32 a9, v1178; GFX908-NEXT: v_accvgpr_write_b32 a10, v2179; GFX908-NEXT: v_mov_b32_e32 v0, s27180; GFX908-NEXT: v_mov_b32_e32 v1, s28181; GFX908-NEXT: v_mov_b32_e32 v2, s29182; GFX908-NEXT: v_accvgpr_write_b32 a11, v0183; GFX908-NEXT: v_accvgpr_write_b32 a12, v1184; GFX908-NEXT: v_accvgpr_write_b32 a13, v2185; GFX908-NEXT: v_mov_b32_e32 v0, s30186; GFX908-NEXT: v_mov_b32_e32 v1, s31187; GFX908-NEXT: v_mov_b32_e32 v2, s0188; GFX908-NEXT: v_accvgpr_write_b32 a14, v0189; GFX908-NEXT: v_accvgpr_write_b32 a15, v1190; GFX908-NEXT: v_accvgpr_write_b32 a16, v2191; GFX908-NEXT: v_mov_b32_e32 v0, s1192; GFX908-NEXT: v_mov_b32_e32 v1, s2193; GFX908-NEXT: v_mov_b32_e32 v2, s3194; GFX908-NEXT: v_accvgpr_write_b32 a17, v0195; GFX908-NEXT: v_accvgpr_write_b32 a18, v1196; GFX908-NEXT: v_accvgpr_write_b32 a19, v2197; GFX908-NEXT: v_mov_b32_e32 v0, s4198; GFX908-NEXT: v_mov_b32_e32 v1, s5199; GFX908-NEXT: v_mov_b32_e32 v2, s6200; GFX908-NEXT: v_accvgpr_write_b32 a20, v0201; GFX908-NEXT: v_accvgpr_write_b32 a21, v1202; GFX908-NEXT: v_accvgpr_write_b32 a22, v2203; GFX908-NEXT: v_mov_b32_e32 v0, s7204; GFX908-NEXT: v_mov_b32_e32 v1, s8205; GFX908-NEXT: v_mov_b32_e32 v2, s9206; GFX908-NEXT: v_mov_b32_e32 v3, s19207; GFX908-NEXT: v_accvgpr_write_b32 a23, v0208; GFX908-NEXT: v_accvgpr_write_b32 a24, v1209; GFX908-NEXT: v_accvgpr_write_b32 a25, v2210; GFX908-NEXT: v_mov_b32_e32 v0, s10211; GFX908-NEXT: v_mov_b32_e32 v1, s11212; GFX908-NEXT: v_mov_b32_e32 v2, s12213; GFX908-NEXT: v_mov_b32_e32 v5, s20214; GFX908-NEXT: v_accvgpr_write_b32 a3, v3215; GFX908-NEXT: v_accvgpr_write_b32 a26, v0216; GFX908-NEXT: v_accvgpr_write_b32 a27, v1217; GFX908-NEXT: v_accvgpr_write_b32 a28, v2218; GFX908-NEXT: v_mov_b32_e32 v0, s13219; GFX908-NEXT: v_mov_b32_e32 v1, s14220; GFX908-NEXT: v_mov_b32_e32 v2, s15221; GFX908-NEXT: v_mov_b32_e32 v3, 1.0222; GFX908-NEXT: v_accvgpr_write_b32 a4, v5223; GFX908-NEXT: v_accvgpr_write_b32 a29, v0224; GFX908-NEXT: v_accvgpr_write_b32 a30, v1225; GFX908-NEXT: v_accvgpr_write_b32 a31, v2226; GFX908-NEXT: v_mov_b32_e32 v0, 2.0227; GFX908-NEXT: s_nop 1228; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v3, v0, a[0:31]229; GFX908-NEXT: s_nop 15230; GFX908-NEXT: s_nop 1231; GFX908-NEXT: v_accvgpr_read_b32 v3, a27232; GFX908-NEXT: v_accvgpr_read_b32 v2, a26233; GFX908-NEXT: v_accvgpr_read_b32 v1, a25234; GFX908-NEXT: v_accvgpr_read_b32 v0, a24235; GFX908-NEXT: s_nop 1236; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:96237; GFX908-NEXT: s_nop 0238; GFX908-NEXT: v_accvgpr_read_b32 v3, a31239; GFX908-NEXT: v_accvgpr_read_b32 v2, a30240; GFX908-NEXT: v_accvgpr_read_b32 v1, a29241; GFX908-NEXT: v_accvgpr_read_b32 v0, a28242; GFX908-NEXT: s_nop 1243; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:112244; GFX908-NEXT: s_nop 0245; GFX908-NEXT: v_accvgpr_read_b32 v3, a19246; GFX908-NEXT: v_accvgpr_read_b32 v2, a18247; GFX908-NEXT: v_accvgpr_read_b32 v1, a17248; GFX908-NEXT: v_accvgpr_read_b32 v0, a16249; GFX908-NEXT: s_nop 1250; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:64251; GFX908-NEXT: s_nop 0252; GFX908-NEXT: v_accvgpr_read_b32 v3, a23253; GFX908-NEXT: v_accvgpr_read_b32 v2, a22254; GFX908-NEXT: v_accvgpr_read_b32 v1, a21255; GFX908-NEXT: v_accvgpr_read_b32 v0, a20256; GFX908-NEXT: s_nop 1257; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:80258; GFX908-NEXT: s_nop 0259; GFX908-NEXT: v_accvgpr_read_b32 v3, a11260; GFX908-NEXT: v_accvgpr_read_b32 v2, a10261; GFX908-NEXT: v_accvgpr_read_b32 v1, a9262; GFX908-NEXT: v_accvgpr_read_b32 v0, a8263; GFX908-NEXT: s_nop 1264; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:32265; GFX908-NEXT: s_nop 0266; GFX908-NEXT: v_accvgpr_read_b32 v3, a15267; GFX908-NEXT: v_accvgpr_read_b32 v2, a14268; GFX908-NEXT: v_accvgpr_read_b32 v1, a13269; GFX908-NEXT: v_accvgpr_read_b32 v0, a12270; GFX908-NEXT: s_nop 1271; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:48272; GFX908-NEXT: s_nop 0273; GFX908-NEXT: v_accvgpr_read_b32 v3, a3274; GFX908-NEXT: v_accvgpr_read_b32 v2, a2275; GFX908-NEXT: v_accvgpr_read_b32 v1, a1276; GFX908-NEXT: v_accvgpr_read_b32 v0, a0277; GFX908-NEXT: s_nop 1278; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35]279; GFX908-NEXT: s_nop 0280; GFX908-NEXT: v_accvgpr_read_b32 v3, a7281; GFX908-NEXT: v_accvgpr_read_b32 v2, a6282; GFX908-NEXT: v_accvgpr_read_b32 v1, a5283; GFX908-NEXT: v_accvgpr_read_b32 v0, a4284; GFX908-NEXT: s_nop 1285; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[34:35] offset:16286; GFX908-NEXT: s_endpgm287bb:288 %in.1 = load <32 x float>, ptr addrspace(1) %arg289 %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)290 store <32 x float> %mai.1, ptr addrspace(1) %arg291 ret void292}293 294define amdgpu_kernel void @test_mfma_f32_32x32x1f32_inline_asm_virtual_agpr(ptr addrspace(1) %arg) {295; GFX908-LABEL: test_mfma_f32_32x32x1f32_inline_asm_virtual_agpr:296; GFX908: ; %bb.0: ; %bb297; GFX908-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24298; GFX908-NEXT: v_mov_b32_e32 v32, 0299; GFX908-NEXT: ;;#ASMSTART300; GFX908-NEXT: ; def a0301; GFX908-NEXT: ;;#ASMEND302; GFX908-NEXT: s_waitcnt lgkmcnt(0)303; GFX908-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:112304; GFX908-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:96305; GFX908-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:80306; GFX908-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:64307; GFX908-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:48308; GFX908-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:32309; GFX908-NEXT: global_load_dwordx4 v[4:7], v32, s[0:1] offset:16310; GFX908-NEXT: global_load_dwordx4 v[0:3], v32, s[0:1]311; GFX908-NEXT: s_waitcnt vmcnt(0)312; GFX908-NEXT: v_accvgpr_write_b32 a0, v0313; GFX908-NEXT: v_accvgpr_write_b32 a1, v1314; GFX908-NEXT: v_accvgpr_write_b32 a2, v2315; GFX908-NEXT: v_accvgpr_write_b32 a3, v3316; GFX908-NEXT: v_accvgpr_write_b32 a4, v4317; GFX908-NEXT: v_accvgpr_write_b32 a5, v5318; GFX908-NEXT: v_accvgpr_write_b32 a6, v6319; GFX908-NEXT: v_accvgpr_write_b32 a7, v7320; GFX908-NEXT: v_accvgpr_write_b32 a8, v8321; GFX908-NEXT: v_accvgpr_write_b32 a9, v9322; GFX908-NEXT: v_accvgpr_write_b32 a10, v10323; GFX908-NEXT: v_accvgpr_write_b32 a11, v11324; GFX908-NEXT: v_accvgpr_write_b32 a12, v12325; GFX908-NEXT: v_accvgpr_write_b32 a13, v13326; GFX908-NEXT: v_accvgpr_write_b32 a14, v14327; GFX908-NEXT: v_accvgpr_write_b32 a15, v15328; GFX908-NEXT: v_accvgpr_write_b32 a16, v16329; GFX908-NEXT: v_accvgpr_write_b32 a17, v17330; GFX908-NEXT: v_accvgpr_write_b32 a18, v18331; GFX908-NEXT: v_accvgpr_write_b32 a19, v19332; GFX908-NEXT: v_accvgpr_write_b32 a20, v20333; GFX908-NEXT: v_accvgpr_write_b32 a21, v21334; GFX908-NEXT: v_accvgpr_write_b32 a22, v22335; GFX908-NEXT: v_accvgpr_write_b32 a23, v23336; GFX908-NEXT: v_accvgpr_write_b32 a24, v24337; GFX908-NEXT: v_accvgpr_write_b32 a25, v25338; GFX908-NEXT: v_accvgpr_write_b32 a26, v26339; GFX908-NEXT: v_accvgpr_write_b32 a27, v27340; GFX908-NEXT: v_accvgpr_write_b32 a28, v28341; GFX908-NEXT: v_accvgpr_write_b32 a29, v29342; GFX908-NEXT: v_accvgpr_write_b32 a30, v30343; GFX908-NEXT: v_accvgpr_write_b32 a31, v31344; GFX908-NEXT: v_mov_b32_e32 v0, 1.0345; GFX908-NEXT: v_mov_b32_e32 v1, 2.0346; GFX908-NEXT: s_nop 1347; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]348; GFX908-NEXT: s_nop 15349; GFX908-NEXT: s_nop 1350; GFX908-NEXT: v_accvgpr_read_b32 v3, a27351; GFX908-NEXT: v_accvgpr_read_b32 v2, a26352; GFX908-NEXT: v_accvgpr_read_b32 v1, a25353; GFX908-NEXT: v_accvgpr_read_b32 v0, a24354; GFX908-NEXT: v_accvgpr_read_b32 v7, a31355; GFX908-NEXT: v_accvgpr_read_b32 v6, a30356; GFX908-NEXT: v_accvgpr_read_b32 v5, a29357; GFX908-NEXT: v_accvgpr_read_b32 v4, a28358; GFX908-NEXT: v_accvgpr_read_b32 v11, a19359; GFX908-NEXT: v_accvgpr_read_b32 v10, a18360; GFX908-NEXT: v_accvgpr_read_b32 v9, a17361; GFX908-NEXT: v_accvgpr_read_b32 v8, a16362; GFX908-NEXT: v_accvgpr_read_b32 v15, a23363; GFX908-NEXT: v_accvgpr_read_b32 v14, a22364; GFX908-NEXT: v_accvgpr_read_b32 v13, a21365; GFX908-NEXT: v_accvgpr_read_b32 v12, a20366; GFX908-NEXT: v_accvgpr_read_b32 v19, a11367; GFX908-NEXT: v_accvgpr_read_b32 v18, a10368; GFX908-NEXT: v_accvgpr_read_b32 v17, a9369; GFX908-NEXT: v_accvgpr_read_b32 v16, a8370; GFX908-NEXT: v_accvgpr_read_b32 v23, a15371; GFX908-NEXT: v_accvgpr_read_b32 v22, a14372; GFX908-NEXT: v_accvgpr_read_b32 v21, a13373; GFX908-NEXT: v_accvgpr_read_b32 v20, a12374; GFX908-NEXT: v_accvgpr_read_b32 v27, a3375; GFX908-NEXT: v_accvgpr_read_b32 v26, a2376; GFX908-NEXT: v_accvgpr_read_b32 v25, a1377; GFX908-NEXT: v_accvgpr_read_b32 v24, a0378; GFX908-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:96379; GFX908-NEXT: s_nop 0380; GFX908-NEXT: v_accvgpr_read_b32 v3, a7381; GFX908-NEXT: v_accvgpr_read_b32 v2, a6382; GFX908-NEXT: v_accvgpr_read_b32 v1, a5383; GFX908-NEXT: v_accvgpr_read_b32 v0, a4384; GFX908-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:112385; GFX908-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:64386; GFX908-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:80387; GFX908-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32388; GFX908-NEXT: global_store_dwordx4 v32, v[20:23], s[0:1] offset:48389; GFX908-NEXT: global_store_dwordx4 v32, v[24:27], s[0:1]390; GFX908-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:16391; GFX908-NEXT: s_endpgm392bb:393 %acc = call i32 asm sideeffect "; def $0", "={a0}"()394 %in.1 = load <32 x float>, ptr addrspace(1) %arg395 %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)396 store <32 x float> %mai.1, ptr addrspace(1) %arg397 ret void398}399 400define amdgpu_kernel void @test_mfma_f32_32x32x1f32_inline_asm_phys_agpr(ptr addrspace(1) %arg) {401; GFX908-LABEL: test_mfma_f32_32x32x1f32_inline_asm_phys_agpr:402; GFX908: ; %bb.0: ; %bb403; GFX908-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24404; GFX908-NEXT: v_mov_b32_e32 v32, 0405; GFX908-NEXT: ;;#ASMSTART406; GFX908-NEXT: ; use a[100:131]407; GFX908-NEXT: ;;#ASMEND408; GFX908-NEXT: s_waitcnt lgkmcnt(0)409; GFX908-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:112410; GFX908-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:96411; GFX908-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:80412; GFX908-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:64413; GFX908-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:48414; GFX908-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:32415; GFX908-NEXT: global_load_dwordx4 v[4:7], v32, s[0:1] offset:16416; GFX908-NEXT: global_load_dwordx4 v[0:3], v32, s[0:1]417; GFX908-NEXT: s_waitcnt vmcnt(0)418; GFX908-NEXT: v_accvgpr_write_b32 a0, v0419; GFX908-NEXT: v_accvgpr_write_b32 a1, v1420; GFX908-NEXT: v_accvgpr_write_b32 a2, v2421; GFX908-NEXT: v_accvgpr_write_b32 a3, v3422; GFX908-NEXT: v_accvgpr_write_b32 a4, v4423; GFX908-NEXT: v_accvgpr_write_b32 a5, v5424; GFX908-NEXT: v_accvgpr_write_b32 a6, v6425; GFX908-NEXT: v_accvgpr_write_b32 a7, v7426; GFX908-NEXT: v_accvgpr_write_b32 a8, v8427; GFX908-NEXT: v_accvgpr_write_b32 a9, v9428; GFX908-NEXT: v_accvgpr_write_b32 a10, v10429; GFX908-NEXT: v_accvgpr_write_b32 a11, v11430; GFX908-NEXT: v_accvgpr_write_b32 a12, v12431; GFX908-NEXT: v_accvgpr_write_b32 a13, v13432; GFX908-NEXT: v_accvgpr_write_b32 a14, v14433; GFX908-NEXT: v_accvgpr_write_b32 a15, v15434; GFX908-NEXT: v_accvgpr_write_b32 a16, v16435; GFX908-NEXT: v_accvgpr_write_b32 a17, v17436; GFX908-NEXT: v_accvgpr_write_b32 a18, v18437; GFX908-NEXT: v_accvgpr_write_b32 a19, v19438; GFX908-NEXT: v_accvgpr_write_b32 a20, v20439; GFX908-NEXT: v_accvgpr_write_b32 a21, v21440; GFX908-NEXT: v_accvgpr_write_b32 a22, v22441; GFX908-NEXT: v_accvgpr_write_b32 a23, v23442; GFX908-NEXT: v_accvgpr_write_b32 a24, v24443; GFX908-NEXT: v_accvgpr_write_b32 a25, v25444; GFX908-NEXT: v_accvgpr_write_b32 a26, v26445; GFX908-NEXT: v_accvgpr_write_b32 a27, v27446; GFX908-NEXT: v_accvgpr_write_b32 a28, v28447; GFX908-NEXT: v_accvgpr_write_b32 a29, v29448; GFX908-NEXT: v_accvgpr_write_b32 a30, v30449; GFX908-NEXT: v_accvgpr_write_b32 a31, v31450; GFX908-NEXT: v_mov_b32_e32 v0, 1.0451; GFX908-NEXT: v_mov_b32_e32 v1, 2.0452; GFX908-NEXT: s_nop 1453; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]454; GFX908-NEXT: s_nop 15455; GFX908-NEXT: s_nop 1456; GFX908-NEXT: v_accvgpr_read_b32 v3, a27457; GFX908-NEXT: v_accvgpr_read_b32 v2, a26458; GFX908-NEXT: v_accvgpr_read_b32 v1, a25459; GFX908-NEXT: v_accvgpr_read_b32 v0, a24460; GFX908-NEXT: v_accvgpr_read_b32 v7, a31461; GFX908-NEXT: v_accvgpr_read_b32 v6, a30462; GFX908-NEXT: v_accvgpr_read_b32 v5, a29463; GFX908-NEXT: v_accvgpr_read_b32 v4, a28464; GFX908-NEXT: v_accvgpr_read_b32 v11, a19465; GFX908-NEXT: v_accvgpr_read_b32 v10, a18466; GFX908-NEXT: v_accvgpr_read_b32 v9, a17467; GFX908-NEXT: v_accvgpr_read_b32 v8, a16468; GFX908-NEXT: v_accvgpr_read_b32 v15, a23469; GFX908-NEXT: v_accvgpr_read_b32 v14, a22470; GFX908-NEXT: v_accvgpr_read_b32 v13, a21471; GFX908-NEXT: v_accvgpr_read_b32 v12, a20472; GFX908-NEXT: v_accvgpr_read_b32 v19, a11473; GFX908-NEXT: v_accvgpr_read_b32 v18, a10474; GFX908-NEXT: v_accvgpr_read_b32 v17, a9475; GFX908-NEXT: v_accvgpr_read_b32 v16, a8476; GFX908-NEXT: v_accvgpr_read_b32 v23, a15477; GFX908-NEXT: v_accvgpr_read_b32 v22, a14478; GFX908-NEXT: v_accvgpr_read_b32 v21, a13479; GFX908-NEXT: v_accvgpr_read_b32 v20, a12480; GFX908-NEXT: v_accvgpr_read_b32 v27, a3481; GFX908-NEXT: v_accvgpr_read_b32 v26, a2482; GFX908-NEXT: v_accvgpr_read_b32 v25, a1483; GFX908-NEXT: v_accvgpr_read_b32 v24, a0484; GFX908-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:96485; GFX908-NEXT: s_nop 0486; GFX908-NEXT: v_accvgpr_read_b32 v3, a7487; GFX908-NEXT: v_accvgpr_read_b32 v2, a6488; GFX908-NEXT: v_accvgpr_read_b32 v1, a5489; GFX908-NEXT: v_accvgpr_read_b32 v0, a4490; GFX908-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:112491; GFX908-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:64492; GFX908-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:80493; GFX908-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32494; GFX908-NEXT: global_store_dwordx4 v32, v[20:23], s[0:1] offset:48495; GFX908-NEXT: global_store_dwordx4 v32, v[24:27], s[0:1]496; GFX908-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:16497; GFX908-NEXT: s_endpgm498bb:499 call void asm sideeffect "; use $0", "{a[100:131]}"(<32 x float> poison)500 %in.1 = load <32 x float>, ptr addrspace(1) %arg501 %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)502 store <32 x float> %mai.1, ptr addrspace(1) %arg503 ret void504}505 506define amdgpu_kernel void @test_mfma_f32_32x32x1f32_inline_asm_no_agprs(ptr addrspace(1) %arg) #0 {507; GFX908-LABEL: test_mfma_f32_32x32x1f32_inline_asm_no_agprs:508; GFX908: ; %bb.0: ; %bb509; GFX908-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24510; GFX908-NEXT: v_mov_b32_e32 v32, 0511; GFX908-NEXT: ;;#ASMSTART512; GFX908-NEXT: ; def v0513; GFX908-NEXT: ;;#ASMEND514; GFX908-NEXT: s_waitcnt lgkmcnt(0)515; GFX908-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:112516; GFX908-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:96517; GFX908-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:80518; GFX908-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:64519; GFX908-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:48520; GFX908-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:32521; GFX908-NEXT: global_load_dwordx4 v[4:7], v32, s[0:1] offset:16522; GFX908-NEXT: global_load_dwordx4 v[0:3], v32, s[0:1]523; GFX908-NEXT: s_waitcnt vmcnt(0)524; GFX908-NEXT: v_accvgpr_write_b32 a0, v0525; GFX908-NEXT: v_accvgpr_write_b32 a1, v1526; GFX908-NEXT: v_accvgpr_write_b32 a2, v2527; GFX908-NEXT: v_accvgpr_write_b32 a3, v3528; GFX908-NEXT: v_accvgpr_write_b32 a4, v4529; GFX908-NEXT: v_accvgpr_write_b32 a5, v5530; GFX908-NEXT: v_accvgpr_write_b32 a6, v6531; GFX908-NEXT: v_accvgpr_write_b32 a7, v7532; GFX908-NEXT: v_accvgpr_write_b32 a8, v8533; GFX908-NEXT: v_accvgpr_write_b32 a9, v9534; GFX908-NEXT: v_accvgpr_write_b32 a10, v10535; GFX908-NEXT: v_accvgpr_write_b32 a11, v11536; GFX908-NEXT: v_accvgpr_write_b32 a12, v12537; GFX908-NEXT: v_accvgpr_write_b32 a13, v13538; GFX908-NEXT: v_accvgpr_write_b32 a14, v14539; GFX908-NEXT: v_accvgpr_write_b32 a15, v15540; GFX908-NEXT: v_accvgpr_write_b32 a16, v16541; GFX908-NEXT: v_accvgpr_write_b32 a17, v17542; GFX908-NEXT: v_accvgpr_write_b32 a18, v18543; GFX908-NEXT: v_accvgpr_write_b32 a19, v19544; GFX908-NEXT: v_accvgpr_write_b32 a20, v20545; GFX908-NEXT: v_accvgpr_write_b32 a21, v21546; GFX908-NEXT: v_accvgpr_write_b32 a22, v22547; GFX908-NEXT: v_accvgpr_write_b32 a23, v23548; GFX908-NEXT: v_accvgpr_write_b32 a24, v24549; GFX908-NEXT: v_accvgpr_write_b32 a25, v25550; GFX908-NEXT: v_accvgpr_write_b32 a26, v26551; GFX908-NEXT: v_accvgpr_write_b32 a27, v27552; GFX908-NEXT: v_accvgpr_write_b32 a28, v28553; GFX908-NEXT: v_accvgpr_write_b32 a29, v29554; GFX908-NEXT: v_accvgpr_write_b32 a30, v30555; GFX908-NEXT: v_accvgpr_write_b32 a31, v31556; GFX908-NEXT: v_mov_b32_e32 v0, 1.0557; GFX908-NEXT: v_mov_b32_e32 v1, 2.0558; GFX908-NEXT: s_nop 1559; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]560; GFX908-NEXT: s_nop 15561; GFX908-NEXT: s_nop 1562; GFX908-NEXT: v_accvgpr_read_b32 v3, a27563; GFX908-NEXT: v_accvgpr_read_b32 v2, a26564; GFX908-NEXT: v_accvgpr_read_b32 v1, a25565; GFX908-NEXT: v_accvgpr_read_b32 v0, a24566; GFX908-NEXT: v_accvgpr_read_b32 v7, a31567; GFX908-NEXT: v_accvgpr_read_b32 v6, a30568; GFX908-NEXT: v_accvgpr_read_b32 v5, a29569; GFX908-NEXT: v_accvgpr_read_b32 v4, a28570; GFX908-NEXT: v_accvgpr_read_b32 v11, a19571; GFX908-NEXT: v_accvgpr_read_b32 v10, a18572; GFX908-NEXT: v_accvgpr_read_b32 v9, a17573; GFX908-NEXT: v_accvgpr_read_b32 v8, a16574; GFX908-NEXT: v_accvgpr_read_b32 v15, a23575; GFX908-NEXT: v_accvgpr_read_b32 v14, a22576; GFX908-NEXT: v_accvgpr_read_b32 v13, a21577; GFX908-NEXT: v_accvgpr_read_b32 v12, a20578; GFX908-NEXT: v_accvgpr_read_b32 v19, a11579; GFX908-NEXT: v_accvgpr_read_b32 v18, a10580; GFX908-NEXT: v_accvgpr_read_b32 v17, a9581; GFX908-NEXT: v_accvgpr_read_b32 v16, a8582; GFX908-NEXT: v_accvgpr_read_b32 v23, a15583; GFX908-NEXT: v_accvgpr_read_b32 v22, a14584; GFX908-NEXT: v_accvgpr_read_b32 v21, a13585; GFX908-NEXT: v_accvgpr_read_b32 v20, a12586; GFX908-NEXT: v_accvgpr_read_b32 v27, a3587; GFX908-NEXT: v_accvgpr_read_b32 v26, a2588; GFX908-NEXT: v_accvgpr_read_b32 v25, a1589; GFX908-NEXT: v_accvgpr_read_b32 v24, a0590; GFX908-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:96591; GFX908-NEXT: s_nop 0592; GFX908-NEXT: v_accvgpr_read_b32 v3, a7593; GFX908-NEXT: v_accvgpr_read_b32 v2, a6594; GFX908-NEXT: v_accvgpr_read_b32 v1, a5595; GFX908-NEXT: v_accvgpr_read_b32 v0, a4596; GFX908-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:112597; GFX908-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:64598; GFX908-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:80599; GFX908-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32600; GFX908-NEXT: global_store_dwordx4 v32, v[20:23], s[0:1] offset:48601; GFX908-NEXT: global_store_dwordx4 v32, v[24:27], s[0:1]602; GFX908-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:16603; GFX908-NEXT: s_endpgm604bb:605 %acc = call i32 asm sideeffect "; def $0", "={v0}"()606 %in.1 = load <32 x float>, ptr addrspace(1) %arg607 %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)608 store <32 x float> %mai.1, ptr addrspace(1) %arg609 ret void610}611 612define amdgpu_kernel void @test_mfma_f32_32x32x1f32_call(ptr addrspace(1) %arg) #1 {613; GFX908-LABEL: test_mfma_f32_32x32x1f32_call:614; GFX908: ; %bb.0: ; %bb615; GFX908-NEXT: s_mov_b32 s36, SCRATCH_RSRC_DWORD0616; GFX908-NEXT: s_mov_b32 s37, SCRATCH_RSRC_DWORD1617; GFX908-NEXT: s_mov_b32 s38, -1618; GFX908-NEXT: s_mov_b32 s39, 0xe00000619; GFX908-NEXT: s_add_u32 s36, s36, s11620; GFX908-NEXT: s_addc_u32 s37, s37, 0621; GFX908-NEXT: s_mov_b32 s12, s8622; GFX908-NEXT: s_add_u32 s8, s4, 44623; GFX908-NEXT: s_mov_b32 s13, s9624; GFX908-NEXT: s_addc_u32 s9, s5, 0625; GFX908-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24626; GFX908-NEXT: s_getpc_b64 s[4:5]627; GFX908-NEXT: s_add_u32 s4, s4, foo@gotpcrel32@lo+4628; GFX908-NEXT: s_addc_u32 s5, s5, foo@gotpcrel32@hi+12629; GFX908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x0630; GFX908-NEXT: s_mov_b32 s14, s10631; GFX908-NEXT: s_mov_b64 s[10:11], s[6:7]632; GFX908-NEXT: v_lshlrev_b32_e32 v2, 20, v2633; GFX908-NEXT: v_lshlrev_b32_e32 v1, 10, v1634; GFX908-NEXT: s_mov_b64 s[4:5], s[0:1]635; GFX908-NEXT: s_mov_b64 s[6:7], s[2:3]636; GFX908-NEXT: s_mov_b64 s[0:1], s[36:37]637; GFX908-NEXT: v_or3_b32 v31, v0, v1, v2638; GFX908-NEXT: s_mov_b64 s[2:3], s[38:39]639; GFX908-NEXT: s_mov_b32 s32, 0640; GFX908-NEXT: v_mov_b32_e32 v40, 0641; GFX908-NEXT: s_waitcnt lgkmcnt(0)642; GFX908-NEXT: s_swappc_b64 s[30:31], s[16:17]643; GFX908-NEXT: global_load_dwordx4 v[28:31], v40, s[34:35] offset:112644; GFX908-NEXT: global_load_dwordx4 v[24:27], v40, s[34:35] offset:96645; GFX908-NEXT: global_load_dwordx4 v[20:23], v40, s[34:35] offset:80646; GFX908-NEXT: global_load_dwordx4 v[16:19], v40, s[34:35] offset:64647; GFX908-NEXT: global_load_dwordx4 v[12:15], v40, s[34:35] offset:48648; GFX908-NEXT: global_load_dwordx4 v[8:11], v40, s[34:35] offset:32649; GFX908-NEXT: global_load_dwordx4 v[4:7], v40, s[34:35] offset:16650; GFX908-NEXT: global_load_dwordx4 v[0:3], v40, s[34:35]651; GFX908-NEXT: s_waitcnt vmcnt(0)652; GFX908-NEXT: v_accvgpr_write_b32 a0, v0653; GFX908-NEXT: v_accvgpr_write_b32 a1, v1654; GFX908-NEXT: v_accvgpr_write_b32 a2, v2655; GFX908-NEXT: v_accvgpr_write_b32 a3, v3656; GFX908-NEXT: v_accvgpr_write_b32 a4, v4657; GFX908-NEXT: v_accvgpr_write_b32 a5, v5658; GFX908-NEXT: v_accvgpr_write_b32 a6, v6659; GFX908-NEXT: v_accvgpr_write_b32 a7, v7660; GFX908-NEXT: v_accvgpr_write_b32 a8, v8661; GFX908-NEXT: v_accvgpr_write_b32 a9, v9662; GFX908-NEXT: v_accvgpr_write_b32 a10, v10663; GFX908-NEXT: v_accvgpr_write_b32 a11, v11664; GFX908-NEXT: v_accvgpr_write_b32 a12, v12665; GFX908-NEXT: v_accvgpr_write_b32 a13, v13666; GFX908-NEXT: v_accvgpr_write_b32 a14, v14667; GFX908-NEXT: v_accvgpr_write_b32 a15, v15668; GFX908-NEXT: v_accvgpr_write_b32 a16, v16669; GFX908-NEXT: v_accvgpr_write_b32 a17, v17670; GFX908-NEXT: v_accvgpr_write_b32 a18, v18671; GFX908-NEXT: v_accvgpr_write_b32 a19, v19672; GFX908-NEXT: v_accvgpr_write_b32 a20, v20673; GFX908-NEXT: v_accvgpr_write_b32 a21, v21674; GFX908-NEXT: v_accvgpr_write_b32 a22, v22675; GFX908-NEXT: v_accvgpr_write_b32 a23, v23676; GFX908-NEXT: v_accvgpr_write_b32 a24, v24677; GFX908-NEXT: v_accvgpr_write_b32 a25, v25678; GFX908-NEXT: v_accvgpr_write_b32 a26, v26679; GFX908-NEXT: v_accvgpr_write_b32 a27, v27680; GFX908-NEXT: v_accvgpr_write_b32 a28, v28681; GFX908-NEXT: v_accvgpr_write_b32 a29, v29682; GFX908-NEXT: v_accvgpr_write_b32 a30, v30683; GFX908-NEXT: v_accvgpr_write_b32 a31, v31684; GFX908-NEXT: v_mov_b32_e32 v0, 1.0685; GFX908-NEXT: v_mov_b32_e32 v1, 2.0686; GFX908-NEXT: s_nop 1687; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]688; GFX908-NEXT: s_nop 15689; GFX908-NEXT: s_nop 1690; GFX908-NEXT: v_accvgpr_read_b32 v3, a27691; GFX908-NEXT: v_accvgpr_read_b32 v2, a26692; GFX908-NEXT: v_accvgpr_read_b32 v1, a25693; GFX908-NEXT: v_accvgpr_read_b32 v0, a24694; GFX908-NEXT: v_accvgpr_read_b32 v7, a31695; GFX908-NEXT: v_accvgpr_read_b32 v6, a30696; GFX908-NEXT: v_accvgpr_read_b32 v5, a29697; GFX908-NEXT: v_accvgpr_read_b32 v4, a28698; GFX908-NEXT: v_accvgpr_read_b32 v11, a19699; GFX908-NEXT: v_accvgpr_read_b32 v10, a18700; GFX908-NEXT: v_accvgpr_read_b32 v9, a17701; GFX908-NEXT: v_accvgpr_read_b32 v8, a16702; GFX908-NEXT: v_accvgpr_read_b32 v15, a23703; GFX908-NEXT: v_accvgpr_read_b32 v14, a22704; GFX908-NEXT: v_accvgpr_read_b32 v13, a21705; GFX908-NEXT: v_accvgpr_read_b32 v12, a20706; GFX908-NEXT: v_accvgpr_read_b32 v19, a11707; GFX908-NEXT: v_accvgpr_read_b32 v18, a10708; GFX908-NEXT: v_accvgpr_read_b32 v17, a9709; GFX908-NEXT: v_accvgpr_read_b32 v16, a8710; GFX908-NEXT: v_accvgpr_read_b32 v23, a15711; GFX908-NEXT: v_accvgpr_read_b32 v22, a14712; GFX908-NEXT: v_accvgpr_read_b32 v21, a13713; GFX908-NEXT: v_accvgpr_read_b32 v20, a12714; GFX908-NEXT: v_accvgpr_read_b32 v27, a3715; GFX908-NEXT: v_accvgpr_read_b32 v26, a2716; GFX908-NEXT: v_accvgpr_read_b32 v25, a1717; GFX908-NEXT: v_accvgpr_read_b32 v24, a0718; GFX908-NEXT: global_store_dwordx4 v40, v[0:3], s[34:35] offset:96719; GFX908-NEXT: s_nop 0720; GFX908-NEXT: v_accvgpr_read_b32 v3, a7721; GFX908-NEXT: v_accvgpr_read_b32 v2, a6722; GFX908-NEXT: v_accvgpr_read_b32 v1, a5723; GFX908-NEXT: v_accvgpr_read_b32 v0, a4724; GFX908-NEXT: global_store_dwordx4 v40, v[4:7], s[34:35] offset:112725; GFX908-NEXT: global_store_dwordx4 v40, v[8:11], s[34:35] offset:64726; GFX908-NEXT: global_store_dwordx4 v40, v[12:15], s[34:35] offset:80727; GFX908-NEXT: global_store_dwordx4 v40, v[16:19], s[34:35] offset:32728; GFX908-NEXT: global_store_dwordx4 v40, v[20:23], s[34:35] offset:48729; GFX908-NEXT: global_store_dwordx4 v40, v[24:27], s[34:35]730; GFX908-NEXT: global_store_dwordx4 v40, v[0:3], s[34:35] offset:16731; GFX908-NEXT: s_endpgm732bb:733 call void @foo()734 %in.1 = load <32 x float>, ptr addrspace(1) %arg735 %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)736 store <32 x float> %mai.1, ptr addrspace(1) %arg737 ret void738}739 740; We could avoid scan to find calls since we see these during lowering before selection.741; However, in SDag lowering and selection is done block by block, so it would only work742; in Global ISel.743define amdgpu_kernel void @test_mfma_f32_32x32x1f32_call_multi_bb(ptr addrspace(1) %arg, i1 %c0) #1 {744; GFX908-LABEL: test_mfma_f32_32x32x1f32_call_multi_bb:745; GFX908: ; %bb.0: ; %bb1746; GFX908-NEXT: s_mov_b32 s52, SCRATCH_RSRC_DWORD0747; GFX908-NEXT: s_mov_b32 s53, SCRATCH_RSRC_DWORD1748; GFX908-NEXT: s_mov_b32 s54, -1749; GFX908-NEXT: s_mov_b32 s55, 0xe00000750; GFX908-NEXT: s_add_u32 s52, s52, s11751; GFX908-NEXT: s_mov_b32 s14, s10752; GFX908-NEXT: s_mov_b32 s12, s8753; GFX908-NEXT: s_mov_b64 s[10:11], s[6:7]754; GFX908-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24755; GFX908-NEXT: s_load_dword s8, s[4:5], 0x2c756; GFX908-NEXT: v_mov_b32_e32 v6, 1.0757; GFX908-NEXT: v_mov_b32_e32 v7, 0758; GFX908-NEXT: s_addc_u32 s53, s53, 0759; GFX908-NEXT: s_waitcnt lgkmcnt(0)760; GFX908-NEXT: s_load_dwordx16 s[36:51], s[6:7], 0x0761; GFX908-NEXT: s_load_dwordx16 s[16:31], s[6:7], 0x40762; GFX908-NEXT: s_bitcmp0_b32 s8, 0763; GFX908-NEXT: s_mov_b32 s32, 0764; GFX908-NEXT: s_waitcnt lgkmcnt(0)765; GFX908-NEXT: v_mov_b32_e32 v3, s36766; GFX908-NEXT: v_mov_b32_e32 v4, s37767; GFX908-NEXT: v_mov_b32_e32 v5, s40768; GFX908-NEXT: v_accvgpr_write_b32 a0, v3769; GFX908-NEXT: v_accvgpr_write_b32 a1, v4770; GFX908-NEXT: v_mov_b32_e32 v3, s38771; GFX908-NEXT: v_mov_b32_e32 v4, s39772; GFX908-NEXT: v_accvgpr_write_b32 a4, v5773; GFX908-NEXT: v_accvgpr_write_b32 a2, v3774; GFX908-NEXT: v_accvgpr_write_b32 a3, v4775; GFX908-NEXT: v_mov_b32_e32 v3, s41776; GFX908-NEXT: v_mov_b32_e32 v4, s42777; GFX908-NEXT: v_mov_b32_e32 v5, s43778; GFX908-NEXT: v_accvgpr_write_b32 a5, v3779; GFX908-NEXT: v_accvgpr_write_b32 a6, v4780; GFX908-NEXT: v_accvgpr_write_b32 a7, v5781; GFX908-NEXT: v_mov_b32_e32 v3, s44782; GFX908-NEXT: v_mov_b32_e32 v4, s45783; GFX908-NEXT: v_mov_b32_e32 v5, s46784; GFX908-NEXT: v_accvgpr_write_b32 a8, v3785; GFX908-NEXT: v_accvgpr_write_b32 a9, v4786; GFX908-NEXT: v_accvgpr_write_b32 a10, v5787; GFX908-NEXT: v_mov_b32_e32 v3, s47788; GFX908-NEXT: v_mov_b32_e32 v4, s48789; GFX908-NEXT: v_mov_b32_e32 v5, s49790; GFX908-NEXT: v_accvgpr_write_b32 a11, v3791; GFX908-NEXT: v_accvgpr_write_b32 a12, v4792; GFX908-NEXT: v_accvgpr_write_b32 a13, v5793; GFX908-NEXT: v_mov_b32_e32 v3, s50794; GFX908-NEXT: v_mov_b32_e32 v4, s51795; GFX908-NEXT: v_mov_b32_e32 v5, s16796; GFX908-NEXT: v_accvgpr_write_b32 a14, v3797; GFX908-NEXT: v_accvgpr_write_b32 a15, v4798; GFX908-NEXT: v_accvgpr_write_b32 a16, v5799; GFX908-NEXT: v_mov_b32_e32 v3, s17800; GFX908-NEXT: v_mov_b32_e32 v4, s18801; GFX908-NEXT: v_mov_b32_e32 v5, s19802; GFX908-NEXT: v_accvgpr_write_b32 a17, v3803; GFX908-NEXT: v_accvgpr_write_b32 a18, v4804; GFX908-NEXT: v_accvgpr_write_b32 a19, v5805; GFX908-NEXT: v_mov_b32_e32 v3, s20806; GFX908-NEXT: v_mov_b32_e32 v4, s21807; GFX908-NEXT: v_mov_b32_e32 v5, s22808; GFX908-NEXT: v_accvgpr_write_b32 a20, v3809; GFX908-NEXT: v_accvgpr_write_b32 a21, v4810; GFX908-NEXT: v_accvgpr_write_b32 a22, v5811; GFX908-NEXT: v_mov_b32_e32 v3, s23812; GFX908-NEXT: v_mov_b32_e32 v4, s24813; GFX908-NEXT: v_mov_b32_e32 v5, s25814; GFX908-NEXT: v_accvgpr_write_b32 a23, v3815; GFX908-NEXT: v_accvgpr_write_b32 a24, v4816; GFX908-NEXT: v_accvgpr_write_b32 a25, v5817; GFX908-NEXT: v_mov_b32_e32 v3, s26818; GFX908-NEXT: v_mov_b32_e32 v4, s27819; GFX908-NEXT: v_mov_b32_e32 v5, s28820; GFX908-NEXT: v_accvgpr_write_b32 a26, v3821; GFX908-NEXT: v_accvgpr_write_b32 a27, v4822; GFX908-NEXT: v_accvgpr_write_b32 a28, v5823; GFX908-NEXT: v_mov_b32_e32 v3, s29824; GFX908-NEXT: v_mov_b32_e32 v4, s30825; GFX908-NEXT: v_mov_b32_e32 v5, s31826; GFX908-NEXT: v_accvgpr_write_b32 a29, v3827; GFX908-NEXT: v_accvgpr_write_b32 a30, v4828; GFX908-NEXT: v_accvgpr_write_b32 a31, v5829; GFX908-NEXT: v_mov_b32_e32 v3, 2.0830; GFX908-NEXT: s_nop 1831; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v6, v3, a[0:31] cbsz:1 abid:2 blgp:3832; GFX908-NEXT: s_nop 15833; GFX908-NEXT: s_nop 1834; GFX908-NEXT: v_accvgpr_read_b32 v6, a27835; GFX908-NEXT: v_accvgpr_read_b32 v5, a26836; GFX908-NEXT: v_accvgpr_read_b32 v4, a25837; GFX908-NEXT: v_accvgpr_read_b32 v3, a24838; GFX908-NEXT: s_nop 1839; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7] offset:96840; GFX908-NEXT: s_nop 0841; GFX908-NEXT: v_accvgpr_read_b32 v6, a31842; GFX908-NEXT: v_accvgpr_read_b32 v5, a30843; GFX908-NEXT: v_accvgpr_read_b32 v4, a29844; GFX908-NEXT: v_accvgpr_read_b32 v3, a28845; GFX908-NEXT: s_nop 1846; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7] offset:112847; GFX908-NEXT: s_nop 0848; GFX908-NEXT: v_accvgpr_read_b32 v6, a19849; GFX908-NEXT: v_accvgpr_read_b32 v5, a18850; GFX908-NEXT: v_accvgpr_read_b32 v4, a17851; GFX908-NEXT: v_accvgpr_read_b32 v3, a16852; GFX908-NEXT: s_nop 1853; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7] offset:64854; GFX908-NEXT: s_nop 0855; GFX908-NEXT: v_accvgpr_read_b32 v6, a23856; GFX908-NEXT: v_accvgpr_read_b32 v5, a22857; GFX908-NEXT: v_accvgpr_read_b32 v4, a21858; GFX908-NEXT: v_accvgpr_read_b32 v3, a20859; GFX908-NEXT: s_nop 1860; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7] offset:80861; GFX908-NEXT: s_nop 0862; GFX908-NEXT: v_accvgpr_read_b32 v6, a11863; GFX908-NEXT: v_accvgpr_read_b32 v5, a10864; GFX908-NEXT: v_accvgpr_read_b32 v4, a9865; GFX908-NEXT: v_accvgpr_read_b32 v3, a8866; GFX908-NEXT: s_nop 1867; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7] offset:32868; GFX908-NEXT: s_nop 0869; GFX908-NEXT: v_accvgpr_read_b32 v6, a15870; GFX908-NEXT: v_accvgpr_read_b32 v5, a14871; GFX908-NEXT: v_accvgpr_read_b32 v4, a13872; GFX908-NEXT: v_accvgpr_read_b32 v3, a12873; GFX908-NEXT: s_nop 1874; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7] offset:48875; GFX908-NEXT: s_nop 0876; GFX908-NEXT: v_accvgpr_read_b32 v6, a3877; GFX908-NEXT: v_accvgpr_read_b32 v5, a2878; GFX908-NEXT: v_accvgpr_read_b32 v4, a1879; GFX908-NEXT: v_accvgpr_read_b32 v3, a0880; GFX908-NEXT: s_nop 1881; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7]882; GFX908-NEXT: s_nop 0883; GFX908-NEXT: v_accvgpr_read_b32 v6, a7884; GFX908-NEXT: v_accvgpr_read_b32 v5, a6885; GFX908-NEXT: v_accvgpr_read_b32 v4, a5886; GFX908-NEXT: v_accvgpr_read_b32 v3, a4887; GFX908-NEXT: s_nop 1888; GFX908-NEXT: global_store_dwordx4 v7, v[3:6], s[6:7] offset:16889; GFX908-NEXT: s_cbranch_scc1 .LBB6_2890; GFX908-NEXT: ; %bb.1: ; %bb2891; GFX908-NEXT: s_add_u32 s8, s4, 48892; GFX908-NEXT: s_mov_b32 s13, s9893; GFX908-NEXT: s_addc_u32 s9, s5, 0894; GFX908-NEXT: s_getpc_b64 s[4:5]895; GFX908-NEXT: s_add_u32 s4, s4, foo@gotpcrel32@lo+4896; GFX908-NEXT: s_addc_u32 s5, s5, foo@gotpcrel32@hi+12897; GFX908-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x0898; GFX908-NEXT: v_lshlrev_b32_e32 v2, 20, v2899; GFX908-NEXT: v_lshlrev_b32_e32 v1, 10, v1900; GFX908-NEXT: s_mov_b64 s[4:5], s[0:1]901; GFX908-NEXT: s_mov_b64 s[6:7], s[2:3]902; GFX908-NEXT: s_mov_b64 s[0:1], s[52:53]903; GFX908-NEXT: v_or3_b32 v31, v0, v1, v2904; GFX908-NEXT: s_mov_b64 s[2:3], s[54:55]905; GFX908-NEXT: s_waitcnt lgkmcnt(0)906; GFX908-NEXT: s_swappc_b64 s[30:31], s[16:17]907; GFX908-NEXT: .LBB6_2: ; %bb3908; GFX908-NEXT: s_endpgm909bb1:910 %in.1 = load <32 x float>, ptr addrspace(1) %arg911 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)912 store <32 x float> %mai.1, ptr addrspace(1) %arg913 br i1 %c0, label %bb2, label %bb3914 br label %bb2915 916bb2:917 call void @foo()918 br label %bb3919 920bb3:921 ret void922}923 924define void @test_mfma_f32_32x32x1f32_nonentry_noagpr(ptr addrspace(1) %arg) #0 {925; GFX908-LABEL: test_mfma_f32_32x32x1f32_nonentry_noagpr:926; GFX908: ; %bb.0: ; %bb927; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)928; GFX908-NEXT: global_load_dwordx4 v[30:33], v[0:1], off offset:112929; GFX908-NEXT: global_load_dwordx4 v[26:29], v[0:1], off offset:96930; GFX908-NEXT: global_load_dwordx4 v[22:25], v[0:1], off offset:80931; GFX908-NEXT: global_load_dwordx4 v[18:21], v[0:1], off offset:64932; GFX908-NEXT: global_load_dwordx4 v[14:17], v[0:1], off offset:48933; GFX908-NEXT: global_load_dwordx4 v[10:13], v[0:1], off offset:32934; GFX908-NEXT: global_load_dwordx4 v[6:9], v[0:1], off offset:16935; GFX908-NEXT: global_load_dwordx4 v[2:5], v[0:1], off936; GFX908-NEXT: s_waitcnt vmcnt(0)937; GFX908-NEXT: v_accvgpr_write_b32 a0, v2938; GFX908-NEXT: v_accvgpr_write_b32 a1, v3939; GFX908-NEXT: v_accvgpr_write_b32 a2, v4940; GFX908-NEXT: v_accvgpr_write_b32 a3, v5941; GFX908-NEXT: v_accvgpr_write_b32 a4, v6942; GFX908-NEXT: v_accvgpr_write_b32 a5, v7943; GFX908-NEXT: v_accvgpr_write_b32 a6, v8944; GFX908-NEXT: v_accvgpr_write_b32 a7, v9945; GFX908-NEXT: v_accvgpr_write_b32 a8, v10946; GFX908-NEXT: v_accvgpr_write_b32 a9, v11947; GFX908-NEXT: v_accvgpr_write_b32 a10, v12948; GFX908-NEXT: v_accvgpr_write_b32 a11, v13949; GFX908-NEXT: v_accvgpr_write_b32 a12, v14950; GFX908-NEXT: v_accvgpr_write_b32 a13, v15951; GFX908-NEXT: v_accvgpr_write_b32 a14, v16952; GFX908-NEXT: v_accvgpr_write_b32 a15, v17953; GFX908-NEXT: v_accvgpr_write_b32 a16, v18954; GFX908-NEXT: v_accvgpr_write_b32 a17, v19955; GFX908-NEXT: v_accvgpr_write_b32 a18, v20956; GFX908-NEXT: v_accvgpr_write_b32 a19, v21957; GFX908-NEXT: v_accvgpr_write_b32 a20, v22958; GFX908-NEXT: v_accvgpr_write_b32 a21, v23959; GFX908-NEXT: v_accvgpr_write_b32 a22, v24960; GFX908-NEXT: v_accvgpr_write_b32 a23, v25961; GFX908-NEXT: v_accvgpr_write_b32 a24, v26962; GFX908-NEXT: v_accvgpr_write_b32 a25, v27963; GFX908-NEXT: v_accvgpr_write_b32 a26, v28964; GFX908-NEXT: v_accvgpr_write_b32 a27, v29965; GFX908-NEXT: v_accvgpr_write_b32 a28, v30966; GFX908-NEXT: v_accvgpr_write_b32 a29, v31967; GFX908-NEXT: v_accvgpr_write_b32 a30, v32968; GFX908-NEXT: v_accvgpr_write_b32 a31, v33969; GFX908-NEXT: v_mov_b32_e32 v2, 1.0970; GFX908-NEXT: v_mov_b32_e32 v3, 2.0971; GFX908-NEXT: s_nop 1972; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]973; GFX908-NEXT: s_nop 15974; GFX908-NEXT: s_nop 1975; GFX908-NEXT: v_accvgpr_read_b32 v5, a27976; GFX908-NEXT: v_accvgpr_read_b32 v4, a26977; GFX908-NEXT: v_accvgpr_read_b32 v3, a25978; GFX908-NEXT: v_accvgpr_read_b32 v2, a24979; GFX908-NEXT: v_accvgpr_read_b32 v9, a31980; GFX908-NEXT: v_accvgpr_read_b32 v8, a30981; GFX908-NEXT: v_accvgpr_read_b32 v7, a29982; GFX908-NEXT: v_accvgpr_read_b32 v6, a28983; GFX908-NEXT: v_accvgpr_read_b32 v13, a19984; GFX908-NEXT: v_accvgpr_read_b32 v12, a18985; GFX908-NEXT: v_accvgpr_read_b32 v11, a17986; GFX908-NEXT: v_accvgpr_read_b32 v10, a16987; GFX908-NEXT: v_accvgpr_read_b32 v17, a23988; GFX908-NEXT: v_accvgpr_read_b32 v16, a22989; GFX908-NEXT: v_accvgpr_read_b32 v15, a21990; GFX908-NEXT: v_accvgpr_read_b32 v14, a20991; GFX908-NEXT: v_accvgpr_read_b32 v21, a11992; GFX908-NEXT: v_accvgpr_read_b32 v20, a10993; GFX908-NEXT: v_accvgpr_read_b32 v19, a9994; GFX908-NEXT: v_accvgpr_read_b32 v18, a8995; GFX908-NEXT: v_accvgpr_read_b32 v25, a15996; GFX908-NEXT: v_accvgpr_read_b32 v24, a14997; GFX908-NEXT: v_accvgpr_read_b32 v23, a13998; GFX908-NEXT: v_accvgpr_read_b32 v22, a12999; GFX908-NEXT: v_accvgpr_read_b32 v29, a31000; GFX908-NEXT: v_accvgpr_read_b32 v28, a21001; GFX908-NEXT: v_accvgpr_read_b32 v27, a11002; GFX908-NEXT: v_accvgpr_read_b32 v26, a01003; GFX908-NEXT: global_store_dwordx4 v[0:1], v[2:5], off offset:961004; GFX908-NEXT: s_nop 01005; GFX908-NEXT: v_accvgpr_read_b32 v5, a71006; GFX908-NEXT: v_accvgpr_read_b32 v4, a61007; GFX908-NEXT: v_accvgpr_read_b32 v3, a51008; GFX908-NEXT: v_accvgpr_read_b32 v2, a41009; GFX908-NEXT: global_store_dwordx4 v[0:1], v[6:9], off offset:1121010; GFX908-NEXT: global_store_dwordx4 v[0:1], v[10:13], off offset:641011; GFX908-NEXT: global_store_dwordx4 v[0:1], v[14:17], off offset:801012; GFX908-NEXT: global_store_dwordx4 v[0:1], v[18:21], off offset:321013; GFX908-NEXT: global_store_dwordx4 v[0:1], v[22:25], off offset:481014; GFX908-NEXT: global_store_dwordx4 v[0:1], v[26:29], off1015; GFX908-NEXT: global_store_dwordx4 v[0:1], v[2:5], off offset:161016; GFX908-NEXT: s_waitcnt vmcnt(0)1017; GFX908-NEXT: s_setpc_b64 s[30:31]1018bb:1019 %in.1 = load <32 x float>, ptr addrspace(1) %arg1020 %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)1021 store <32 x float> %mai.1, ptr addrspace(1) %arg1022 ret void1023}1024 1025define void @test_mfma_f32_32x32x1f32_nonentry_with_agpr(ptr addrspace(1) %arg) #3 {1026; GFX908-LABEL: test_mfma_f32_32x32x1f32_nonentry_with_agpr:1027; GFX908: ; %bb.0: ; %bb1028; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1029; GFX908-NEXT: global_load_dwordx4 v[30:33], v[0:1], off offset:1121030; GFX908-NEXT: global_load_dwordx4 v[26:29], v[0:1], off offset:961031; GFX908-NEXT: global_load_dwordx4 v[22:25], v[0:1], off offset:801032; GFX908-NEXT: global_load_dwordx4 v[18:21], v[0:1], off offset:641033; GFX908-NEXT: global_load_dwordx4 v[14:17], v[0:1], off offset:481034; GFX908-NEXT: global_load_dwordx4 v[10:13], v[0:1], off offset:321035; GFX908-NEXT: global_load_dwordx4 v[6:9], v[0:1], off offset:161036; GFX908-NEXT: global_load_dwordx4 v[2:5], v[0:1], off1037; GFX908-NEXT: s_waitcnt vmcnt(0)1038; GFX908-NEXT: v_accvgpr_write_b32 a0, v21039; GFX908-NEXT: v_accvgpr_write_b32 a1, v31040; GFX908-NEXT: v_accvgpr_write_b32 a2, v41041; GFX908-NEXT: v_accvgpr_write_b32 a3, v51042; GFX908-NEXT: v_accvgpr_write_b32 a4, v61043; GFX908-NEXT: v_accvgpr_write_b32 a5, v71044; GFX908-NEXT: v_accvgpr_write_b32 a6, v81045; GFX908-NEXT: v_accvgpr_write_b32 a7, v91046; GFX908-NEXT: v_accvgpr_write_b32 a8, v101047; GFX908-NEXT: v_accvgpr_write_b32 a9, v111048; GFX908-NEXT: v_accvgpr_write_b32 a10, v121049; GFX908-NEXT: v_accvgpr_write_b32 a11, v131050; GFX908-NEXT: v_accvgpr_write_b32 a12, v141051; GFX908-NEXT: v_accvgpr_write_b32 a13, v151052; GFX908-NEXT: v_accvgpr_write_b32 a14, v161053; GFX908-NEXT: v_accvgpr_write_b32 a15, v171054; GFX908-NEXT: v_accvgpr_write_b32 a16, v181055; GFX908-NEXT: v_accvgpr_write_b32 a17, v191056; GFX908-NEXT: v_accvgpr_write_b32 a18, v201057; GFX908-NEXT: v_accvgpr_write_b32 a19, v211058; GFX908-NEXT: v_accvgpr_write_b32 a20, v221059; GFX908-NEXT: v_accvgpr_write_b32 a21, v231060; GFX908-NEXT: v_accvgpr_write_b32 a22, v241061; GFX908-NEXT: v_accvgpr_write_b32 a23, v251062; GFX908-NEXT: v_accvgpr_write_b32 a24, v261063; GFX908-NEXT: v_accvgpr_write_b32 a25, v271064; GFX908-NEXT: v_accvgpr_write_b32 a26, v281065; GFX908-NEXT: v_accvgpr_write_b32 a27, v291066; GFX908-NEXT: v_accvgpr_write_b32 a28, v301067; GFX908-NEXT: v_accvgpr_write_b32 a29, v311068; GFX908-NEXT: v_accvgpr_write_b32 a30, v321069; GFX908-NEXT: v_accvgpr_write_b32 a31, v331070; GFX908-NEXT: v_mov_b32_e32 v2, 1.01071; GFX908-NEXT: v_mov_b32_e32 v3, 2.01072; GFX908-NEXT: s_nop 11073; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v3, a[0:31]1074; GFX908-NEXT: s_nop 151075; GFX908-NEXT: s_nop 11076; GFX908-NEXT: v_accvgpr_read_b32 v5, a271077; GFX908-NEXT: v_accvgpr_read_b32 v4, a261078; GFX908-NEXT: v_accvgpr_read_b32 v3, a251079; GFX908-NEXT: v_accvgpr_read_b32 v2, a241080; GFX908-NEXT: v_accvgpr_read_b32 v9, a311081; GFX908-NEXT: v_accvgpr_read_b32 v8, a301082; GFX908-NEXT: v_accvgpr_read_b32 v7, a291083; GFX908-NEXT: v_accvgpr_read_b32 v6, a281084; GFX908-NEXT: v_accvgpr_read_b32 v13, a191085; GFX908-NEXT: v_accvgpr_read_b32 v12, a181086; GFX908-NEXT: v_accvgpr_read_b32 v11, a171087; GFX908-NEXT: v_accvgpr_read_b32 v10, a161088; GFX908-NEXT: v_accvgpr_read_b32 v17, a231089; GFX908-NEXT: v_accvgpr_read_b32 v16, a221090; GFX908-NEXT: v_accvgpr_read_b32 v15, a211091; GFX908-NEXT: v_accvgpr_read_b32 v14, a201092; GFX908-NEXT: v_accvgpr_read_b32 v21, a111093; GFX908-NEXT: v_accvgpr_read_b32 v20, a101094; GFX908-NEXT: v_accvgpr_read_b32 v19, a91095; GFX908-NEXT: v_accvgpr_read_b32 v18, a81096; GFX908-NEXT: v_accvgpr_read_b32 v25, a151097; GFX908-NEXT: v_accvgpr_read_b32 v24, a141098; GFX908-NEXT: v_accvgpr_read_b32 v23, a131099; GFX908-NEXT: v_accvgpr_read_b32 v22, a121100; GFX908-NEXT: v_accvgpr_read_b32 v29, a31101; GFX908-NEXT: v_accvgpr_read_b32 v28, a21102; GFX908-NEXT: v_accvgpr_read_b32 v27, a11103; GFX908-NEXT: v_accvgpr_read_b32 v26, a01104; GFX908-NEXT: global_store_dwordx4 v[0:1], v[2:5], off offset:961105; GFX908-NEXT: s_nop 01106; GFX908-NEXT: v_accvgpr_read_b32 v5, a71107; GFX908-NEXT: v_accvgpr_read_b32 v4, a61108; GFX908-NEXT: v_accvgpr_read_b32 v3, a51109; GFX908-NEXT: v_accvgpr_read_b32 v2, a41110; GFX908-NEXT: global_store_dwordx4 v[0:1], v[6:9], off offset:1121111; GFX908-NEXT: global_store_dwordx4 v[0:1], v[10:13], off offset:641112; GFX908-NEXT: global_store_dwordx4 v[0:1], v[14:17], off offset:801113; GFX908-NEXT: global_store_dwordx4 v[0:1], v[18:21], off offset:321114; GFX908-NEXT: global_store_dwordx4 v[0:1], v[22:25], off offset:481115; GFX908-NEXT: global_store_dwordx4 v[0:1], v[26:29], off1116; GFX908-NEXT: global_store_dwordx4 v[0:1], v[2:5], off offset:161117; GFX908-NEXT: s_waitcnt vmcnt(0)1118; GFX908-NEXT: s_setpc_b64 s[30:31]1119bb:1120 %in.1 = load <32 x float>, ptr addrspace(1) %arg1121 %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)1122 store <32 x float> %mai.1, ptr addrspace(1) %arg1123 ret void1124}1125 1126declare void @foo()1127 1128attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-waves-per-eu"="2" "amdgpu-agpr-alloc"="0" }1129attributes #1 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-waves-per-eu"="2" }1130attributes #2 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="0" }1131attributes #3 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-waves-per-eu"="2" }1132;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:1133; GCN: {{.*}}1134; GFX90A: {{.*}}1135