diff options
| -rw-r--r-- | llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 56 |
1 files changed, 32 insertions, 24 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index bad4216173d..4c67cdea4d5 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -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>]>; |
