maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

AMDGPU: Fix some missing GCCBuiltin declarations 

llvm-svn: 366250
This commit is contained in:
Matt Arsenault 2019-07-16 19:44:14 +00:00
parent 2d10407719
commit afdf6b3c37
1 changed files with 32 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

56
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -296,29 +296,33 @@ def int_amdgcn_fract : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_cvt_pkrtz : Intrinsic<
[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
def int_amdgcn_cvt_pkrtz : GCCBuiltin<"__builtin_amdgcn_cvt_pkrtz">,
Intrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_cvt_pknorm_i16 : Intrinsic<
[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
def int_amdgcn_cvt_pknorm_i16 :
GCCBuiltin<"__builtin_amdgcn_cvt_pknorm_i16">,
Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_cvt_pknorm_u16 : Intrinsic<
[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
def int_amdgcn_cvt_pknorm_u16 :
GCCBuiltin<"__builtin_amdgcn_cvt_pknorm_u16">,
Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_cvt_pk_i16 : Intrinsic<
def int_amdgcn_cvt_pk_i16 :
GCCBuiltin<"__builtin_amdgcn_cvt_pk_i16">,
Intrinsic<
[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_cvt_pk_u16 : Intrinsic<
[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
def int_amdgcn_cvt_pk_u16 : GCCBuiltin<"__builtin_amdgcn_cvt_pk_u16">,
Intrinsic<[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_class : Intrinsic<
@ -1245,14 +1249,17 @@ def int_amdgcn_ds_swizzle :
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrConvergent, ImmArg<1>]>;
def int_amdgcn_ubfe : Intrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
def int_amdgcn_ubfe :
GCCBuiltin<"__builtin_amdgcn_ubfe">,
Intrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_sbfe : Intrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
def int_amdgcn_sbfe : GCCBuiltin<"__builtin_amdgcn_sbfe">,
Intrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_lerp :
@ -1340,13 +1347,14 @@ def int_amdgcn_writelane :
[IntrNoMem, IntrConvergent]
>;
def int_amdgcn_alignbit : Intrinsic<[llvm_i32_ty],
def int_amdgcn_alignbit :
GCCBuiltin<"__builtin_amdgcn_alignbit">, Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_alignbyte : Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
def int_amdgcn_alignbyte : GCCBuiltin<"__builtin_amdgcn_alignbyte">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
@ -1515,13 +1523,13 @@ def int_amdgcn_ds_bpermute :
//===----------------------------------------------------------------------===//
// llvm.amdgcn.permlane16 <old> <src0> <src1> <src2> <fi> <bound_control>
def int_amdgcn_permlane16 :
def int_amdgcn_permlane16 : GCCBuiltin<"__builtin_amdgcn_permlane16">,
Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
[IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
// llvm.amdgcn.permlanex16 <old> <src0> <src1> <src2> <fi> <bound_control>
def int_amdgcn_permlanex16 :
def int_amdgcn_permlanex16 : GCCBuiltin<"__builtin_amdgcn_permlanex16">,
Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
[IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;