Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -6981,6 +6981,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: lib/Sema/SemaCast.cpp =================================================================== --- lib/Sema/SemaCast.cpp +++ lib/Sema/SemaCast.cpp @@ -2282,19 +2282,42 @@ // 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. + + // FIXME: C++ might want to emit different errors here. 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(); + SrcExpr = ExprError(); + return; + } + + DestPtr = DestPPtr->getPointeeType().getTypePtr(); + SrcPtr = SrcPPtr->getPointeeType().getTypePtr(); + Nested = true; + DiagID = diag::err_typecheck_incompatible_nested_address_space; } } } Index: lib/Sema/SemaExpr.cpp =================================================================== --- lib/Sema/SemaExpr.cpp +++ lib/Sema/SemaExpr.cpp @@ -7703,9 +7703,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*. @@ -7778,8 +7778,16 @@ // 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. + if (lhq.getAddressSpace() != rhq.getAddressSpace()) + return Sema::IncompatibleNestedPointerQualifiers; + } while (isa(lhptee) && isa(rhptee)); if (lhptee == rhptee) @@ -14179,9 +14187,32 @@ 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()) { + DiagKind = diag::err_typecheck_incompatible_nested_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 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-error {{casting '__global float **' to type '__local int **' changes address space of nested 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 nested 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 nested 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'}} }