494 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 -check-prefixes=GCN,GFX908 %s3; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -check-prefixes=GCN,GFX90A %s4 5define amdgpu_kernel void @max_12regs_13a_used(i32 %cond, ptr addrspace(1) %arg, ptr addrspace(1) %out) #2 {6; GFX908-LABEL: max_12regs_13a_used:7; GFX908: ; %bb.0: ; %bb8; GFX908-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x2c9; GFX908-NEXT: v_mov_b32_e32 v0, 1.010; GFX908-NEXT: s_load_dword s0, s[0:1], 0x2411; GFX908-NEXT: s_waitcnt lgkmcnt(0)12; GFX908-NEXT: s_cmp_lg_u32 s0, 013; GFX908-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x014; GFX908-NEXT: s_waitcnt lgkmcnt(0)15; GFX908-NEXT: v_mov_b32_e32 v5, s816; GFX908-NEXT: v_mov_b32_e32 v1, s917; GFX908-NEXT: v_mov_b32_e32 v2, s1018; GFX908-NEXT: v_accvgpr_write_b32 a0, v519; GFX908-NEXT: v_mov_b32_e32 v5, s1120; GFX908-NEXT: v_accvgpr_write_b32 a1, v121; GFX908-NEXT: v_accvgpr_write_b32 a2, v222; GFX908-NEXT: v_accvgpr_write_b32 a3, v523; GFX908-NEXT: s_nop 024; GFX908-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v0, a[0:3]25; GFX908-NEXT: v_mfma_f32_4x4x1f32 a[4:7], v0, v0, a[0:3]26; GFX908-NEXT: s_cbranch_scc0 .LBB0_227; GFX908-NEXT: ; %bb.1: ; %st28; GFX908-NEXT: ;;#ASMSTART29; GFX908-NEXT: ;;#ASMEND30; GFX908-NEXT: s_endpgm31; GFX908-NEXT: .LBB0_2: ; %use32; GFX908-NEXT: s_nop 233; GFX908-NEXT: v_accvgpr_read_b32 v1, a434; GFX908-NEXT: v_accvgpr_read_b32 v4, a735; GFX908-NEXT: v_accvgpr_read_b32 v2, a536; GFX908-NEXT: v_accvgpr_read_b32 v3, a637; GFX908-NEXT: v_accvgpr_write_b32 a4, 438; GFX908-NEXT: v_accvgpr_write_b32 a8, 539; GFX908-NEXT: v_accvgpr_write_b32 a9, 140; GFX908-NEXT: v_accvgpr_write_b32 a10, 241; GFX908-NEXT: v_accvgpr_write_b32 a11, 342; GFX908-NEXT: ;;#ASMSTART43; GFX908-NEXT: ;;#ASMEND44; GFX908-NEXT: v_accvgpr_write_b32 a7, v445; GFX908-NEXT: v_accvgpr_write_b32 a6, v346; GFX908-NEXT: v_accvgpr_write_b32 a5, v247; GFX908-NEXT: v_accvgpr_write_b32 a4, v148; GFX908-NEXT: v_mov_b32_e32 v1, v049; GFX908-NEXT: v_mov_b32_e32 v2, v050; GFX908-NEXT: v_mov_b32_e32 v3, v051; GFX908-NEXT: v_mov_b32_e32 v4, 052; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]53; GFX908-NEXT: s_waitcnt vmcnt(0)54; GFX908-NEXT: ;;#ASMSTART55; GFX908-NEXT: ;;#ASMEND56; GFX908-NEXT: s_endpgm57;58; GFX90A-LABEL: max_12regs_13a_used:59; GFX90A: ; %bb.0: ; %bb60; GFX90A-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x2c61; GFX90A-NEXT: v_mov_b32_e32 v0, 1.062; GFX90A-NEXT: s_load_dword s0, s[0:1], 0x2463; GFX90A-NEXT: s_waitcnt lgkmcnt(0)64; GFX90A-NEXT: s_cmp_lg_u32 s0, 065; GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x066; GFX90A-NEXT: s_waitcnt lgkmcnt(0)67; GFX90A-NEXT: v_accvgpr_write_b32 a0, s868; GFX90A-NEXT: v_accvgpr_write_b32 a1, s969; GFX90A-NEXT: v_accvgpr_write_b32 a2, s1070; GFX90A-NEXT: v_accvgpr_write_b32 a3, s1171; GFX90A-NEXT: s_nop 172; GFX90A-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v0, a[0:3]73; GFX90A-NEXT: v_mfma_f32_4x4x1f32 a[4:7], v0, v0, a[0:3]74; GFX90A-NEXT: s_cbranch_scc0 .LBB0_275; GFX90A-NEXT: ; %bb.1: ; %st76; GFX90A-NEXT: ;;#ASMSTART77; GFX90A-NEXT: ;;#ASMEND78; GFX90A-NEXT: s_endpgm79; GFX90A-NEXT: .LBB0_2: ; %use80; GFX90A-NEXT: s_nop 381; GFX90A-NEXT: v_accvgpr_read_b32 v9, a782; GFX90A-NEXT: v_accvgpr_read_b32 v8, a683; GFX90A-NEXT: v_accvgpr_read_b32 v7, a584; GFX90A-NEXT: v_accvgpr_read_b32 v6, a485; GFX90A-NEXT: v_accvgpr_write_b32 a4, 486; GFX90A-NEXT: v_accvgpr_write_b32 a8, 587; GFX90A-NEXT: v_accvgpr_write_b32 a9, 188; GFX90A-NEXT: v_accvgpr_write_b32 a10, 289; GFX90A-NEXT: v_accvgpr_write_b32 a11, 390; GFX90A-NEXT: ;;#ASMSTART91; GFX90A-NEXT: ;;#ASMEND92; GFX90A-NEXT: v_accvgpr_write_b32 a4, v693; GFX90A-NEXT: v_mov_b32_e32 v4, 094; GFX90A-NEXT: v_accvgpr_write_b32 a5, v795; GFX90A-NEXT: v_accvgpr_write_b32 a6, v896; GFX90A-NEXT: v_accvgpr_write_b32 a7, v997; GFX90A-NEXT: v_mov_b32_e32 v1, v098; GFX90A-NEXT: v_mov_b32_e32 v2, v099; GFX90A-NEXT: v_mov_b32_e32 v3, v0100; GFX90A-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]101; GFX90A-NEXT: s_waitcnt vmcnt(0)102; GFX90A-NEXT: ;;#ASMSTART103; GFX90A-NEXT: ;;#ASMEND104; GFX90A-NEXT: s_endpgm105bb:106 %in.1 = load <4 x float>, ptr addrspace(1) %arg107 %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %in.1, i32 0, i32 0, i32 0)108 %mai.2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai.1, i32 0, i32 0, i32 0)109 %cmp = icmp eq i32 %cond, 0110 br i1 %cmp, label %use, label %st111 112use:113 call void asm sideeffect "", "a,a,a,a,a"(i32 1, i32 2, i32 3, i32 4, i32 5)114 store volatile <4 x float> <float 1.0, float 1.0, float 1.0, float 1.0>, ptr addrspace(1) %out115 br label %st116 117st:118 %gep1 = getelementptr <4 x float>, ptr addrspace(1) %out, i64 16119 %gep2 = getelementptr <4 x float>, ptr addrspace(1) %out, i64 32120 call void asm sideeffect "", "a,a"(<4 x float> %mai.1, <4 x float> %mai.2)121 ret void122}123; GCN: ScratchSize: 0124 125define amdgpu_kernel void @max_10_vgprs_used_9a() #1 {126; GFX908-LABEL: max_10_vgprs_used_9a:127; GFX908: ; %bb.0:128; GFX908-NEXT: ;;#ASMSTART129; GFX908-NEXT: ;;#ASMEND130; GFX908-NEXT: v_accvgpr_read_b32 v0, a0131; GFX908-NEXT: v_accvgpr_read_b32 v3, a3132; GFX908-NEXT: v_accvgpr_read_b32 v1, a1133; GFX908-NEXT: v_accvgpr_read_b32 v2, a2134; GFX908-NEXT: ;;#ASMSTART135; GFX908-NEXT: ;;#ASMEND136; GFX908-NEXT: ;;#ASMSTART137; GFX908-NEXT: ;;#ASMEND138; GFX908-NEXT: ;;#ASMSTART139; GFX908-NEXT: ;;#ASMEND140; GFX908-NEXT: v_accvgpr_read_b32 v5, a2141; GFX908-NEXT: v_accvgpr_read_b32 v4, a1142; GFX908-NEXT: v_accvgpr_write_b32 a4, v3143; GFX908-NEXT: v_accvgpr_write_b32 a1, v0144; GFX908-NEXT: v_accvgpr_write_b32 a3, v2145; GFX908-NEXT: v_accvgpr_write_b32 a2, v1146; GFX908-NEXT: ;;#ASMSTART147; GFX908-NEXT: ;;#ASMEND148; GFX908-NEXT: v_accvgpr_write_b32 a0, v4149; GFX908-NEXT: v_accvgpr_write_b32 a1, v5150; GFX908-NEXT: ;;#ASMSTART151; GFX908-NEXT: ;;#ASMEND152; GFX908-NEXT: s_endpgm153;154; GFX90A-LABEL: max_10_vgprs_used_9a:155; GFX90A: ; %bb.0:156; GFX90A-NEXT: ;;#ASMSTART157; GFX90A-NEXT: ;;#ASMEND158; GFX90A-NEXT: v_accvgpr_read_b32 v0, a0159; GFX90A-NEXT: v_accvgpr_read_b32 v1, a1160; GFX90A-NEXT: v_accvgpr_read_b32 v2, a2161; GFX90A-NEXT: v_accvgpr_read_b32 v3, a3162; GFX90A-NEXT: ;;#ASMSTART163; GFX90A-NEXT: ;;#ASMEND164; GFX90A-NEXT: ;;#ASMSTART165; GFX90A-NEXT: ;;#ASMEND166; GFX90A-NEXT: ;;#ASMSTART167; GFX90A-NEXT: ;;#ASMEND168; GFX90A-NEXT: v_accvgpr_read_b32 v5, a3169; GFX90A-NEXT: v_accvgpr_read_b32 v4, a2170; GFX90A-NEXT: v_accvgpr_write_b32 a5, v3171; GFX90A-NEXT: v_accvgpr_write_b32 a4, v2172; GFX90A-NEXT: v_accvgpr_write_b32 a3, v1173; GFX90A-NEXT: v_accvgpr_write_b32 a2, v0174; GFX90A-NEXT: ;;#ASMSTART175; GFX90A-NEXT: ;;#ASMEND176; GFX90A-NEXT: v_accvgpr_write_b32 a0, v4177; GFX90A-NEXT: v_accvgpr_write_b32 a1, v5178; GFX90A-NEXT: ;;#ASMSTART179; GFX90A-NEXT: ;;#ASMEND180; GFX90A-NEXT: s_endpgm181 %a1 = call <4 x i32> asm sideeffect "", "=a"()182 %a2 = call <4 x i32> asm sideeffect "", "=a"()183 %a3 = call i32 asm sideeffect "", "=a"()184 %a4 = call <2 x i32> asm sideeffect "", "=a"()185 call void asm sideeffect "", "a,a,a"(<4 x i32> %a1, <4 x i32> %a2, i32 %a3)186 call void asm sideeffect "", "a"(<2 x i32> %a4)187 ret void188}189; GCN: ScratchSize: 0190 191define amdgpu_kernel void @max_32regs_mfma32(ptr addrspace(1) %arg) #3 {192; GFX908-LABEL: max_32regs_mfma32:193; GFX908: ; %bb.0: ; %bb194; GFX908-NEXT: v_mov_b32_e32 v1, 0x40400000195; GFX908-NEXT: v_mov_b32_e32 v2, 0x40c00000196; GFX908-NEXT: v_mov_b32_e32 v3, 0x40e00000197; GFX908-NEXT: v_accvgpr_write_b32 a2, v1198; GFX908-NEXT: v_mov_b32_e32 v1, 0x40a00000199; GFX908-NEXT: v_accvgpr_write_b32 a5, v2200; GFX908-NEXT: v_accvgpr_write_b32 a6, v3201; GFX908-NEXT: v_accvgpr_write_b32 a4, v1202; GFX908-NEXT: v_mov_b32_e32 v1, 0x41000000203; GFX908-NEXT: v_mov_b32_e32 v2, 0x41100000204; GFX908-NEXT: v_mov_b32_e32 v3, 0x41200000205; GFX908-NEXT: v_accvgpr_write_b32 a7, v1206; GFX908-NEXT: v_accvgpr_write_b32 a8, v2207; GFX908-NEXT: v_accvgpr_write_b32 a9, v3208; GFX908-NEXT: v_mov_b32_e32 v1, 0x41300000209; GFX908-NEXT: v_mov_b32_e32 v2, 0x41400000210; GFX908-NEXT: v_mov_b32_e32 v3, 0x41500000211; GFX908-NEXT: v_accvgpr_write_b32 a10, v1212; GFX908-NEXT: v_accvgpr_write_b32 a11, v2213; GFX908-NEXT: v_accvgpr_write_b32 a12, v3214; GFX908-NEXT: v_mov_b32_e32 v1, 0x41600000215; GFX908-NEXT: v_mov_b32_e32 v2, 0x41700000216; GFX908-NEXT: v_mov_b32_e32 v3, 0x41800000217; GFX908-NEXT: v_accvgpr_write_b32 a13, v1218; GFX908-NEXT: v_accvgpr_write_b32 a14, v2219; GFX908-NEXT: v_accvgpr_write_b32 a15, v3220; GFX908-NEXT: v_mov_b32_e32 v1, 0x41880000221; GFX908-NEXT: v_mov_b32_e32 v2, 0x41900000222; GFX908-NEXT: v_mov_b32_e32 v3, 0x41980000223; GFX908-NEXT: v_accvgpr_write_b32 a16, v1224; GFX908-NEXT: v_accvgpr_write_b32 a17, v2225; GFX908-NEXT: v_accvgpr_write_b32 a18, v3226; GFX908-NEXT: v_mov_b32_e32 v1, 0x41a00000227; GFX908-NEXT: v_mov_b32_e32 v2, 0x41a80000228; GFX908-NEXT: v_mov_b32_e32 v3, 0x41b00000229; GFX908-NEXT: v_accvgpr_write_b32 a19, v1230; GFX908-NEXT: v_accvgpr_write_b32 a20, v2231; GFX908-NEXT: v_accvgpr_write_b32 a21, v3232; GFX908-NEXT: v_mov_b32_e32 v1, 0x41b80000233; GFX908-NEXT: v_mov_b32_e32 v2, 0x41c00000234; GFX908-NEXT: v_mov_b32_e32 v3, 0x41c80000235; GFX908-NEXT: v_accvgpr_write_b32 a22, v1236; GFX908-NEXT: v_accvgpr_write_b32 a23, v2237; GFX908-NEXT: v_accvgpr_write_b32 a24, v3238; GFX908-NEXT: v_mov_b32_e32 v1, 0x41d00000239; GFX908-NEXT: v_mov_b32_e32 v2, 0x41d80000240; GFX908-NEXT: v_mov_b32_e32 v3, 0x41e00000241; GFX908-NEXT: v_accvgpr_write_b32 a25, v1242; GFX908-NEXT: v_accvgpr_write_b32 a26, v2243; GFX908-NEXT: v_accvgpr_write_b32 a27, v3244; GFX908-NEXT: v_mov_b32_e32 v1, 0x41e80000245; GFX908-NEXT: v_mov_b32_e32 v2, 0x41f00000246; GFX908-NEXT: v_mov_b32_e32 v3, 0x41f80000247; GFX908-NEXT: v_mov_b32_e32 v4, 1.0248; GFX908-NEXT: ;;#ASMSTART249; GFX908-NEXT: ;;#ASMEND250; GFX908-NEXT: v_accvgpr_read_b32 v5, a0251; GFX908-NEXT: v_accvgpr_write_b32 a0, 1.0252; GFX908-NEXT: v_accvgpr_write_b32 a1, 2.0253; GFX908-NEXT: v_accvgpr_write_b32 a3, 4.0254; GFX908-NEXT: v_accvgpr_write_b32 a28, v1255; GFX908-NEXT: v_accvgpr_write_b32 a29, v2256; GFX908-NEXT: v_accvgpr_write_b32 a30, v3257; GFX908-NEXT: v_accvgpr_write_b32 a31, 2.0258; GFX908-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24259; GFX908-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v4, v4, a[0:31]260; GFX908-NEXT: v_mov_b32_e32 v0, 0261; GFX908-NEXT: s_nop 13262; GFX908-NEXT: v_accvgpr_write_b32 a1, v5263; GFX908-NEXT: ;;#ASMSTART264; GFX908-NEXT: ;;#ASMEND265; GFX908-NEXT: s_nop 1266; GFX908-NEXT: v_accvgpr_read_b32 v1, a0267; GFX908-NEXT: s_waitcnt lgkmcnt(0)268; GFX908-NEXT: s_nop 0269; GFX908-NEXT: global_store_dword v0, v1, s[2:3]270; GFX908-NEXT: s_endpgm271;272; GFX90A-LABEL: max_32regs_mfma32:273; GFX90A: ; %bb.0: ; %bb274; GFX90A-NEXT: v_mov_b32_e32 v1, 0x40400000275; GFX90A-NEXT: v_accvgpr_write_b32 a2, v1276; GFX90A-NEXT: v_mov_b32_e32 v1, 0x40a00000277; GFX90A-NEXT: v_accvgpr_write_b32 a4, v1278; GFX90A-NEXT: v_mov_b32_e32 v1, 0x40c00000279; GFX90A-NEXT: v_accvgpr_write_b32 a5, v1280; GFX90A-NEXT: v_mov_b32_e32 v1, 0x40e00000281; GFX90A-NEXT: v_accvgpr_write_b32 a6, v1282; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41000000283; GFX90A-NEXT: v_accvgpr_write_b32 a7, v1284; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41100000285; GFX90A-NEXT: v_accvgpr_write_b32 a8, v1286; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41200000287; GFX90A-NEXT: v_accvgpr_write_b32 a9, v1288; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41300000289; GFX90A-NEXT: v_accvgpr_write_b32 a10, v1290; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41400000291; GFX90A-NEXT: v_accvgpr_write_b32 a11, v1292; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41500000293; GFX90A-NEXT: v_accvgpr_write_b32 a12, v1294; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41600000295; GFX90A-NEXT: v_accvgpr_write_b32 a13, v1296; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41700000297; GFX90A-NEXT: v_accvgpr_write_b32 a14, v1298; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41800000299; GFX90A-NEXT: v_accvgpr_write_b32 a15, v1300; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41880000301; GFX90A-NEXT: v_accvgpr_write_b32 a16, v1302; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41900000303; GFX90A-NEXT: v_accvgpr_write_b32 a17, v1304; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41980000305; GFX90A-NEXT: v_accvgpr_write_b32 a18, v1306; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41a00000307; GFX90A-NEXT: v_accvgpr_write_b32 a19, v1308; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41a80000309; GFX90A-NEXT: v_accvgpr_write_b32 a20, v1310; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41b00000311; GFX90A-NEXT: v_accvgpr_write_b32 a21, v1312; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41b80000313; GFX90A-NEXT: v_accvgpr_write_b32 a22, v1314; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41c00000315; GFX90A-NEXT: v_accvgpr_write_b32 a23, v1316; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41c80000317; GFX90A-NEXT: v_accvgpr_write_b32 a24, v1318; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41d00000319; GFX90A-NEXT: v_accvgpr_write_b32 a25, v1320; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41d80000321; GFX90A-NEXT: v_accvgpr_write_b32 a26, v1322; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41e00000323; GFX90A-NEXT: v_mov_b32_e32 v2, 1.0324; GFX90A-NEXT: v_accvgpr_write_b32 a27, v1325; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41e80000326; GFX90A-NEXT: v_accvgpr_write_b32 a28, v1327; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41f00000328; GFX90A-NEXT: ;;#ASMSTART329; GFX90A-NEXT: ;;#ASMEND330; GFX90A-NEXT: v_accvgpr_write_b32 a1, 2.0331; GFX90A-NEXT: v_accvgpr_write_b32 a29, v1332; GFX90A-NEXT: v_mov_b32_e32 v1, 0x41f80000333; GFX90A-NEXT: v_accvgpr_read_b32 v3, a0334; GFX90A-NEXT: v_accvgpr_write_b32 a0, 1.0335; GFX90A-NEXT: v_accvgpr_write_b32 a3, 4.0336; GFX90A-NEXT: v_accvgpr_write_b32 a30, v1337; GFX90A-NEXT: v_accvgpr_mov_b32 a31, a1338; GFX90A-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24339; GFX90A-NEXT: v_mov_b32_e32 v0, 0340; GFX90A-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v2, v2, a[0:31]341; GFX90A-NEXT: s_nop 15342; GFX90A-NEXT: s_nop 2343; GFX90A-NEXT: v_accvgpr_write_b32 a1, v3344; GFX90A-NEXT: ;;#ASMSTART345; GFX90A-NEXT: ;;#ASMEND346; GFX90A-NEXT: s_waitcnt lgkmcnt(0)347; GFX90A-NEXT: global_store_dword v0, a0, s[2:3]348; GFX90A-NEXT: s_endpgm349bb:350 %v = call i32 asm sideeffect "", "=a"()351 br label %use352 353use:354 %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 1.0, <32 x float> <float 1.0, float 2.0, float 3.0, float 4.0, float 5.0, float 6.0, float 7.0, float 8.0, float 9.0, float 10.0, float 11.0, float 12.0, float 13.0, float 14.0, float 15.0, float 16.0, float 17.0, float 18.0, float 19.0, float 20.0, float 21.0, float 22.0, float 23.0, float 24.0, float 25.0, float 26.0, float 27.0, float 28.0, float 29.0, float 30.0, float 31.0, float 2.0>, i32 0, i32 0, i32 0)355 call void asm sideeffect "", "a"(i32 %v)356 %elt1 = extractelement <32 x float> %mai.1, i32 0357 store float %elt1, ptr addrspace(1) %arg358 ret void359}360; GCN: ScratchSize: 0361 362; Should spill agprs to memory for both gfx908 and gfx90a.363define amdgpu_kernel void @max_6regs_used_8a(ptr addrspace(1) %arg) #4 {364; GFX908-LABEL: max_6regs_used_8a:365; GFX908: ; %bb.0:366; GFX908-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24367; GFX908-NEXT: ;;#ASMSTART368; GFX908-NEXT: ; def v1369; GFX908-NEXT: ;;#ASMEND370; GFX908-NEXT: v_lshlrev_b32_e32 v4, 4, v0371; GFX908-NEXT: ;;#ASMSTART372; GFX908-NEXT: ; def a[0:3]373; GFX908-NEXT: ;;#ASMEND374; GFX908-NEXT: s_mov_b32 s8, SCRATCH_RSRC_DWORD0375; GFX908-NEXT: v_accvgpr_write_b32 a4, v1376; GFX908-NEXT: s_waitcnt lgkmcnt(0)377; GFX908-NEXT: global_load_dwordx4 v[0:3], v4, s[2:3]378; GFX908-NEXT: s_mov_b32 s9, SCRATCH_RSRC_DWORD1379; GFX908-NEXT: s_mov_b32 s10, -1380; GFX908-NEXT: s_mov_b32 s11, 0xe00000381; GFX908-NEXT: s_add_u32 s8, s8, s5382; GFX908-NEXT: s_addc_u32 s9, s9, 0383; GFX908-NEXT: v_accvgpr_read_b32 v5, a0 ; Reload Reuse384; GFX908-NEXT: s_nop 1385; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 ; 4-byte Folded Spill386; GFX908-NEXT: v_accvgpr_read_b32 v5, a1 ; Reload Reuse387; GFX908-NEXT: s_nop 1388; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 offset:4 ; 4-byte Folded Spill389; GFX908-NEXT: v_accvgpr_read_b32 v5, a2 ; Reload Reuse390; GFX908-NEXT: s_nop 1391; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 offset:8 ; 4-byte Folded Spill392; GFX908-NEXT: v_accvgpr_read_b32 v5, a3 ; Reload Reuse393; GFX908-NEXT: s_nop 1394; GFX908-NEXT: buffer_store_dword v5, off, s[8:11], 0 offset:12 ; 4-byte Folded Spill395; GFX908-NEXT: s_waitcnt vmcnt(4)396; GFX908-NEXT: v_accvgpr_write_b32 a0, v0397; GFX908-NEXT: v_accvgpr_write_b32 a1, v1398; GFX908-NEXT: v_accvgpr_write_b32 a2, v2399; GFX908-NEXT: v_accvgpr_write_b32 a3, v3400; GFX908-NEXT: v_mov_b32_e32 v0, 1.0401; GFX908-NEXT: s_nop 1402; GFX908-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v0, a[0:3]403; GFX908-NEXT: s_nop 3404; GFX908-NEXT: v_accvgpr_read_b32 v0, a0405; GFX908-NEXT: v_accvgpr_read_b32 v1, a1406; GFX908-NEXT: v_accvgpr_read_b32 v2, a2407; GFX908-NEXT: v_accvgpr_read_b32 v3, a3408; GFX908-NEXT: s_nop 1409; GFX908-NEXT: global_store_dwordx4 v4, v[0:3], s[2:3]410; GFX908-NEXT: buffer_load_dword v0, off, s[8:11], 0 ; 4-byte Folded Reload411; GFX908-NEXT: s_nop 0412; GFX908-NEXT: buffer_load_dword v1, off, s[8:11], 0 offset:4 ; 4-byte Folded Reload413; GFX908-NEXT: buffer_load_dword v2, off, s[8:11], 0 offset:8 ; 4-byte Folded Reload414; GFX908-NEXT: buffer_load_dword v3, off, s[8:11], 0 offset:12 ; 4-byte Folded Reload415; GFX908-NEXT: s_waitcnt vmcnt(0)416; GFX908-NEXT: global_store_dwordx4 v[0:1], v[0:3], off417; GFX908-NEXT: s_waitcnt vmcnt(0)418; GFX908-NEXT: v_accvgpr_read_b32 v0, a4419; GFX908-NEXT: ;;#ASMSTART420; GFX908-NEXT: ; use v0421; GFX908-NEXT: ;;#ASMEND422; GFX908-NEXT: s_endpgm423;424; GFX90A-LABEL: max_6regs_used_8a:425; GFX90A: ; %bb.0:426; GFX90A-NEXT: s_mov_b32 s8, SCRATCH_RSRC_DWORD0427; GFX90A-NEXT: s_mov_b32 s9, SCRATCH_RSRC_DWORD1428; GFX90A-NEXT: s_mov_b32 s10, -1429; GFX90A-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x24430; GFX90A-NEXT: s_mov_b32 s11, 0xe00000431; GFX90A-NEXT: s_add_u32 s8, s8, s5432; GFX90A-NEXT: s_addc_u32 s9, s9, 0433; GFX90A-NEXT: ;;#ASMSTART434; GFX90A-NEXT: ; def v1435; GFX90A-NEXT: ;;#ASMEND436; GFX90A-NEXT: ;;#ASMSTART437; GFX90A-NEXT: ; def a[0:3]438; GFX90A-NEXT: ;;#ASMEND439; GFX90A-NEXT: buffer_store_dword a0, off, s[8:11], 0 ; 4-byte Folded Spill440; GFX90A-NEXT: s_nop 0441; GFX90A-NEXT: buffer_store_dword a1, off, s[8:11], 0 offset:4 ; 4-byte Folded Spill442; GFX90A-NEXT: buffer_store_dword a2, off, s[8:11], 0 offset:8 ; 4-byte Folded Spill443; GFX90A-NEXT: buffer_store_dword a3, off, s[8:11], 0 offset:12 ; 4-byte Folded Spill444; GFX90A-NEXT: v_lshlrev_b32_e32 v0, 4, v0445; GFX90A-NEXT: s_waitcnt lgkmcnt(0)446; GFX90A-NEXT: global_load_dwordx4 a[0:3], v0, s[2:3]447; GFX90A-NEXT: v_mov_b32_e32 v2, 1.0448; GFX90A-NEXT: s_waitcnt vmcnt(0)449; GFX90A-NEXT: s_nop 0450; GFX90A-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v2, v2, a[0:3]451; GFX90A-NEXT: buffer_load_dword v2, off, s[8:11], 0 ; 4-byte Folded Reload452; GFX90A-NEXT: buffer_load_dword v3, off, s[8:11], 0 offset:4 ; 4-byte Folded Reload453; GFX90A-NEXT: buffer_load_dword v4, off, s[8:11], 0 offset:8 ; 4-byte Folded Reload454; GFX90A-NEXT: buffer_load_dword v5, off, s[8:11], 0 offset:12 ; 4-byte Folded Reload455; GFX90A-NEXT: s_nop 0456; GFX90A-NEXT: global_store_dwordx4 v0, a[0:3], s[2:3]457; GFX90A-NEXT: s_waitcnt vmcnt(1)458; GFX90A-NEXT: global_store_dwordx4 v[0:1], v[2:5], off459; GFX90A-NEXT: s_waitcnt vmcnt(0)460; GFX90A-NEXT: ;;#ASMSTART461; GFX90A-NEXT: ; use v1462; GFX90A-NEXT: ;;#ASMEND463; GFX90A-NEXT: s_endpgm464 %tid = call i32 @llvm.amdgcn.workitem.id.x()465 %v0 = call float asm sideeffect "; def $0", "=v"()466 %a4 = call <4 x float> asm sideeffect "; def $0", "=a"()467 %gep = getelementptr inbounds <4 x float>, ptr addrspace(1) %arg, i32 %tid468 %mai.in = load <4 x float>, ptr addrspace(1) %gep469 %mai.out = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai.in, i32 0, i32 0, i32 0)470 store <4 x float> %mai.out, ptr addrspace(1) %gep471 store volatile <4 x float> %a4, ptr addrspace(1) poison472 call void asm sideeffect "; use $0", "v"(float %v0);473 ret void474}475; GCN: ScratchSize: 20476 477declare i32 @llvm.amdgcn.workitem.id.x()478declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32, i32, i32)479declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)480declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)481 482 483attributes #1 = { nounwind "amdgpu-num-vgpr"="10" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }484attributes #2 = { nounwind "amdgpu-num-vgpr"="12" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }485attributes #3 = { nounwind "amdgpu-num-vgpr"="32" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }486attributes #4 = { nounwind "amdgpu-num-vgpr"="6" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }487 488attributes #1 = { nounwind "amdgpu-num-vgpr"="10" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }489attributes #2 = { nounwind "amdgpu-num-vgpr"="12" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }490attributes #3 = { nounwind "amdgpu-num-vgpr"="32" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }491attributes #4 = { nounwind "amdgpu-num-vgpr"="6" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }492;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:493; GCN: {{.*}}494