r199208 - [OpenCL] Disallow casts between address spaces.

[Alp Toker]

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

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