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

[svenvh]

Summary

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.

[Anastasia]

clang/lib/Sema/SemaDecl.cpp

8638

Btw I am surprised that we recurse to check the underlying type of an array but not the pointer. I guess this spec wording implies that pointed types should be checked too?

The size in bytes of these types areimplementation-defined and in addition can also be different for the OpenCL device and the host processor making it difficult to allocate buffer objects to be passed as arguments to a kernel declared as pointer to these types.

Also I can't find anything specific to the arrays...

Do you think we need to chase this with the spec?

8669

I feel this breaks the original flow a little bit. It would have been nicer if this could be part of getOpenCLKernelParameterType, but it might become more complex then... We would probably need to introduce something like ValidPtrPtrKernelParam and InvalidAddrSpacePtrPtrKernelParam.

[svenvh]

Recursively call getOpenCLKernelParameterType to check address spaces.

[svenvh]

clang/lib/Sema/SemaDecl.cpp

8638

I have raised https://github.com/KhronosGroup/OpenCL-Docs/issues/504

8669

Agreed. I have simplified the patch a bit, by not distinguishing between the outer pointer-to-pointer type and any inner pointer-to-pointer types for the diagnostic.

[Anastasia]

LGTM! Thanks!

The testing can be improved before committing!

clang/lib/Sema/SemaDecl.cpp

8638

Great! Thanks!

clang/test/SemaOpenCL/invalid-kernel-parameters.cl

13

Btw this was missing in the original testing, could we add a line with __constant and __local? Or perhaps just replace the first two global in this line.

[svenvh]

clang/test/SemaOpenCL/invalid-kernel-parameters.cl

13

Sure, will do!

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_static_kernel : Error<
   "kernel functions cannot be declared static">;
 def err_method_kernel : Error<
   "kernel functions cannot be class members">;
 def err_template_kernel : Error<
   "kernel functions cannot be used in a template declaration, instantiation or specialization">;
 def err_opencl_ptrptr_kernel_param : Error<
   "kernel parameter cannot be declared as a pointer to a pointer">;
 def err_kernel_arg_address_space : Error<
-  "pointer arguments to kernel functions must reside in '__global', "
-  "'__constant' or '__local' address space">;
+  "%select{|nested }0pointer arguments to kernel functions must reside in "
+  "'__global', '__constant' or '__local' address space">;
 def err_opencl_ext_vector_component_invalid_length : Error<
   "vector component access has invalid length %0.  Supported: 1,2,3,4,8,16.">;
 def err_opencl_function_variable : Error<
   "%select{non-kernel function|function scope}0 variable cannot be declared in %1 address space">;
 def err_opencl_addrspace_scope : Error<
   "variables in the %0 address space can only be declared in the outermost "
   "scope of a kernel function">;
 def err_static_function_scope : Error<

clang/lib/Sema/SemaDecl.cpp

do {
 
     Ty = DesugaredTy;
     DesugaredTy = Ty.getSingleStepDesugaredType(C);
   } while (DesugaredTy != Ty);
 
   return false;
 }
 
+static bool hasOpenCLValidPtrParameterAddrSpace(QualType Ty) {
+  LangAS AS = Ty.getAddressSpace();
+  return AS != LangAS::opencl_generic && AS != LangAS::opencl_private &&
+         AS != LangAS::Default;
+}
+
 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)
+    if (!hasOpenCLValidPtrParameterAddrSpace(PointeeType))
       return InvalidAddrSpacePtrKernelParam;
     return PtrKernelParam;
   }
 
   // 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
   // uintptr_t or a struct and/or union that contain fields declared to be one
if (PT->isRecordType())
     return RecordKernelParam;
 
   // Look into an array argument to check if it has a forbidden type.
   if (PT->isArrayType()) {
     const Type *UnderlyingTy = PT->getPointeeOrArrayElementType();
     // Call ourself to check an underlying type of an array. Since the
     // getPointeeOrArrayElementType returns an innermost type which is not an
     // array, this recursive call only happens once.
     return getOpenCLKernelParameterType(S, QualType(UnderlyingTy, 0));
   }
 
   return ValidKernelParam;
 }
 
 static void checkIsValidOpenCLKernelParameter(
   Sema &S,
   Declarator &D,
   ParmVarDecl *Param,
   llvm::SmallPtrSetImpl<const Type *> &ValidTypes) {
   QualType PT = Param->getType();
 
   // Cache the valid types we encounter to avoid rechecking structs that are
   // used again
   if (ValidTypes.count(PT.getTypePtr()))
     return;
 
   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.
+  case PtrPtrKernelParam: {
+    // 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;
+    }
+
+    // Loop through inner pointer types to ensure address spaces are valid.
+    QualType NestedPT = PT->getPointeeType();
+    do {
+      if (!hasOpenCLValidPtrParameterAddrSpace(NestedPT->getPointeeType())) {
+        S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space) << 1;
+        D.setInvalidType();
+        return;
+      }
+      NestedPT = NestedPT->getPointeeType();
+    } while (NestedPT->isPointerType());
+
+    ValidTypes.insert(PT.getTypePtr());
+    return;
+  }
 
   case InvalidAddrSpacePtrKernelParam:
     // OpenCL v1.0 s6.5:
     // __kernel function arguments declared to be a pointer of a type can point
     // to one of the following address spaces only : __global, __local or
     // __constant.
-    S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space);
+    S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space) << 0;
     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
     // uintptr_t or a struct and/or union that contain fields declared to be
     // one of these built-in scalar types.

clang/test/SemaOpenCL/invalid-kernel-parameters.cl

 // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown
 // RUN: %clang_cc1 -fsyntax-only -verify %s -triple spir-unknown-unknown -cl-std=CL2.0
 
 kernel void half_arg(half x) { } // expected-error{{declaring function parameter of type '__private half' is not allowed; did you forget * ?}}
 
 #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@+3{{kernel parameter cannot be declared as a pointer to a pointer}}
+// expected-error@+3{{kernel parameter cannot be declared as a pointer to a pointer}}
+#endif
 kernel void no_ptrptr(global int * global *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) { }
 
 // expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 __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{{nested 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{{nested 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{{nested 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
 
 typedef __SIZE_TYPE__ size_t; // expected-note{{'__private size_t' (aka '__private unsigned int') declared here}}
                               // expected-note@-1{{'size_t' (aka 'unsigned int') declared here}}
 typedef __PTRDIFF_TYPE__ ptrdiff_t; // expected-note{{'__private ptrdiff_t' (aka '__private int') declared here}}
 typedef __INTPTR_TYPE__ intptr_t; // expected-note{{'__private intptr_t' (aka '__private int') declared here}}