1150 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py2; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 < %s | FileCheck -check-prefix=GFX908 %s3; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefix=GFX90A %s4 5; This testcase would fail on GFX908 due to not having a free VGPR available to6; copy between AGPRs.7define void @no_free_vgprs_at_agpr_to_agpr_copy(float %v0, float %v1) #0 {8; GFX908-LABEL: no_free_vgprs_at_agpr_to_agpr_copy:9; GFX908: ; %bb.0:10; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)11; GFX908-NEXT: v_mov_b32_e32 v32, v112; GFX908-NEXT: v_mov_b32_e32 v33, v013; GFX908-NEXT: ;;#ASMSTART14; GFX908-NEXT: ; def v[0:31] a[0:15]15; GFX908-NEXT: ;;#ASMEND16; GFX908-NEXT: v_accvgpr_read_b32 v39, a1517; GFX908-NEXT: s_nop 118; GFX908-NEXT: v_accvgpr_write_b32 a31, v3919; GFX908-NEXT: v_accvgpr_read_b32 v39, a1420; GFX908-NEXT: s_nop 121; GFX908-NEXT: v_accvgpr_write_b32 a30, v3922; GFX908-NEXT: v_accvgpr_read_b32 v39, a1323; GFX908-NEXT: s_nop 124; GFX908-NEXT: v_accvgpr_write_b32 a29, v3925; GFX908-NEXT: v_accvgpr_read_b32 v39, a1226; GFX908-NEXT: s_nop 127; GFX908-NEXT: v_accvgpr_write_b32 a28, v3928; GFX908-NEXT: v_accvgpr_read_b32 v39, a1129; GFX908-NEXT: s_nop 130; GFX908-NEXT: v_accvgpr_write_b32 a27, v3931; GFX908-NEXT: v_accvgpr_read_b32 v39, a1032; GFX908-NEXT: s_nop 133; GFX908-NEXT: v_accvgpr_write_b32 a26, v3934; GFX908-NEXT: v_accvgpr_read_b32 v39, a935; GFX908-NEXT: s_nop 136; GFX908-NEXT: v_accvgpr_write_b32 a25, v3937; GFX908-NEXT: v_accvgpr_read_b32 v39, a838; GFX908-NEXT: s_nop 139; GFX908-NEXT: v_accvgpr_write_b32 a24, v3940; GFX908-NEXT: v_accvgpr_read_b32 v39, a741; GFX908-NEXT: s_nop 142; GFX908-NEXT: v_accvgpr_write_b32 a23, v3943; GFX908-NEXT: v_accvgpr_read_b32 v39, a644; GFX908-NEXT: s_nop 145; GFX908-NEXT: v_accvgpr_write_b32 a22, v3946; GFX908-NEXT: v_accvgpr_read_b32 v39, a547; GFX908-NEXT: s_nop 148; GFX908-NEXT: v_accvgpr_write_b32 a21, v3949; GFX908-NEXT: v_accvgpr_read_b32 v39, a450; GFX908-NEXT: s_nop 151; GFX908-NEXT: v_accvgpr_write_b32 a20, v3952; GFX908-NEXT: v_accvgpr_read_b32 v39, a353; GFX908-NEXT: s_nop 154; GFX908-NEXT: v_accvgpr_write_b32 a19, v3955; GFX908-NEXT: v_accvgpr_read_b32 v39, a256; GFX908-NEXT: s_nop 157; GFX908-NEXT: v_accvgpr_write_b32 a18, v3958; GFX908-NEXT: v_accvgpr_read_b32 v39, a159; GFX908-NEXT: s_nop 160; GFX908-NEXT: v_accvgpr_write_b32 a17, v3961; GFX908-NEXT: v_accvgpr_read_b32 v39, a062; GFX908-NEXT: s_nop 163; GFX908-NEXT: v_accvgpr_write_b32 a16, v3964; GFX908-NEXT: s_nop 065; GFX908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]66; GFX908-NEXT: s_nop 967; GFX908-NEXT: v_accvgpr_read_b32 v39, a0 ; Reload Reuse68; GFX908-NEXT: v_accvgpr_read_b32 v38, a11 ; Reload Reuse69; GFX908-NEXT: v_accvgpr_read_b32 v37, a12 ; Reload Reuse70; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 ; 4-byte Folded Spill71; GFX908-NEXT: v_accvgpr_read_b32 v39, a1 ; Reload Reuse72; GFX908-NEXT: v_accvgpr_read_b32 v36, a13 ; Reload Reuse73; GFX908-NEXT: v_accvgpr_read_b32 v35, a14 ; Reload Reuse74; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill75; GFX908-NEXT: v_accvgpr_read_b32 v39, a2 ; Reload Reuse76; GFX908-NEXT: v_accvgpr_read_b32 v34, a15 ; Reload Reuse77; GFX908-NEXT: s_nop 078; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill79; GFX908-NEXT: v_accvgpr_read_b32 v39, a3 ; Reload Reuse80; GFX908-NEXT: s_nop 181; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill82; GFX908-NEXT: v_accvgpr_read_b32 v39, a4 ; Reload Reuse83; GFX908-NEXT: s_nop 184; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill85; GFX908-NEXT: v_accvgpr_read_b32 v39, a5 ; Reload Reuse86; GFX908-NEXT: s_nop 187; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill88; GFX908-NEXT: v_accvgpr_read_b32 v39, a6 ; Reload Reuse89; GFX908-NEXT: s_nop 190; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill91; GFX908-NEXT: v_accvgpr_read_b32 v39, a7 ; Reload Reuse92; GFX908-NEXT: s_nop 193; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill94; GFX908-NEXT: v_accvgpr_read_b32 v39, a8 ; Reload Reuse95; GFX908-NEXT: s_nop 196; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill97; GFX908-NEXT: v_accvgpr_read_b32 v39, a9 ; Reload Reuse98; GFX908-NEXT: s_nop 199; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill100; GFX908-NEXT: v_accvgpr_read_b32 v39, a10 ; Reload Reuse101; GFX908-NEXT: s_nop 1102; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill103; GFX908-NEXT: ;;#ASMSTART104; GFX908-NEXT: ; copy105; GFX908-NEXT: ;;#ASMEND106; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 ; 4-byte Folded Reload107; GFX908-NEXT: v_accvgpr_read_b32 v32, a1108; GFX908-NEXT: s_waitcnt vmcnt(0)109; GFX908-NEXT: v_accvgpr_write_b32 a0, v39 ; Reload Reuse110; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload111; GFX908-NEXT: v_accvgpr_write_b32 a16, v32112; GFX908-NEXT: v_accvgpr_write_b32 a11, v38 ; Reload Reuse113; GFX908-NEXT: v_accvgpr_write_b32 a12, v37 ; Reload Reuse114; GFX908-NEXT: v_accvgpr_write_b32 a13, v36 ; Reload Reuse115; GFX908-NEXT: v_accvgpr_write_b32 a14, v35 ; Reload Reuse116; GFX908-NEXT: s_waitcnt vmcnt(0)117; GFX908-NEXT: v_accvgpr_write_b32 a1, v39 ; Reload Reuse118; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload119; GFX908-NEXT: s_waitcnt vmcnt(0)120; GFX908-NEXT: v_accvgpr_write_b32 a2, v39 ; Reload Reuse121; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload122; GFX908-NEXT: s_waitcnt vmcnt(0)123; GFX908-NEXT: v_accvgpr_write_b32 a3, v39 ; Reload Reuse124; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload125; GFX908-NEXT: s_waitcnt vmcnt(0)126; GFX908-NEXT: v_accvgpr_write_b32 a4, v39 ; Reload Reuse127; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload128; GFX908-NEXT: s_waitcnt vmcnt(0)129; GFX908-NEXT: v_accvgpr_write_b32 a5, v39 ; Reload Reuse130; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload131; GFX908-NEXT: s_waitcnt vmcnt(0)132; GFX908-NEXT: v_accvgpr_write_b32 a6, v39 ; Reload Reuse133; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload134; GFX908-NEXT: s_waitcnt vmcnt(0)135; GFX908-NEXT: v_accvgpr_write_b32 a7, v39 ; Reload Reuse136; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload137; GFX908-NEXT: s_waitcnt vmcnt(0)138; GFX908-NEXT: v_accvgpr_write_b32 a8, v39 ; Reload Reuse139; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload140; GFX908-NEXT: s_waitcnt vmcnt(0)141; GFX908-NEXT: v_accvgpr_write_b32 a9, v39 ; Reload Reuse142; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload143; GFX908-NEXT: s_waitcnt vmcnt(0)144; GFX908-NEXT: v_accvgpr_write_b32 a10, v39 ; Reload Reuse145; GFX908-NEXT: v_accvgpr_write_b32 a15, v34 ; Reload Reuse146; GFX908-NEXT: ;;#ASMSTART147; GFX908-NEXT: ; copy148; GFX908-NEXT: ;;#ASMEND149; GFX908-NEXT: v_accvgpr_read_b32 v39, a2150; GFX908-NEXT: s_nop 1151; GFX908-NEXT: v_accvgpr_write_b32 a3, v39152; GFX908-NEXT: ;;#ASMSTART153; GFX908-NEXT: ; use a3 v[0:31]154; GFX908-NEXT: ;;#ASMEND155; GFX908-NEXT: s_setpc_b64 s[30:31]156;157; GFX90A-LABEL: no_free_vgprs_at_agpr_to_agpr_copy:158; GFX90A: ; %bb.0:159; GFX90A-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)160; GFX90A-NEXT: v_mov_b32_e32 v33, v0161; GFX90A-NEXT: v_mov_b32_e32 v32, v1162; GFX90A-NEXT: ;;#ASMSTART163; GFX90A-NEXT: ; def v[0:31] a[0:15]164; GFX90A-NEXT: ;;#ASMEND165; GFX90A-NEXT: v_accvgpr_mov_b32 a31, a15166; GFX90A-NEXT: v_accvgpr_mov_b32 a30, a14167; GFX90A-NEXT: v_accvgpr_mov_b32 a29, a13168; GFX90A-NEXT: v_accvgpr_mov_b32 a28, a12169; GFX90A-NEXT: v_accvgpr_mov_b32 a27, a11170; GFX90A-NEXT: v_accvgpr_mov_b32 a26, a10171; GFX90A-NEXT: v_accvgpr_mov_b32 a25, a9172; GFX90A-NEXT: v_accvgpr_mov_b32 a24, a8173; GFX90A-NEXT: v_accvgpr_mov_b32 a23, a7174; GFX90A-NEXT: v_accvgpr_mov_b32 a22, a6175; GFX90A-NEXT: v_accvgpr_mov_b32 a21, a5176; GFX90A-NEXT: v_accvgpr_mov_b32 a20, a4177; GFX90A-NEXT: v_accvgpr_mov_b32 a19, a3178; GFX90A-NEXT: v_accvgpr_mov_b32 a18, a2179; GFX90A-NEXT: v_accvgpr_mov_b32 a17, a1180; GFX90A-NEXT: v_accvgpr_mov_b32 a16, a0181; GFX90A-NEXT: s_nop 1182; GFX90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]183; GFX90A-NEXT: s_nop 10184; GFX90A-NEXT: buffer_store_dword a0, off, s[0:3], s32 ; 4-byte Folded Spill185; GFX90A-NEXT: s_nop 0186; GFX90A-NEXT: buffer_store_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill187; GFX90A-NEXT: buffer_store_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill188; GFX90A-NEXT: buffer_store_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill189; GFX90A-NEXT: buffer_store_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill190; GFX90A-NEXT: buffer_store_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill191; GFX90A-NEXT: buffer_store_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill192; GFX90A-NEXT: buffer_store_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill193; GFX90A-NEXT: buffer_store_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill194; GFX90A-NEXT: buffer_store_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill195; GFX90A-NEXT: v_accvgpr_read_b32 v39, a10 ; Reload Reuse196; GFX90A-NEXT: v_accvgpr_read_b32 v38, a11 ; Reload Reuse197; GFX90A-NEXT: v_accvgpr_read_b32 v37, a12 ; Reload Reuse198; GFX90A-NEXT: v_accvgpr_read_b32 v36, a13 ; Reload Reuse199; GFX90A-NEXT: v_accvgpr_read_b32 v35, a14 ; Reload Reuse200; GFX90A-NEXT: v_accvgpr_read_b32 v34, a15 ; Reload Reuse201; GFX90A-NEXT: ;;#ASMSTART202; GFX90A-NEXT: ; copy203; GFX90A-NEXT: ;;#ASMEND204; GFX90A-NEXT: v_accvgpr_mov_b32 a16, a1205; GFX90A-NEXT: buffer_load_dword a0, off, s[0:3], s32 ; 4-byte Folded Reload206; GFX90A-NEXT: buffer_load_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload207; GFX90A-NEXT: buffer_load_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload208; GFX90A-NEXT: buffer_load_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload209; GFX90A-NEXT: buffer_load_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload210; GFX90A-NEXT: buffer_load_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload211; GFX90A-NEXT: buffer_load_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload212; GFX90A-NEXT: buffer_load_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload213; GFX90A-NEXT: buffer_load_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload214; GFX90A-NEXT: buffer_load_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload215; GFX90A-NEXT: v_accvgpr_write_b32 a10, v39 ; Reload Reuse216; GFX90A-NEXT: v_accvgpr_write_b32 a11, v38 ; Reload Reuse217; GFX90A-NEXT: v_accvgpr_write_b32 a12, v37 ; Reload Reuse218; GFX90A-NEXT: v_accvgpr_write_b32 a13, v36 ; Reload Reuse219; GFX90A-NEXT: v_accvgpr_write_b32 a14, v35 ; Reload Reuse220; GFX90A-NEXT: s_waitcnt vmcnt(0)221; GFX90A-NEXT: v_accvgpr_write_b32 a15, v34 ; Reload Reuse222; GFX90A-NEXT: ;;#ASMSTART223; GFX90A-NEXT: ; copy224; GFX90A-NEXT: ;;#ASMEND225; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a2226; GFX90A-NEXT: ;;#ASMSTART227; GFX90A-NEXT: ; use a3 v[0:31]228; GFX90A-NEXT: ;;#ASMEND229; GFX90A-NEXT: s_setpc_b64 s[30:31]230 %asm = call { <32 x i32>, <16 x float> } asm sideeffect "; def $0 $1","=${v[0:31]},=${a[0:15]}"()231 %vgpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 0232 %agpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 1233 %mfma = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %v0, float %v1, <16 x float> %agpr0, i32 0, i32 0, i32 0)234 %agpr1 = call i32 asm sideeffect "; copy ", "={a1},a,~{a[0:15]}"(<16 x float> %agpr0)235 %agpr2 = call i32 asm sideeffect "; copy ", "={a2},a,{a[0:15]}"(i32 %agpr1, <16 x float> %mfma)236 call void asm sideeffect "; use $0 $1","{a3},{v[0:31]}"(i32 %agpr2, <32 x i32> %vgpr0)237 ret void238}239 240; Check that we do make use of v32 if there are no AGPRs present in the function241define amdgpu_kernel void @no_agpr_no_reserve(ptr addrspace(1) %arg) #5 {242; GFX908-LABEL: no_agpr_no_reserve:243; GFX908: ; %bb.0:244; GFX908-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0245; GFX908-NEXT: v_lshlrev_b32_e32 v0, 7, v0246; GFX908-NEXT: s_waitcnt lgkmcnt(0)247; GFX908-NEXT: global_load_dwordx4 v[1:4], v0, s[0:1] offset:16248; GFX908-NEXT: global_load_dwordx4 v[5:8], v0, s[0:1]249; GFX908-NEXT: global_load_dwordx4 v[9:12], v0, s[0:1] offset:48250; GFX908-NEXT: global_load_dwordx4 v[13:16], v0, s[0:1] offset:32251; GFX908-NEXT: global_load_dwordx4 v[17:20], v0, s[0:1] offset:80252; GFX908-NEXT: global_load_dwordx4 v[21:24], v0, s[0:1] offset:64253; GFX908-NEXT: global_load_dwordx4 v[25:28], v0, s[0:1] offset:112254; GFX908-NEXT: global_load_dwordx4 v[29:32], v0, s[0:1] offset:96255; GFX908-NEXT: s_waitcnt vmcnt(7)256; GFX908-NEXT: v_add_u32_e32 v4, v4, v4257; GFX908-NEXT: v_add_u32_e32 v3, v3, v3258; GFX908-NEXT: v_add_u32_e32 v2, v2, v2259; GFX908-NEXT: v_add_u32_e32 v1, v1, v1260; GFX908-NEXT: s_waitcnt vmcnt(6)261; GFX908-NEXT: v_add_u32_e32 v8, v8, v8262; GFX908-NEXT: v_add_u32_e32 v7, v7, v7263; GFX908-NEXT: v_add_u32_e32 v6, v6, v6264; GFX908-NEXT: s_waitcnt vmcnt(0)265; GFX908-NEXT: v_add_u32_e32 v32, v32, v32266; GFX908-NEXT: v_add_u32_e32 v31, v31, v31267; GFX908-NEXT: v_add_u32_e32 v30, v30, v30268; GFX908-NEXT: v_add_u32_e32 v29, v29, v29269; GFX908-NEXT: v_add_u32_e32 v5, v5, v5270; GFX908-NEXT: v_add_u32_e32 v12, v12, v12271; GFX908-NEXT: v_add_u32_e32 v11, v11, v11272; GFX908-NEXT: v_add_u32_e32 v10, v10, v10273; GFX908-NEXT: v_add_u32_e32 v9, v9, v9274; GFX908-NEXT: v_add_u32_e32 v16, v16, v16275; GFX908-NEXT: v_add_u32_e32 v15, v15, v15276; GFX908-NEXT: v_add_u32_e32 v14, v14, v14277; GFX908-NEXT: v_add_u32_e32 v13, v13, v13278; GFX908-NEXT: v_add_u32_e32 v20, v20, v20279; GFX908-NEXT: v_add_u32_e32 v19, v19, v19280; GFX908-NEXT: v_add_u32_e32 v18, v18, v18281; GFX908-NEXT: v_add_u32_e32 v17, v17, v17282; GFX908-NEXT: v_add_u32_e32 v24, v24, v24283; GFX908-NEXT: v_add_u32_e32 v23, v23, v23284; GFX908-NEXT: v_add_u32_e32 v22, v22, v22285; GFX908-NEXT: v_add_u32_e32 v21, v21, v21286; GFX908-NEXT: v_add_u32_e32 v28, v28, v28287; GFX908-NEXT: v_add_u32_e32 v27, v27, v27288; GFX908-NEXT: v_add_u32_e32 v26, v26, v26289; GFX908-NEXT: v_add_u32_e32 v25, v25, v25290; GFX908-NEXT: global_store_dwordx4 v0, v[29:32], s[0:1] offset:96291; GFX908-NEXT: global_store_dwordx4 v0, v[25:28], s[0:1] offset:112292; GFX908-NEXT: global_store_dwordx4 v0, v[21:24], s[0:1] offset:64293; GFX908-NEXT: global_store_dwordx4 v0, v[17:20], s[0:1] offset:80294; GFX908-NEXT: global_store_dwordx4 v0, v[13:16], s[0:1] offset:32295; GFX908-NEXT: global_store_dwordx4 v0, v[9:12], s[0:1] offset:48296; GFX908-NEXT: global_store_dwordx4 v0, v[5:8], s[0:1]297; GFX908-NEXT: global_store_dwordx4 v0, v[1:4], s[0:1] offset:16298; GFX908-NEXT: s_endpgm299;300; GFX90A-LABEL: no_agpr_no_reserve:301; GFX90A: ; %bb.0:302; GFX90A-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0303; GFX90A-NEXT: v_and_b32_e32 v0, 0x3ff, v0304; GFX90A-NEXT: v_lshlrev_b32_e32 v32, 7, v0305; GFX90A-NEXT: s_waitcnt lgkmcnt(0)306; GFX90A-NEXT: global_load_dwordx4 v[0:3], v32, s[0:1] offset:16307; GFX90A-NEXT: global_load_dwordx4 v[4:7], v32, s[0:1]308; GFX90A-NEXT: global_load_dwordx4 v[8:11], v32, s[0:1] offset:48309; GFX90A-NEXT: global_load_dwordx4 v[12:15], v32, s[0:1] offset:32310; GFX90A-NEXT: global_load_dwordx4 v[16:19], v32, s[0:1] offset:80311; GFX90A-NEXT: global_load_dwordx4 v[20:23], v32, s[0:1] offset:64312; GFX90A-NEXT: global_load_dwordx4 v[24:27], v32, s[0:1] offset:112313; GFX90A-NEXT: global_load_dwordx4 v[28:31], v32, s[0:1] offset:96314; GFX90A-NEXT: s_waitcnt vmcnt(7)315; GFX90A-NEXT: v_add_u32_e32 v3, v3, v3316; GFX90A-NEXT: v_add_u32_e32 v2, v2, v2317; GFX90A-NEXT: v_add_u32_e32 v1, v1, v1318; GFX90A-NEXT: v_add_u32_e32 v0, v0, v0319; GFX90A-NEXT: s_waitcnt vmcnt(6)320; GFX90A-NEXT: v_add_u32_e32 v7, v7, v7321; GFX90A-NEXT: v_add_u32_e32 v6, v6, v6322; GFX90A-NEXT: v_add_u32_e32 v5, v5, v5323; GFX90A-NEXT: s_waitcnt vmcnt(0)324; GFX90A-NEXT: v_add_u32_e32 v31, v31, v31325; GFX90A-NEXT: v_add_u32_e32 v30, v30, v30326; GFX90A-NEXT: v_add_u32_e32 v29, v29, v29327; GFX90A-NEXT: v_add_u32_e32 v28, v28, v28328; GFX90A-NEXT: v_add_u32_e32 v4, v4, v4329; GFX90A-NEXT: v_add_u32_e32 v11, v11, v11330; GFX90A-NEXT: v_add_u32_e32 v10, v10, v10331; GFX90A-NEXT: v_add_u32_e32 v9, v9, v9332; GFX90A-NEXT: v_add_u32_e32 v8, v8, v8333; GFX90A-NEXT: v_add_u32_e32 v15, v15, v15334; GFX90A-NEXT: v_add_u32_e32 v14, v14, v14335; GFX90A-NEXT: v_add_u32_e32 v13, v13, v13336; GFX90A-NEXT: v_add_u32_e32 v12, v12, v12337; GFX90A-NEXT: v_add_u32_e32 v19, v19, v19338; GFX90A-NEXT: v_add_u32_e32 v18, v18, v18339; GFX90A-NEXT: v_add_u32_e32 v17, v17, v17340; GFX90A-NEXT: v_add_u32_e32 v16, v16, v16341; GFX90A-NEXT: v_add_u32_e32 v23, v23, v23342; GFX90A-NEXT: v_add_u32_e32 v22, v22, v22343; GFX90A-NEXT: v_add_u32_e32 v21, v21, v21344; GFX90A-NEXT: v_add_u32_e32 v20, v20, v20345; GFX90A-NEXT: v_add_u32_e32 v27, v27, v27346; GFX90A-NEXT: v_add_u32_e32 v26, v26, v26347; GFX90A-NEXT: v_add_u32_e32 v25, v25, v25348; GFX90A-NEXT: v_add_u32_e32 v24, v24, v24349; GFX90A-NEXT: global_store_dwordx4 v32, v[28:31], s[0:1] offset:96350; GFX90A-NEXT: global_store_dwordx4 v32, v[24:27], s[0:1] offset:112351; GFX90A-NEXT: global_store_dwordx4 v32, v[20:23], s[0:1] offset:64352; GFX90A-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:80353; GFX90A-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:32354; GFX90A-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:48355; GFX90A-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1]356; GFX90A-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] offset:16357; GFX90A-NEXT: s_endpgm358 %id = call i32 @llvm.amdgcn.workitem.id.x()359 %gep = getelementptr inbounds <32 x i32>, ptr addrspace(1) %arg, i32 %id360 %load = load <32 x i32>, ptr addrspace(1) %gep361 %add = add <32 x i32> %load, %load362 store <32 x i32> %add, ptr addrspace(1) %gep363 ret void364}365 366; FIXME: This case is broken. The asm value passed in v32 is live367; through the range where the reserved def for the copy is introduced,368; clobbering the user value.369define void @v32_asm_def_use(float %v0, float %v1) #4 {370; GFX908-LABEL: v32_asm_def_use:371; GFX908: ; %bb.0:372; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)373; GFX908-NEXT: v_mov_b32_e32 v33, v1374; GFX908-NEXT: v_mov_b32_e32 v34, v0375; GFX908-NEXT: ;;#ASMSTART376; GFX908-NEXT: ; def v[0:31] a[0:15]377; GFX908-NEXT: ;;#ASMEND378; GFX908-NEXT: v_accvgpr_read_b32 v35, a15379; GFX908-NEXT: ;;#ASMSTART380; GFX908-NEXT: ; def v32381; GFX908-NEXT: ;;#ASMEND382; GFX908-NEXT: s_nop 1383; GFX908-NEXT: v_accvgpr_write_b32 a31, v35384; GFX908-NEXT: v_accvgpr_read_b32 v35, a14385; GFX908-NEXT: s_nop 1386; GFX908-NEXT: v_accvgpr_write_b32 a30, v35387; GFX908-NEXT: v_accvgpr_read_b32 v35, a13388; GFX908-NEXT: s_nop 1389; GFX908-NEXT: v_accvgpr_write_b32 a29, v35390; GFX908-NEXT: v_accvgpr_read_b32 v35, a12391; GFX908-NEXT: s_nop 1392; GFX908-NEXT: v_accvgpr_write_b32 a28, v35393; GFX908-NEXT: v_accvgpr_read_b32 v35, a11394; GFX908-NEXT: s_nop 1395; GFX908-NEXT: v_accvgpr_write_b32 a27, v35396; GFX908-NEXT: v_accvgpr_read_b32 v35, a10397; GFX908-NEXT: s_nop 1398; GFX908-NEXT: v_accvgpr_write_b32 a26, v35399; GFX908-NEXT: v_accvgpr_read_b32 v35, a9400; GFX908-NEXT: s_nop 1401; GFX908-NEXT: v_accvgpr_write_b32 a25, v35402; GFX908-NEXT: v_accvgpr_read_b32 v35, a8403; GFX908-NEXT: s_nop 1404; GFX908-NEXT: v_accvgpr_write_b32 a24, v35405; GFX908-NEXT: v_accvgpr_read_b32 v35, a7406; GFX908-NEXT: s_nop 1407; GFX908-NEXT: v_accvgpr_write_b32 a23, v35408; GFX908-NEXT: v_accvgpr_read_b32 v35, a6409; GFX908-NEXT: s_nop 1410; GFX908-NEXT: v_accvgpr_write_b32 a22, v35411; GFX908-NEXT: v_accvgpr_read_b32 v35, a5412; GFX908-NEXT: s_nop 1413; GFX908-NEXT: v_accvgpr_write_b32 a21, v35414; GFX908-NEXT: v_accvgpr_read_b32 v35, a4415; GFX908-NEXT: s_nop 1416; GFX908-NEXT: v_accvgpr_write_b32 a20, v35417; GFX908-NEXT: v_accvgpr_read_b32 v35, a3418; GFX908-NEXT: s_nop 1419; GFX908-NEXT: v_accvgpr_write_b32 a19, v35420; GFX908-NEXT: v_accvgpr_read_b32 v35, a2421; GFX908-NEXT: s_nop 1422; GFX908-NEXT: v_accvgpr_write_b32 a18, v35423; GFX908-NEXT: v_accvgpr_read_b32 v35, a1424; GFX908-NEXT: s_nop 1425; GFX908-NEXT: v_accvgpr_write_b32 a17, v35426; GFX908-NEXT: v_accvgpr_read_b32 v35, a0427; GFX908-NEXT: s_nop 1428; GFX908-NEXT: v_accvgpr_write_b32 a16, v35429; GFX908-NEXT: ;;#ASMSTART430; GFX908-NEXT: ; copy431; GFX908-NEXT: ;;#ASMEND432; GFX908-NEXT: v_accvgpr_read_b32 v35, a1433; GFX908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v34, v33, a[16:31]434; GFX908-NEXT: s_nop 0435; GFX908-NEXT: v_accvgpr_write_b32 a32, v35436; GFX908-NEXT: ;;#ASMSTART437; GFX908-NEXT: ; copy438; GFX908-NEXT: ;;#ASMEND439; GFX908-NEXT: s_nop 7440; GFX908-NEXT: v_accvgpr_read_b32 v35, a2441; GFX908-NEXT: s_nop 1442; GFX908-NEXT: v_accvgpr_write_b32 a3, v35443; GFX908-NEXT: ;;#ASMSTART444; GFX908-NEXT: ; use a3 v[0:31]445; GFX908-NEXT: ;;#ASMEND446; GFX908-NEXT: ;;#ASMSTART447; GFX908-NEXT: ; use v32448; GFX908-NEXT: ;;#ASMEND449; GFX908-NEXT: s_setpc_b64 s[30:31]450;451; GFX90A-LABEL: v32_asm_def_use:452; GFX90A: ; %bb.0:453; GFX90A-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)454; GFX90A-NEXT: v_mov_b32_e32 v34, v0455; GFX90A-NEXT: v_mov_b32_e32 v33, v1456; GFX90A-NEXT: ;;#ASMSTART457; GFX90A-NEXT: ; def v[0:31] a[0:15]458; GFX90A-NEXT: ;;#ASMEND459; GFX90A-NEXT: v_accvgpr_mov_b32 a31, a15460; GFX90A-NEXT: v_accvgpr_mov_b32 a30, a14461; GFX90A-NEXT: v_accvgpr_mov_b32 a29, a13462; GFX90A-NEXT: v_accvgpr_mov_b32 a28, a12463; GFX90A-NEXT: v_accvgpr_mov_b32 a27, a11464; GFX90A-NEXT: v_accvgpr_mov_b32 a26, a10465; GFX90A-NEXT: v_accvgpr_mov_b32 a25, a9466; GFX90A-NEXT: v_accvgpr_mov_b32 a24, a8467; GFX90A-NEXT: v_accvgpr_mov_b32 a23, a7468; GFX90A-NEXT: v_accvgpr_mov_b32 a22, a6469; GFX90A-NEXT: v_accvgpr_mov_b32 a21, a5470; GFX90A-NEXT: v_accvgpr_mov_b32 a20, a4471; GFX90A-NEXT: v_accvgpr_mov_b32 a19, a3472; GFX90A-NEXT: v_accvgpr_mov_b32 a18, a2473; GFX90A-NEXT: v_accvgpr_mov_b32 a17, a1474; GFX90A-NEXT: v_accvgpr_mov_b32 a16, a0475; GFX90A-NEXT: ;;#ASMSTART476; GFX90A-NEXT: ; def v32477; GFX90A-NEXT: ;;#ASMEND478; GFX90A-NEXT: ;;#ASMSTART479; GFX90A-NEXT: ; copy480; GFX90A-NEXT: ;;#ASMEND481; GFX90A-NEXT: v_accvgpr_read_b32 v35, a32 ; Reload Reuse482; GFX90A-NEXT: v_accvgpr_mov_b32 a32, a1483; GFX90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v34, v33, a[16:31]484; GFX90A-NEXT: ;;#ASMSTART485; GFX90A-NEXT: ; copy486; GFX90A-NEXT: ;;#ASMEND487; GFX90A-NEXT: v_accvgpr_write_b32 a32, v35 ; Reload Reuse488; GFX90A-NEXT: s_nop 9489; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a2490; GFX90A-NEXT: ;;#ASMSTART491; GFX90A-NEXT: ; use a3 v[0:31]492; GFX90A-NEXT: ;;#ASMEND493; GFX90A-NEXT: ;;#ASMSTART494; GFX90A-NEXT: ; use v32495; GFX90A-NEXT: ;;#ASMEND496; GFX90A-NEXT: s_setpc_b64 s[30:31]497 %asm = call { <32 x i32>, <16 x float> } asm sideeffect "; def $0 $1","=${v[0:31]},=${a[0:15]}"()498 %vgpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 0499 %agpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 1500 %mfma = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %v0, float %v1, <16 x float> %agpr0, i32 0, i32 0, i32 0)501 %v32 = call i32 asm sideeffect "; def $0","=${v32}"()502 %agpr1 = call i32 asm sideeffect "; copy ", "={a1},a,~{a[0:15]}"(<16 x float> %agpr0)503 %agpr2 = call i32 asm sideeffect "; copy ", "={a2},a,{a[0:15]}"(i32 %agpr1, <16 x float> %mfma)504 call void asm sideeffect "; use $0 $1","{a3},{v[0:31]}"(i32 %agpr2, <32 x i32> %vgpr0)505 call void asm sideeffect "; use $0","${v32}"(i32 %v32)506 ret void507}508 509define amdgpu_kernel void @introduced_copy_to_sgpr(i64 %arg, i32 %arg1, i32 %arg2, i64 %arg3, <2 x half> %arg4, <2 x half> %arg5) #3 {510; GFX908-LABEL: introduced_copy_to_sgpr:511; GFX908: ; %bb.0: ; %bb512; GFX908-NEXT: global_load_ushort v16, v[0:1], off glc513; GFX908-NEXT: s_load_dwordx4 s[4:7], s[8:9], 0x0514; GFX908-NEXT: s_load_dwordx2 s[10:11], s[8:9], 0x10515; GFX908-NEXT: s_load_dword s0, s[8:9], 0x18516; GFX908-NEXT: s_mov_b32 s12, 0517; GFX908-NEXT: s_mov_b32 s9, s12518; GFX908-NEXT: s_waitcnt lgkmcnt(0)519; GFX908-NEXT: v_cvt_f32_u32_e32 v0, s7520; GFX908-NEXT: s_sub_i32 s1, 0, s7521; GFX908-NEXT: v_cvt_f32_f16_e32 v18, s0522; GFX908-NEXT: v_mov_b32_e32 v17, 0523; GFX908-NEXT: v_rcp_iflag_f32_e32 v0, v0524; GFX908-NEXT: v_mul_f32_e32 v0, 0x4f7ffffe, v0525; GFX908-NEXT: v_cvt_u32_f32_e32 v0, v0526; GFX908-NEXT: v_readfirstlane_b32 s2, v0527; GFX908-NEXT: s_mul_i32 s1, s1, s2528; GFX908-NEXT: s_mul_hi_u32 s1, s2, s1529; GFX908-NEXT: s_add_i32 s2, s2, s1530; GFX908-NEXT: s_mul_hi_u32 s1, s6, s2531; GFX908-NEXT: s_mul_i32 s2, s1, s7532; GFX908-NEXT: s_sub_i32 s2, s6, s2533; GFX908-NEXT: s_add_i32 s3, s1, 1534; GFX908-NEXT: s_sub_i32 s6, s2, s7535; GFX908-NEXT: s_cmp_ge_u32 s2, s7536; GFX908-NEXT: s_cselect_b32 s1, s3, s1537; GFX908-NEXT: s_cselect_b32 s2, s6, s2538; GFX908-NEXT: s_add_i32 s3, s1, 1539; GFX908-NEXT: s_cmp_ge_u32 s2, s7540; GFX908-NEXT: s_cselect_b32 s8, s3, s1541; GFX908-NEXT: s_lshr_b32 s2, s0, 16542; GFX908-NEXT: v_cvt_f32_f16_e32 v19, s2543; GFX908-NEXT: s_lshl_b64 s[6:7], s[4:5], 5544; GFX908-NEXT: s_lshl_b64 s[14:15], s[10:11], 5545; GFX908-NEXT: v_mov_b32_e32 v0, 0546; GFX908-NEXT: s_and_b64 s[0:1], exec, s[0:1]547; GFX908-NEXT: s_or_b32 s14, s14, 28548; GFX908-NEXT: s_lshl_b64 s[16:17], s[8:9], 5549; GFX908-NEXT: v_mov_b32_e32 v1, 0550; GFX908-NEXT: s_waitcnt vmcnt(0)551; GFX908-NEXT: v_readfirstlane_b32 s2, v16552; GFX908-NEXT: s_and_b32 s2, 0xffff, s2553; GFX908-NEXT: s_mul_i32 s3, s5, s2554; GFX908-NEXT: s_mul_hi_u32 s5, s4, s2555; GFX908-NEXT: s_mul_i32 s2, s4, s2556; GFX908-NEXT: s_add_i32 s3, s5, s3557; GFX908-NEXT: s_lshl_b64 s[4:5], s[2:3], 5558; GFX908-NEXT: s_branch .LBB3_2559; GFX908-NEXT: .LBB3_1: ; %Flow20560; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1561; GFX908-NEXT: s_andn2_b64 vcc, exec, s[2:3]562; GFX908-NEXT: s_cbranch_vccz .LBB3_12563; GFX908-NEXT: .LBB3_2: ; %bb9564; GFX908-NEXT: ; =>This Loop Header: Depth=1565; GFX908-NEXT: ; Child Loop BB3_5 Depth 2566; GFX908-NEXT: s_mov_b64 s[18:19], -1567; GFX908-NEXT: s_mov_b64 vcc, s[0:1]568; GFX908-NEXT: s_cbranch_vccz .LBB3_10569; GFX908-NEXT: ; %bb.3: ; %bb14570; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1571; GFX908-NEXT: global_load_dwordx2 v[2:3], v[0:1], off572; GFX908-NEXT: v_cmp_gt_i64_e64 s[2:3], s[10:11], -1573; GFX908-NEXT: s_mov_b32 s13, s12574; GFX908-NEXT: v_cndmask_b32_e64 v6, 0, 1, s[2:3]575; GFX908-NEXT: v_mov_b32_e32 v4, s12576; GFX908-NEXT: v_cmp_ne_u32_e64 s[2:3], 1, v6577; GFX908-NEXT: v_mov_b32_e32 v6, s12578; GFX908-NEXT: v_mov_b32_e32 v8, s12579; GFX908-NEXT: v_mov_b32_e32 v5, s13580; GFX908-NEXT: v_mov_b32_e32 v7, s13581; GFX908-NEXT: v_mov_b32_e32 v9, s13582; GFX908-NEXT: v_cmp_lt_i64_e64 s[18:19], s[10:11], 0583; GFX908-NEXT: v_mov_b32_e32 v11, v5584; GFX908-NEXT: s_mov_b64 s[20:21], s[14:15]585; GFX908-NEXT: v_mov_b32_e32 v10, v4586; GFX908-NEXT: s_waitcnt vmcnt(0)587; GFX908-NEXT: v_readfirstlane_b32 s9, v2588; GFX908-NEXT: v_readfirstlane_b32 s13, v3589; GFX908-NEXT: s_add_u32 s9, s9, 1590; GFX908-NEXT: s_addc_u32 s13, s13, 0591; GFX908-NEXT: s_mul_hi_u32 s22, s6, s9592; GFX908-NEXT: s_mul_i32 s13, s6, s13593; GFX908-NEXT: s_mul_i32 s23, s7, s9594; GFX908-NEXT: s_add_i32 s13, s22, s13595; GFX908-NEXT: s_mul_i32 s9, s6, s9596; GFX908-NEXT: s_add_i32 s13, s13, s23597; GFX908-NEXT: s_branch .LBB3_5598; GFX908-NEXT: .LBB3_4: ; %bb58599; GFX908-NEXT: ; in Loop: Header=BB3_5 Depth=2600; GFX908-NEXT: v_add_co_u32_sdwa v2, vcc, v2, v16 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:WORD_0601; GFX908-NEXT: v_addc_co_u32_e32 v3, vcc, 0, v3, vcc602; GFX908-NEXT: s_add_u32 s20, s20, s4603; GFX908-NEXT: v_cmp_lt_i64_e64 s[24:25], -1, v[2:3]604; GFX908-NEXT: s_addc_u32 s21, s21, s5605; GFX908-NEXT: s_mov_b64 s[22:23], 0606; GFX908-NEXT: s_andn2_b64 vcc, exec, s[24:25]607; GFX908-NEXT: s_cbranch_vccz .LBB3_9608; GFX908-NEXT: .LBB3_5: ; %bb16609; GFX908-NEXT: ; Parent Loop BB3_2 Depth=1610; GFX908-NEXT: ; => This Inner Loop Header: Depth=2611; GFX908-NEXT: s_add_u32 s22, s20, s9612; GFX908-NEXT: s_addc_u32 s23, s21, s13613; GFX908-NEXT: global_load_dword v21, v17, s[22:23] offset:-12 glc614; GFX908-NEXT: s_waitcnt vmcnt(0)615; GFX908-NEXT: global_load_dword v20, v17, s[22:23] offset:-8 glc616; GFX908-NEXT: s_waitcnt vmcnt(0)617; GFX908-NEXT: global_load_dword v12, v17, s[22:23] offset:-4 glc618; GFX908-NEXT: s_waitcnt vmcnt(0)619; GFX908-NEXT: global_load_dword v12, v17, s[22:23] glc620; GFX908-NEXT: s_waitcnt vmcnt(0)621; GFX908-NEXT: ds_read_b64 v[12:13], v17622; GFX908-NEXT: ds_read_b64 v[14:15], v0623; GFX908-NEXT: s_and_b64 vcc, exec, s[2:3]624; GFX908-NEXT: s_waitcnt lgkmcnt(0)625; GFX908-NEXT: s_cbranch_vccnz .LBB3_7626; GFX908-NEXT: ; %bb.6: ; %bb51627; GFX908-NEXT: ; in Loop: Header=BB3_5 Depth=2628; GFX908-NEXT: v_cvt_f32_f16_sdwa v22, v21 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1629; GFX908-NEXT: v_cvt_f32_f16_e32 v21, v21630; GFX908-NEXT: v_cvt_f32_f16_sdwa v23, v20 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1631; GFX908-NEXT: v_cvt_f32_f16_e32 v20, v20632; GFX908-NEXT: v_add_f32_e32 v24, v18, v12633; GFX908-NEXT: v_add_f32_e32 v25, v19, v13634; GFX908-NEXT: v_add_f32_e32 v26, 0, v12635; GFX908-NEXT: v_add_f32_e32 v27, 0, v13636; GFX908-NEXT: v_add_f32_e32 v15, v22, v15637; GFX908-NEXT: v_add_f32_e32 v14, v21, v14638; GFX908-NEXT: v_add_f32_e32 v13, v23, v13639; GFX908-NEXT: v_add_f32_e32 v12, v20, v12640; GFX908-NEXT: v_add_f32_e32 v5, v5, v25641; GFX908-NEXT: v_add_f32_e32 v4, v4, v24642; GFX908-NEXT: v_add_f32_e32 v7, v7, v27643; GFX908-NEXT: v_add_f32_e32 v6, v6, v26644; GFX908-NEXT: v_add_f32_e32 v8, v8, v14645; GFX908-NEXT: v_add_f32_e32 v9, v9, v15646; GFX908-NEXT: v_add_f32_e32 v10, v10, v12647; GFX908-NEXT: v_add_f32_e32 v11, v11, v13648; GFX908-NEXT: s_branch .LBB3_4649; GFX908-NEXT: .LBB3_7: ; in Loop: Header=BB3_5 Depth=2650; GFX908-NEXT: s_mov_b64 s[22:23], s[18:19]651; GFX908-NEXT: s_andn2_b64 vcc, exec, s[22:23]652; GFX908-NEXT: s_cbranch_vccz .LBB3_4653; GFX908-NEXT: ; %bb.8: ; in Loop: Header=BB3_2 Depth=1654; GFX908-NEXT: s_mov_b64 s[22:23], -1655; GFX908-NEXT: ; implicit-def: $vgpr2_vgpr3656; GFX908-NEXT: ; implicit-def: $sgpr20_sgpr21657; GFX908-NEXT: .LBB3_9: ; %loop.exit.guard658; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1659; GFX908-NEXT: s_xor_b64 s[18:19], s[22:23], -1660; GFX908-NEXT: .LBB3_10: ; %Flow19661; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1662; GFX908-NEXT: s_mov_b64 s[2:3], -1663; GFX908-NEXT: s_and_b64 vcc, exec, s[18:19]664; GFX908-NEXT: s_cbranch_vccz .LBB3_1665; GFX908-NEXT: ; %bb.11: ; %bb12666; GFX908-NEXT: ; in Loop: Header=BB3_2 Depth=1667; GFX908-NEXT: s_add_u32 s10, s10, s8668; GFX908-NEXT: s_addc_u32 s11, s11, 0669; GFX908-NEXT: s_add_u32 s14, s14, s16670; GFX908-NEXT: s_addc_u32 s15, s15, s17671; GFX908-NEXT: s_mov_b64 s[2:3], 0672; GFX908-NEXT: s_branch .LBB3_1673; GFX908-NEXT: .LBB3_12: ; %DummyReturnBlock674; GFX908-NEXT: s_endpgm675;676; GFX90A-LABEL: introduced_copy_to_sgpr:677; GFX90A: ; %bb.0: ; %bb678; GFX90A-NEXT: global_load_ushort v18, v[0:1], off glc679; GFX90A-NEXT: s_load_dwordx4 s[4:7], s[8:9], 0x0680; GFX90A-NEXT: s_load_dwordx2 s[10:11], s[8:9], 0x10681; GFX90A-NEXT: s_load_dword s0, s[8:9], 0x18682; GFX90A-NEXT: s_mov_b32 s12, 0683; GFX90A-NEXT: s_mov_b32 s9, s12684; GFX90A-NEXT: s_waitcnt lgkmcnt(0)685; GFX90A-NEXT: v_cvt_f32_u32_e32 v0, s7686; GFX90A-NEXT: s_sub_i32 s1, 0, s7687; GFX90A-NEXT: v_mov_b32_e32 v19, 0688; GFX90A-NEXT: v_pk_mov_b32 v[2:3], 0, 0689; GFX90A-NEXT: v_rcp_iflag_f32_e32 v0, v0690; GFX90A-NEXT: v_mul_f32_e32 v0, 0x4f7ffffe, v0691; GFX90A-NEXT: v_cvt_u32_f32_e32 v1, v0692; GFX90A-NEXT: v_cvt_f32_f16_e32 v0, s0693; GFX90A-NEXT: v_readfirstlane_b32 s2, v1694; GFX90A-NEXT: s_mul_i32 s1, s1, s2695; GFX90A-NEXT: s_mul_hi_u32 s1, s2, s1696; GFX90A-NEXT: s_add_i32 s2, s2, s1697; GFX90A-NEXT: s_mul_hi_u32 s1, s6, s2698; GFX90A-NEXT: s_mul_i32 s2, s1, s7699; GFX90A-NEXT: s_sub_i32 s2, s6, s2700; GFX90A-NEXT: s_add_i32 s3, s1, 1701; GFX90A-NEXT: s_sub_i32 s6, s2, s7702; GFX90A-NEXT: s_cmp_ge_u32 s2, s7703; GFX90A-NEXT: s_cselect_b32 s1, s3, s1704; GFX90A-NEXT: s_cselect_b32 s2, s6, s2705; GFX90A-NEXT: s_add_i32 s3, s1, 1706; GFX90A-NEXT: s_cmp_ge_u32 s2, s7707; GFX90A-NEXT: s_cselect_b32 s8, s3, s1708; GFX90A-NEXT: s_lshr_b32 s2, s0, 16709; GFX90A-NEXT: v_cvt_f32_f16_e32 v1, s2710; GFX90A-NEXT: s_lshl_b64 s[6:7], s[4:5], 5711; GFX90A-NEXT: s_lshl_b64 s[14:15], s[10:11], 5712; GFX90A-NEXT: s_and_b64 s[0:1], exec, s[0:1]713; GFX90A-NEXT: s_or_b32 s14, s14, 28714; GFX90A-NEXT: s_lshl_b64 s[16:17], s[8:9], 5715; GFX90A-NEXT: s_waitcnt vmcnt(0)716; GFX90A-NEXT: v_readfirstlane_b32 s2, v18717; GFX90A-NEXT: s_and_b32 s2, 0xffff, s2718; GFX90A-NEXT: s_mul_i32 s3, s5, s2719; GFX90A-NEXT: s_mul_hi_u32 s5, s4, s2720; GFX90A-NEXT: s_mul_i32 s2, s4, s2721; GFX90A-NEXT: s_add_i32 s3, s5, s3722; GFX90A-NEXT: s_lshl_b64 s[4:5], s[2:3], 5723; GFX90A-NEXT: s_branch .LBB3_2724; GFX90A-NEXT: .LBB3_1: ; %Flow20725; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1726; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[2:3]727; GFX90A-NEXT: s_cbranch_vccz .LBB3_12728; GFX90A-NEXT: .LBB3_2: ; %bb9729; GFX90A-NEXT: ; =>This Loop Header: Depth=1730; GFX90A-NEXT: ; Child Loop BB3_5 Depth 2731; GFX90A-NEXT: s_mov_b64 s[18:19], -1732; GFX90A-NEXT: s_mov_b64 vcc, s[0:1]733; GFX90A-NEXT: s_cbranch_vccz .LBB3_10734; GFX90A-NEXT: ; %bb.3: ; %bb14735; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1736; GFX90A-NEXT: global_load_dwordx2 v[4:5], v[2:3], off737; GFX90A-NEXT: v_cmp_gt_i64_e64 s[2:3], s[10:11], -1738; GFX90A-NEXT: s_mov_b32 s13, s12739; GFX90A-NEXT: v_cndmask_b32_e64 v8, 0, 1, s[2:3]740; GFX90A-NEXT: v_pk_mov_b32 v[6:7], s[12:13], s[12:13] op_sel:[0,1]741; GFX90A-NEXT: v_cmp_ne_u32_e64 s[2:3], 1, v8742; GFX90A-NEXT: v_pk_mov_b32 v[8:9], s[12:13], s[12:13] op_sel:[0,1]743; GFX90A-NEXT: v_pk_mov_b32 v[10:11], s[12:13], s[12:13] op_sel:[0,1]744; GFX90A-NEXT: v_cmp_lt_i64_e64 s[18:19], s[10:11], 0745; GFX90A-NEXT: s_mov_b64 s[20:21], s[14:15]746; GFX90A-NEXT: v_pk_mov_b32 v[12:13], v[6:7], v[6:7] op_sel:[0,1]747; GFX90A-NEXT: s_waitcnt vmcnt(0)748; GFX90A-NEXT: v_readfirstlane_b32 s9, v4749; GFX90A-NEXT: v_readfirstlane_b32 s13, v5750; GFX90A-NEXT: s_add_u32 s9, s9, 1751; GFX90A-NEXT: s_addc_u32 s13, s13, 0752; GFX90A-NEXT: s_mul_hi_u32 s22, s6, s9753; GFX90A-NEXT: s_mul_i32 s13, s6, s13754; GFX90A-NEXT: s_mul_i32 s23, s7, s9755; GFX90A-NEXT: s_add_i32 s13, s22, s13756; GFX90A-NEXT: s_mul_i32 s9, s6, s9757; GFX90A-NEXT: s_add_i32 s13, s13, s23758; GFX90A-NEXT: s_branch .LBB3_5759; GFX90A-NEXT: .LBB3_4: ; %bb58760; GFX90A-NEXT: ; in Loop: Header=BB3_5 Depth=2761; GFX90A-NEXT: v_add_co_u32_sdwa v4, vcc, v4, v18 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:WORD_0762; GFX90A-NEXT: v_addc_co_u32_e32 v5, vcc, 0, v5, vcc763; GFX90A-NEXT: s_add_u32 s20, s20, s4764; GFX90A-NEXT: s_addc_u32 s21, s21, s5765; GFX90A-NEXT: v_cmp_lt_i64_e64 s[24:25], -1, v[4:5]766; GFX90A-NEXT: s_mov_b64 s[22:23], 0767; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[24:25]768; GFX90A-NEXT: s_cbranch_vccz .LBB3_9769; GFX90A-NEXT: .LBB3_5: ; %bb16770; GFX90A-NEXT: ; Parent Loop BB3_2 Depth=1771; GFX90A-NEXT: ; => This Inner Loop Header: Depth=2772; GFX90A-NEXT: s_add_u32 s22, s20, s9773; GFX90A-NEXT: s_addc_u32 s23, s21, s13774; GFX90A-NEXT: global_load_dword v21, v19, s[22:23] offset:-12 glc775; GFX90A-NEXT: s_waitcnt vmcnt(0)776; GFX90A-NEXT: global_load_dword v20, v19, s[22:23] offset:-8 glc777; GFX90A-NEXT: s_waitcnt vmcnt(0)778; GFX90A-NEXT: global_load_dword v14, v19, s[22:23] offset:-4 glc779; GFX90A-NEXT: s_waitcnt vmcnt(0)780; GFX90A-NEXT: global_load_dword v14, v19, s[22:23] glc781; GFX90A-NEXT: s_waitcnt vmcnt(0)782; GFX90A-NEXT: ds_read_b64 v[14:15], v19783; GFX90A-NEXT: ds_read_b64 v[16:17], v0784; GFX90A-NEXT: s_and_b64 vcc, exec, s[2:3]785; GFX90A-NEXT: ; kill: killed $sgpr22 killed $sgpr23786; GFX90A-NEXT: s_waitcnt lgkmcnt(0)787; GFX90A-NEXT: s_cbranch_vccnz .LBB3_7788; GFX90A-NEXT: ; %bb.6: ; %bb51789; GFX90A-NEXT: ; in Loop: Header=BB3_5 Depth=2790; GFX90A-NEXT: v_cvt_f32_f16_sdwa v23, v21 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1791; GFX90A-NEXT: v_cvt_f32_f16_e32 v22, v21792; GFX90A-NEXT: v_cvt_f32_f16_sdwa v21, v20 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1793; GFX90A-NEXT: v_cvt_f32_f16_e32 v20, v20794; GFX90A-NEXT: v_pk_add_f32 v[24:25], v[0:1], v[14:15]795; GFX90A-NEXT: v_pk_add_f32 v[26:27], v[14:15], 0 op_sel_hi:[1,0]796; GFX90A-NEXT: v_pk_add_f32 v[16:17], v[22:23], v[16:17]797; GFX90A-NEXT: v_pk_add_f32 v[14:15], v[20:21], v[14:15]798; GFX90A-NEXT: v_pk_add_f32 v[6:7], v[6:7], v[24:25]799; GFX90A-NEXT: v_pk_add_f32 v[8:9], v[8:9], v[26:27]800; GFX90A-NEXT: v_pk_add_f32 v[10:11], v[10:11], v[16:17]801; GFX90A-NEXT: v_pk_add_f32 v[12:13], v[12:13], v[14:15]802; GFX90A-NEXT: s_branch .LBB3_4803; GFX90A-NEXT: .LBB3_7: ; in Loop: Header=BB3_5 Depth=2804; GFX90A-NEXT: s_mov_b64 s[22:23], s[18:19]805; GFX90A-NEXT: s_andn2_b64 vcc, exec, s[22:23]806; GFX90A-NEXT: s_cbranch_vccz .LBB3_4807; GFX90A-NEXT: ; %bb.8: ; in Loop: Header=BB3_2 Depth=1808; GFX90A-NEXT: s_mov_b64 s[22:23], -1809; GFX90A-NEXT: ; implicit-def: $vgpr4_vgpr5810; GFX90A-NEXT: ; implicit-def: $sgpr20_sgpr21811; GFX90A-NEXT: .LBB3_9: ; %loop.exit.guard812; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1813; GFX90A-NEXT: s_xor_b64 s[18:19], s[22:23], -1814; GFX90A-NEXT: .LBB3_10: ; %Flow19815; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1816; GFX90A-NEXT: s_mov_b64 s[2:3], -1817; GFX90A-NEXT: s_and_b64 vcc, exec, s[18:19]818; GFX90A-NEXT: s_cbranch_vccz .LBB3_1819; GFX90A-NEXT: ; %bb.11: ; %bb12820; GFX90A-NEXT: ; in Loop: Header=BB3_2 Depth=1821; GFX90A-NEXT: s_add_u32 s10, s10, s8822; GFX90A-NEXT: s_addc_u32 s11, s11, 0823; GFX90A-NEXT: s_add_u32 s14, s14, s16824; GFX90A-NEXT: s_addc_u32 s15, s15, s17825; GFX90A-NEXT: s_mov_b64 s[2:3], 0826; GFX90A-NEXT: s_branch .LBB3_1827; GFX90A-NEXT: .LBB3_12: ; %DummyReturnBlock828; GFX90A-NEXT: s_endpgm829bb:830 %i = load volatile i16, ptr addrspace(4) poison, align 2831 %i6 = zext i16 %i to i64832 %i7 = udiv i32 %arg1, %arg2833 %i8 = zext i32 %i7 to i64834 br label %bb9835 836bb9: ; preds = %bb12, %bb837 %i10 = phi i64 [ %arg3, %bb ], [ %i13, %bb12 ]838 %undef = freeze i1 poison839 br i1 %undef, label %bb14, label %bb12840 841bb12: ; preds = %bb58, %bb9842 %i13 = add nuw nsw i64 %i10, %i8843 br label %bb9844 845bb14: ; preds = %bb9846 %i11 = icmp slt i64 %i10, 0847 %i15 = load i64, ptr addrspace(1) null, align 8848 br label %bb16849 850bb16: ; preds = %bb58, %bb14851 %i17 = phi i64 [ %i65, %bb58 ], [ %i15, %bb14 ]852 %i18 = phi <2 x float> [ %i59, %bb58 ], [ zeroinitializer, %bb14 ]853 %i19 = phi <2 x float> [ %i60, %bb58 ], [ zeroinitializer, %bb14 ]854 %i20 = phi <2 x float> [ %i61, %bb58 ], [ zeroinitializer, %bb14 ]855 %i21 = phi <2 x float> [ %i62, %bb58 ], [ zeroinitializer, %bb14 ]856 %i22 = add nsw i64 %i17, 1857 %i23 = mul nsw i64 %i22, %arg858 %i24 = add nsw i64 %i23, %i10859 %i25 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 8860 %i27 = load volatile <2 x half>, ptr addrspace(1) %i25, align 16861 %i28 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 10862 %i30 = load volatile <2 x half>, ptr addrspace(1) %i28, align 4863 %i31 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 12864 %i33 = load volatile <2 x half>, ptr addrspace(1) %i31, align 8865 %i34 = getelementptr inbounds [16 x half], ptr addrspace(1) null, i64 %i24, i64 14866 %i36 = load volatile <2 x half>, ptr addrspace(1) %i34, align 4867 %i43 = load volatile <2 x float>, ptr addrspace(3) null, align 8868 %i46 = load volatile <2 x float>, ptr addrspace(3) poison, align 32869 fence syncscope("workgroup") acquire870 br i1 %i11, label %bb58, label %bb51871 872bb51: ; preds = %bb16873 %i37 = fpext <2 x half> %arg4 to <2 x float>874 %i39 = fpext <2 x half> %i27 to <2 x float>875 %i40 = fpext <2 x half> %i30 to <2 x float>876 %i41 = fpext <2 x half> %i33 to <2 x float>877 %i42 = fpext <2 x half> %i36 to <2 x float>878 %i44 = fadd contract <2 x float> %i37, %i43879 %i45 = fadd contract <2 x float> %i43, zeroinitializer880 %i47 = fadd contract <2 x float> %i39, %i46881 %i48 = fadd contract <2 x float> %i40, %i43882 %i49 = fadd contract <2 x float> %i41, zeroinitializer883 %i50 = fadd contract <2 x float> %i42, zeroinitializer884 %i52 = fadd contract <2 x float> %i18, %i44885 %i53 = fadd contract <2 x float> %i19, %i45886 %i54 = fadd contract <2 x float> %i20, %i47887 %i55 = fadd contract <2 x float> %i21, %i48888 %i56 = fadd contract <2 x float> %i49, zeroinitializer889 %i57 = fadd contract <2 x float> %i50, zeroinitializer890 br label %bb58891 892bb58: ; preds = %bb51, %bb16893 %i59 = phi <2 x float> [ %i18, %bb16 ], [ %i52, %bb51 ]894 %i60 = phi <2 x float> [ %i19, %bb16 ], [ %i53, %bb51 ]895 %i61 = phi <2 x float> [ %i20, %bb16 ], [ %i54, %bb51 ]896 %i62 = phi <2 x float> [ %i21, %bb16 ], [ %i55, %bb51 ]897 %i63 = phi <2 x float> [ zeroinitializer, %bb16 ], [ %i56, %bb51 ]898 %i64 = phi <2 x float> [ zeroinitializer, %bb16 ], [ %i57, %bb51 ]899 %i65 = add nsw i64 %i17, %i6900 %i66 = icmp slt i64 %i65, 0901 br i1 %i66, label %bb16, label %bb12902}903 904; This testcase would fail on GFX908 due to not having a free VGPR available to905; copy SGPR to AGPR.906define void @no_free_vgprs_at_sgpr_to_agpr_copy(float %v0, float %v1) #0 {907; GFX908-LABEL: no_free_vgprs_at_sgpr_to_agpr_copy:908; GFX908: ; %bb.0:909; GFX908-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)910; GFX908-NEXT: v_mov_b32_e32 v32, v1911; GFX908-NEXT: v_mov_b32_e32 v33, v0912; GFX908-NEXT: ;;#ASMSTART913; GFX908-NEXT: ; def v[0:31] s[0:15]914; GFX908-NEXT: ;;#ASMEND915; GFX908-NEXT: v_mov_b32_e32 v39, s15916; GFX908-NEXT: s_nop 1917; GFX908-NEXT: v_accvgpr_write_b32 a31, v39918; GFX908-NEXT: v_mov_b32_e32 v39, s14919; GFX908-NEXT: s_nop 1920; GFX908-NEXT: v_accvgpr_write_b32 a30, v39921; GFX908-NEXT: v_mov_b32_e32 v39, s13922; GFX908-NEXT: s_nop 1923; GFX908-NEXT: v_accvgpr_write_b32 a29, v39924; GFX908-NEXT: v_mov_b32_e32 v39, s12925; GFX908-NEXT: s_nop 1926; GFX908-NEXT: v_accvgpr_write_b32 a28, v39927; GFX908-NEXT: v_mov_b32_e32 v39, s11928; GFX908-NEXT: s_nop 1929; GFX908-NEXT: v_accvgpr_write_b32 a27, v39930; GFX908-NEXT: v_mov_b32_e32 v39, s10931; GFX908-NEXT: s_nop 1932; GFX908-NEXT: v_accvgpr_write_b32 a26, v39933; GFX908-NEXT: v_mov_b32_e32 v39, s9934; GFX908-NEXT: s_nop 1935; GFX908-NEXT: v_accvgpr_write_b32 a25, v39936; GFX908-NEXT: v_mov_b32_e32 v39, s8937; GFX908-NEXT: s_nop 1938; GFX908-NEXT: v_accvgpr_write_b32 a24, v39939; GFX908-NEXT: v_mov_b32_e32 v39, s7940; GFX908-NEXT: s_nop 1941; GFX908-NEXT: v_accvgpr_write_b32 a23, v39942; GFX908-NEXT: v_mov_b32_e32 v39, s6943; GFX908-NEXT: s_nop 1944; GFX908-NEXT: v_accvgpr_write_b32 a22, v39945; GFX908-NEXT: v_mov_b32_e32 v39, s5946; GFX908-NEXT: s_nop 1947; GFX908-NEXT: v_accvgpr_write_b32 a21, v39948; GFX908-NEXT: v_mov_b32_e32 v39, s4949; GFX908-NEXT: s_nop 1950; GFX908-NEXT: v_accvgpr_write_b32 a20, v39951; GFX908-NEXT: v_mov_b32_e32 v39, s3952; GFX908-NEXT: s_nop 1953; GFX908-NEXT: v_accvgpr_write_b32 a19, v39954; GFX908-NEXT: v_mov_b32_e32 v39, s2955; GFX908-NEXT: s_nop 1956; GFX908-NEXT: v_accvgpr_write_b32 a18, v39957; GFX908-NEXT: v_mov_b32_e32 v39, s1958; GFX908-NEXT: s_nop 1959; GFX908-NEXT: v_accvgpr_write_b32 a17, v39960; GFX908-NEXT: v_mov_b32_e32 v39, s0961; GFX908-NEXT: s_nop 1962; GFX908-NEXT: v_accvgpr_write_b32 a16, v39963; GFX908-NEXT: s_nop 0964; GFX908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]965; GFX908-NEXT: s_nop 9966; GFX908-NEXT: v_accvgpr_read_b32 v39, a0 ; Reload Reuse967; GFX908-NEXT: v_accvgpr_read_b32 v38, a11 ; Reload Reuse968; GFX908-NEXT: v_accvgpr_read_b32 v37, a12 ; Reload Reuse969; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 ; 4-byte Folded Spill970; GFX908-NEXT: v_accvgpr_read_b32 v39, a1 ; Reload Reuse971; GFX908-NEXT: v_accvgpr_read_b32 v36, a13 ; Reload Reuse972; GFX908-NEXT: v_accvgpr_read_b32 v35, a14 ; Reload Reuse973; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill974; GFX908-NEXT: v_accvgpr_read_b32 v39, a2 ; Reload Reuse975; GFX908-NEXT: v_accvgpr_read_b32 v34, a15 ; Reload Reuse976; GFX908-NEXT: s_nop 0977; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill978; GFX908-NEXT: v_accvgpr_read_b32 v39, a3 ; Reload Reuse979; GFX908-NEXT: s_nop 1980; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill981; GFX908-NEXT: v_accvgpr_read_b32 v39, a4 ; Reload Reuse982; GFX908-NEXT: s_nop 1983; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill984; GFX908-NEXT: v_accvgpr_read_b32 v39, a5 ; Reload Reuse985; GFX908-NEXT: s_nop 1986; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill987; GFX908-NEXT: v_accvgpr_read_b32 v39, a6 ; Reload Reuse988; GFX908-NEXT: s_nop 1989; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill990; GFX908-NEXT: v_accvgpr_read_b32 v39, a7 ; Reload Reuse991; GFX908-NEXT: s_nop 1992; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill993; GFX908-NEXT: v_accvgpr_read_b32 v39, a8 ; Reload Reuse994; GFX908-NEXT: s_nop 1995; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill996; GFX908-NEXT: v_accvgpr_read_b32 v39, a9 ; Reload Reuse997; GFX908-NEXT: s_nop 1998; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill999; GFX908-NEXT: v_accvgpr_read_b32 v39, a10 ; Reload Reuse1000; GFX908-NEXT: s_nop 11001; GFX908-NEXT: buffer_store_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill1002; GFX908-NEXT: ;;#ASMSTART1003; GFX908-NEXT: ; copy1004; GFX908-NEXT: ;;#ASMEND1005; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 ; 4-byte Folded Reload1006; GFX908-NEXT: v_accvgpr_read_b32 v33, a11007; GFX908-NEXT: s_waitcnt vmcnt(0)1008; GFX908-NEXT: v_accvgpr_write_b32 a0, v39 ; Reload Reuse1009; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload1010; GFX908-NEXT: v_accvgpr_write_b32 a32, v331011; GFX908-NEXT: v_accvgpr_write_b32 a11, v38 ; Reload Reuse1012; GFX908-NEXT: v_accvgpr_write_b32 a12, v37 ; Reload Reuse1013; GFX908-NEXT: v_accvgpr_write_b32 a13, v36 ; Reload Reuse1014; GFX908-NEXT: v_accvgpr_write_b32 a14, v35 ; Reload Reuse1015; GFX908-NEXT: s_waitcnt vmcnt(0)1016; GFX908-NEXT: v_accvgpr_write_b32 a1, v39 ; Reload Reuse1017; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload1018; GFX908-NEXT: s_waitcnt vmcnt(0)1019; GFX908-NEXT: v_accvgpr_write_b32 a2, v39 ; Reload Reuse1020; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload1021; GFX908-NEXT: s_waitcnt vmcnt(0)1022; GFX908-NEXT: v_accvgpr_write_b32 a3, v39 ; Reload Reuse1023; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload1024; GFX908-NEXT: s_waitcnt vmcnt(0)1025; GFX908-NEXT: v_accvgpr_write_b32 a4, v39 ; Reload Reuse1026; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload1027; GFX908-NEXT: s_waitcnt vmcnt(0)1028; GFX908-NEXT: v_accvgpr_write_b32 a5, v39 ; Reload Reuse1029; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload1030; GFX908-NEXT: s_waitcnt vmcnt(0)1031; GFX908-NEXT: v_accvgpr_write_b32 a6, v39 ; Reload Reuse1032; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload1033; GFX908-NEXT: s_waitcnt vmcnt(0)1034; GFX908-NEXT: v_accvgpr_write_b32 a7, v39 ; Reload Reuse1035; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload1036; GFX908-NEXT: s_waitcnt vmcnt(0)1037; GFX908-NEXT: v_accvgpr_write_b32 a8, v39 ; Reload Reuse1038; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload1039; GFX908-NEXT: s_waitcnt vmcnt(0)1040; GFX908-NEXT: v_accvgpr_write_b32 a9, v39 ; Reload Reuse1041; GFX908-NEXT: buffer_load_dword v39, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload1042; GFX908-NEXT: s_waitcnt vmcnt(0)1043; GFX908-NEXT: v_accvgpr_write_b32 a10, v39 ; Reload Reuse1044; GFX908-NEXT: v_accvgpr_write_b32 a15, v34 ; Reload Reuse1045; GFX908-NEXT: ;;#ASMSTART1046; GFX908-NEXT: ; copy1047; GFX908-NEXT: ;;#ASMEND1048; GFX908-NEXT: v_accvgpr_read_b32 v39, a21049; GFX908-NEXT: s_nop 11050; GFX908-NEXT: v_accvgpr_write_b32 a3, v391051; GFX908-NEXT: ;;#ASMSTART1052; GFX908-NEXT: ; use a3 v[0:31]1053; GFX908-NEXT: ;;#ASMEND1054; GFX908-NEXT: s_setpc_b64 s[30:31]1055;1056; GFX90A-LABEL: no_free_vgprs_at_sgpr_to_agpr_copy:1057; GFX90A: ; %bb.0:1058; GFX90A-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)1059; GFX90A-NEXT: v_mov_b32_e32 v33, v01060; GFX90A-NEXT: v_mov_b32_e32 v32, v11061; GFX90A-NEXT: ;;#ASMSTART1062; GFX90A-NEXT: ; def v[0:31] s[0:15]1063; GFX90A-NEXT: ;;#ASMEND1064; GFX90A-NEXT: v_accvgpr_write_b32 a31, s151065; GFX90A-NEXT: v_accvgpr_write_b32 a30, s141066; GFX90A-NEXT: v_accvgpr_write_b32 a29, s131067; GFX90A-NEXT: v_accvgpr_write_b32 a28, s121068; GFX90A-NEXT: v_accvgpr_write_b32 a27, s111069; GFX90A-NEXT: v_accvgpr_write_b32 a26, s101070; GFX90A-NEXT: v_accvgpr_write_b32 a25, s91071; GFX90A-NEXT: v_accvgpr_write_b32 a24, s81072; GFX90A-NEXT: v_accvgpr_write_b32 a23, s71073; GFX90A-NEXT: v_accvgpr_write_b32 a22, s61074; GFX90A-NEXT: v_accvgpr_write_b32 a21, s51075; GFX90A-NEXT: v_accvgpr_write_b32 a20, s41076; GFX90A-NEXT: v_accvgpr_write_b32 a19, s31077; GFX90A-NEXT: v_accvgpr_write_b32 a18, s21078; GFX90A-NEXT: v_accvgpr_write_b32 a17, s11079; GFX90A-NEXT: v_accvgpr_write_b32 a16, s01080; GFX90A-NEXT: v_accvgpr_read_b32 v34, a32 ; Reload Reuse1081; GFX90A-NEXT: s_nop 01082; GFX90A-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v33, v32, a[16:31]1083; GFX90A-NEXT: s_nop 101084; GFX90A-NEXT: buffer_store_dword a0, off, s[0:3], s32 ; 4-byte Folded Spill1085; GFX90A-NEXT: s_nop 01086; GFX90A-NEXT: buffer_store_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Spill1087; GFX90A-NEXT: buffer_store_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Spill1088; GFX90A-NEXT: buffer_store_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Spill1089; GFX90A-NEXT: buffer_store_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Spill1090; GFX90A-NEXT: buffer_store_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Spill1091; GFX90A-NEXT: buffer_store_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Spill1092; GFX90A-NEXT: buffer_store_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Spill1093; GFX90A-NEXT: buffer_store_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Spill1094; GFX90A-NEXT: buffer_store_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill1095; GFX90A-NEXT: buffer_store_dword a10, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill1096; GFX90A-NEXT: v_accvgpr_read_b32 v39, a11 ; Reload Reuse1097; GFX90A-NEXT: v_accvgpr_read_b32 v38, a12 ; Reload Reuse1098; GFX90A-NEXT: v_accvgpr_read_b32 v37, a13 ; Reload Reuse1099; GFX90A-NEXT: v_accvgpr_read_b32 v36, a14 ; Reload Reuse1100; GFX90A-NEXT: v_accvgpr_read_b32 v35, a15 ; Reload Reuse1101; GFX90A-NEXT: ;;#ASMSTART1102; GFX90A-NEXT: ; copy1103; GFX90A-NEXT: ;;#ASMEND1104; GFX90A-NEXT: v_accvgpr_mov_b32 a32, a11105; GFX90A-NEXT: buffer_load_dword a0, off, s[0:3], s32 ; 4-byte Folded Reload1106; GFX90A-NEXT: buffer_load_dword a1, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload1107; GFX90A-NEXT: buffer_load_dword a2, off, s[0:3], s32 offset:8 ; 4-byte Folded Reload1108; GFX90A-NEXT: buffer_load_dword a3, off, s[0:3], s32 offset:12 ; 4-byte Folded Reload1109; GFX90A-NEXT: buffer_load_dword a4, off, s[0:3], s32 offset:16 ; 4-byte Folded Reload1110; GFX90A-NEXT: buffer_load_dword a5, off, s[0:3], s32 offset:20 ; 4-byte Folded Reload1111; GFX90A-NEXT: buffer_load_dword a6, off, s[0:3], s32 offset:24 ; 4-byte Folded Reload1112; GFX90A-NEXT: buffer_load_dword a7, off, s[0:3], s32 offset:28 ; 4-byte Folded Reload1113; GFX90A-NEXT: buffer_load_dword a8, off, s[0:3], s32 offset:32 ; 4-byte Folded Reload1114; GFX90A-NEXT: buffer_load_dword a9, off, s[0:3], s32 offset:36 ; 4-byte Folded Reload1115; GFX90A-NEXT: buffer_load_dword a10, off, s[0:3], s32 offset:40 ; 4-byte Folded Reload1116; GFX90A-NEXT: v_accvgpr_write_b32 a11, v39 ; Reload Reuse1117; GFX90A-NEXT: v_accvgpr_write_b32 a12, v38 ; Reload Reuse1118; GFX90A-NEXT: v_accvgpr_write_b32 a13, v37 ; Reload Reuse1119; GFX90A-NEXT: v_accvgpr_write_b32 a14, v36 ; Reload Reuse1120; GFX90A-NEXT: s_waitcnt vmcnt(0)1121; GFX90A-NEXT: v_accvgpr_write_b32 a15, v35 ; Reload Reuse1122; GFX90A-NEXT: ;;#ASMSTART1123; GFX90A-NEXT: ; copy1124; GFX90A-NEXT: ;;#ASMEND1125; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a21126; GFX90A-NEXT: ;;#ASMSTART1127; GFX90A-NEXT: ; use a3 v[0:31]1128; GFX90A-NEXT: ;;#ASMEND1129; GFX90A-NEXT: v_accvgpr_write_b32 a32, v34 ; Reload Reuse1130; GFX90A-NEXT: s_setpc_b64 s[30:31]1131 %asm = call { <32 x i32>, <16 x float> } asm sideeffect "; def $0 $1","=${v[0:31]},=${s[0:15]}"()1132 %vgpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 01133 %agpr0 = extractvalue { <32 x i32>, <16 x float> } %asm, 11134 %mfma = call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float %v0, float %v1, <16 x float> %agpr0, i32 0, i32 0, i32 0)1135 %agpr1 = call i32 asm sideeffect "; copy ", "={a1},a,~{a[0:15]}"(<16 x float> %agpr0)1136 %agpr2 = call i32 asm sideeffect "; copy ", "={a2},a,{a[0:15]}"(i32 %agpr1, <16 x float> %mfma)1137 call void asm sideeffect "; use $0 $1","{a3},{v[0:31]}"(i32 %agpr2, <32 x i32> %vgpr0)1138 ret void1139}1140 1141declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32 immarg, i32 immarg, i32 immarg) #11142declare i32 @llvm.amdgcn.workitem.id.x() #21143 1144attributes #0 = { "amdgpu-waves-per-eu"="6,6" }1145attributes #1 = { convergent nounwind readnone willreturn }1146attributes #2 = { nounwind readnone willreturn }1147attributes #3 = { "amdgpu-waves-per-eu"="7,7" "amdgpu-agpr-alloc"="0" }1148attributes #4 = { "amdgpu-waves-per-eu"="6,6" "amdgpu-flat-work-group-size"="1024,1024" }1149attributes #5 = { "amdgpu-waves-per-eu"="6,6" "amdgpu-agpr-alloc"="0" }1150