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

[Sven van Haastregt via Phabricator via cfe-commits]

svenvh updated this revision to Diff 308075.
svenvh added a comment.

Recursively call `getOpenCLKernelParameterType` to check address spaces.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D92091/new/

https://reviews.llvm.org/D92091

Files:
  clang/lib/Sema/SemaDecl.cpp
  clang/test/SemaOpenCL/invalid-kernel-parameters.cl


Index: clang/test/SemaOpenCL/invalid-kernel-parameters.cl
===================================================================
--- clang/test/SemaOpenCL/invalid-kernel-parameters.cl
+++ clang/test/SemaOpenCL/invalid-kernel-parameters.cl
@@ -5,8 +5,12 @@
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
-// expected-error at +1{{kernel parameter cannot be declared as a pointer to a pointer}}
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+// expected-error at +3{{kernel parameter cannot be declared as a pointer to a pointer}}
+// expected-error at +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 at +1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 __kernel void no_privateptr(__private int *i) { }
@@ -17,6 +21,15 @@
 // expected-error at +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 at -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 at -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 at -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
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -8593,12 +8593,21 @@
 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;
   }
 
@@ -8651,11 +8660,17 @@
 
   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:


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D92091.308075.patch
Type: text/x-patch
Size: 3816 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20201127/3a68600c/attachment-0001.bin>