svenvh created this revision.
svenvh added a reviewer: Anastasia.
Herald added a subscriber: yaxunl.
Herald added a project: clang.
svenvh requested review of this revision.
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.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D92091
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
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@+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) { }
@@ -17,6 +21,15 @@
// 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
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -8590,14 +8590,18 @@
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;
}
@@ -8650,20 +8654,38 @@
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.
- S.Diag(Param->getLocation(), diag::err_opencl_ptrptr_kernel_param);
- D.setInvalidType();
+ 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;
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9749,8 +9749,8 @@
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<
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
