Hi,
This patch enforces the restriction in OpenCL 1.2 that a kernel
parameter cannot be declared as a pointer to the private address space.
See OpenCL 1.2 spec section 6.9a, first item.
Cheers,
Fraser
--
Fraser Cormack
Compiler Developer
Codeplay Software Ltd
45 York Place, Edinburgh, EH1 3HP
Tel: 0131 466 0503
Fax: 0131 557 6600
Website: http://www.codeplay.com
Twitter: https://twitter.com/codeplaysoft
This email and any attachments may contain confidential and /or privileged
information and is for use by the addressee only. If you are not the intended
recipient, please notify Codeplay Software Ltd immediately and delete the
message from your computer. You may not copy or forward it,or use or disclose
its contents to any other person. Any views or other information in this
message which do not relate to our business are not authorized by Codeplay
software Ltd, nor does this message form part of any contract unless so stated.
As internet communications are capable of data corruption Codeplay Software Ltd
does not accept any responsibility for any changes made to this message after
it was sent. Please note that Codeplay Software Ltd does not accept any
liability or responsibility for viruses and it is your responsibility to scan
any attachments.
Company registered in England and Wales, number: 04567874
Registered office: 81 Linkfield Street, Redhill RH1 6BY
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td (revision 204905)
+++ include/clang/Basic/DiagnosticSemaKinds.td (working copy)
@@ -6813,6 +6813,8 @@
"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<
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp (revision 204905)
+++ lib/Sema/SemaDecl.cpp (working copy)
@@ -6378,6 +6378,7 @@
ValidKernelParam,
PtrPtrKernelParam,
PtrKernelParam,
+ PrivatePtrKernelParam,
InvalidKernelParam,
RecordKernelParam
};
@@ -6385,7 +6386,10 @@
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 @@
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 @@
// 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()
Index: test/SemaOpenCL/invalid-kernel.cl
===================================================================
--- test/SemaOpenCL/invalid-kernel.cl (revision 204905)
+++ test/SemaOpenCL/invalid-kernel.cl (working copy)
@@ -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
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
