On 14/01/2014 13:15, David Tweed wrote:
So this is an absolutely fine commit, but I guess this does highlight
that, once people get seriously stuck into implementing OpenCL 2.0
(and so we'll acquire the generic address space, to/from which one can
cast pointers in other address spaces) we'll have subtly different
rules for the two versions, at which point "getLangOpts.OpenCL" is
probably not going to be fine grained enough. What's the thought on
how to handle that in general in clang?
Hi David,
There's precedent for this in LangOptions::ObjC1 / LangOptions::ObjC2.
Opts.ObjC1 = Opts.ObjC2 = 1;
I think it'd be best to tweak that approach when it comes to OpenCL
though, leaving the current 'OpenCL' flag as-is and introducing
'OpenCL2' as needed to make adjustments for the newer standard.
My reasoning is that "if (ObjC1) ..." has turned out in hindsight to be
unnatural given that it really means "any Objective-C version".
Alp.
On Tue, Jan 14, 2014 at 12:47 PM, Joey Gouly <[email protected]> wrote:
Author: joey
Date: Tue Jan 14 06:47:29 2014
New Revision: 199208
URL: http://llvm.org/viewvc/llvm-project?rev=199208&view=rev
Log:
[OpenCL] Disallow casts between address spaces.
Modified:
cfe/trunk/lib/Sema/SemaCast.cpp
cfe/trunk/test/SemaOpenCL/address-spaces.cl
Modified: cfe/trunk/lib/Sema/SemaCast.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCast.cpp?rev=199208&r1=199207&r2=199208&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCast.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCast.cpp Tue Jan 14 06:47:29 2014
@@ -2170,6 +2170,21 @@ void CastOperation::CheckCStyleCast() {
assert(!SrcType->isPlaceholderType());
+ // OpenCL v1 s6.5: Casting a pointer to address space A to a pointer to
+ // address space B is illegal.
+ if (Self.getLangOpts().OpenCL && DestType->isPointerType() &&
+ SrcType->isPointerType()) {
+ if (DestType->getPointeeType().getAddressSpace() !=
+ SrcType->getPointeeType().getAddressSpace()) {
+ Self.Diag(OpRange.getBegin(),
+ diag::err_typecheck_incompatible_address_space)
+ << SrcType << DestType << Sema::AA_Casting
+ << SrcExpr.get()->getSourceRange();
+ SrcExpr = ExprError();
+ return;
+ }
+ }
+
if (Self.RequireCompleteType(OpRange.getBegin(), DestType,
diag::err_typecheck_cast_to_incomplete)) {
SrcExpr = ExprError();
Modified: cfe/trunk/test/SemaOpenCL/address-spaces.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/address-spaces.cl?rev=199208&r1=199207&r2=199208&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/address-spaces.cl (original)
+++ cfe/trunk/test/SemaOpenCL/address-spaces.cl Tue Jan 14 06:47:29 2014
@@ -11,3 +11,29 @@ __kernel void foo(__global int *gip) {
ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes
address space of pointer}}
ip = &ci; // expected-error {{assigning '__constant int *' to 'int *'
changes address space of pointer}}
}
+
+void explicit_cast(global int* g, local int* l, constant int* c, private int*
p)
+{
+ g = (global int*) l; // expected-error {{casting '__local int *' to type
'__global int *' changes address space of pointer}}
+ g = (global int*) c; // expected-error {{casting '__constant int *' to
type '__global int *' changes address space of pointer}}
+ g = (global int*) p; // expected-error {{casting 'int *' to type
'__global int *' changes address space of pointer}}
+
+ l = (local int*) g; // expected-error {{casting '__global int *' to type
'__local int *' changes address space of pointer}}
+ l = (local int*) c; // expected-error {{casting '__constant int *' to
type '__local int *' changes address space of pointer}}
+ l = (local int*) p; // expected-error {{casting 'int *' to type '__local
int *' changes address space of pointer}}
+
+ c = (constant int*) g; // expected-error {{casting '__global int *' to type
'__constant int *' changes address space of pointer}}
+ c = (constant int*) l; // expected-error {{casting '__local int *' to type
'__constant int *' changes address space of pointer}}
+ c = (constant int*) p; // expected-error {{casting 'int *' to type
'__constant int *' changes address space of pointer}}
+
+ p = (private int*) g; // expected-error {{casting '__global int *' to type
'int *' changes address space of pointer}}
+ p = (private int*) l; // expected-error {{casting '__local int *' to type
'int *' changes address space of pointer}}
+ p = (private int*) c; // expected-error {{casting '__constant int *' to
type 'int *' changes address space of pointer}}
+}
+
+void ok_explicit_casts(global int *g, global int* g2, local int* l, local int*
l2, private int* p, private int* p2)
+{
+ g = (global int*) g2;
+ l = (local int*) l2;
+ p = (private int*) p2;
+}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
--
http://www.nuanti.com
the browser experts
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
