171 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -o - %s | FileCheck -check-prefix=GCN %s3 4define amdgpu_kernel void @MFMAExpInterleave(ptr addrspace(1) %out0, ptr addrspace(1) %out1, float %in0, <4 x float> %in1) {5; GCN-LABEL: MFMAExpInterleave:6; GCN: ; %bb.0:7; GCN-NEXT: s_load_dword s6, s[4:5], 0x108; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x209; GCN-NEXT: v_mov_b32_e32 v1, 0x3fb8aa3b10; GCN-NEXT: v_mov_b32_e32 v0, 1.011; GCN-NEXT: s_mov_b32 s7, 0x42b1721812; GCN-NEXT: s_waitcnt lgkmcnt(0)13; GCN-NEXT: v_mul_f32_e32 v2, s6, v114; GCN-NEXT: v_rndne_f32_e32 v3, v215; GCN-NEXT: v_sub_f32_e32 v4, v2, v316; GCN-NEXT: v_fma_f32 v1, s6, v1, -v217; GCN-NEXT: v_mov_b32_e32 v2, 0x32a5705f18; GCN-NEXT: v_accvgpr_write_b32 a0, s019; GCN-NEXT: v_fmac_f32_e32 v1, s6, v220; GCN-NEXT: v_accvgpr_write_b32 a1, s121; GCN-NEXT: v_accvgpr_write_b32 a2, s222; GCN-NEXT: v_accvgpr_write_b32 a3, s323; GCN-NEXT: v_add_f32_e32 v1, v4, v124; GCN-NEXT: v_cvt_i32_f32_e32 v2, v325; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]26; GCN-NEXT: v_exp_f32_e32 v1, v127; GCN-NEXT: s_mov_b32 s0, 0x3fb8aa3b28; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]29; GCN-NEXT: ; iglp_opt mask(0x00000003)30; GCN-NEXT: v_ldexp_f32 v1, v1, v231; GCN-NEXT: v_mov_b32_e32 v2, 0xc2ce8ed032; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s6, v233; GCN-NEXT: v_mov_b32_e32 v2, 0x42b1721834; GCN-NEXT: s_nop 035; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc36; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v237; GCN-NEXT: v_mov_b32_e32 v2, 0x7f80000038; GCN-NEXT: s_mov_b32 s6, 0xc2ce8ed039; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v1, vcc40; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v141; GCN-NEXT: v_fma_f32 v4, v1, s0, -v342; GCN-NEXT: v_rndne_f32_e32 v5, v343; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v144; GCN-NEXT: v_sub_f32_e32 v3, v3, v545; GCN-NEXT: v_add_f32_e32 v3, v3, v446; GCN-NEXT: v_exp_f32_e32 v3, v347; GCN-NEXT: v_cvt_i32_f32_e32 v4, v548; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v149; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]50; GCN-NEXT: v_ldexp_f32 v3, v3, v451; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc52; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v153; GCN-NEXT: s_nop 154; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc55; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v156; GCN-NEXT: v_fma_f32 v4, v1, s0, -v357; GCN-NEXT: v_rndne_f32_e32 v5, v358; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v159; GCN-NEXT: v_sub_f32_e32 v3, v3, v560; GCN-NEXT: v_add_f32_e32 v3, v3, v461; GCN-NEXT: v_exp_f32_e32 v3, v362; GCN-NEXT: v_cvt_i32_f32_e32 v4, v563; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v164; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]65; GCN-NEXT: v_ldexp_f32 v3, v3, v466; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc67; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v168; GCN-NEXT: s_nop 169; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc70; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v171; GCN-NEXT: v_fma_f32 v4, v1, s0, -v372; GCN-NEXT: v_rndne_f32_e32 v5, v373; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v174; GCN-NEXT: v_sub_f32_e32 v3, v3, v575; GCN-NEXT: v_add_f32_e32 v3, v3, v476; GCN-NEXT: v_exp_f32_e32 v3, v377; GCN-NEXT: v_cvt_i32_f32_e32 v4, v578; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v179; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]80; GCN-NEXT: v_ldexp_f32 v3, v3, v481; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc82; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v183; GCN-NEXT: s_nop 184; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc85; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v186; GCN-NEXT: v_fma_f32 v4, v1, s0, -v387; GCN-NEXT: v_rndne_f32_e32 v5, v388; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v189; GCN-NEXT: v_sub_f32_e32 v3, v3, v590; GCN-NEXT: v_add_f32_e32 v3, v3, v491; GCN-NEXT: v_exp_f32_e32 v3, v392; GCN-NEXT: v_cvt_i32_f32_e32 v4, v593; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v194; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]95; GCN-NEXT: v_ldexp_f32 v3, v3, v496; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc97; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v198; GCN-NEXT: s_nop 199; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc100; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1101; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3102; GCN-NEXT: v_rndne_f32_e32 v5, v3103; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1104; GCN-NEXT: v_sub_f32_e32 v3, v3, v5105; GCN-NEXT: v_add_f32_e32 v3, v3, v4106; GCN-NEXT: v_exp_f32_e32 v3, v3107; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5108; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1109; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]110; GCN-NEXT: v_ldexp_f32 v3, v3, v4111; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc112; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1113; GCN-NEXT: s_nop 1114; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc115; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1116; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3117; GCN-NEXT: v_rndne_f32_e32 v5, v3118; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1119; GCN-NEXT: v_sub_f32_e32 v3, v3, v5120; GCN-NEXT: v_add_f32_e32 v3, v3, v4121; GCN-NEXT: v_exp_f32_e32 v3, v3122; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5123; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]124; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1125; GCN-NEXT: v_ldexp_f32 v0, v3, v4126; GCN-NEXT: s_nop 0127; GCN-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc128; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1129; GCN-NEXT: s_nop 1130; GCN-NEXT: v_cndmask_b32_e32 v0, v2, v0, vcc131; GCN-NEXT: v_mul_f32_e32 v1, 0x3fb8aa3b, v0132; GCN-NEXT: v_fma_f32 v3, v0, s0, -v1133; GCN-NEXT: v_rndne_f32_e32 v4, v1134; GCN-NEXT: v_fmac_f32_e32 v3, 0x32a5705f, v0135; GCN-NEXT: v_sub_f32_e32 v1, v1, v4136; GCN-NEXT: v_add_f32_e32 v1, v1, v3137; GCN-NEXT: v_exp_f32_e32 v1, v1138; GCN-NEXT: v_cvt_i32_f32_e32 v3, v4139; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0140; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v0141; GCN-NEXT: v_mov_b32_e32 v4, 0142; GCN-NEXT: v_ldexp_f32 v1, v1, v3143; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc144; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v0145; GCN-NEXT: s_waitcnt lgkmcnt(0)146; GCN-NEXT: global_store_dwordx4 v4, a[0:3], s[0:1]147; GCN-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc148; GCN-NEXT: global_store_dword v4, v0, s[2:3]149; GCN-NEXT: s_endpgm150 %mai0 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %in1, i32 0, i32 0, i32 0)151 %mai1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai0, i32 0, i32 0, i32 0)152 %mai2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai1, i32 0, i32 0, i32 0)153 %mai3 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai2, i32 0, i32 0, i32 0)154 %mai4 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai3, i32 0, i32 0, i32 0)155 %mai5 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai4, i32 0, i32 0, i32 0)156 %mai6 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai5, i32 0, i32 0, i32 0)157 %mai7 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 1.0, <4 x float> %mai6, i32 0, i32 0, i32 0)158 %exp0 = call float @llvm.exp.f32(float %in0)159 %exp1 = call float @llvm.exp.f32(float %exp0)160 %exp2 = call float @llvm.exp.f32(float %exp1)161 %exp3 = call float @llvm.exp.f32(float %exp2)162 %exp4 = call float @llvm.exp.f32(float %exp3)163 %exp5 = call float @llvm.exp.f32(float %exp4)164 %exp6 = call float @llvm.exp.f32(float %exp5)165 %exp7 = call float @llvm.exp.f32(float %exp6)166 store <4 x float> %mai7, ptr addrspace(1) %out0167 store float %exp7, ptr addrspace(1) %out1168 tail call void @llvm.amdgcn.iglp.opt(i32 3)169 ret void170}171