caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
6a01ac7d06df875206f746fc982f58c161249285
clang-p2996/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.simple.ll
Jeffrey Byrnes 9ac52ce8d6 [AMDGPU] Add iglp_opt(3) for simple mfma / exp interleaving (#117269) 
Adds a minimal iglp_opt to do simple exp / mfma interleaving.
2024-12-06 15:19:07 -08:00

171 lines
7.8 KiB
LLVM
Raw Blame History

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -o - %s | FileCheck -check-prefix=GCN %s
define amdgpu_kernel void @MFMAExpInterleave(ptr addrspace(1) %out0, ptr addrspace(1) %out1, float %in0, <4 x float> %in1) {
; GCN-LABEL: MFMAExpInterleave:
; GCN: ; %bb.0:
; GCN-NEXT: s_load_dword s6, s[4:5], 0x10
; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x20
; GCN-NEXT: v_mov_b32_e32 v1, 0x3fb8aa3b
; GCN-NEXT: v_mov_b32_e32 v0, 1.0
; GCN-NEXT: s_mov_b32 s7, 0x42b17218
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: v_mul_f32_e32 v2, s6, v1
; GCN-NEXT: v_rndne_f32_e32 v3, v2
; GCN-NEXT: v_sub_f32_e32 v4, v2, v3
; GCN-NEXT: v_fma_f32 v1, s6, v1, -v2
; GCN-NEXT: v_mov_b32_e32 v2, 0x32a5705f
; GCN-NEXT: v_accvgpr_write_b32 a0, s0
; GCN-NEXT: v_fmac_f32_e32 v1, s6, v2
; GCN-NEXT: v_accvgpr_write_b32 a1, s1
; GCN-NEXT: v_accvgpr_write_b32 a2, s2
; GCN-NEXT: v_accvgpr_write_b32 a3, s3
; GCN-NEXT: v_add_f32_e32 v1, v4, v1
; GCN-NEXT: v_cvt_i32_f32_e32 v2, v3
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: v_exp_f32_e32 v1, v1
; GCN-NEXT: s_mov_b32 s0, 0x3fb8aa3b
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: ; iglp_opt mask(0x00000003)
; GCN-NEXT: v_ldexp_f32 v1, v1, v2
; GCN-NEXT: v_mov_b32_e32 v2, 0xc2ce8ed0
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s6, v2
; GCN-NEXT: v_mov_b32_e32 v2, 0x42b17218
; GCN-NEXT: s_nop 0
; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v2
; GCN-NEXT: v_mov_b32_e32 v2, 0x7f800000
; GCN-NEXT: s_mov_b32 s6, 0xc2ce8ed0
; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v1, vcc
; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1
; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3
; GCN-NEXT: v_rndne_f32_e32 v5, v3
; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1
; GCN-NEXT: v_sub_f32_e32 v3, v3, v5
; GCN-NEXT: v_add_f32_e32 v3, v3, v4
; GCN-NEXT: v_exp_f32_e32 v3, v3
; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: v_ldexp_f32 v3, v3, v4
; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1
; GCN-NEXT: s_nop 1
; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc
; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1
; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3
; GCN-NEXT: v_rndne_f32_e32 v5, v3
; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1
; GCN-NEXT: v_sub_f32_e32 v3, v3, v5
; GCN-NEXT: v_add_f32_e32 v3, v3, v4
; GCN-NEXT: v_exp_f32_e32 v3, v3
; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: v_ldexp_f32 v3, v3, v4
; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1
; GCN-NEXT: s_nop 1
; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc
; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1
; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3
; GCN-NEXT: v_rndne_f32_e32 v5, v3
; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1
; GCN-NEXT: v_sub_f32_e32 v3, v3, v5
; GCN-NEXT: v_add_f32_e32 v3, v3, v4
; GCN-NEXT: v_exp_f32_e32 v3, v3
; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: v_ldexp_f32 v3, v3, v4
; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1
; GCN-NEXT: s_nop 1
; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc
; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1
; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3
; GCN-NEXT: v_rndne_f32_e32 v5, v3
; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1
; GCN-NEXT: v_sub_f32_e32 v3, v3, v5
; GCN-NEXT: v_add_f32_e32 v3, v3, v4
; GCN-NEXT: v_exp_f32_e32 v3, v3
; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: v_ldexp_f32 v3, v3, v4
; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1
; GCN-NEXT: s_nop 1
; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc
; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1
; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3
; GCN-NEXT: v_rndne_f32_e32 v5, v3
; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1
; GCN-NEXT: v_sub_f32_e32 v3, v3, v5
; GCN-NEXT: v_add_f32_e32 v3, v3, v4
; GCN-NEXT: v_exp_f32_e32 v3, v3
; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: v_ldexp_f32 v3, v3, v4
; GCN-NEXT: v_cndmask_b32_e32 v3, 0, v3, vcc
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1
; GCN-NEXT: s_nop 1
; GCN-NEXT: v_cndmask_b32_e32 v1, v2, v3, vcc
; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v1
; GCN-NEXT: v_fma_f32 v4, v1, s0, -v3
; GCN-NEXT: v_rndne_f32_e32 v5, v3
; GCN-NEXT: v_fmac_f32_e32 v4, 0x32a5705f, v1
; GCN-NEXT: v_sub_f32_e32 v3, v3, v5
; GCN-NEXT: v_add_f32_e32 v3, v3, v4
; GCN-NEXT: v_exp_f32_e32 v3, v3
; GCN-NEXT: v_cvt_i32_f32_e32 v4, v5
; GCN-NEXT: v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v0, a[0:3]
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v1
; GCN-NEXT: v_ldexp_f32 v0, v3, v4
; GCN-NEXT: s_nop 0
; GCN-NEXT: v_cndmask_b32_e32 v0, 0, v0, vcc
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v1
; GCN-NEXT: s_nop 1
; GCN-NEXT: v_cndmask_b32_e32 v0, v2, v0, vcc
; GCN-NEXT: v_mul_f32_e32 v1, 0x3fb8aa3b, v0
; GCN-NEXT: v_fma_f32 v3, v0, s0, -v1
; GCN-NEXT: v_rndne_f32_e32 v4, v1
; GCN-NEXT: v_fmac_f32_e32 v3, 0x32a5705f, v0
; GCN-NEXT: v_sub_f32_e32 v1, v1, v4
; GCN-NEXT: v_add_f32_e32 v1, v1, v3
; GCN-NEXT: v_exp_f32_e32 v1, v1
; GCN-NEXT: v_cvt_i32_f32_e32 v3, v4
; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0
; GCN-NEXT: v_cmp_ngt_f32_e32 vcc, s6, v0
; GCN-NEXT: v_mov_b32_e32 v4, 0
; GCN-NEXT: v_ldexp_f32 v1, v1, v3
; GCN-NEXT: v_cndmask_b32_e32 v1, 0, v1, vcc
; GCN-NEXT: v_cmp_nlt_f32_e32 vcc, s7, v0
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: global_store_dwordx4 v4, a[0:3], s[0:1]
; GCN-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc
; GCN-NEXT: global_store_dword v4, v0, s[2:3]
; GCN-NEXT: s_endpgm
%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)
%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)
%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)
%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)
%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)
%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)
%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)
%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)
%exp0 = call float @llvm.exp.f32(float %in0)
%exp1 = call float @llvm.exp.f32(float %exp0)
%exp2 = call float @llvm.exp.f32(float %exp1)
%exp3 = call float @llvm.exp.f32(float %exp2)
%exp4 = call float @llvm.exp.f32(float %exp3)
%exp5 = call float @llvm.exp.f32(float %exp4)
%exp6 = call float @llvm.exp.f32(float %exp5)
%exp7 = call float @llvm.exp.f32(float %exp6)
store <4 x float> %mai7, ptr addrspace(1) %out0
store float %exp7, ptr addrspace(1) %out1
tail call void @llvm.amdgcn.iglp.opt(i32 3)
ret void
}
Reference in New Issue View Git Blame Copy Permalink