278 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 42; RUN: llc -mtriple=amdgcn -mcpu=gfx950 --amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefixes=HEURRC %s3; RUN: llc -mtriple=amdgcn -mcpu=gfx950 --amdgpu-mfma-vgpr-form=1 < %s | FileCheck -enable-var-scope --check-prefixes=VGPRRC %s4 5declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half>, <8 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)6 7define <4 x float> @default(<8 x half> %arg0, <8 x half> %arg1, <4 x float> %arg2) {8; HEURRC-LABEL: default:9; HEURRC: ; %bb.0:10; HEURRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)11; HEURRC-NEXT: v_accvgpr_write_b32 a0, v812; HEURRC-NEXT: v_accvgpr_write_b32 a1, v913; HEURRC-NEXT: v_accvgpr_write_b32 a2, v1014; HEURRC-NEXT: v_accvgpr_write_b32 a3, v1115; HEURRC-NEXT: s_nop 116; HEURRC-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], v[4:7], a[0:3]17; HEURRC-NEXT: s_nop 718; HEURRC-NEXT: v_accvgpr_read_b32 v0, a019; HEURRC-NEXT: v_accvgpr_read_b32 v1, a120; HEURRC-NEXT: v_accvgpr_read_b32 v2, a221; HEURRC-NEXT: v_accvgpr_read_b32 v3, a322; HEURRC-NEXT: s_setpc_b64 s[30:31]23;24; VGPRRC-LABEL: default:25; VGPRRC: ; %bb.0:26; VGPRRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)27; VGPRRC-NEXT: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[4:7], v[8:11]28; VGPRRC-NEXT: s_setpc_b64 s[30:31]29 %result = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %arg0, <8 x half> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0)30 ret <4 x float> %result31}32 33define <4 x float> @request_agpr(<8 x half> %arg0, <8 x half> %arg1, <4 x float> %arg2) #0 {34; HEURRC-LABEL: request_agpr:35; HEURRC: ; %bb.0:36; HEURRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)37; HEURRC-NEXT: v_accvgpr_write_b32 a0, v838; HEURRC-NEXT: v_accvgpr_write_b32 a1, v939; HEURRC-NEXT: v_accvgpr_write_b32 a2, v1040; HEURRC-NEXT: v_accvgpr_write_b32 a3, v1141; HEURRC-NEXT: s_nop 142; HEURRC-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], v[0:3], v[4:7], a[0:3]43; HEURRC-NEXT: s_nop 744; HEURRC-NEXT: v_accvgpr_read_b32 v0, a045; HEURRC-NEXT: v_accvgpr_read_b32 v1, a146; HEURRC-NEXT: v_accvgpr_read_b32 v2, a247; HEURRC-NEXT: v_accvgpr_read_b32 v3, a348; HEURRC-NEXT: s_setpc_b64 s[30:31]49;50; VGPRRC-LABEL: request_agpr:51; VGPRRC: ; %bb.0:52; VGPRRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)53; VGPRRC-NEXT: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[4:7], v[8:11]54; VGPRRC-NEXT: s_setpc_b64 s[30:31]55 %result = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %arg0, <8 x half> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0)56 ret <4 x float> %result57}58 59define <4 x float> @request_no_agpr(<8 x half> %arg0, <8 x half> %arg1, <4 x float> %arg2) #1 {60; HEURRC-LABEL: request_no_agpr:61; HEURRC: ; %bb.0:62; HEURRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)63; HEURRC-NEXT: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[4:7], v[8:11]64; HEURRC-NEXT: s_setpc_b64 s[30:31]65;66; VGPRRC-LABEL: request_no_agpr:67; VGPRRC: ; %bb.0:68; VGPRRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)69; VGPRRC-NEXT: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[4:7], v[8:11]70; VGPRRC-NEXT: s_setpc_b64 s[30:31]71 %result = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %arg0, <8 x half> %arg1, <4 x float> %arg2, i32 0, i32 0, i32 0)72 ret <4 x float> %result73}74 75; Make sure this selects the VGPR form, if AGPRs available, but not76; enough.77define amdgpu_kernel void @not_enough_agprs(ptr addrspace(1) %arg) #2 {78; HEURRC-LABEL: not_enough_agprs:79; HEURRC: ; %bb.0: ; %bb80; HEURRC-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x2481; HEURRC-NEXT: v_mov_b32_e32 v33, 1.082; HEURRC-NEXT: v_mov_b32_e32 v34, 2.083; HEURRC-NEXT: v_mov_b32_e32 v32, 084; HEURRC-NEXT: s_waitcnt lgkmcnt(0)85; HEURRC-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x086; HEURRC-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x4087; HEURRC-NEXT: s_waitcnt lgkmcnt(0)88; HEURRC-NEXT: v_mov_b32_e32 v0, s1689; HEURRC-NEXT: v_mov_b32_e32 v1, s1790; HEURRC-NEXT: v_mov_b32_e32 v2, s1891; HEURRC-NEXT: v_mov_b32_e32 v3, s1992; HEURRC-NEXT: v_mov_b32_e32 v4, s2093; HEURRC-NEXT: v_mov_b32_e32 v5, s2194; HEURRC-NEXT: v_mov_b32_e32 v6, s2295; HEURRC-NEXT: v_mov_b32_e32 v7, s2396; HEURRC-NEXT: v_mov_b32_e32 v8, s2497; HEURRC-NEXT: v_mov_b32_e32 v9, s2598; HEURRC-NEXT: v_mov_b32_e32 v10, s2699; HEURRC-NEXT: v_mov_b32_e32 v11, s27100; HEURRC-NEXT: v_mov_b32_e32 v12, s28101; HEURRC-NEXT: v_mov_b32_e32 v13, s29102; HEURRC-NEXT: v_mov_b32_e32 v14, s30103; HEURRC-NEXT: v_mov_b32_e32 v15, s31104; HEURRC-NEXT: v_mov_b32_e32 v16, s0105; HEURRC-NEXT: v_mov_b32_e32 v17, s1106; HEURRC-NEXT: v_mov_b32_e32 v18, s2107; HEURRC-NEXT: v_mov_b32_e32 v19, s3108; HEURRC-NEXT: v_mov_b32_e32 v20, s4109; HEURRC-NEXT: v_mov_b32_e32 v21, s5110; HEURRC-NEXT: v_mov_b32_e32 v22, s6111; HEURRC-NEXT: v_mov_b32_e32 v23, s7112; HEURRC-NEXT: v_mov_b32_e32 v24, s8113; HEURRC-NEXT: v_mov_b32_e32 v25, s9114; HEURRC-NEXT: v_mov_b32_e32 v26, s10115; HEURRC-NEXT: v_mov_b32_e32 v27, s11116; HEURRC-NEXT: v_mov_b32_e32 v28, s12117; HEURRC-NEXT: v_mov_b32_e32 v29, s13118; HEURRC-NEXT: v_mov_b32_e32 v30, s14119; HEURRC-NEXT: v_mov_b32_e32 v31, s15120; HEURRC-NEXT: s_nop 1121; HEURRC-NEXT: v_mfma_f32_32x32x1_2b_f32 v[0:31], v33, v34, v[0:31] cbsz:1 abid:2 blgp:3122; HEURRC-NEXT: s_nop 15123; HEURRC-NEXT: s_nop 1124; HEURRC-NEXT: global_store_dwordx4 v32, v[24:27], s[34:35] offset:96125; HEURRC-NEXT: global_store_dwordx4 v32, v[28:31], s[34:35] offset:112126; HEURRC-NEXT: global_store_dwordx4 v32, v[16:19], s[34:35] offset:64127; HEURRC-NEXT: global_store_dwordx4 v32, v[20:23], s[34:35] offset:80128; HEURRC-NEXT: global_store_dwordx4 v32, v[8:11], s[34:35] offset:32129; HEURRC-NEXT: global_store_dwordx4 v32, v[12:15], s[34:35] offset:48130; HEURRC-NEXT: global_store_dwordx4 v32, v[0:3], s[34:35]131; HEURRC-NEXT: global_store_dwordx4 v32, v[4:7], s[34:35] offset:16132; HEURRC-NEXT: s_endpgm133;134; VGPRRC-LABEL: not_enough_agprs:135; VGPRRC: ; %bb.0: ; %bb136; VGPRRC-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24137; VGPRRC-NEXT: v_mov_b32_e32 v33, 1.0138; VGPRRC-NEXT: v_mov_b32_e32 v34, 2.0139; VGPRRC-NEXT: v_mov_b32_e32 v32, 0140; VGPRRC-NEXT: s_waitcnt lgkmcnt(0)141; VGPRRC-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0142; VGPRRC-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40143; VGPRRC-NEXT: s_waitcnt lgkmcnt(0)144; VGPRRC-NEXT: v_mov_b32_e32 v0, s16145; VGPRRC-NEXT: v_mov_b32_e32 v1, s17146; VGPRRC-NEXT: v_mov_b32_e32 v2, s18147; VGPRRC-NEXT: v_mov_b32_e32 v3, s19148; VGPRRC-NEXT: v_mov_b32_e32 v4, s20149; VGPRRC-NEXT: v_mov_b32_e32 v5, s21150; VGPRRC-NEXT: v_mov_b32_e32 v6, s22151; VGPRRC-NEXT: v_mov_b32_e32 v7, s23152; VGPRRC-NEXT: v_mov_b32_e32 v8, s24153; VGPRRC-NEXT: v_mov_b32_e32 v9, s25154; VGPRRC-NEXT: v_mov_b32_e32 v10, s26155; VGPRRC-NEXT: v_mov_b32_e32 v11, s27156; VGPRRC-NEXT: v_mov_b32_e32 v12, s28157; VGPRRC-NEXT: v_mov_b32_e32 v13, s29158; VGPRRC-NEXT: v_mov_b32_e32 v14, s30159; VGPRRC-NEXT: v_mov_b32_e32 v15, s31160; VGPRRC-NEXT: v_mov_b32_e32 v16, s0161; VGPRRC-NEXT: v_mov_b32_e32 v17, s1162; VGPRRC-NEXT: v_mov_b32_e32 v18, s2163; VGPRRC-NEXT: v_mov_b32_e32 v19, s3164; VGPRRC-NEXT: v_mov_b32_e32 v20, s4165; VGPRRC-NEXT: v_mov_b32_e32 v21, s5166; VGPRRC-NEXT: v_mov_b32_e32 v22, s6167; VGPRRC-NEXT: v_mov_b32_e32 v23, s7168; VGPRRC-NEXT: v_mov_b32_e32 v24, s8169; VGPRRC-NEXT: v_mov_b32_e32 v25, s9170; VGPRRC-NEXT: v_mov_b32_e32 v26, s10171; VGPRRC-NEXT: v_mov_b32_e32 v27, s11172; VGPRRC-NEXT: v_mov_b32_e32 v28, s12173; VGPRRC-NEXT: v_mov_b32_e32 v29, s13174; VGPRRC-NEXT: v_mov_b32_e32 v30, s14175; VGPRRC-NEXT: v_mov_b32_e32 v31, s15176; VGPRRC-NEXT: s_nop 1177; VGPRRC-NEXT: v_mfma_f32_32x32x1_2b_f32 v[0:31], v33, v34, v[0:31] cbsz:1 abid:2 blgp:3178; VGPRRC-NEXT: s_nop 15179; VGPRRC-NEXT: s_nop 1180; VGPRRC-NEXT: global_store_dwordx4 v32, v[24:27], s[34:35] offset:96181; VGPRRC-NEXT: global_store_dwordx4 v32, v[28:31], s[34:35] offset:112182; VGPRRC-NEXT: global_store_dwordx4 v32, v[16:19], s[34:35] offset:64183; VGPRRC-NEXT: global_store_dwordx4 v32, v[20:23], s[34:35] offset:80184; VGPRRC-NEXT: global_store_dwordx4 v32, v[8:11], s[34:35] offset:32185; VGPRRC-NEXT: global_store_dwordx4 v32, v[12:15], s[34:35] offset:48186; VGPRRC-NEXT: global_store_dwordx4 v32, v[0:3], s[34:35]187; VGPRRC-NEXT: global_store_dwordx4 v32, v[4:7], s[34:35] offset:16188; VGPRRC-NEXT: s_endpgm189bb:190 %in.1 = load <32 x float>, ptr addrspace(1) %arg, align 128191 %mai.1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %in.1, i32 1, i32 2, i32 3)192 store <32 x float> %mai.1, ptr addrspace(1) %arg, align 128193 ret void194}195 196define <16 x float> @mfma_scale_respect_flag(<8 x i32> %arg0, <8 x i32> %arg1, <16 x float> %arg2, i32 %scale0, i32 %scale1) #2 {197; HEURRC-LABEL: mfma_scale_respect_flag:198; HEURRC: ; %bb.0:199; HEURRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)200; HEURRC-NEXT: scratch_load_dword a15, off, s32201; HEURRC-NEXT: scratch_load_dword v31, off, s32 offset:8202; HEURRC-NEXT: scratch_load_dword v32, off, s32 offset:4203; HEURRC-NEXT: v_accvgpr_write_b32 a0, v16204; HEURRC-NEXT: v_accvgpr_write_b32 a1, v17205; HEURRC-NEXT: v_accvgpr_write_b32 a2, v18206; HEURRC-NEXT: v_accvgpr_write_b32 a3, v19207; HEURRC-NEXT: v_accvgpr_write_b32 a4, v20208; HEURRC-NEXT: v_accvgpr_write_b32 a5, v21209; HEURRC-NEXT: v_accvgpr_write_b32 a6, v22210; HEURRC-NEXT: v_accvgpr_write_b32 a7, v23211; HEURRC-NEXT: v_accvgpr_write_b32 a8, v24212; HEURRC-NEXT: v_accvgpr_write_b32 a9, v25213; HEURRC-NEXT: v_accvgpr_write_b32 a10, v26214; HEURRC-NEXT: v_accvgpr_write_b32 a11, v27215; HEURRC-NEXT: v_accvgpr_write_b32 a12, v28216; HEURRC-NEXT: v_accvgpr_write_b32 a13, v29217; HEURRC-NEXT: v_accvgpr_write_b32 a14, v30218; HEURRC-NEXT: s_waitcnt vmcnt(0)219; HEURRC-NEXT: s_nop 0220; HEURRC-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15], v32, v31 op_sel_hi:[0,0,0]221; HEURRC-NEXT: s_nop 15222; HEURRC-NEXT: s_nop 3223; HEURRC-NEXT: v_accvgpr_read_b32 v0, a0224; HEURRC-NEXT: v_accvgpr_read_b32 v1, a1225; HEURRC-NEXT: v_accvgpr_read_b32 v2, a2226; HEURRC-NEXT: v_accvgpr_read_b32 v3, a3227; HEURRC-NEXT: v_accvgpr_read_b32 v4, a4228; HEURRC-NEXT: v_accvgpr_read_b32 v5, a5229; HEURRC-NEXT: v_accvgpr_read_b32 v6, a6230; HEURRC-NEXT: v_accvgpr_read_b32 v7, a7231; HEURRC-NEXT: v_accvgpr_read_b32 v8, a8232; HEURRC-NEXT: v_accvgpr_read_b32 v9, a9233; HEURRC-NEXT: v_accvgpr_read_b32 v10, a10234; HEURRC-NEXT: v_accvgpr_read_b32 v11, a11235; HEURRC-NEXT: v_accvgpr_read_b32 v12, a12236; HEURRC-NEXT: v_accvgpr_read_b32 v13, a13237; HEURRC-NEXT: v_accvgpr_read_b32 v14, a14238; HEURRC-NEXT: v_accvgpr_read_b32 v15, a15239; HEURRC-NEXT: s_setpc_b64 s[30:31]240;241; VGPRRC-LABEL: mfma_scale_respect_flag:242; VGPRRC: ; %bb.0:243; VGPRRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)244; VGPRRC-NEXT: scratch_load_dword v31, off, s32245; VGPRRC-NEXT: scratch_load_dword v32, off, s32 offset:8246; VGPRRC-NEXT: scratch_load_dword v33, off, s32 offset:4247; VGPRRC-NEXT: s_waitcnt vmcnt(0)248; VGPRRC-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 v[16:31], v[0:7], v[8:15], v[16:31], v33, v32 op_sel_hi:[0,0,0]249; VGPRRC-NEXT: s_nop 15250; VGPRRC-NEXT: s_nop 3251; VGPRRC-NEXT: v_mov_b32_e32 v0, v16252; VGPRRC-NEXT: v_mov_b32_e32 v1, v17253; VGPRRC-NEXT: v_mov_b32_e32 v2, v18254; VGPRRC-NEXT: v_mov_b32_e32 v3, v19255; VGPRRC-NEXT: v_mov_b32_e32 v4, v20256; VGPRRC-NEXT: v_mov_b32_e32 v5, v21257; VGPRRC-NEXT: v_mov_b32_e32 v6, v22258; VGPRRC-NEXT: v_mov_b32_e32 v7, v23259; VGPRRC-NEXT: v_mov_b32_e32 v8, v24260; VGPRRC-NEXT: v_mov_b32_e32 v9, v25261; VGPRRC-NEXT: v_mov_b32_e32 v10, v26262; VGPRRC-NEXT: v_mov_b32_e32 v11, v27263; VGPRRC-NEXT: v_mov_b32_e32 v12, v28264; VGPRRC-NEXT: v_mov_b32_e32 v13, v29265; VGPRRC-NEXT: v_mov_b32_e32 v14, v30266; VGPRRC-NEXT: v_mov_b32_e32 v15, v31267; VGPRRC-NEXT: s_setpc_b64 s[30:31]268 %result = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <16 x float> %arg2,269 i32 0, ; cbsz270 i32 0, ; blgp271 i32 0, i32 %scale0, i32 0, i32 %scale1)272 ret <16 x float> %result273}274 275attributes #0 = { "amdgpu-agpr-alloc"="32,256" }276attributes #1 = { "amdgpu-agpr-alloc"="0,0" }277attributes #2 = { nounwind "amdgpu-agpr-alloc"="20" }278