1109 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 < %s | FileCheck -check-prefixes=GFX942 %s3; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefixes=GFX90a %s4; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx1250 < %s | FileCheck -check-prefixes=GFX1250 %s5 6define amdgpu_kernel void @preload_block_count_x(ptr addrspace(1) inreg %out) #0 {7; GFX942-LABEL: preload_block_count_x:8; GFX942: ; %bb.1:9; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x010; GFX942-NEXT: s_load_dword s4, s[0:1], 0x811; GFX942-NEXT: s_waitcnt lgkmcnt(0)12; GFX942-NEXT: s_branch .LBB0_013; GFX942-NEXT: .p2align 814; GFX942-NEXT: ; %bb.2:15; GFX942-NEXT: .LBB0_0:16; GFX942-NEXT: v_mov_b32_e32 v0, 017; GFX942-NEXT: v_mov_b32_e32 v1, s418; GFX942-NEXT: global_store_dword v0, v1, s[2:3]19; GFX942-NEXT: s_endpgm20;21; GFX90a-LABEL: preload_block_count_x:22; GFX90a: ; %bb.1:23; GFX90a-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x024; GFX90a-NEXT: s_load_dword s10, s[4:5], 0x825; GFX90a-NEXT: s_waitcnt lgkmcnt(0)26; GFX90a-NEXT: s_branch .LBB0_027; GFX90a-NEXT: .p2align 828; GFX90a-NEXT: ; %bb.2:29; GFX90a-NEXT: .LBB0_0:30; GFX90a-NEXT: v_mov_b32_e32 v0, 031; GFX90a-NEXT: v_mov_b32_e32 v1, s1032; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]33; GFX90a-NEXT: s_endpgm34;35; GFX1250-LABEL: preload_block_count_x:36; GFX1250: ; %bb.0:37; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 138; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s439; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]40; GFX1250-NEXT: s_endpgm41 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()42 %load = load i32, ptr addrspace(4) %imp_arg_ptr43 store i32 %load, ptr addrspace(1) %out44 ret void45}46 47define amdgpu_kernel void @preload_unused_arg_block_count_x(ptr addrspace(1) inreg %out, i32 inreg) #0 {48; GFX942-LABEL: preload_unused_arg_block_count_x:49; GFX942: ; %bb.1:50; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x051; GFX942-NEXT: s_load_dwordx2 s[4:5], s[0:1], 0x852; GFX942-NEXT: s_load_dword s6, s[0:1], 0x1053; GFX942-NEXT: s_waitcnt lgkmcnt(0)54; GFX942-NEXT: s_branch .LBB1_055; GFX942-NEXT: .p2align 856; GFX942-NEXT: ; %bb.2:57; GFX942-NEXT: .LBB1_0:58; GFX942-NEXT: v_mov_b32_e32 v0, 059; GFX942-NEXT: v_mov_b32_e32 v1, s660; GFX942-NEXT: global_store_dword v0, v1, s[2:3]61; GFX942-NEXT: s_endpgm62;63; GFX90a-LABEL: preload_unused_arg_block_count_x:64; GFX90a: ; %bb.1:65; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x066; GFX90a-NEXT: s_load_dword s12, s[4:5], 0x1067; GFX90a-NEXT: s_waitcnt lgkmcnt(0)68; GFX90a-NEXT: s_branch .LBB1_069; GFX90a-NEXT: .p2align 870; GFX90a-NEXT: ; %bb.2:71; GFX90a-NEXT: .LBB1_0:72; GFX90a-NEXT: v_mov_b32_e32 v0, 073; GFX90a-NEXT: v_mov_b32_e32 v1, s1274; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]75; GFX90a-NEXT: s_endpgm76;77; GFX1250-LABEL: preload_unused_arg_block_count_x:78; GFX1250: ; %bb.0:79; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 180; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s681; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]82; GFX1250-NEXT: s_endpgm83 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()84 %load = load i32, ptr addrspace(4) %imp_arg_ptr85 store i32 %load, ptr addrspace(1) %out86 ret void87}88 89define amdgpu_kernel void @no_free_sgprs_block_count_x(ptr addrspace(1) inreg %out, i256 inreg) {90; GFX942-LABEL: no_free_sgprs_block_count_x:91; GFX942: ; %bb.1:92; GFX942-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x093; GFX942-NEXT: s_waitcnt lgkmcnt(0)94; GFX942-NEXT: s_branch .LBB2_095; GFX942-NEXT: .p2align 896; GFX942-NEXT: ; %bb.2:97; GFX942-NEXT: .LBB2_0:98; GFX942-NEXT: s_load_dword s0, s[4:5], 0x2899; GFX942-NEXT: v_mov_b32_e32 v0, 0100; GFX942-NEXT: s_waitcnt lgkmcnt(0)101; GFX942-NEXT: v_mov_b32_e32 v1, s0102; GFX942-NEXT: global_store_dword v0, v1, s[8:9]103; GFX942-NEXT: s_endpgm104;105; GFX90a-LABEL: no_free_sgprs_block_count_x:106; GFX90a: ; %bb.1:107; GFX90a-NEXT: s_load_dwordx2 s[14:15], s[8:9], 0x0108; GFX90a-NEXT: s_waitcnt lgkmcnt(0)109; GFX90a-NEXT: s_branch .LBB2_0110; GFX90a-NEXT: .p2align 8111; GFX90a-NEXT: ; %bb.2:112; GFX90a-NEXT: .LBB2_0:113; GFX90a-NEXT: s_load_dword s0, s[8:9], 0x28114; GFX90a-NEXT: v_mov_b32_e32 v0, 0115; GFX90a-NEXT: s_waitcnt lgkmcnt(0)116; GFX90a-NEXT: v_mov_b32_e32 v1, s0117; GFX90a-NEXT: global_store_dword v0, v1, s[14:15]118; GFX90a-NEXT: s_endpgm119;120; GFX1250-LABEL: no_free_sgprs_block_count_x:121; GFX1250: ; %bb.0:122; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1123; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s18124; GFX1250-NEXT: global_store_b32 v0, v1, s[8:9]125; GFX1250-NEXT: s_endpgm126 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()127 %load = load i32, ptr addrspace(4) %imp_arg_ptr128 store i32 %load, ptr addrspace(1) %out129 ret void130}131 132define amdgpu_kernel void @no_inreg_block_count_x(ptr addrspace(1) %out) #0 {133; GFX942-LABEL: no_inreg_block_count_x:134; GFX942: ; %bb.0:135; GFX942-NEXT: s_load_dword s4, s[0:1], 0x8136; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0137; GFX942-NEXT: v_mov_b32_e32 v0, 0138; GFX942-NEXT: s_waitcnt lgkmcnt(0)139; GFX942-NEXT: v_mov_b32_e32 v1, s4140; GFX942-NEXT: global_store_dword v0, v1, s[2:3]141; GFX942-NEXT: s_endpgm142;143; GFX90a-LABEL: no_inreg_block_count_x:144; GFX90a: ; %bb.0:145; GFX90a-NEXT: s_load_dword s2, s[4:5], 0x8146; GFX90a-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0147; GFX90a-NEXT: v_mov_b32_e32 v0, 0148; GFX90a-NEXT: s_waitcnt lgkmcnt(0)149; GFX90a-NEXT: v_mov_b32_e32 v1, s2150; GFX90a-NEXT: global_store_dword v0, v1, s[0:1]151; GFX90a-NEXT: s_endpgm152;153; GFX1250-LABEL: no_inreg_block_count_x:154; GFX1250: ; %bb.0:155; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1156; GFX1250-NEXT: s_load_b96 s[0:2], s[0:1], 0x0157; GFX1250-NEXT: s_wait_kmcnt 0x0158; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2159; GFX1250-NEXT: global_store_b32 v0, v1, s[0:1]160; GFX1250-NEXT: s_endpgm161 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()162 %load = load i32, ptr addrspace(4) %imp_arg_ptr163 store i32 %load, ptr addrspace(1) %out164 ret void165}166 167; Implicit arg preloading is currently restricted to cases where all explicit168; args are inreg (preloaded).169 170define amdgpu_kernel void @mixed_inreg_block_count_x(ptr addrspace(1) %out, i32 inreg) #0 {171; GFX942-LABEL: mixed_inreg_block_count_x:172; GFX942: ; %bb.0:173; GFX942-NEXT: s_load_dword s4, s[0:1], 0x10174; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0175; GFX942-NEXT: v_mov_b32_e32 v0, 0176; GFX942-NEXT: s_waitcnt lgkmcnt(0)177; GFX942-NEXT: v_mov_b32_e32 v1, s4178; GFX942-NEXT: global_store_dword v0, v1, s[2:3]179; GFX942-NEXT: s_endpgm180;181; GFX90a-LABEL: mixed_inreg_block_count_x:182; GFX90a: ; %bb.0:183; GFX90a-NEXT: s_load_dword s2, s[4:5], 0x10184; GFX90a-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0185; GFX90a-NEXT: v_mov_b32_e32 v0, 0186; GFX90a-NEXT: s_waitcnt lgkmcnt(0)187; GFX90a-NEXT: v_mov_b32_e32 v1, s2188; GFX90a-NEXT: global_store_dword v0, v1, s[0:1]189; GFX90a-NEXT: s_endpgm190;191; GFX1250-LABEL: mixed_inreg_block_count_x:192; GFX1250: ; %bb.0:193; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1194; GFX1250-NEXT: s_clause 0x1195; GFX1250-NEXT: s_load_b32 s4, s[0:1], 0x10196; GFX1250-NEXT: s_load_b64 s[2:3], s[0:1], 0x0197; GFX1250-NEXT: s_wait_kmcnt 0x0198; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s4199; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]200; GFX1250-NEXT: s_endpgm201 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()202 %load = load i32, ptr addrspace(4) %imp_arg_ptr203 store i32 %load, ptr addrspace(1) %out204 ret void205}206 207define amdgpu_kernel void @incorrect_type_i64_block_count_x(ptr addrspace(1) inreg %out) #0 {208; GFX942-LABEL: incorrect_type_i64_block_count_x:209; GFX942: ; %bb.1:210; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0211; GFX942-NEXT: s_waitcnt lgkmcnt(0)212; GFX942-NEXT: s_branch .LBB5_0213; GFX942-NEXT: .p2align 8214; GFX942-NEXT: ; %bb.2:215; GFX942-NEXT: .LBB5_0:216; GFX942-NEXT: s_load_dwordx2 s[0:1], s[0:1], 0x8217; GFX942-NEXT: v_mov_b32_e32 v0, 0218; GFX942-NEXT: s_waitcnt lgkmcnt(0)219; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[0:1]220; GFX942-NEXT: global_store_dwordx2 v0, v[2:3], s[2:3]221; GFX942-NEXT: s_endpgm222;223; GFX90a-LABEL: incorrect_type_i64_block_count_x:224; GFX90a: ; %bb.1:225; GFX90a-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0226; GFX90a-NEXT: s_waitcnt lgkmcnt(0)227; GFX90a-NEXT: s_branch .LBB5_0228; GFX90a-NEXT: .p2align 8229; GFX90a-NEXT: ; %bb.2:230; GFX90a-NEXT: .LBB5_0:231; GFX90a-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x8232; GFX90a-NEXT: v_mov_b32_e32 v0, 0233; GFX90a-NEXT: s_waitcnt lgkmcnt(0)234; GFX90a-NEXT: v_pk_mov_b32 v[2:3], s[0:1], s[0:1] op_sel:[0,1]235; GFX90a-NEXT: global_store_dwordx2 v0, v[2:3], s[8:9]236; GFX90a-NEXT: s_endpgm237;238; GFX1250-LABEL: incorrect_type_i64_block_count_x:239; GFX1250: ; %bb.0:240; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1241; GFX1250-NEXT: s_load_b64 s[0:1], s[0:1], 0x8242; GFX1250-NEXT: v_mov_b32_e32 v2, 0243; GFX1250-NEXT: s_wait_kmcnt 0x0244; GFX1250-NEXT: v_mov_b64_e32 v[0:1], s[0:1]245; GFX1250-NEXT: global_store_b64 v2, v[0:1], s[2:3]246; GFX1250-NEXT: s_endpgm247 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()248 %load = load i64, ptr addrspace(4) %imp_arg_ptr249 store i64 %load, ptr addrspace(1) %out250 ret void251}252 253define amdgpu_kernel void @incorrect_type_i16_block_count_x(ptr addrspace(1) inreg %out) #0 {254; GFX942-LABEL: incorrect_type_i16_block_count_x:255; GFX942: ; %bb.1:256; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0257; GFX942-NEXT: s_waitcnt lgkmcnt(0)258; GFX942-NEXT: s_branch .LBB6_0259; GFX942-NEXT: .p2align 8260; GFX942-NEXT: ; %bb.2:261; GFX942-NEXT: .LBB6_0:262; GFX942-NEXT: s_load_dword s0, s[0:1], 0x8263; GFX942-NEXT: v_mov_b32_e32 v0, 0264; GFX942-NEXT: s_waitcnt lgkmcnt(0)265; GFX942-NEXT: v_mov_b32_e32 v1, s0266; GFX942-NEXT: global_store_short v0, v1, s[2:3]267; GFX942-NEXT: s_endpgm268;269; GFX90a-LABEL: incorrect_type_i16_block_count_x:270; GFX90a: ; %bb.1:271; GFX90a-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0272; GFX90a-NEXT: s_waitcnt lgkmcnt(0)273; GFX90a-NEXT: s_branch .LBB6_0274; GFX90a-NEXT: .p2align 8275; GFX90a-NEXT: ; %bb.2:276; GFX90a-NEXT: .LBB6_0:277; GFX90a-NEXT: s_load_dword s0, s[4:5], 0x8278; GFX90a-NEXT: v_mov_b32_e32 v0, 0279; GFX90a-NEXT: s_waitcnt lgkmcnt(0)280; GFX90a-NEXT: v_mov_b32_e32 v1, s0281; GFX90a-NEXT: global_store_short v0, v1, s[8:9]282; GFX90a-NEXT: s_endpgm283;284; GFX1250-LABEL: incorrect_type_i16_block_count_x:285; GFX1250: ; %bb.0:286; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1287; GFX1250-NEXT: v_mov_b32_e32 v0, 0288; GFX1250-NEXT: global_load_u16 v1, v0, s[0:1] offset:8289; GFX1250-NEXT: s_wait_loadcnt 0x0290; GFX1250-NEXT: global_store_b16 v0, v1, s[2:3]291; GFX1250-NEXT: s_endpgm292 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()293 %load = load i16, ptr addrspace(4) %imp_arg_ptr294 store i16 %load, ptr addrspace(1) %out295 ret void296}297 298define amdgpu_kernel void @preload_block_count_y(ptr addrspace(1) inreg %out) #0 {299; GFX942-LABEL: preload_block_count_y:300; GFX942: ; %bb.1:301; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0302; GFX942-NEXT: s_load_dwordx2 s[4:5], s[0:1], 0x8303; GFX942-NEXT: s_waitcnt lgkmcnt(0)304; GFX942-NEXT: s_branch .LBB7_0305; GFX942-NEXT: .p2align 8306; GFX942-NEXT: ; %bb.2:307; GFX942-NEXT: .LBB7_0:308; GFX942-NEXT: v_mov_b32_e32 v0, 0309; GFX942-NEXT: v_mov_b32_e32 v1, s5310; GFX942-NEXT: global_store_dword v0, v1, s[2:3]311; GFX942-NEXT: s_endpgm312;313; GFX90a-LABEL: preload_block_count_y:314; GFX90a: ; %bb.1:315; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0316; GFX90a-NEXT: s_waitcnt lgkmcnt(0)317; GFX90a-NEXT: s_branch .LBB7_0318; GFX90a-NEXT: .p2align 8319; GFX90a-NEXT: ; %bb.2:320; GFX90a-NEXT: .LBB7_0:321; GFX90a-NEXT: v_mov_b32_e32 v0, 0322; GFX90a-NEXT: v_mov_b32_e32 v1, s11323; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]324; GFX90a-NEXT: s_endpgm325;326; GFX1250-LABEL: preload_block_count_y:327; GFX1250: ; %bb.0:328; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1329; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s5330; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]331; GFX1250-NEXT: s_endpgm332 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()333 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 4334 %load = load i32, ptr addrspace(4) %gep335 store i32 %load, ptr addrspace(1) %out336 ret void337}338 339define amdgpu_kernel void @random_incorrect_offset(ptr addrspace(1) inreg %out) #0 {340; GFX942-LABEL: random_incorrect_offset:341; GFX942: ; %bb.1:342; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0343; GFX942-NEXT: s_waitcnt lgkmcnt(0)344; GFX942-NEXT: s_branch .LBB8_0345; GFX942-NEXT: .p2align 8346; GFX942-NEXT: ; %bb.2:347; GFX942-NEXT: .LBB8_0:348; GFX942-NEXT: s_load_dword s0, s[0:1], 0xa349; GFX942-NEXT: v_mov_b32_e32 v0, 0350; GFX942-NEXT: s_waitcnt lgkmcnt(0)351; GFX942-NEXT: v_mov_b32_e32 v1, s0352; GFX942-NEXT: global_store_dword v0, v1, s[2:3]353; GFX942-NEXT: s_endpgm354;355; GFX90a-LABEL: random_incorrect_offset:356; GFX90a: ; %bb.1:357; GFX90a-NEXT: s_load_dwordx2 s[8:9], s[4:5], 0x0358; GFX90a-NEXT: s_waitcnt lgkmcnt(0)359; GFX90a-NEXT: s_branch .LBB8_0360; GFX90a-NEXT: .p2align 8361; GFX90a-NEXT: ; %bb.2:362; GFX90a-NEXT: .LBB8_0:363; GFX90a-NEXT: s_load_dword s0, s[4:5], 0xa364; GFX90a-NEXT: v_mov_b32_e32 v0, 0365; GFX90a-NEXT: s_waitcnt lgkmcnt(0)366; GFX90a-NEXT: v_mov_b32_e32 v1, s0367; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]368; GFX90a-NEXT: s_endpgm369;370; GFX1250-LABEL: random_incorrect_offset:371; GFX1250: ; %bb.0:372; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1373; GFX1250-NEXT: s_load_b32 s0, s[0:1], 0xa374; GFX1250-NEXT: s_wait_kmcnt 0x0375; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0376; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]377; GFX1250-NEXT: s_endpgm378 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()379 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 2380 %load = load i32, ptr addrspace(4) %gep381 store i32 %load, ptr addrspace(1) %out382 ret void383}384 385define amdgpu_kernel void @preload_block_count_z(ptr addrspace(1) inreg %out) #0 {386; GFX942-LABEL: preload_block_count_z:387; GFX942: ; %bb.1:388; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0389; GFX942-NEXT: s_load_dwordx2 s[4:5], s[0:1], 0x8390; GFX942-NEXT: s_load_dword s6, s[0:1], 0x10391; GFX942-NEXT: s_waitcnt lgkmcnt(0)392; GFX942-NEXT: s_branch .LBB9_0393; GFX942-NEXT: .p2align 8394; GFX942-NEXT: ; %bb.2:395; GFX942-NEXT: .LBB9_0:396; GFX942-NEXT: v_mov_b32_e32 v0, 0397; GFX942-NEXT: v_mov_b32_e32 v1, s6398; GFX942-NEXT: global_store_dword v0, v1, s[2:3]399; GFX942-NEXT: s_endpgm400;401; GFX90a-LABEL: preload_block_count_z:402; GFX90a: ; %bb.1:403; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0404; GFX90a-NEXT: s_load_dword s12, s[4:5], 0x10405; GFX90a-NEXT: s_waitcnt lgkmcnt(0)406; GFX90a-NEXT: s_branch .LBB9_0407; GFX90a-NEXT: .p2align 8408; GFX90a-NEXT: ; %bb.2:409; GFX90a-NEXT: .LBB9_0:410; GFX90a-NEXT: v_mov_b32_e32 v0, 0411; GFX90a-NEXT: v_mov_b32_e32 v1, s12412; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]413; GFX90a-NEXT: s_endpgm414;415; GFX1250-LABEL: preload_block_count_z:416; GFX1250: ; %bb.0:417; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1418; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s6419; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]420; GFX1250-NEXT: s_endpgm421 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()422 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 8423 %load = load i32, ptr addrspace(4) %gep424 store i32 %load, ptr addrspace(1) %out425 ret void426}427 428define amdgpu_kernel void @preload_block_count_x_imparg_align_ptr_i8(ptr addrspace(1) inreg %out, i8 inreg %val) #0 {429; GFX942-LABEL: preload_block_count_x_imparg_align_ptr_i8:430; GFX942: ; %bb.1:431; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0432; GFX942-NEXT: s_load_dwordx2 s[4:5], s[0:1], 0x8433; GFX942-NEXT: s_load_dword s6, s[0:1], 0x10434; GFX942-NEXT: s_waitcnt lgkmcnt(0)435; GFX942-NEXT: s_branch .LBB10_0436; GFX942-NEXT: .p2align 8437; GFX942-NEXT: ; %bb.2:438; GFX942-NEXT: .LBB10_0:439; GFX942-NEXT: s_and_b32 s0, s4, 0xff440; GFX942-NEXT: s_add_i32 s0, s6, s0441; GFX942-NEXT: v_mov_b32_e32 v0, 0442; GFX942-NEXT: v_mov_b32_e32 v1, s0443; GFX942-NEXT: global_store_dword v0, v1, s[2:3]444; GFX942-NEXT: s_endpgm445;446; GFX90a-LABEL: preload_block_count_x_imparg_align_ptr_i8:447; GFX90a: ; %bb.1:448; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0449; GFX90a-NEXT: s_load_dword s12, s[4:5], 0x10450; GFX90a-NEXT: s_waitcnt lgkmcnt(0)451; GFX90a-NEXT: s_branch .LBB10_0452; GFX90a-NEXT: .p2align 8453; GFX90a-NEXT: ; %bb.2:454; GFX90a-NEXT: .LBB10_0:455; GFX90a-NEXT: s_and_b32 s0, s10, 0xff456; GFX90a-NEXT: s_add_i32 s0, s12, s0457; GFX90a-NEXT: v_mov_b32_e32 v0, 0458; GFX90a-NEXT: v_mov_b32_e32 v1, s0459; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]460; GFX90a-NEXT: s_endpgm461;462; GFX1250-LABEL: preload_block_count_x_imparg_align_ptr_i8:463; GFX1250: ; %bb.0:464; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1465; GFX1250-NEXT: s_and_b32 s0, s4, 0xff466; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)467; GFX1250-NEXT: s_add_co_i32 s0, s6, s0468; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0469; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]470; GFX1250-NEXT: s_endpgm471 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()472 %load = load i32, ptr addrspace(4) %imp_arg_ptr473 %ext = zext i8 %val to i32474 %add = add i32 %load, %ext475 store i32 %add, ptr addrspace(1) %out476 ret void477}478 479define amdgpu_kernel void @preload_block_count_xyz(ptr addrspace(1) inreg %out) #0 {480; GFX942-LABEL: preload_block_count_xyz:481; GFX942: ; %bb.1:482; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0483; GFX942-NEXT: s_load_dwordx2 s[4:5], s[0:1], 0x8484; GFX942-NEXT: s_load_dword s6, s[0:1], 0x10485; GFX942-NEXT: s_waitcnt lgkmcnt(0)486; GFX942-NEXT: s_branch .LBB11_0487; GFX942-NEXT: .p2align 8488; GFX942-NEXT: ; %bb.2:489; GFX942-NEXT: .LBB11_0:490; GFX942-NEXT: v_mov_b32_e32 v3, 0491; GFX942-NEXT: v_mov_b32_e32 v0, s4492; GFX942-NEXT: v_mov_b32_e32 v1, s5493; GFX942-NEXT: v_mov_b32_e32 v2, s6494; GFX942-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3]495; GFX942-NEXT: s_endpgm496;497; GFX90a-LABEL: preload_block_count_xyz:498; GFX90a: ; %bb.1:499; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0500; GFX90a-NEXT: s_load_dword s12, s[4:5], 0x10501; GFX90a-NEXT: s_waitcnt lgkmcnt(0)502; GFX90a-NEXT: s_branch .LBB11_0503; GFX90a-NEXT: .p2align 8504; GFX90a-NEXT: ; %bb.2:505; GFX90a-NEXT: .LBB11_0:506; GFX90a-NEXT: v_mov_b32_e32 v3, 0507; GFX90a-NEXT: v_mov_b32_e32 v0, s10508; GFX90a-NEXT: v_mov_b32_e32 v1, s11509; GFX90a-NEXT: v_mov_b32_e32 v2, s12510; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]511; GFX90a-NEXT: s_endpgm512;513; GFX1250-LABEL: preload_block_count_xyz:514; GFX1250: ; %bb.0:515; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1516; GFX1250-NEXT: v_dual_mov_b32 v3, 0 :: v_dual_mov_b32 v0, s4517; GFX1250-NEXT: v_dual_mov_b32 v1, s5 :: v_dual_mov_b32 v2, s6518; GFX1250-NEXT: global_store_b96 v3, v[0:2], s[2:3]519; GFX1250-NEXT: s_endpgm520 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()521 %gep_x = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 0522 %load_x = load i32, ptr addrspace(4) %gep_x523 %gep_y = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 4524 %load_y = load i32, ptr addrspace(4) %gep_y525 %gep_z = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 8526 %load_z = load i32, ptr addrspace(4) %gep_z527 %ins.0 = insertelement <3 x i32> poison, i32 %load_x, i32 0528 %ins.1 = insertelement <3 x i32> %ins.0, i32 %load_y, i32 1529 %ins.2 = insertelement <3 x i32> %ins.1, i32 %load_z, i32 2530 store <3 x i32> %ins.2, ptr addrspace(1) %out531 ret void532}533 534define amdgpu_kernel void @preload_workgroup_size_x(ptr addrspace(1) inreg %out) #0 {535; GFX942-LABEL: preload_workgroup_size_x:536; GFX942: ; %bb.1:537; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0538; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8539; GFX942-NEXT: s_waitcnt lgkmcnt(0)540; GFX942-NEXT: s_branch .LBB12_0541; GFX942-NEXT: .p2align 8542; GFX942-NEXT: ; %bb.2:543; GFX942-NEXT: .LBB12_0:544; GFX942-NEXT: s_and_b32 s0, s7, 0xffff545; GFX942-NEXT: v_mov_b32_e32 v0, 0546; GFX942-NEXT: v_mov_b32_e32 v1, s0547; GFX942-NEXT: global_store_dword v0, v1, s[2:3]548; GFX942-NEXT: s_endpgm549;550; GFX90a-LABEL: preload_workgroup_size_x:551; GFX90a: ; %bb.1:552; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0553; GFX90a-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10554; GFX90a-NEXT: s_waitcnt lgkmcnt(0)555; GFX90a-NEXT: s_branch .LBB12_0556; GFX90a-NEXT: .p2align 8557; GFX90a-NEXT: ; %bb.2:558; GFX90a-NEXT: .LBB12_0:559; GFX90a-NEXT: s_and_b32 s0, s13, 0xffff560; GFX90a-NEXT: v_mov_b32_e32 v0, 0561; GFX90a-NEXT: v_mov_b32_e32 v1, s0562; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]563; GFX90a-NEXT: s_endpgm564;565; GFX1250-LABEL: preload_workgroup_size_x:566; GFX1250: ; %bb.0:567; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1568; GFX1250-NEXT: s_and_b32 s0, s7, 0xffff569; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)570; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0571; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]572; GFX1250-NEXT: s_endpgm573 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()574 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 12575 %load = load i16, ptr addrspace(4) %gep576 %conv = zext i16 %load to i32577 store i32 %conv, ptr addrspace(1) %out578 ret void579}580 581define amdgpu_kernel void @preload_workgroup_size_y(ptr addrspace(1) inreg %out) #0 {582; GFX942-LABEL: preload_workgroup_size_y:583; GFX942: ; %bb.1:584; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0585; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8586; GFX942-NEXT: s_waitcnt lgkmcnt(0)587; GFX942-NEXT: s_branch .LBB13_0588; GFX942-NEXT: .p2align 8589; GFX942-NEXT: ; %bb.2:590; GFX942-NEXT: .LBB13_0:591; GFX942-NEXT: s_lshr_b32 s0, s7, 16592; GFX942-NEXT: v_mov_b32_e32 v0, 0593; GFX942-NEXT: v_mov_b32_e32 v1, s0594; GFX942-NEXT: global_store_dword v0, v1, s[2:3]595; GFX942-NEXT: s_endpgm596;597; GFX90a-LABEL: preload_workgroup_size_y:598; GFX90a: ; %bb.1:599; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0600; GFX90a-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10601; GFX90a-NEXT: s_waitcnt lgkmcnt(0)602; GFX90a-NEXT: s_branch .LBB13_0603; GFX90a-NEXT: .p2align 8604; GFX90a-NEXT: ; %bb.2:605; GFX90a-NEXT: .LBB13_0:606; GFX90a-NEXT: s_lshr_b32 s0, s13, 16607; GFX90a-NEXT: v_mov_b32_e32 v0, 0608; GFX90a-NEXT: v_mov_b32_e32 v1, s0609; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]610; GFX90a-NEXT: s_endpgm611;612; GFX1250-LABEL: preload_workgroup_size_y:613; GFX1250: ; %bb.0:614; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1615; GFX1250-NEXT: s_lshr_b32 s0, s7, 16616; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)617; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0618; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]619; GFX1250-NEXT: s_endpgm620 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()621 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 14622 %load = load i16, ptr addrspace(4) %gep623 %conv = zext i16 %load to i32624 store i32 %conv, ptr addrspace(1) %out625 ret void626}627 628define amdgpu_kernel void @preload_workgroup_size_z(ptr addrspace(1) inreg %out) #0 {629; GFX942-LABEL: preload_workgroup_size_z:630; GFX942: ; %bb.1:631; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0632; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8633; GFX942-NEXT: s_load_dword s8, s[0:1], 0x18634; GFX942-NEXT: s_waitcnt lgkmcnt(0)635; GFX942-NEXT: s_branch .LBB14_0636; GFX942-NEXT: .p2align 8637; GFX942-NEXT: ; %bb.2:638; GFX942-NEXT: .LBB14_0:639; GFX942-NEXT: s_and_b32 s0, s8, 0xffff640; GFX942-NEXT: v_mov_b32_e32 v0, 0641; GFX942-NEXT: v_mov_b32_e32 v1, s0642; GFX942-NEXT: global_store_dword v0, v1, s[2:3]643; GFX942-NEXT: s_endpgm644;645; GFX90a-LABEL: preload_workgroup_size_z:646; GFX90a: ; %bb.1:647; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0648; GFX90a-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10649; GFX90a-NEXT: s_load_dword s14, s[4:5], 0x18650; GFX90a-NEXT: s_waitcnt lgkmcnt(0)651; GFX90a-NEXT: s_branch .LBB14_0652; GFX90a-NEXT: .p2align 8653; GFX90a-NEXT: ; %bb.2:654; GFX90a-NEXT: .LBB14_0:655; GFX90a-NEXT: s_and_b32 s0, s14, 0xffff656; GFX90a-NEXT: v_mov_b32_e32 v0, 0657; GFX90a-NEXT: v_mov_b32_e32 v1, s0658; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]659; GFX90a-NEXT: s_endpgm660;661; GFX1250-LABEL: preload_workgroup_size_z:662; GFX1250: ; %bb.0:663; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1664; GFX1250-NEXT: s_and_b32 s0, s8, 0xffff665; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)666; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0667; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]668; GFX1250-NEXT: s_endpgm669 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()670 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 16671 %load = load i16, ptr addrspace(4) %gep672 %conv = zext i16 %load to i32673 store i32 %conv, ptr addrspace(1) %out674 ret void675}676 677define amdgpu_kernel void @preload_workgroup_size_xyz(ptr addrspace(1) inreg %out) #0 {678; GFX942-LABEL: preload_workgroup_size_xyz:679; GFX942: ; %bb.1:680; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0681; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8682; GFX942-NEXT: s_load_dword s8, s[0:1], 0x18683; GFX942-NEXT: s_waitcnt lgkmcnt(0)684; GFX942-NEXT: s_branch .LBB15_0685; GFX942-NEXT: .p2align 8686; GFX942-NEXT: ; %bb.2:687; GFX942-NEXT: .LBB15_0:688; GFX942-NEXT: s_lshr_b32 s0, s7, 16689; GFX942-NEXT: s_and_b32 s1, s7, 0xffff690; GFX942-NEXT: s_and_b32 s4, s8, 0xffff691; GFX942-NEXT: v_mov_b32_e32 v3, 0692; GFX942-NEXT: v_mov_b32_e32 v0, s1693; GFX942-NEXT: v_mov_b32_e32 v1, s0694; GFX942-NEXT: v_mov_b32_e32 v2, s4695; GFX942-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3]696; GFX942-NEXT: s_endpgm697;698; GFX90a-LABEL: preload_workgroup_size_xyz:699; GFX90a: ; %bb.1:700; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0701; GFX90a-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10702; GFX90a-NEXT: s_load_dword s14, s[4:5], 0x18703; GFX90a-NEXT: s_waitcnt lgkmcnt(0)704; GFX90a-NEXT: s_branch .LBB15_0705; GFX90a-NEXT: .p2align 8706; GFX90a-NEXT: ; %bb.2:707; GFX90a-NEXT: .LBB15_0:708; GFX90a-NEXT: s_lshr_b32 s0, s13, 16709; GFX90a-NEXT: s_and_b32 s1, s13, 0xffff710; GFX90a-NEXT: s_and_b32 s2, s14, 0xffff711; GFX90a-NEXT: v_mov_b32_e32 v3, 0712; GFX90a-NEXT: v_mov_b32_e32 v0, s1713; GFX90a-NEXT: v_mov_b32_e32 v1, s0714; GFX90a-NEXT: v_mov_b32_e32 v2, s2715; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]716; GFX90a-NEXT: s_endpgm717;718; GFX1250-LABEL: preload_workgroup_size_xyz:719; GFX1250: ; %bb.0:720; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1721; GFX1250-NEXT: s_lshr_b32 s0, s7, 16722; GFX1250-NEXT: s_and_b32 s1, s7, 0xffff723; GFX1250-NEXT: s_and_b32 s4, s8, 0xffff724; GFX1250-NEXT: v_dual_mov_b32 v3, 0 :: v_dual_mov_b32 v0, s1725; GFX1250-NEXT: v_dual_mov_b32 v1, s0 :: v_dual_mov_b32 v2, s4726; GFX1250-NEXT: global_store_b96 v3, v[0:2], s[2:3]727; GFX1250-NEXT: s_endpgm728 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()729 %gep_x = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 12730 %load_x = load i16, ptr addrspace(4) %gep_x731 %conv_x = zext i16 %load_x to i32732 %gep_y = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 14733 %load_y = load i16, ptr addrspace(4) %gep_y734 %conv_y = zext i16 %load_y to i32735 %gep_z = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 16736 %load_z = load i16, ptr addrspace(4) %gep_z737 %conv_z = zext i16 %load_z to i32738 %ins.0 = insertelement <3 x i32> poison, i32 %conv_x, i32 0739 %ins.1 = insertelement <3 x i32> %ins.0, i32 %conv_y, i32 1740 %ins.2 = insertelement <3 x i32> %ins.1, i32 %conv_z, i32 2741 store <3 x i32> %ins.2, ptr addrspace(1) %out742 ret void743}744 745define amdgpu_kernel void @preload_remainder_x(ptr addrspace(1) inreg %out) #0 {746; GFX942-LABEL: preload_remainder_x:747; GFX942: ; %bb.1:748; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0749; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8750; GFX942-NEXT: s_load_dword s8, s[0:1], 0x18751; GFX942-NEXT: s_waitcnt lgkmcnt(0)752; GFX942-NEXT: s_branch .LBB16_0753; GFX942-NEXT: .p2align 8754; GFX942-NEXT: ; %bb.2:755; GFX942-NEXT: .LBB16_0:756; GFX942-NEXT: s_lshr_b32 s0, s8, 16757; GFX942-NEXT: v_mov_b32_e32 v0, 0758; GFX942-NEXT: v_mov_b32_e32 v1, s0759; GFX942-NEXT: global_store_dword v0, v1, s[2:3]760; GFX942-NEXT: s_endpgm761;762; GFX90a-LABEL: preload_remainder_x:763; GFX90a: ; %bb.1:764; GFX90a-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x0765; GFX90a-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x10766; GFX90a-NEXT: s_load_dword s14, s[4:5], 0x18767; GFX90a-NEXT: s_waitcnt lgkmcnt(0)768; GFX90a-NEXT: s_branch .LBB16_0769; GFX90a-NEXT: .p2align 8770; GFX90a-NEXT: ; %bb.2:771; GFX90a-NEXT: .LBB16_0:772; GFX90a-NEXT: s_lshr_b32 s0, s14, 16773; GFX90a-NEXT: v_mov_b32_e32 v0, 0774; GFX90a-NEXT: v_mov_b32_e32 v1, s0775; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]776; GFX90a-NEXT: s_endpgm777;778; GFX1250-LABEL: preload_remainder_x:779; GFX1250: ; %bb.0:780; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1781; GFX1250-NEXT: s_lshr_b32 s0, s8, 16782; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)783; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0784; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]785; GFX1250-NEXT: s_endpgm786 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()787 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 18788 %load = load i16, ptr addrspace(4) %gep789 %conv = zext i16 %load to i32790 store i32 %conv, ptr addrspace(1) %out791 ret void792}793 794define amdgpu_kernel void @preloadremainder_y(ptr addrspace(1) inreg %out) #0 {795; GFX942-LABEL: preloadremainder_y:796; GFX942: ; %bb.1:797; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0798; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8799; GFX942-NEXT: s_load_dwordx2 s[8:9], s[0:1], 0x18800; GFX942-NEXT: s_waitcnt lgkmcnt(0)801; GFX942-NEXT: s_branch .LBB17_0802; GFX942-NEXT: .p2align 8803; GFX942-NEXT: ; %bb.2:804; GFX942-NEXT: .LBB17_0:805; GFX942-NEXT: s_and_b32 s0, s9, 0xffff806; GFX942-NEXT: v_mov_b32_e32 v0, 0807; GFX942-NEXT: v_mov_b32_e32 v1, s0808; GFX942-NEXT: global_store_dword v0, v1, s[2:3]809; GFX942-NEXT: s_endpgm810;811; GFX90a-LABEL: preloadremainder_y:812; GFX90a: ; %bb.1:813; GFX90a-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0814; GFX90a-NEXT: s_waitcnt lgkmcnt(0)815; GFX90a-NEXT: s_branch .LBB17_0816; GFX90a-NEXT: .p2align 8817; GFX90a-NEXT: ; %bb.2:818; GFX90a-NEXT: .LBB17_0:819; GFX90a-NEXT: s_and_b32 s0, s15, 0xffff820; GFX90a-NEXT: v_mov_b32_e32 v0, 0821; GFX90a-NEXT: v_mov_b32_e32 v1, s0822; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]823; GFX90a-NEXT: s_endpgm824;825; GFX1250-LABEL: preloadremainder_y:826; GFX1250: ; %bb.0:827; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1828; GFX1250-NEXT: s_and_b32 s0, s9, 0xffff829; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)830; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0831; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]832; GFX1250-NEXT: s_endpgm833 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()834 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 20835 %load = load i16, ptr addrspace(4) %gep836 %conv = zext i16 %load to i32837 store i32 %conv, ptr addrspace(1) %out838 ret void839}840 841define amdgpu_kernel void @preloadremainder_z(ptr addrspace(1) inreg %out) #0 {842; GFX942-LABEL: preloadremainder_z:843; GFX942: ; %bb.1:844; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0845; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8846; GFX942-NEXT: s_load_dwordx2 s[8:9], s[0:1], 0x18847; GFX942-NEXT: s_waitcnt lgkmcnt(0)848; GFX942-NEXT: s_branch .LBB18_0849; GFX942-NEXT: .p2align 8850; GFX942-NEXT: ; %bb.2:851; GFX942-NEXT: .LBB18_0:852; GFX942-NEXT: s_lshr_b32 s0, s9, 16853; GFX942-NEXT: v_mov_b32_e32 v0, 0854; GFX942-NEXT: v_mov_b32_e32 v1, s0855; GFX942-NEXT: global_store_dword v0, v1, s[2:3]856; GFX942-NEXT: s_endpgm857;858; GFX90a-LABEL: preloadremainder_z:859; GFX90a: ; %bb.1:860; GFX90a-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0861; GFX90a-NEXT: s_waitcnt lgkmcnt(0)862; GFX90a-NEXT: s_branch .LBB18_0863; GFX90a-NEXT: .p2align 8864; GFX90a-NEXT: ; %bb.2:865; GFX90a-NEXT: .LBB18_0:866; GFX90a-NEXT: s_lshr_b32 s0, s15, 16867; GFX90a-NEXT: v_mov_b32_e32 v0, 0868; GFX90a-NEXT: v_mov_b32_e32 v1, s0869; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]870; GFX90a-NEXT: s_endpgm871;872; GFX1250-LABEL: preloadremainder_z:873; GFX1250: ; %bb.0:874; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1875; GFX1250-NEXT: s_lshr_b32 s0, s9, 16876; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)877; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0878; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]879; GFX1250-NEXT: s_endpgm880 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()881 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 22882 %load = load i16, ptr addrspace(4) %gep883 %conv = zext i16 %load to i32884 store i32 %conv, ptr addrspace(1) %out885 ret void886}887 888define amdgpu_kernel void @preloadremainder_xyz(ptr addrspace(1) inreg %out) #0 {889; GFX942-LABEL: preloadremainder_xyz:890; GFX942: ; %bb.1:891; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x0892; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x8893; GFX942-NEXT: s_load_dwordx2 s[8:9], s[0:1], 0x18894; GFX942-NEXT: s_waitcnt lgkmcnt(0)895; GFX942-NEXT: s_branch .LBB19_0896; GFX942-NEXT: .p2align 8897; GFX942-NEXT: ; %bb.2:898; GFX942-NEXT: .LBB19_0:899; GFX942-NEXT: s_lshr_b32 s0, s9, 16900; GFX942-NEXT: s_lshr_b32 s1, s8, 16901; GFX942-NEXT: s_and_b32 s4, s9, 0xffff902; GFX942-NEXT: v_mov_b32_e32 v3, 0903; GFX942-NEXT: v_mov_b32_e32 v0, s1904; GFX942-NEXT: v_mov_b32_e32 v1, s4905; GFX942-NEXT: v_mov_b32_e32 v2, s0906; GFX942-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3]907; GFX942-NEXT: s_endpgm908;909; GFX90a-LABEL: preloadremainder_xyz:910; GFX90a: ; %bb.1:911; GFX90a-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0912; GFX90a-NEXT: s_waitcnt lgkmcnt(0)913; GFX90a-NEXT: s_branch .LBB19_0914; GFX90a-NEXT: .p2align 8915; GFX90a-NEXT: ; %bb.2:916; GFX90a-NEXT: .LBB19_0:917; GFX90a-NEXT: s_lshr_b32 s0, s15, 16918; GFX90a-NEXT: s_lshr_b32 s1, s14, 16919; GFX90a-NEXT: s_and_b32 s2, s15, 0xffff920; GFX90a-NEXT: v_mov_b32_e32 v3, 0921; GFX90a-NEXT: v_mov_b32_e32 v0, s1922; GFX90a-NEXT: v_mov_b32_e32 v1, s2923; GFX90a-NEXT: v_mov_b32_e32 v2, s0924; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]925; GFX90a-NEXT: s_endpgm926;927; GFX1250-LABEL: preloadremainder_xyz:928; GFX1250: ; %bb.0:929; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1930; GFX1250-NEXT: s_lshr_b32 s0, s9, 16931; GFX1250-NEXT: s_lshr_b32 s1, s8, 16932; GFX1250-NEXT: s_and_b32 s4, s9, 0xffff933; GFX1250-NEXT: v_dual_mov_b32 v3, 0 :: v_dual_mov_b32 v0, s1934; GFX1250-NEXT: v_dual_mov_b32 v1, s4 :: v_dual_mov_b32 v2, s0935; GFX1250-NEXT: global_store_b96 v3, v[0:2], s[2:3]936; GFX1250-NEXT: s_endpgm937 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()938 %gep_x = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 18939 %load_x = load i16, ptr addrspace(4) %gep_x940 %conv_x = zext i16 %load_x to i32941 %gep_y = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 20942 %load_y = load i16, ptr addrspace(4) %gep_y943 %conv_y = zext i16 %load_y to i32944 %gep_z = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 22945 %load_z = load i16, ptr addrspace(4) %gep_z946 %conv_z = zext i16 %load_z to i32947 %ins.0 = insertelement <3 x i32> poison, i32 %conv_x, i32 0948 %ins.1 = insertelement <3 x i32> %ins.0, i32 %conv_y, i32 1949 %ins.2 = insertelement <3 x i32> %ins.1, i32 %conv_z, i32 2950 store <3 x i32> %ins.2, ptr addrspace(1) %out951 ret void952}953 954define amdgpu_kernel void @no_free_sgprs_preloadremainder_z(ptr addrspace(1) inreg %out) {955; GFX942-LABEL: no_free_sgprs_preloadremainder_z:956; GFX942: ; %bb.1:957; GFX942-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0958; GFX942-NEXT: s_waitcnt lgkmcnt(0)959; GFX942-NEXT: s_branch .LBB20_0960; GFX942-NEXT: .p2align 8961; GFX942-NEXT: ; %bb.2:962; GFX942-NEXT: .LBB20_0:963; GFX942-NEXT: s_lshr_b32 s0, s15, 16964; GFX942-NEXT: v_mov_b32_e32 v0, 0965; GFX942-NEXT: v_mov_b32_e32 v1, s0966; GFX942-NEXT: global_store_dword v0, v1, s[8:9]967; GFX942-NEXT: s_endpgm968;969; GFX90a-LABEL: no_free_sgprs_preloadremainder_z:970; GFX90a: ; %bb.1:971; GFX90a-NEXT: s_load_dwordx2 s[14:15], s[8:9], 0x0972; GFX90a-NEXT: s_waitcnt lgkmcnt(0)973; GFX90a-NEXT: s_branch .LBB20_0974; GFX90a-NEXT: .p2align 8975; GFX90a-NEXT: ; %bb.2:976; GFX90a-NEXT: .LBB20_0:977; GFX90a-NEXT: s_load_dword s0, s[8:9], 0x1c978; GFX90a-NEXT: v_mov_b32_e32 v0, 0979; GFX90a-NEXT: s_waitcnt lgkmcnt(0)980; GFX90a-NEXT: s_lshr_b32 s0, s0, 16981; GFX90a-NEXT: v_mov_b32_e32 v1, s0982; GFX90a-NEXT: global_store_dword v0, v1, s[14:15]983; GFX90a-NEXT: s_endpgm984;985; GFX1250-LABEL: no_free_sgprs_preloadremainder_z:986; GFX1250: ; %bb.0:987; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1988; GFX1250-NEXT: s_lshr_b32 s0, s15, 16989; GFX1250-NEXT: s_delay_alu instid0(SALU_CYCLE_1)990; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0991; GFX1250-NEXT: global_store_b32 v0, v1, s[8:9]992; GFX1250-NEXT: s_endpgm993 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()994 %gep = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 22995 %load = load i16, ptr addrspace(4) %gep996 %conv = zext i16 %load to i32997 store i32 %conv, ptr addrspace(1) %out998 ret void999}1000 1001; Check for consistency between isel and earlier passes preload SGPR accounting with max preload SGPRs.1002 1003define amdgpu_kernel void @preload_block_max_user_sgprs(ptr addrspace(1) inreg %out, i192 inreg %t0, i32 inreg %t1) #0 {1004; GFX942-LABEL: preload_block_max_user_sgprs:1005; GFX942: ; %bb.1:1006; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x01007; GFX942-NEXT: s_load_dwordx8 s[4:11], s[0:1], 0x81008; GFX942-NEXT: s_load_dword s12, s[0:1], 0x281009; GFX942-NEXT: s_waitcnt lgkmcnt(0)1010; GFX942-NEXT: s_branch .LBB21_01011; GFX942-NEXT: .p2align 81012; GFX942-NEXT: ; %bb.2:1013; GFX942-NEXT: .LBB21_0:1014; GFX942-NEXT: v_mov_b32_e32 v0, 01015; GFX942-NEXT: v_mov_b32_e32 v1, s121016; GFX942-NEXT: global_store_dword v0, v1, s[2:3]1017; GFX942-NEXT: s_endpgm1018;1019; GFX90a-LABEL: preload_block_max_user_sgprs:1020; GFX90a: ; %bb.1:1021; GFX90a-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x01022; GFX90a-NEXT: s_waitcnt lgkmcnt(0)1023; GFX90a-NEXT: s_branch .LBB21_01024; GFX90a-NEXT: .p2align 81025; GFX90a-NEXT: ; %bb.2:1026; GFX90a-NEXT: .LBB21_0:1027; GFX90a-NEXT: s_load_dword s0, s[4:5], 0x281028; GFX90a-NEXT: v_mov_b32_e32 v0, 01029; GFX90a-NEXT: s_waitcnt lgkmcnt(0)1030; GFX90a-NEXT: v_mov_b32_e32 v1, s01031; GFX90a-NEXT: global_store_dword v0, v1, s[8:9]1032; GFX90a-NEXT: s_endpgm1033;1034; GFX1250-LABEL: preload_block_max_user_sgprs:1035; GFX1250: ; %bb.0:1036; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 11037; GFX1250-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s121038; GFX1250-NEXT: global_store_b32 v0, v1, s[2:3]1039; GFX1250-NEXT: s_endpgm1040 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()1041 %load = load i32, ptr addrspace(4) %imp_arg_ptr1042 store i32 %load, ptr addrspace(1) %out1043 ret void1044}1045 1046define amdgpu_kernel void @preload_block_count_z_workgroup_size_z_remainder_z(ptr addrspace(1) inreg %out) #0 {1047; GFX942-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:1048; GFX942: ; %bb.1:1049; GFX942-NEXT: s_load_dwordx2 s[2:3], s[0:1], 0x01050; GFX942-NEXT: s_load_dwordx4 s[4:7], s[0:1], 0x81051; GFX942-NEXT: s_load_dwordx2 s[8:9], s[0:1], 0x181052; GFX942-NEXT: s_waitcnt lgkmcnt(0)1053; GFX942-NEXT: s_branch .LBB22_01054; GFX942-NEXT: .p2align 81055; GFX942-NEXT: ; %bb.2:1056; GFX942-NEXT: .LBB22_0:1057; GFX942-NEXT: s_lshr_b32 s0, s9, 161058; GFX942-NEXT: s_and_b32 s1, s8, 0xffff1059; GFX942-NEXT: v_mov_b32_e32 v3, 01060; GFX942-NEXT: v_mov_b32_e32 v0, s61061; GFX942-NEXT: v_mov_b32_e32 v1, s11062; GFX942-NEXT: v_mov_b32_e32 v2, s01063; GFX942-NEXT: global_store_dwordx3 v3, v[0:2], s[2:3]1064; GFX942-NEXT: s_endpgm1065;1066; GFX90a-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:1067; GFX90a: ; %bb.1:1068; GFX90a-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x01069; GFX90a-NEXT: s_waitcnt lgkmcnt(0)1070; GFX90a-NEXT: s_branch .LBB22_01071; GFX90a-NEXT: .p2align 81072; GFX90a-NEXT: ; %bb.2:1073; GFX90a-NEXT: .LBB22_0:1074; GFX90a-NEXT: s_lshr_b32 s0, s15, 161075; GFX90a-NEXT: s_and_b32 s1, s14, 0xffff1076; GFX90a-NEXT: v_mov_b32_e32 v3, 01077; GFX90a-NEXT: v_mov_b32_e32 v0, s121078; GFX90a-NEXT: v_mov_b32_e32 v1, s11079; GFX90a-NEXT: v_mov_b32_e32 v2, s01080; GFX90a-NEXT: global_store_dwordx3 v3, v[0:2], s[8:9]1081; GFX90a-NEXT: s_endpgm1082;1083; GFX1250-LABEL: preload_block_count_z_workgroup_size_z_remainder_z:1084; GFX1250: ; %bb.0:1085; GFX1250-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 11086; GFX1250-NEXT: s_lshr_b32 s0, s9, 161087; GFX1250-NEXT: s_and_b32 s1, s8, 0xffff1088; GFX1250-NEXT: v_dual_mov_b32 v3, 0 :: v_dual_mov_b32 v0, s61089; GFX1250-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s01090; GFX1250-NEXT: global_store_b96 v3, v[0:2], s[2:3]1091; GFX1250-NEXT: s_endpgm1092 %imp_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()1093 %gep0 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 81094 %gep1 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 161095 %gep2 = getelementptr i8, ptr addrspace(4) %imp_arg_ptr, i32 221096 %load0 = load i32, ptr addrspace(4) %gep01097 %load1 = load i16, ptr addrspace(4) %gep11098 %load2 = load i16, ptr addrspace(4) %gep21099 %conv1 = zext i16 %load1 to i321100 %conv2 = zext i16 %load2 to i321101 %ins.0 = insertelement <3 x i32> poison, i32 %load0, i32 01102 %ins.1 = insertelement <3 x i32> %ins.0, i32 %conv1, i32 11103 %ins.2 = insertelement <3 x i32> %ins.1, i32 %conv2, i32 21104 store <3 x i32> %ins.2, ptr addrspace(1) %out1105 ret void1106}1107 1108attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }1109