> 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".
We actually have LangOpts::OpenCLVersion, I could guard this on "<= 120", however that's not right because OpenCL 2.0 only allows casts to the generic address space, and since we don't have any support for the generic address space, I can't really do much. I'm going to leave this as-is, and when we (someone) starts work on OpenCL 2.0, they'll have to add support for the generic address space and tweak this a little. Make sense? Joey > > 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=1992 08&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 _______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
