Index: lib/Sema/SemaCast.cpp =================================================================== --- lib/Sema/SemaCast.cpp +++ lib/Sema/SemaCast.cpp @@ -2282,19 +2282,39 @@ // 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 consider it illegal + // regardless of addrspace compatibility. This makes + // local int ** p; + // return (generic int **) p; + // illegal 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; + 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(), + diag::err_typecheck_incompatible_address_space) + << SrcType << DestType << Sema::AA_Casting + << SrcExpr.get()->getSourceRange(); + SrcExpr = ExprError(); + return; + } + + DestPtr = DestPPtr->getPointeeType().getTypePtr(); + SrcPtr = SrcPPtr->getPointeeType().getTypePtr(); + Nested = true; } } } Index: lib/Sema/SemaExpr.cpp =================================================================== --- lib/Sema/SemaExpr.cpp +++ lib/Sema/SemaExpr.cpp @@ -7705,7 +7705,7 @@ if (!lhq.compatiblyIncludes(rhq)) { // Treat address-space mismatches as fatal. TODO: address subspaces 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*. @@ -7778,6 +7778,12 @@ // level of indirection, this must be the issue. if (isa(lhptee) && isa(rhptee)) { do { + // Inconsistent address spaces at this point is invalid, even if the + // address spaces would be compatible. + if (cast(lhptee)->getPointeeType().getAddressSpace() != + cast(rhptee)->getPointeeType().getAddressSpace()) + return Sema::IncompatibleNestedPointerQualifiers; + lhptee = cast(lhptee)->getPointeeType().getTypePtr(); rhptee = cast(rhptee)->getPointeeType().getTypePtr(); } while (isa(lhptee) && isa(rhptee)); @@ -14179,9 +14185,34 @@ return false; DiagKind = diag::ext_typecheck_convert_discards_qualifiers; break; - case IncompatibleNestedPointerQualifiers: + case IncompatibleNestedPointerQualifiers: { + // Perform array-to-pointer decay if necessary. + if (SrcType->isArrayType()) SrcType = Context.getArrayDecayedType(SrcType); + + // Address space qualifier mismatch in nested pointers is always an error. + // Strip the pointers until we either hit a mismatch or hit the end of the + // nesting. + // Use pointee type here since this should only care about nested + // qualifiers. + // XXX: Canonical types? + const Type *SrcPtr = SrcType->getPointeeType().getTypePtr(); + const Type *DstPtr = DstType->getPointeeType().getTypePtr(); + DiagKind = diag::ext_nested_pointer_qualifier_mismatch; + while (SrcPtr->isPointerType() && DstPtr->isPointerType()) { + if (SrcPtr->getPointeeType().getAddressSpace() != + DstPtr->getPointeeType().getAddressSpace()) { + // XXX: Should this be a different variation of the error, like + // 'changes address space in nested pointer qualifiers'? + DiagKind = diag::err_typecheck_incompatible_address_space; + isInvalid = true; + break; + } + SrcPtr = SrcPtr->getPointeeType().getTypePtr(); + DstPtr = DstPtr->getPointeeType().getTypePtr(); + } break; + } case IntToBlockPointer: DiagKind = diag::err_int_to_block_pointer; break; Index: test/CodeGenOpenCL/numbered-address-space.cl =================================================================== --- test/CodeGenOpenCL/numbered-address-space.cl +++ 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: test/SemaOpenCL/address-spaces.cl =================================================================== --- test/SemaOpenCL/address-spaces.cl +++ test/SemaOpenCL/address-spaces.cl @@ -52,6 +52,76 @@ 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 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 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 pointer}} + ll = (__local int * __private *)gg_f; // expected-error {{casting '__global float **' to type '__local int **' changes address space of pointer}} + + 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 pointer}} + gg_f = (__global float * __private *)gg; + + // FIXME: This doesn't seem right. + __local int * __global * __private * lll; + lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}} +} +#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 {{casting '__global float **' to type '__global int *' changes address space of pointer}} + + 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 {{casting '__global float **' to type '__local int *' changes address space of pointer}} + + 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 **'}} + ll = (__local int * __private *)gg_f; // expected-error {{casting '__global float **' to type '__local int **' changes address space of pointer}} + + 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; +} +#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: test/SemaOpenCL/event_t_overload.cl =================================================================== --- test/SemaOpenCL/event_t_overload.cl +++ 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: test/SemaOpenCL/numbered-address-space.cl =================================================================== --- test/SemaOpenCL/numbered-address-space.cl +++ 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: test/SemaOpenCL/queue_t_overload.cl =================================================================== --- test/SemaOpenCL/queue_t_overload.cl +++ 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'}} }