]> granicus.if.org Git - clang/commitdiff
AMDGPU: Fix missing declaration of queue ptr builtin
authorMatt Arsenault <Matthew.Arsenault@amd.com>
Thu, 2 Aug 2018 18:24:55 +0000 (18:24 +0000)
committerMatt Arsenault <Matthew.Arsenault@amd.com>
Thu, 2 Aug 2018 18:24:55 +0000 (18:24 +0000)
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@338754 91177308-0d34-0410-b5e6-96231b3b80d8

include/clang/Basic/BuiltinsAMDGPU.def
test/CodeGenOpenCL/builtins-amdgcn.cl

index 3e60eb2a0afb68069906a875637995c0f4735a0f..a8036a7004c7385bd4079c34ba8a50978ecc81cf 100644 (file)
@@ -24,6 +24,7 @@
 BUILTIN(__builtin_amdgcn_dispatch_ptr, "v*4", "nc")
 BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "v*4", "nc")
 BUILTIN(__builtin_amdgcn_implicitarg_ptr, "v*4", "nc")
+BUILTIN(__builtin_amdgcn_queue_ptr, "v*4", "nc")
 
 BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui", "nc")
index e92cf42dd1c0afab1f4898d247ef7da22bbd1caf..dc7f480209af33fa7eea3f9dfd798adda739665e 100644 (file)
@@ -462,6 +462,13 @@ void test_dispatch_ptr(__constant unsigned char ** out)
   *out = __builtin_amdgcn_dispatch_ptr();
 }
 
+// CHECK-LABEL: @test_queue_ptr
+// CHECK: call i8 addrspace(4)* @llvm.amdgcn.queue.ptr()
+void test_queue_ptr(__constant unsigned char ** out)
+{
+  *out = __builtin_amdgcn_queue_ptr();
+}
+
 // CHECK-LABEL: @test_kernarg_segment_ptr
 // CHECK: call i8 addrspace(4)* @llvm.amdgcn.kernarg.segment.ptr()
 void test_kernarg_segment_ptr(__constant unsigned char ** out)