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

[Matt Arsenault via llvm-commits]

arsenm added inline comments.

================
Comment at: include/llvm/IR/IntrinsicsAMDGPU.td:131
@@ +130,3 @@
+def int_amdgcn_dispatch_ptr :
+  GCCBuiltin<"__builtin_amdgcn_disaptch_ptr">,
+  Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
----------------
Typo: disaptch

================
Comment at: include/llvm/IR/IntrinsicsAMDGPU.td:132
@@ +131,3 @@
+  GCCBuiltin<"__builtin_amdgcn_disaptch_ptr">,
+  Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
+
----------------
Is there a way to mark noalias on the return value?

================
Comment at: lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp:106-109
@@ -105,5 +105,6 @@
 
     { "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" }
   };
----------------
Later this should probably be cleaned up to only check the dispatch.ptr one for HSA, and the read.global.size.x should not set this


http://reviews.llvm.org/D14898