r199208 - [OpenCL] Disallow casts between address spaces.

[Joey Gouly]

> 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 <joey.gouly at gmail.com> 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
>> cfe-commits at cs.uiuc.edu
>> http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
>
>

-- 
http://www.nuanti.com
the browser experts

_______________________________________________
cfe-commits mailing list
cfe-commits at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits