Enforce the restriction that a parameter to a kernel function

cannot be a pointer to the private address space (as clarified
in the OpenCL 1.2 specification).

Patch by Fraser Cormack!

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

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index e3ac4c2ef774a89d514affeea5aa098e34de571f..9094e9b323f385627f7d94d4a01da8ebe6c40c6a 100644 (file)
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6816,6 +6816,8 @@ def err_static_kernel : Error<
   "kernel functions cannot be declared static">;
 def err_opencl_ptrptr_kernel_param : Error<
   "kernel parameter cannot be declared as a pointer to a pointer">;
+def err_opencl_private_ptr_kernel_param : Error<
+  "kernel parameter cannot be declared as a pointer to the __private address space">;
 def err_static_function_scope : Error<
   "variables in function scope cannot be declared static">;
 def err_opencl_bitfields : Error<

diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index bfbd870eefc6072b3135f9356f86d951d993b74a..82579d89cd2c8c5fceb32088832e6d831965e070 100644 (file)
--- a/lib/Sema/SemaDecl.cpp
+++ b/lib/Sema/SemaDecl.cpp
@@ -6378,6 +6378,7 @@ enum OpenCLParamType {
   ValidKernelParam,
   PtrPtrKernelParam,
   PtrKernelParam,
+  PrivatePtrKernelParam,
   InvalidKernelParam,
   RecordKernelParam
 };
@@ -6385,7 +6386,10 @@ enum OpenCLParamType {
 static OpenCLParamType getOpenCLKernelParameterType(QualType PT) {
   if (PT->isPointerType()) {
     QualType PointeeType = PT->getPointeeType();
-    return PointeeType->isPointerType() ? PtrPtrKernelParam : PtrKernelParam;
+    if (PointeeType->isPointerType())
+      return PtrPtrKernelParam;
+    return PointeeType.getAddressSpace() == 0 ? PrivatePtrKernelParam
+                                              : PtrKernelParam;
   }
 
   // TODO: Forbid the other integer types (size_t, ptrdiff_t...) when they can
@@ -6430,6 +6434,14 @@ static void checkIsValidOpenCLKernelParameter(
     D.setInvalidType();
     return;
 
+  case PrivatePtrKernelParam:
+    // OpenCL v1.2 s6.9.a:
+    // A kernel function argument cannot be declared as a
+    // pointer to the private address space.
+    S.Diag(Param->getLocation(), diag::err_opencl_private_ptr_kernel_param);
+    D.setInvalidType();
+    return;
+
     // OpenCL v1.2 s6.9.k:
     // Arguments to kernel functions in a program cannot be declared with the
     // built-in scalar types bool, half, size_t, ptrdiff_t, intptr_t, and
@@ -6509,7 +6521,8 @@ static void checkIsValidOpenCLKernelParameter(
       // Arguments to kernel functions that are declared to be a struct or union
       // do not allow OpenCL objects to be passed as elements of the struct or
       // union.
-      if (ParamType == PtrKernelParam || ParamType == PtrPtrKernelParam) {
+      if (ParamType == PtrKernelParam || ParamType == PtrPtrKernelParam ||
+          ParamType == PrivatePtrKernelParam) {
         S.Diag(Param->getLocation(),
                diag::err_record_with_pointers_kernel_param)
           << PT->isUnionType()

diff --git a/test/SemaOpenCL/invalid-kernel.cl b/test/SemaOpenCL/invalid-kernel.cl
index 62991d4642018af19bc58658db0783784c4699ee..df73eddc3b5dee1ce81fb2bac24541a32147b7c1 100644 (file)
--- a/test/SemaOpenCL/invalid-kernel.cl
+++ b/test/SemaOpenCL/invalid-kernel.cl
@@ -2,6 +2,8 @@
 
 kernel void no_ptrptr(global int **i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}}
 
+__kernel void no_privateptr(__private int *i) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}}
+
 kernel int bar()  { // expected-error {{kernel must have void return type}}
   return 6;
 }