[OpenCL] Allow pointer-to-pointer kernel args beyond CL 1.2

The restriction on pointer-to-pointer kernel arguments has been
relaxed in OpenCL 2.0.  Apply the same address space restrictions for
pointer argument types to the inner pointer types.

Differential Revision: https://reviews.llvm.org/D92091

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 9c282a7..2116b3f 100644 (file)
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8591,12 +8591,21 @@ static bool isOpenCLSizeDependentType(ASTContext &C, QualType Ty) {
 static OpenCLParamType getOpenCLKernelParameterType(Sema &S, QualType PT) {
   if (PT->isPointerType()) {
     QualType PointeeType = PT->getPointeeType();
-    if (PointeeType->isPointerType())
-      return PtrPtrKernelParam;
     if (PointeeType.getAddressSpace() == LangAS::opencl_generic ||
         PointeeType.getAddressSpace() == LangAS::opencl_private ||
         PointeeType.getAddressSpace() == LangAS::Default)
       return InvalidAddrSpacePtrKernelParam;
+
+    if (PointeeType->isPointerType()) {
+      // This is a pointer to pointer parameter.
+      // Recursively check inner type.
+      OpenCLParamType ParamKind = getOpenCLKernelParameterType(S, PointeeType);
+      if (ParamKind == InvalidAddrSpacePtrKernelParam ||
+          ParamKind == InvalidKernelParam)
+        return ParamKind;
+
+      return PtrPtrKernelParam;
+    }
     return PtrKernelParam;
   }
 
@@ -8649,11 +8658,17 @@ static void checkIsValidOpenCLKernelParameter(
 
   switch (getOpenCLKernelParameterType(S, PT)) {
   case PtrPtrKernelParam:
-    // OpenCL v1.2 s6.9.a:
-    // A kernel function argument cannot be declared as a
-    // pointer to a pointer type.
-    S.Diag(Param->getLocation(), diag::err_opencl_ptrptr_kernel_param);
-    D.setInvalidType();
+    // OpenCL v3.0 s6.11.a:
+    // A kernel function argument cannot be declared as a pointer to a pointer
+    // type. [...] This restriction only applies to OpenCL C 1.2 or below.
+    if (S.getLangOpts().OpenCLVersion < 120 &&
+        !S.getLangOpts().OpenCLCPlusPlus) {
+      S.Diag(Param->getLocation(), diag::err_opencl_ptrptr_kernel_param);
+      D.setInvalidType();
+      return;
+    }
+
+    ValidTypes.insert(PT.getTypePtr());
     return;
 
   case InvalidAddrSpacePtrKernelParam:

diff --git a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
index 26859ee..9fce92d 100644 (file)
--- a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
+++ b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
@@ -5,8 +5,14 @@ kernel void half_arg(half x) { } // expected-error{{declaring function parameter
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
-// expected-error@+1{{kernel parameter cannot be declared as a pointer to a pointer}}
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+// expected-error@+4{{kernel parameter cannot be declared as a pointer to a pointer}}
+// expected-error@+4{{kernel parameter cannot be declared as a pointer to a pointer}}
+// expected-error@+4{{kernel parameter cannot be declared as a pointer to a pointer}}
+#endif
 kernel void no_ptrptr(global int * global *i) { }
+kernel void no_lptrcptr(constant int * local *i) { }
+kernel void no_ptrptrptr(global int * global * global *i) { }
 
 // expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 __kernel void no_privateptr(__private int *i) { }
@@ -17,6 +23,15 @@ __kernel void no_privatearray(__private int i[]) { }
 // expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 __kernel void no_addrsp_ptr(int *ptr) { }
 
+#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+kernel void no_ptr_private_ptr(private int * global *i) { }
+// expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
+kernel void no_ptr_ptr_private_ptr(private int * global * global *i) { }
+// expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
+kernel void no_ptr_private_ptr_ptr(global int * private * global *i) { }
+// expected-error@-1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
+#endif
+
 // Disallowed: parameters with type
 // bool, half, size_t, ptrdiff_t, intptr_t, and uintptr_t
 // or a struct / union with any of these types in them