Index: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td @@ -7005,6 +7005,19 @@ "sending to parameter of different type}0,1" "|%diff{casting $ to type $|casting between types}0,1}2" " changes address space of pointer">; +def err_typecheck_incompatible_nested_address_space : Error< + "%select{%diff{assigning $ to $|assigning to different types}1,0" + "|%diff{passing $ to parameter of type $|" + "passing to parameter of different type}0,1" + "|%diff{returning $ from a function with result type $|" + "returning from function with different return type}0,1" + "|%diff{converting $ to type $|converting between types}0,1" + "|%diff{initializing $ with an expression of type $|" + "initializing with expression of different type}0,1" + "|%diff{sending $ to parameter of type $|" + "sending to parameter of different type}0,1" + "|%diff{casting $ to type $|casting between types}0,1}2" + " changes address space of nested pointer">; def err_typecheck_incompatible_ownership : Error< "%select{%diff{assigning $ to $|assigning to different types}1,0" "|%diff{passing $ to parameter of type $|" Index: cfe/trunk/include/clang/Sema/Sema.h =================================================================== --- cfe/trunk/include/clang/Sema/Sema.h +++ cfe/trunk/include/clang/Sema/Sema.h @@ -9726,6 +9726,12 @@ /// like address spaces. IncompatiblePointerDiscardsQualifiers, + /// IncompatibleNestedPointerAddressSpaceMismatch - The assignment + /// changes address spaces in nested pointer types which is not allowed. + /// For instance, converting __private int ** to __generic int ** is + /// illegal even though __private could be converted to __generic. + IncompatibleNestedPointerAddressSpaceMismatch, + /// IncompatibleNestedPointerQualifiers - The assignment is between two /// nested pointer types, and the qualifiers other than the first two /// levels differ e.g. char ** -> const char **, but we accept them as an Index: cfe/trunk/lib/Sema/SemaCast.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaCast.cpp +++ cfe/trunk/lib/Sema/SemaCast.cpp @@ -2323,19 +2323,41 @@ // In OpenCL only conversions between pointers to objects in overlapping // addr spaces are allowed. v2.0 s6.5.5 - Generic addr space overlaps // with any named one, except for constant. + + // Converting the top level pointee addrspace is permitted for compatible + // addrspaces (such as 'generic int *' to 'local int *' or vice versa), but + // if any of the nested pointee addrspaces differ, we emit a warning + // regardless of addrspace compatibility. This makes + // local int ** p; + // return (generic int **) p; + // warn even though local -> generic is permitted. if (Self.getLangOpts().OpenCL) { - auto SrcPtrType = SrcType->getAs(); - if (!SrcPtrType) - return; - auto DestPtrType = DestType->getAs(); - if (!DestPtrType) - return; - if (!DestPtrType->isAddressSpaceOverlapping(*SrcPtrType)) { - Self.Diag(OpRange.getBegin(), - diag::err_typecheck_incompatible_address_space) - << SrcType << DestType << Sema::AA_Casting - << SrcExpr.get()->getSourceRange(); - SrcExpr = ExprError(); + const Type *DestPtr, *SrcPtr; + bool Nested = false; + unsigned DiagID = diag::err_typecheck_incompatible_address_space; + DestPtr = Self.getASTContext().getCanonicalType(DestType.getTypePtr()), + SrcPtr = Self.getASTContext().getCanonicalType(SrcType.getTypePtr()); + + while (isa(DestPtr) && isa(SrcPtr)) { + const PointerType *DestPPtr = cast(DestPtr); + const PointerType *SrcPPtr = cast(SrcPtr); + QualType DestPPointee = DestPPtr->getPointeeType(); + QualType SrcPPointee = SrcPPtr->getPointeeType(); + if (Nested ? DestPPointee.getAddressSpace() != + SrcPPointee.getAddressSpace() + : !DestPPtr->isAddressSpaceOverlapping(*SrcPPtr)) { + Self.Diag(OpRange.getBegin(), DiagID) + << SrcType << DestType << Sema::AA_Casting + << SrcExpr.get()->getSourceRange(); + if (!Nested) + SrcExpr = ExprError(); + return; + } + + DestPtr = DestPPtr->getPointeeType().getTypePtr(); + SrcPtr = SrcPPtr->getPointeeType().getTypePtr(); + Nested = true; + DiagID = diag::ext_nested_pointer_qualifier_mismatch; } } } Index: cfe/trunk/lib/Sema/SemaExpr.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaExpr.cpp +++ cfe/trunk/lib/Sema/SemaExpr.cpp @@ -7724,9 +7724,9 @@ } if (!lhq.compatiblyIncludes(rhq)) { - // Treat address-space mismatches as fatal. TODO: address subspaces + // Treat address-space mismatches as fatal. if (!lhq.isAddressSpaceSupersetOf(rhq)) - ConvTy = Sema::IncompatiblePointerDiscardsQualifiers; + return Sema::IncompatiblePointerDiscardsQualifiers; // It's okay to add or remove GC or lifetime qualifiers when converting to // and from void*. @@ -7799,8 +7799,22 @@ // level of indirection, this must be the issue. if (isa(lhptee) && isa(rhptee)) { do { - lhptee = cast(lhptee)->getPointeeType().getTypePtr(); - rhptee = cast(rhptee)->getPointeeType().getTypePtr(); + std::tie(lhptee, lhq) = + cast(lhptee)->getPointeeType().split().asPair(); + std::tie(rhptee, rhq) = + cast(rhptee)->getPointeeType().split().asPair(); + + // Inconsistent address spaces at this point is invalid, even if the + // address spaces would be compatible. + // FIXME: This doesn't catch address space mismatches for pointers of + // different nesting levels, like: + // __local int *** a; + // int ** b = a; + // It's not clear how to actually determine when such pointers are + // invalidly incompatible. + if (lhq.getAddressSpace() != rhq.getAddressSpace()) + return Sema::IncompatibleNestedPointerAddressSpaceMismatch; + } while (isa(lhptee) && isa(rhptee)); if (lhptee == rhptee) @@ -14213,6 +14227,9 @@ case IncompatibleNestedPointerQualifiers: DiagKind = diag::ext_nested_pointer_qualifier_mismatch; break; + case IncompatibleNestedPointerAddressSpaceMismatch: + DiagKind = diag::err_typecheck_incompatible_nested_address_space; + break; case IntToBlockPointer: DiagKind = diag::err_int_to_block_pointer; break; Index: cfe/trunk/test/CodeGenOpenCL/numbered-address-space.cl =================================================================== --- cfe/trunk/test/CodeGenOpenCL/numbered-address-space.cl +++ cfe/trunk/test/CodeGenOpenCL/numbered-address-space.cl @@ -11,12 +11,6 @@ *generic_ptr = 4; } -// CHECK-LABEL: @test_numbered_as_to_builtin( -// CHECK: addrspacecast i32 addrspace(42)* %0 to float addrspace(3)* -void test_numbered_as_to_builtin(__attribute__((address_space(42))) int *arbitary_numbered_ptr, float src) { - volatile float result = __builtin_amdgcn_ds_fmaxf(arbitary_numbered_ptr, src, 0, 0, false); -} - // CHECK-LABEL: @test_generic_as_to_builtin_parameter_explicit_cast( // CHECK: addrspacecast i32 addrspace(3)* %0 to i32* void test_generic_as_to_builtin_parameter_explicit_cast(__local int *local_ptr, float src) { @@ -25,10 +19,7 @@ } // CHECK-LABEL: @test_generic_as_to_builtin_parameter_implicit_cast( -// CHECK: addrspacecast i32* %2 to float addrspace(3)* +// CHECK: bitcast i32 addrspace(3)* %0 to float addrspace(3)* void test_generic_as_to_builtin_parameter_implicit_cast(__local int *local_ptr, float src) { - generic int* generic_ptr = local_ptr; - - volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); + volatile float result = __builtin_amdgcn_ds_fmaxf(local_ptr, src, 0, 0, false); } - Index: cfe/trunk/test/SemaOpenCL/address-spaces.cl =================================================================== --- cfe/trunk/test/SemaOpenCL/address-spaces.cl +++ cfe/trunk/test/SemaOpenCL/address-spaces.cl @@ -124,6 +124,106 @@ p = (__private int *)p2; } +#if !__OPENCL_CPP_VERSION__ +void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) { + g = gg; // expected-error {{assigning '__global int **' to '__global int *' changes address space of pointer}} + g = l; // expected-error {{assigning '__local int *' to '__global int *' changes address space of pointer}} + g = ll; // expected-error {{assigning '__local int **' to '__global int *' changes address space of pointer}} + g = gg_f; // expected-error {{assigning '__global float **' to '__global int *' changes address space of pointer}} + g = (__global int *)gg_f; // expected-error {{casting '__global float **' to type '__global int *' changes address space of pointer}} + + gg = g; // expected-error {{assigning '__global int *' to '__global int **' changes address space of pointer}} + gg = l; // expected-error {{assigning '__local int *' to '__global int **' changes address space of pointer}} + gg = ll; // expected-error {{assigning '__local int **' to '__global int **' changes address space of nested pointer}} + gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int **' from '__global float **'}} + gg = (__global int * __private *)gg_f; + + l = g; // expected-error {{assigning '__global int *' to '__local int *' changes address space of pointer}} + l = gg; // expected-error {{assigning '__global int **' to '__local int *' changes address space of pointer}} + l = ll; // expected-error {{assigning '__local int **' to '__local int *' changes address space of pointer}} + l = gg_f; // expected-error {{assigning '__global float **' to '__local int *' changes address space of pointer}} + l = (__local int *)gg_f; // expected-error {{casting '__global float **' to type '__local int *' changes address space of pointer}} + + ll = g; // expected-error {{assigning '__global int *' to '__local int **' changes address space of pointer}} + ll = gg; // expected-error {{assigning '__global int **' to '__local int **' changes address space of nested pointer}} + ll = l; // expected-error {{assigning '__local int *' to '__local int **' changes address space of pointer}} + ll = gg_f; // expected-error {{assigning '__global float **' to '__local int **' changes address space of nested pointer}} + ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float **' to type '__local int **' discards qualifiers in nested pointer types}} + + gg_f = g; // expected-error {{assigning '__global int *' to '__global float **' changes address space of pointer}} + gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float **' from '__global int **'}} + gg_f = l; // expected-error {{assigning '__local int *' to '__global float **' changes address space of pointer}} + gg_f = ll; // expected-error {{assigning '__local int **' to '__global float **' changes address space of nested pointer}} + gg_f = (__global float * __private *)gg; + + // FIXME: This doesn't seem right. This should be an error, not a warning. + __local int * __global * __private * lll; + lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}} + + typedef __local int * l_t; + typedef __global int * g_t; + __private l_t * pl; + __private g_t * pg; + gg = pl; // expected-error {{assigning 'l_t *' (aka '__local int **') to '__global int **' changes address space of nested pointer}} + pl = gg; // expected-error {{assigning '__global int **' to 'l_t *' (aka '__local int **') changes address space of nested pointer}} + gg = pg; + pg = gg; + pg = pl; // expected-error {{assigning 'l_t *' (aka '__local int **') to 'g_t *' (aka '__global int **') changes address space of nested pointer}} + pl = pg; // expected-error {{assigning 'g_t *' (aka '__global int **') to 'l_t *' (aka '__local int **') changes address space of nested pointer}} + + ll = (__local int * __private *)(void *)gg; + void *vp = ll; +} +#else +void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) { + g = gg; // expected-error {{assigning to '__global int *' from incompatible type '__global int **'}} + g = l; // expected-error {{assigning to '__global int *' from incompatible type '__local int *'}} + g = ll; // expected-error {{assigning to '__global int *' from incompatible type '__local int **'}} + g = gg_f; // expected-error {{assigning to '__global int *' from incompatible type '__global float **'}} + g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__global int *' converts between mismatching address spaces}} + + gg = g; // expected-error {{assigning to '__global int **' from incompatible type '__global int *'; take the address with &}} + gg = l; // expected-error {{assigning to '__global int **' from incompatible type '__local int *'}} + gg = ll; // expected-error {{assigning to '__global int **' from incompatible type '__local int **'}} + gg = gg_f; // expected-error {{assigning to '__global int **' from incompatible type '__global float **'}} + gg = (__global int * __private *)gg_f; + + l = g; // expected-error {{assigning to '__local int *' from incompatible type '__global int *'}} + l = gg; // expected-error {{assigning to '__local int *' from incompatible type '__global int **'}} + l = ll; // expected-error {{assigning to '__local int *' from incompatible type '__local int **'}} + l = gg_f; // expected-error {{assigning to '__local int *' from incompatible type '__global float **'}} + l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float **' to '__local int *' converts between mismatching address spaces}} + + ll = g; // expected-error {{assigning to '__local int **' from incompatible type '__global int *'}} + ll = gg; // expected-error {{assigning to '__local int **' from incompatible type '__global int **'}} + ll = l; // expected-error {{assigning to '__local int **' from incompatible type '__local int *'; take the address with &}} + ll = gg_f; // expected-error {{assigning to '__local int **' from incompatible type '__global float **'}} + // FIXME: The below becomes a reinterpret_cast, and therefore does not emit an error + // even though the address space mismatches in the nested pointers. + ll = (__local int * __private *)gg; + + gg_f = g; // expected-error {{assigning to '__global float **' from incompatible type '__global int *'}} + gg_f = gg; // expected-error {{assigning to '__global float **' from incompatible type '__global int **'}} + gg_f = l; // expected-error {{assigning to '__global float **' from incompatible type '__local int *'}} + gg_f = ll; // expected-error {{assigning to '__global float **' from incompatible type '__local int **'}} + gg_f = (__global float * __private *)gg; + + typedef __local int * l_t; + typedef __global int * g_t; + __private l_t * pl; + __private g_t * pg; + gg = pl; // expected-error {{assigning to '__global int **' from incompatible type 'l_t *' (aka '__local int **')}} + pl = gg; // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type '__global int **'}} + gg = pg; + pg = gg; + pg = pl; // expected-error {{assigning to 'g_t *' (aka '__global int **') from incompatible type 'l_t *' (aka '__local int **')}} + pl = pg; // expected-error {{assigning to 'l_t *' (aka '__local int **') from incompatible type 'g_t *' (aka '__global int **')}} + + ll = (__local int * __private *)(void *)gg; + void *vp = ll; +} +#endif + __private int func_return_priv(void); //expected-error {{return value cannot be qualified with address space}} __global int func_return_global(void); //expected-error {{return value cannot be qualified with address space}} __local int func_return_local(void); //expected-error {{return value cannot be qualified with address space}} Index: cfe/trunk/test/SemaOpenCL/event_t_overload.cl =================================================================== --- cfe/trunk/test/SemaOpenCL/event_t_overload.cl +++ cfe/trunk/test/SemaOpenCL/event_t_overload.cl @@ -7,5 +7,5 @@ event_t evt; foo(evt, src1); foo(0, src2); - foo(evt, src3); // expected-error {{call to 'foo' is ambiguous}} + foo(evt, src3); // expected-error {{no matching function for call to 'foo'}} } Index: cfe/trunk/test/SemaOpenCL/numbered-address-space.cl =================================================================== --- cfe/trunk/test/SemaOpenCL/numbered-address-space.cl +++ cfe/trunk/test/SemaOpenCL/numbered-address-space.cl @@ -26,6 +26,6 @@ void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) { generic int* generic_ptr = as3_ptr; - volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__local float *'}} + volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-error {{passing '__generic int *' to parameter of type '__local float *' changes address space of pointer}} } Index: cfe/trunk/test/SemaOpenCL/queue_t_overload.cl =================================================================== --- cfe/trunk/test/SemaOpenCL/queue_t_overload.cl +++ cfe/trunk/test/SemaOpenCL/queue_t_overload.cl @@ -7,6 +7,6 @@ queue_t q; foo(q, src1); foo(0, src2); - foo(q, src3); // expected-error {{call to 'foo' is ambiguous}} + foo(q, src3); // expected-error {{no matching function for call to 'foo'}} foo(1, src3); // expected-error {{no matching function for call to 'foo'}} }