[PATCH] D14898: AMDGPU: Add llvm.amdgcn.dispatch.ptr intrinsic

[Tom Stellard via llvm-commits]

tstellarAMD created this revision.
tstellarAMD added a reviewer: arsenm.
tstellarAMD added a subscriber: llvm-commits.
Herald added a subscriber: arsenm.

This returns a pointer to the dispatch packet, which can be used to load
information about the kernel dispach.

http://reviews.llvm.org/D14898

Files:
  include/llvm/IR/IntrinsicsAMDGPU.td
  lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp
  lib/Target/AMDGPU/SIISelLowering.cpp

Index: lib/Target/AMDGPU/SIISelLowering.cpp
===================================================================
--- lib/Target/AMDGPU/SIISelLowering.cpp
+++ lib/Target/AMDGPU/SIISelLowering.cpp
@@ -1083,6 +1083,10 @@
   // TODO: Should this propagate fast-math-flags?
 
   switch (IntrinsicID) {
+  case Intrinsic::amdgcn_dispatch_ptr:
+    return CreateLiveInRegister(DAG, &AMDGPU::SReg_64RegClass,
+      TRI->getPreloadedValue(MF, SIRegisterInfo::DISPATCH_PTR), VT);
+
   case Intrinsic::r600_read_ngroups_x:
     return LowerParameter(DAG, VT, VT, DL, DAG.getEntryNode(),
                           SI::KernelInputOffsets::NGROUPS_X,
Index: lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp
===================================================================
--- lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp
+++ lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp
@@ -105,7 +105,8 @@
 
     { "llvm.r600.read.global.size.x", "amdgpu-dispatch-ptr" },
     { "llvm.r600.read.global.size.y", "amdgpu-dispatch-ptr" },
-    { "llvm.r600.read.global.size.z", "amdgpu-dispatch-ptr" }
+    { "llvm.r600.read.global.size.z", "amdgpu-dispatch-ptr" },
+    { "llvm.amdgcn.dispatch.ptr",     "amdgpu-dispatch-ptr" }
   };
 
   // TODO: Intrinsics that require queue ptr.
Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -127,4 +127,8 @@
   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
   Intrinsic<[], [], []>;
 
+def int_amdgcn_dispatch_ptr :
+  GCCBuiltin<"__builtin_amdgcn_disaptch_ptr">,
+  Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
+
 }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D14898.40855.patch
Type: text/x-patch
Size: 1712 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20151121/2d60a9e0/attachment.bin>