diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-07-22 17:01:30 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-07-22 17:01:30 +0000 |
| commit | 8d718dcfdae7c4ce7fae4ed51ac8a931142dc89d (patch) | |
| tree | 15837044b6336d9e7cf1f7e10ac1874a83c73373 /llvm/include | |
| parent | f9245b75c013653a59ced45a1340e39d9a6bbcb9 (diff) | |
| download | bcm5719-llvm-8d718dcfdae7c4ce7fae4ed51ac8a931142dc89d.tar.gz bcm5719-llvm-8d718dcfdae7c4ce7fae4ed51ac8a931142dc89d.zip | |
AMDGPU: Add HSA dispatch id intrinsic
llvm-svn: 276437
Diffstat (limited to 'llvm/include')
| -rw-r--r-- | llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 52 |
1 files changed, 32 insertions, 20 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 387335ce2aa..b8a7c257b54 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -70,10 +70,42 @@ def int_r600_recipsqrt_clamped : Intrinsic< let TargetPrefix = "amdgcn" in { +//===----------------------------------------------------------------------===// +// ABI Special Intrinsics +//===----------------------------------------------------------------------===// + defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz; defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named <"__builtin_amdgcn_workgroup_id">; +def int_amdgcn_dispatch_ptr : + GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">, + Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; + +def int_amdgcn_queue_ptr : + GCCBuiltin<"__builtin_amdgcn_queue_ptr">, + Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; + +def int_amdgcn_kernarg_segment_ptr : + GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">, + Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; + +def int_amdgcn_implicitarg_ptr : + GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">, + Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; + +def int_amdgcn_groupstaticsize : + GCCBuiltin<"__builtin_amdgcn_groupstaticsize">, + Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; + +def int_amdgcn_dispatch_id : + GCCBuiltin<"__builtin_amdgcn_dispatch_id">, + Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>; + +//===----------------------------------------------------------------------===// +// Instruction Intrinsics +//===----------------------------------------------------------------------===// + def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">, Intrinsic<[], [], [IntrConvergent]>; @@ -331,26 +363,6 @@ def int_amdgcn_s_getreg : GCCBuiltin<"__builtin_amdgcn_s_getreg">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>; -def int_amdgcn_groupstaticsize : - GCCBuiltin<"__builtin_amdgcn_groupstaticsize">, - Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; - -def int_amdgcn_dispatch_ptr : - GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">, - Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; - -def int_amdgcn_queue_ptr : - GCCBuiltin<"__builtin_amdgcn_queue_ptr">, - Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; - -def int_amdgcn_kernarg_segment_ptr : - GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">, - Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; - -def int_amdgcn_implicitarg_ptr : - GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">, - Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; - // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0> def int_amdgcn_interp_p1 : GCCBuiltin<"__builtin_amdgcn_interp_p1">, |
