Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -6679,6 +6679,10 @@ def err_bad_cxx_cast_vector_to_vector_different_size : Error< "%select{||reinterpret_cast||C-style cast|}0 from vector %1 " "to vector %2 of different size">; +def warn_bad_cxx_cast_nested_pointer_addr_space : Warning< + "%select{reinterpret_cast|C-style cast}0 from %1 to %2 " + "changes address space of nested pointers">, + InGroup; def err_bad_lvalue_to_rvalue_cast : Error< "cannot cast from lvalue of type %1 to rvalue reference type %2; types are " "not compatible">; Index: clang/lib/Sema/SemaCast.cpp =================================================================== --- clang/lib/Sema/SemaCast.cpp +++ clang/lib/Sema/SemaCast.cpp @@ -2311,6 +2311,24 @@ return SuccessResult; } + // Diagnose address space conversion in nested pointers. + QualType DestPtee = DestType->getPointeeType().isNull() + ? DestType->getPointeeType() + : DestType->getPointeeType()->getPointeeType(); + QualType SrcPtee = SrcType->getPointeeType().isNull() + ? SrcType->getPointeeType() + : SrcType->getPointeeType()->getPointeeType(); + while (!DestPtee.isNull() && !SrcPtee.isNull()) { + if (DestPtee.getAddressSpace() != SrcPtee.getAddressSpace()) { + Self.Diag(OpRange.getBegin(), + diag::warn_bad_cxx_cast_nested_pointer_addr_space) + << CStyle << SrcType << DestType << SrcExpr.get()->getSourceRange(); + break; + } + DestPtee = DestPtee->getPointeeType(); + SrcPtee = SrcPtee->getPointeeType(); + } + // C++ 5.2.10p7: A pointer to an object can be explicitly converted to // a pointer to an object of different type. // Void pointers are not specified, but supported by every compiler out there. Index: clang/lib/Sema/SemaOverload.cpp =================================================================== --- clang/lib/Sema/SemaOverload.cpp +++ clang/lib/Sema/SemaOverload.cpp @@ -3176,7 +3176,7 @@ /// FromType and \p ToType is permissible, given knowledge about whether every /// outer layer is const-qualified. static bool isQualificationConversionStep(QualType FromType, QualType ToType, - bool CStyle, + bool CStyle, bool IsTopLevel, bool &PreviousToQualsIncludeConst, bool &ObjCLifetimeConversion) { Qualifiers FromQuals = FromType.getQualifiers(); @@ -3213,11 +3213,15 @@ if (!CStyle && !ToQuals.compatiblyIncludes(FromQuals)) return false; - // For a C-style cast, just require the address spaces to overlap. - // FIXME: Does "superset" also imply the representation of a pointer is the - // same? We're assuming that it does here and in compatiblyIncludes. - if (CStyle && !ToQuals.isAddressSpaceSupersetOf(FromQuals) && - !FromQuals.isAddressSpaceSupersetOf(ToQuals)) + // If address spaces mismatch: + // - in top level it is only valid to convert to addr space that is a + // superset in all cases apart from C-style casts where we allow + // conversions between overlapping address spaces. + // - in non-top levels it is not a valid conversion. + if (ToQuals.getAddressSpace() != FromQuals.getAddressSpace() && + (!IsTopLevel || + !(ToQuals.isAddressSpaceSupersetOf(FromQuals) || + (CStyle && FromQuals.isAddressSpaceSupersetOf(ToQuals))))) return false; // -- if the cv 1,j and cv 2,j are different, then const is in @@ -3258,9 +3262,9 @@ bool PreviousToQualsIncludeConst = true; bool UnwrappedAnyPointer = false; while (Context.UnwrapSimilarTypes(FromType, ToType)) { - if (!isQualificationConversionStep(FromType, ToType, CStyle, - PreviousToQualsIncludeConst, - ObjCLifetimeConversion)) + if (!isQualificationConversionStep( + FromType, ToType, CStyle, !UnwrappedAnyPointer, + PreviousToQualsIncludeConst, ObjCLifetimeConversion)) return false; UnwrappedAnyPointer = true; } @@ -4504,7 +4508,7 @@ // If we find a qualifier mismatch, the types are not reference-compatible, // but are still be reference-related if they're similar. bool ObjCLifetimeConversion = false; - if (!isQualificationConversionStep(T2, T1, /*CStyle=*/false, + if (!isQualificationConversionStep(T2, T1, /*CStyle=*/false, TopLevel, PreviousToQualsIncludeConst, ObjCLifetimeConversion)) return (ConvertedReferent || Context.hasSimilarType(T1, T2)) Index: clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl =================================================================== --- clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl +++ clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl @@ -513,17 +513,37 @@ #endif // Case 3: Corresponded inner pointees has overlapping but not equivalent address spaces. - // FIXME: Should this really be allowed in C++ mode? var_as_as_int = var_asc_asc_int; -#if !__OPENCL_CPP_VERSION__ #ifdef GENERIC +#if !__OPENCL_CPP_VERSION__ // expected-error@-3 {{assigning '__local int *__local *__private' to '__generic int *__generic *__private' changes address space of nested pointer}} +#else +// expected-error@-5 {{assigning to '__generic int *__generic *' from incompatible type '__local int *__local *__private'}} #endif #endif - var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int; + + var_as_as_int = (AS int *AS *)var_asc_asc_int; +#ifdef GENERIC #if !__OPENCL_CPP_VERSION__ +// expected-warning@-3 {{casting '__local int *__local *' to type '__generic int *__generic *' discards qualifiers in nested pointer types}} +#else +// expected-warning@-5 {{C-style cast from '__local int *__local *' to '__generic int *__generic *' changes address space of nested pointers}} +#endif +#endif + + var_as_as_int = (AS int *AS *)var_asc_asn_int; +#if !__OPENCL_CPP_VERSION__ +// expected-warning-re@-2 {{casting '__{{global|local|constant}} int *__{{local|constant|global}} *' to type '__{{global|constant|generic}} int *__{{global|constant|generic}} *' discards qualifiers in nested pointer types}} +#else +// expected-warning-re@-4 {{C-style cast from '__{{global|local|constant}} int *__{{local|constant|global}} *' to '__{{global|constant|generic}} int *__{{global|constant|generic}} *' changes address space of nested pointers}} +#endif + + var_as_as_int = 0 ? var_as_as_int : var_asc_asc_int; #ifdef GENERIC +#if !__OPENCL_CPP_VERSION__ // expected-warning@-3{{pointer type mismatch ('__generic int *__generic *' and '__local int *__local *')}} +#else +// expected-error@-5 {{incompatible operand types ('__generic int *__generic *' and '__local int *__local *')}} #endif #endif } Index: clang/test/SemaOpenCL/address-spaces.cl =================================================================== --- clang/test/SemaOpenCL/address-spaces.cl +++ clang/test/SemaOpenCL/address-spaces.cl @@ -198,9 +198,7 @@ ll = gg; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global int *__private *__private'}} ll = l; // expected-error {{assigning to '__local int *__private *' from incompatible type '__local int *__private'; take the address with &}} ll = gg_f; // expected-error {{assigning to '__local int *__private *' from incompatible type '__global float *__private *__private'}} - // 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; + ll = (__local int *__private *)gg; //expected-warning{{C-style cast from '__global int *__private *' to '__local int *__private *' changes address space of nested pointers}} gg_f = g; // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private'}} gg_f = gg; // expected-error {{assigning to '__global float *__private *' from incompatible type '__global int *__private *__private'}} Index: clang/test/SemaOpenCLCXX/address-space-castoperators.cl =================================================================== --- /dev/null +++ clang/test/SemaOpenCLCXX/address-space-castoperators.cl @@ -0,0 +1,12 @@ +//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s + +void nester_ptr() { + local int * * locgen; + constant int * * congen; + int * * gengen; + + gengen = const_cast(locgen); //expected-error{{const_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}} + gengen = static_cast(locgen); //expected-error{{static_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}} +// CHECK-NOT: AddressSpaceConversion + gengen = reinterpret_cast(locgen); //expected-warning{{reinterpret_cast from '__local int *__generic *' to '__generic int *__generic *' changes address space of nested pointers}} +} Index: clang/test/SemaOpenCLCXX/address-space-deduction.cl =================================================================== --- clang/test/SemaOpenCLCXX/address-space-deduction.cl +++ clang/test/SemaOpenCLCXX/address-space-deduction.cl @@ -105,7 +105,7 @@ __kernel void k() { __local float x[2]; - __local float(*p)[2]; + float(*p)[2]; t1(x); t2(&x); t3(&x); Index: clang/test/SemaOpenCLCXX/address-space-references.cl =================================================================== --- clang/test/SemaOpenCLCXX/address-space-references.cl +++ clang/test/SemaOpenCLCXX/address-space-references.cl @@ -13,3 +13,16 @@ void foo() { bar(1) // expected-error{{binding reference of type 'const __global unsigned int' to value of type 'int' changes address space}} } + +// Test addr space conversion with nested pointers + +extern void nestptr(int *&); // expected-note {{candidate function not viable: no known conversion from '__global int *__private' to '__generic int *__generic &__private' for 1st argument}} +extern void nestptr_const(int * const &); // expected-note {{candidate function not viable: cannot pass pointer to address space '__constant' as a pointer to address space '__generic' in 1st argument}} +int test_nestptr(__global int *glob, __constant int *cons, int* gen) { + nestptr(glob); // expected-error{{no matching function for call to 'nestptr'}} + // Addr space conversion first occurs on a temporary. + nestptr_const(glob); + // No legal conversion between disjoint addr spaces. + nestptr_const(cons); // expected-error{{no matching function for call to 'nestptr_const'}} + return *(*cons ? glob : gen); +}