[OpenCL] Restrict address space conversions in nested pointers

[Anastasia]

Summary

This patch is fixing the issue reported in:
http://lists.llvm.org/pipermail/cfe-dev/2020-January/064273.html
and in the bug:
https://bugs.llvm.org/show_bug.cgi?id=39674

Since address space conversion changes pointer representation it is not legal to do it for the nested pointers even with compatible address spaces. Because the address space conversion in the nested levels can't be generated. The behavior implemented by this patch is as follows:

• Reject any implicit conversions of nested pointers with address spaces.
• Reject address space of nested pointers conversions in safe casts e.g. const_cast or static_cast.
• Allow conversion in low level C-style or reinterpret_cast but with a warning (this aligns with OpenCL C behavior).

[rjmccall]

clang/test/Misc/warning-flags.c

41 ↗

(On Diff #240205)

Please give this warning a warning group rather than changing this test.

clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl

521

Can we specialize the diagnostic here so that we get the good diagnostic in both language modes?

[jeroen.dobbelaere]

This patch looks good to me.

I do agree with John that it would be good to add the warning to a warning group.

[Anastasia]

clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl

521

I am not sure of the approach so I am wondering if you have any good suggestions.

It seems in C++ Clang exits with quite generic Incompatible in the implicit conversion case because the functionality for doing the semantic checks doesn't generally set Sema::AssignConvertType.

https://clang.llvm.org/doxygen/SemaExpr_8cpp_source.html

Sema::CheckSingleAssignmentConstraints:


  if (getLangOpts().CPlusPlus) {
      if (!LHSType->isRecordType() && !LHSType->isAtomicType()) {
        // C++ 5.17p3: If the left operand is not of class type, the
        // expression is implicitly converted (C++ 4) to the
        // cv-unqualified type of the left operand.
        QualType RHSType = RHS.get()->getType();
        if (Diagnose) {
          RHS = PerformImplicitConversion(RHS.get(), LHSType.getUnqualifiedType(),
                                          AA_Assigning);
        } else {
          ...
        }
        if (RHS.isInvalid())
          return Incompatible;

However for example in IsStandardConversion I can see that Sema::AssignConvertType is being set but only for classifying the conversion sequence. It is never propagated outside.

http://clang.llvm.org/doxygen/SemaOverload_8cpp_source.html

IsStandardConversion:


    ExprResult ER = ExprResult{From};
    Sema::AssignConvertType Conv =
        S.CheckSingleAssignmentConstraints(ToType, ER,
                                           /*Diagnose=*/false,
                                           /*DiagnoseCFAudited=*/false,
                                           /*ConvertRHS=*/false);
    ImplicitConversionKind SecondConv;
    switch (Conv) {
    case Sema::Compatible:
      SecondConv = ICK_C_Only_Conversion;
      break;
    // For our purposes, discarding qualifiers is just as bad as using an
    // incompatible pointer. Note that an IncompatiblePointer conversion can drop
    // qualifiers, as well.
    case Sema::CompatiblePointerDiscardsQualifiers:
    case Sema::IncompatiblePointer:
    case Sema::IncompatiblePointerSign:
      SecondConv = ICK_Incompatible_Pointer_Conversion;
      break;
    default:
      return false;
    }

In most of other places it's not being used for C++ mode. I feel like C++ flow was not written to set the conversion type failure. Maybe there is something else I could use in C++ mode to specialize this diagnostic? Otherwise, it seems like quite a lot of Sema would have to be modified for C++ flow because implicit conversions are used in many different places.

[rjmccall]

Is there no follow-up code when actually emitting the failure diagnostic which tries to figure out a more specific cause of failure?

[Anastasia]

In D73360#1850001, @rjmccall wrote:

Is there no follow-up code when actually emitting the failure diagnostic which tries to figure out a more specific cause of failure?

After looking at this a bit more it seems C++ part calls the diagnostic code earlier:

Sema::PerformImplicitConversion:

   case ImplicitConversionSequence::BadConversion:
     bool Diagnosed =
         DiagnoseAssignmentResult(Incompatible, From->getExprLoc(), ToType,
                                  From->getType(), From, Action);

It just seems to always pass Incompatible. We could add follow up code before? This would be easier than modifying the existing flow completely. However, looking at this holistically we would need to repeat all checks classifying the failure types...

[rjmccall]

If it's reasonable to persist the failure kind, that would probably be good, but it might be a lot of work.

[Anastasia]

In D73360#1852147, @rjmccall wrote:

If it's reasonable to persist the failure kind, that would probably be good, but it might be a lot of work.

I have looked at it earlier and it indeed affects a lot of Sema in C++ mode. Besides amount of work I am not sure I understand the flow enough to come up with a good approach.

Maybe we can finalize this review first to fix the regression. Because I would like it to be ported to the release 10 branch too. Then I can spend some time to see if I can come up with something better for the diagnostic too. :)

[Anastasia]

Add warning into IncompatiblePointerTypesDiscardsQualifiers group.

[rjmccall]

Okay, we can go with this for now, since it does fix a bug.

[Anastasia]

In D73360#1857073, @rjmccall wrote:

Okay, we can go with this for now, since it does fix a bug.

Thanks! Btw I uploaded straw-man patch for improving the diagnostics: https://reviews.llvm.org/D74116.

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_bad_cxx_cast_vector_to_scalar_different_size : Error<
   "%select{||reinterpret_cast||C-style cast|}0 from vector %1 "
   "to scalar %2 of different size">;
 def err_bad_cxx_cast_scalar_to_vector_different_size : Error<
   "%select{||reinterpret_cast||C-style cast|}0 from scalar %1 "
   "to vector %2 of different size">;
 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">;
 def err_bad_rvalue_to_rvalue_cast : Error<
   "cannot cast from rvalue of type %1 to rvalue reference type %2; types are "
   "not compatible">;
 def err_bad_static_cast_pointer_nonpointer : Error<
   "cannot cast from type %1 to pointer type %2">;

clang/lib/Sema/SemaCast.cpp

if (DestType->isFunctionPointerType()) {
     // See above.
     Self.Diag(OpRange.getBegin(),
               Self.getLangOpts().CPlusPlus11 ?
                 diag::warn_cxx98_compat_cast_fn_obj : diag::ext_cast_fn_obj)
       << OpRange;
     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.
   // So we finish by allowing everything that remains - it's got to be two
   // object pointers.
   return SuccessResult;
 }
 

clang/lib/Sema/SemaOverload.cpp

 
 /// Perform a single iteration of the loop for checking if a qualification
 /// conversion is valid.
 ///
 /// Specifically, check whether any change between the qualifiers of \p
 /// 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();
   Qualifiers ToQuals = ToType.getQualifiers();
 
   // Ignore __unaligned qualifier if this type is void.
   if (ToType.getUnqualifiedType()->isVoidType())
     FromQuals.removeUnaligned();
if (FromQuals.getObjCGCAttr() != ToQuals.getObjCGCAttr() &&
     ToQuals.removeObjCGCAttr();
   }
 
   //   -- for every j > 0, if const is in cv 1,j then const is in cv
   //      2,j, and similarly for volatile.
   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
   //      every cv for 0 < k < j.
   if (!CStyle && FromQuals.getCVRQualifiers() != ToQuals.getCVRQualifiers() &&
       !PreviousToQualsIncludeConst)
     return false;
 
if (FromType.getUnqualifiedType() == ToType.getUnqualifiedType())
     return false;
 
   // (C++ 4.4p4):
   //   A conversion can add cv-qualifiers at levels other than the first
   //   in multi-level pointers, subject to the following rules: [...]
   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;
   }
 
   // We are left with FromType and ToType being the pointee types
   // after unwrapping the original FromType and ToType the same number
   // of times. If we unwrapped any pointers, and if FromType and
   // ToType have the same unqualified type (since we checked
do {
 
     // MS compiler ignores __unaligned qualifier for references; do the same.
     T1 = withoutUnaligned(Context, T1);
     T2 = withoutUnaligned(Context, T2);
 
     // 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))
                  ? Ref_Related
                  : Ref_Incompatible;
 
     // FIXME: Should we track this for any level other than the first?
     if (ObjCLifetimeConversion)

clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl

void test_pointer_chains() {
   var_as_as_int = 0 ? var_as_as_int : var_asc_asn_int;
 #if !__OPENCL_CPP_VERSION__
 // expected-warning-re@-2{{pointer type mismatch ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
 #else
 // expected-error-re@-4{{incompatible operand types ('__{{(generic|global|constant)}} int *__{{(generic|global|constant)}} *' and '__{{(local|global|constant)}} int *__{{(constant|local|global)}} *')}}
 #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
 }

clang/test/SemaOpenCL/address-spaces.cl

void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
   l = ll;    // expected-error {{assigning to '__local int *' from incompatible type '__local int *__private *__private'}}
   l = gg_f;  // expected-error {{assigning to '__local int *' from incompatible type '__global float *__private *__private'}}
   l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__local int *' converts between mismatching address spaces}}
 
   ll = g;    // expected-error {{assigning to '__local int *__private *' from incompatible type '__global int *__private'}}
   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'}}
   gg_f = l;  // expected-error {{assigning to '__global float *__private *' from incompatible type '__local int *__private'}}
   gg_f = ll; // expected-error {{assigning to '__global float *__private *' from incompatible type '__local int *__private *__private'}}
   gg_f = (__global float * __private *)gg;
 
   typedef __local int * l_t;

clang/test/SemaOpenCLCXX/address-space-castoperators.cl

+//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}}
+}

clang/test/SemaOpenCLCXX/address-space-deduction.cl

 void t3(float(((*fYZ)))[2]);
 //CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *__private)))[2])'
 void t4(float(((**fYZ)))[2]);
 //CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *__private))[2])'
 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);
   t4(&p);
   t5(&p);
 }

clang/test/SemaOpenCLCXX/address-space-references.cl

 // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify -fsyntax-only
 
 __global const int& f(__global float &ref) {
   return ref; // expected-error{{reference of type 'const __global int &' cannot bind to a temporary object because of address space mismatch}}
 }
 
 int bar(const __global unsigned int &i); // expected-note{{passing argument to parameter 'i' here}}
 //FIXME: With the overload below the call should be resolved
 // successfully. However, current overload resolution logic
 // can't detect this case and therefore fails.
 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);
+}