From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Mon, 9 Oct 2017 17:44:18 +0000 (+0000)
Subject: AMDGPU: Fix missing declaration for __builtin_amdgcn_dispatch_ptr
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=bdc2045617344adaffd47aae222d66f252b89c70;p=clang

AMDGPU: Fix missing declaration for __builtin_amdgcn_dispatch_ptr

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@315219 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 6542acafe4..3e5b0d13ae 100644
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -21,6 +21,7 @@
 // SI+ only builtins.
 //===----------------------------------------------------------------------===//
 
+BUILTIN(__builtin_amdgcn_dispatch_ptr, "Uc*2", "nc")
 BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "Uc*2", "nc")
 BUILTIN(__builtin_amdgcn_implicitarg_ptr, "Uc*2", "nc")
 
diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index f75620ba60..65b0666ad1 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -421,6 +421,13 @@ void test_read_exec(global ulong* out) {
 
 // CHECK: declare i64 @llvm.read_register.i64(metadata) #[[NOUNWIND_READONLY:[0-9]+]]
 
+// CHECK-LABEL: @test_dispatch_ptr
+// CHECK: call i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
+void test_dispatch_ptr(__attribute__((address_space(2))) unsigned char ** out)
+{
+  *out = __builtin_amdgcn_dispatch_ptr();
+}
+
 // CHECK-LABEL: @test_kernarg_segment_ptr
 // CHECK: call i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr()
 void test_kernarg_segment_ptr(__attribute__((address_space(2))) unsigned char ** out)