maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Sven van Haastregt 2020-12-01 11:33:10 +00:00
parent cba4accda0
commit 523775f967
2 changed files with 38 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

29
clang/lib/Sema/SemaDecl.cpp
View File

@ -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:

17
clang/test/SemaOpenCL/invalid-kernel-parameters.cl
View File

@ -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