diff options
| -rw-r--r-- | clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaCast.cpp | 18 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaOverload.cpp | 24 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl | 26 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/address-spaces.cl | 4 | ||||
| -rw-r--r-- | clang/test/SemaOpenCLCXX/address-space-castoperators.cl | 12 | ||||
| -rw-r--r-- | clang/test/SemaOpenCLCXX/address-space-deduction.cl | 2 | ||||
| -rw-r--r-- | clang/test/SemaOpenCLCXX/address-space-references.cl | 13 |
8 files changed, 86 insertions, 17 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 504c905..104a2a5 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -6760,6 +6760,10 @@ def err_bad_cxx_cast_scalar_to_vector_different_size : Error< 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<IncompatiblePointerTypesDiscardsQualifiers>; def err_bad_lvalue_to_rvalue_cast : Error< "cannot cast from lvalue of type %1 to rvalue reference type %2; types are " "not compatible">; diff --git a/clang/lib/Sema/SemaCast.cpp b/clang/lib/Sema/SemaCast.cpp index a905ebc..7a8cbca 100644 --- a/clang/lib/Sema/SemaCast.cpp +++ b/clang/lib/Sema/SemaCast.cpp @@ -2311,6 +2311,24 @@ static TryCastResult TryReinterpretCast(Sema &Self, ExprResult &SrcExpr, 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. diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp index 9b89bac..858e7ae 100644 --- a/clang/lib/Sema/SemaOverload.cpp +++ b/clang/lib/Sema/SemaOverload.cpp @@ -3180,7 +3180,7 @@ static bool isNonTrivialObjCLifetimeConversion(Qualifiers FromQuals, /// 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(); @@ -3217,11 +3217,15 @@ static bool isQualificationConversionStep(QualType FromType, QualType ToType, 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 @@ -3262,9 +3266,9 @@ Sema::IsQualificationConversion(QualType FromType, QualType ToType, 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; } @@ -4508,7 +4512,7 @@ Sema::CompareReferenceRelationship(SourceLocation Loc, // 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)) diff --git a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl index 14f7cf3a..5efea21 100644 --- a/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl +++ b/clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl @@ -513,17 +513,37 @@ void test_pointer_chains() { #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 } diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl index bc9a07d8..07547ea 100644 --- a/clang/test/SemaOpenCL/address-spaces.cl +++ b/clang/test/SemaOpenCL/address-spaces.cl @@ -198,9 +198,7 @@ void nested(__global int *g, __global int * __private *gg, __local int *l, __loc 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'}} diff --git a/clang/test/SemaOpenCLCXX/address-space-castoperators.cl b/clang/test/SemaOpenCLCXX/address-space-castoperators.cl new file mode 100644 index 0000000..d61a9a7 --- /dev/null +++ b/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<int**>(locgen); //expected-error{{const_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}} + gengen = static_cast<int**>(locgen); //expected-error{{static_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}} +// CHECK-NOT: AddressSpaceConversion + gengen = reinterpret_cast<int**>(locgen); //expected-warning{{reinterpret_cast from '__local int *__generic *' to '__generic int *__generic *' changes address space of nested pointers}} +} diff --git a/clang/test/SemaOpenCLCXX/address-space-deduction.cl b/clang/test/SemaOpenCLCXX/address-space-deduction.cl index 0275ae5..6a81a8b 100644 --- a/clang/test/SemaOpenCLCXX/address-space-deduction.cl +++ b/clang/test/SemaOpenCLCXX/address-space-deduction.cl @@ -105,7 +105,7 @@ void t5(float (*(*fYZ))[2]); __kernel void k() { __local float x[2]; - __local float(*p)[2]; + float(*p)[2]; t1(x); t2(&x); t3(&x); diff --git a/clang/test/SemaOpenCLCXX/address-space-references.cl b/clang/test/SemaOpenCLCXX/address-space-references.cl index a6c7c84..068318d 100644 --- a/clang/test/SemaOpenCLCXX/address-space-references.cl +++ b/clang/test/SemaOpenCLCXX/address-space-references.cl @@ -13,3 +13,16 @@ int bar(const unsigned int &i); 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); +} |