llvm.org GIT mirror llvm / 2daf966
AMDGPU : Add intrinsic for instruction v_cvt_pk_u8_f32 Differential Revision: http://reviews.llvm.org/D23336 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@278403 91177308-0d34-0410-b5e6-96231b3b80d8 Wei Ding 3 years ago
4 changed file(s) with 70 addition(s) and 0 deletion(s). Raw diff Collapse all Expand all
529529 GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
530530 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
531531
532 def int_amdgcn_cvt_pk_u8_f32 :
533 GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
534 Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
535
532536 def int_amdgcn_icmp :
533537 Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
534538 [IntrNoMem, IntrConvergent]>;
16541654 def VOP_F64_F64_F64_F64 : VOPProfile <[f64, f64, f64, f64]>;
16551655 def VOP_I32_I32_I32_I32 : VOPProfile <[i32, i32, i32, i32]>;
16561656 def VOP_I64_I32_I32_I64 : VOPProfile <[i64, i32, i32, i64]>;
1657 def VOP_I32_F32_I32_I32 : VOPProfile <[i32, f32, i32, i32]>;
16571658
16581659 // This class is used only with VOPC instructions. Use $sdst for out operand
16591660 class SIInstAlias :
16141614 defm V_SAD_U32 : VOP3Inst , "v_sad_u32",
16151615 VOP_I32_I32_I32_I32
16161616 >;
1617
1618 defm V_CVT_PK_U8_F32 : VOP3Inst, "v_cvt_pk_u8_f32",
1619 VOP_I32_F32_I32_I32, int_amdgcn_cvt_pk_u8_f32
1620 >;
1621
16171622 //def V_CVT_PK_U8_F32 : VOP3_U8 <0x0000015e, "v_cvt_pk_u8_f32", []>;
16181623 defm V_DIV_FIXUP_F32 : VOP3Inst <
16191624 vop3<0x15f, 0x1de>, "v_div_fixup_f32", VOP_F32_F32_F32_F32, AMDGPUdiv_fixup
0 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
1 ; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
2
3 declare i32 @llvm.amdgcn.cvt.pk.u8.f32(float, i32, i32) #0
4
5 ; GCN-LABEL: {{^}}v_cvt_pk_u8_f32_idx_0:
6 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 0, v{{[0-9]+}}
7 define void @v_cvt_pk_u8_f32_idx_0(i32 addrspace(1)* %out, float %src, i32 %reg) {
8 %result = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 0, i32 %reg) #0
9 store i32 %result, i32 addrspace(1)* %out, align 4
10 ret void
11 }
12
13 ; GCN-LABEL: {{^}}v_cvt_pk_u8_f32_idx_1:
14 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 1, v{{[0-9]+}}
15 define void @v_cvt_pk_u8_f32_idx_1(i32 addrspace(1)* %out, float %src, i32 %reg) {
16 %result = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 1, i32 %reg) #0
17 store i32 %result, i32 addrspace(1)* %out, align 4
18 ret void
19 }
20
21 ; GCN-LABEL: {{^}}v_cvt_pk_u8_f32_idx_2:
22 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 2, v{{[0-9]+}}
23 define void @v_cvt_pk_u8_f32_idx_2(i32 addrspace(1)* %out, float %src, i32 %reg) {
24 %result = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 2, i32 %reg) #0
25 store i32 %result, i32 addrspace(1)* %out, align 4
26 ret void
27 }
28
29 ; GCN-LABEL: {{^}}v_cvt_pk_u8_f32_idx_3:
30 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 3, v{{[0-9]+}}
31 define void @v_cvt_pk_u8_f32_idx_3(i32 addrspace(1)* %out, float %src, i32 %reg) {
32 %result = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 3, i32 %reg) #0
33 store i32 %result, i32 addrspace(1)* %out, align 4
34 ret void
35 }
36
37 ; GCN-LABEL: {{^}}v_cvt_pk_u8_f32_combine:
38 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 0, v{{[0-9]+}}
39 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 1, v{{[0-9]+}}
40 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 2, v{{[0-9]+}}
41 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, 3, v{{[0-9]+}}
42 define void @v_cvt_pk_u8_f32_combine(i32 addrspace(1)* %out, float %src, i32 %reg) {
43 %result0 = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 0, i32 %reg) #0
44 %result1 = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 1, i32 %result0) #0
45 %result2 = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 2, i32 %result1) #0
46 %result3 = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 3, i32 %result2) #0
47 store i32 %result3, i32 addrspace(1)* %out, align 4
48 ret void
49 }
50
51 ; GCN-LABEL: {{^}}v_cvt_pk_u8_f32_idx:
52 ; GCN: v_cvt_pk_u8_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
53 define void @v_cvt_pk_u8_f32_idx(i32 addrspace(1)* %out, float %src, i32 %idx, i32 %reg) {
54 %result = call i32 @llvm.amdgcn.cvt.pk.u8.f32(float %src, i32 %idx, i32 %reg) #0
55 store i32 %result, i32 addrspace(1)* %out, align 4
56 ret void
57 }
58
59 attributes #0 = { nounwind readnone }