r199208 - [OpenCL] Disallow casts between address spaces.

[David Tweed]

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?

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



-- 
cheers, dave tweed__________________________
high-performance computing and machine vision expert: david.tweed at gmail.com
"while having code so boring anyone can maintain it, use Python." --
attempted insult seen on slashdot