Author: davidtweed Date: Thu Mar 27 11:34:11 2014 New Revision: 204941 URL: http://llvm.org/viewvc/llvm-project?rev=204941&view=rev Log: Enforce the restriction that a parameter to a kernel function cannot be a pointer to the private address space (as clarified in the OpenCL 1.2 specification).
Patch by Fraser Cormack! Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/lib/Sema/SemaDecl.cpp cfe/trunk/test/SemaOpenCL/invalid-kernel.cl Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=204941&r1=204940&r2=204941&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Mar 27 11:34:11 2014 @@ -6816,6 +6816,8 @@ def err_static_kernel : Error< "kernel functions cannot be declared static">; def err_opencl_ptrptr_kernel_param : Error< "kernel parameter cannot be declared as a pointer to a pointer">; +def err_opencl_private_ptr_kernel_param : Error< + "kernel parameter cannot be declared as a pointer to the __private address space">; def err_static_function_scope : Error< "variables in function scope cannot be declared static">; def err_opencl_bitfields : Error< Modified: cfe/trunk/lib/Sema/SemaDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=204941&r1=204940&r2=204941&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaDecl.cpp Thu Mar 27 11:34:11 2014 @@ -6378,6 +6378,7 @@ enum OpenCLParamType { ValidKernelParam, PtrPtrKernelParam, PtrKernelParam, + PrivatePtrKernelParam, InvalidKernelParam, RecordKernelParam }; @@ -6385,7 +6386,10 @@ enum OpenCLParamType { static OpenCLParamType getOpenCLKernelParameterType(QualType PT) { if (PT->isPointerType()) { QualType PointeeType = PT->getPointeeType(); - return PointeeType->isPointerType() ? PtrPtrKernelParam : PtrKernelParam; + if (PointeeType->isPointerType()) + return PtrPtrKernelParam; + return PointeeType.getAddressSpace() == 0 ? PrivatePtrKernelParam + : PtrKernelParam; } // TODO: Forbid the other integer types (size_t, ptrdiff_t...) when they can @@ -6430,6 +6434,14 @@ static void checkIsValidOpenCLKernelPara D.setInvalidType(); return; + case PrivatePtrKernelParam: + // OpenCL v1.2 s6.9.a: + // A kernel function argument cannot be declared as a + // pointer to the private address space. + S.Diag(Param->getLocation(), diag::err_opencl_private_ptr_kernel_param); + 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 @@ -6509,7 +6521,8 @@ static void checkIsValidOpenCLKernelPara // Arguments to kernel functions that are declared to be a struct or union // do not allow OpenCL objects to be passed as elements of the struct or // union. - if (ParamType == PtrKernelParam || ParamType == PtrPtrKernelParam) { + if (ParamType == PtrKernelParam || ParamType == PtrPtrKernelParam || + ParamType == PrivatePtrKernelParam) { S.Diag(Param->getLocation(), diag::err_record_with_pointers_kernel_param) << PT->isUnionType() Modified: cfe/trunk/test/SemaOpenCL/invalid-kernel.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/invalid-kernel.cl?rev=204941&r1=204940&r2=204941&view=diff ============================================================================== --- cfe/trunk/test/SemaOpenCL/invalid-kernel.cl (original) +++ cfe/trunk/test/SemaOpenCL/invalid-kernel.cl Thu Mar 27 11:34:11 2014 @@ -2,6 +2,8 @@ kernel void no_ptrptr(global int **i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}} +__kernel void no_privateptr(__private int *i) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}} + kernel int bar() { // expected-error {{kernel must have void return type}} return 6; } _______________________________________________ cfe-commits mailing list cfe-commits@cs.uiuc.edu http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits