[OpenCL] Enable address spaces for references in C++

[Anastasia]

Summary

I first enabled AS deduction for references that allowed to inherit the right conversion diagnostics based on qualification conversion rules implemented earlier for the pointer type.

Then in order to tests the deduction rules fully, I had to enable some extra features from OpenCL 2.0 that are also valid in C++.

A number of ICEs fired in the CodeGen due to missing addrspacecast. Not convinced the current solution is good though. May be it would be cleaner to add a separate CastKind here - CK_LValueAddressSpaceConversion? Although I am not entirely clear about the benefits yet.

[Anastasia]

lib/AST/Expr.cpp

1609

I don't like this assert now. Would adding extra variable be cleaner here?

[rjmccall]

lib/AST/Expr.cpp

1609

Yeah, this assertion doesn't make any sense like this. It should be checking whether the cast is a gl-value and, if so, requiring the subexpression to also be a gl-value and then asserting the difference between the type. But you can certainly do an address-space conversion on l-values that just happen to be of pointer or block-pointer type.

lib/CodeGen/CGExpr.cpp

4252

Please use the performAddrSpaceCast target hook instead of directly constructing an LLVM addrspacecast.

lib/Sema/DeclSpec.cpp

576

Please update the comment above this.

lib/Sema/SemaDecl.cpp

7366

Please update the comment above this.

lib/Sema/SemaInit.cpp

7614

Please extract a function to do an l-value qualification conversion just in case we add more non-trivial conversions that we need to represent.

[Anastasia]

Addressed comments from John.

[rjmccall]

lib/AST/Expr.cpp

1609

No, if this is a gl-value cast, the assertion must ignore whether there's a pointee type, or it will be messed up on gl-values of pointer types.

That is, if I have a gl-value of type char * __private, I should be able to do an address-space promotion to get a gl-value of type char * __generic. It's okay that the pointers are into the same address space here — in fact, it's more than okay, it's necessary.

[Anastasia]

Changed the assert for address space conversion.

[rjmccall]

lib/AST/Expr.cpp

1609

Thanks, that's right now. Although please assert that the base has the same value kind; I've seen bugs before where ICEs tried to implicitly materialize their arguments, and it's really frustrating to root out.

lib/Sema/SemaExprCXX.cpp

4285

If ToType is a reference type, the address space will be on its pointee type.

[Anastasia]

• Extended assert
  • Handled AS of ToType

[rjmccall]

lib/Sema/SemaExprCXX.cpp

4289

Okay. But if ToType *isn't* a reference type, this will never be an address-space conversion. I feel like this code could be written more clearly to express what it's trying to do.

[Anastasia]

Rewrite how CastKind is set for reference and pointer type.

[Anastasia]

lib/Sema/SemaExprCXX.cpp

4289

I hope it makes more sense now. Btw, it also applies to pointer type.

[rjmccall]

lib/Sema/SemaExprCXX.cpp

4289

The logic is wrong for pointer types; if you're converting pointers, you need to be checking the address space of the pointee type of the from type.

It sounds like this is totally inadequately tested; please flesh out the test with all of these cases. While you're at it, please ensure that there are tests verifying that we don't allowing address-space changes in nested positions.

[Anastasia]

Fixed check for AS mismatch of pointer type and added missing test case

[Anastasia]

lib/Sema/SemaExprCXX.cpp

4289

Thanks for spotting this bug! The generated IR for the test was still correct because AS of FromType happened to correctly mismatch AS of pointee of ToType.

I failed to construct the test case where it would miss classifying addrspacecast due to OpenCL or C++ sema rules but I managed to add a case in which addrspacecast was incorrectly added for pointers where it wasn't needed (see line 36 of the test). I think this code is covered now.

As for the address space position in pointers, the following test checks the address spaces of pointers in addrspacecast. For the other program paths we also have a test with similar checks in test/CodeGenOpenCL/address-spaces-conversions.cl that we now run for C++ mode too.

BTW, while trying to construct a test case for the bug, I have discovered that multiple pointer indirection casting isn't working correctly. I.e. for the following program:

kernel void foo(){
   __private int** loc;
   int** loc_p = loc;
   **loc_p = 1;
}

We generate:

bitcast i32* addrspace(4)* %0 to i32 addrspace(4)* addrspace(4)*

in OpenCL C and then perform store over pointer in AS 4 (generic). We have now lost the information that the original pointer was in private AS and that the adjustment of AS segment has to be performed before accessing memory pointed by the pointer. Based on the current specification of addrspacecast in https://llvm.org/docs/LangRef.html#addrspacecast-to-instruction I am not very clear whether it can be used for this case without any modifications or clarifications and also what would happen if there are multiple AS mismatches. I am going to look at this issue separately in more details. In OpenCL C++ an ICE is triggered for this though. Let me know if you have any thoughts on this.

[rjmccall]

lib/Sema/SemaExprCXX.cpp

4289

Thanks, the check looks good now.

BTW, while trying to construct a test case for the bug, I have discovered that multiple pointer indirection casting isn't working correctly.

This needs to be an error in Sema. The only qualification conversions that should be allowed in general on nested pointers (i.e. on T in T** or T*&) are the basic C qualifiers: const, volatile, and restrict; any other qualification change there is unsound.

[Anastasia]

lib/Sema/SemaExprCXX.cpp

4289

I see. I guess it's because C++ rules don't cover address spaces.

It feels like it would be a regression for OpenCL C++ vs OpenCL C to reject nested pointers with address spaces because it was allowed before. :(

However, the generation for OpenCL C and C are incorrect currently. I will try to sort that all out as a separate patch though, if it makes sense?

[rjmccall]

lib/Sema/SemaExprCXX.cpp

4289

C++'s rules assume that qualifiers don't introduce real representation differences and that operations on qualified types are compatible with operations on unqualified types. That's not true of qualifiers in general: address space qualifiers can change representations, ARC qualifiers can have incompatible semantics, etc. There is no way to soundly implement a conversion from __private int ** to __generic int **, just there's no way to soundly implement a conversion from Derived ** to Base **.

If you want to allow this conversion anyway for source-compatibility reasons (and I don't think that's a good idea), it should be a bitcast.

[Anastasia]

Do you think there is anything else to do for this patch?

lib/Sema/SemaExprCXX.cpp

4289

Ok, then bitcast is not a good solution because it has an issue of loosing address space information. Perhaps disallowing it completely is a better approach in this case. I have created a bug to investigate it further and may be request some feedback from other OpenCL developers:
https://bugs.llvm.org/show_bug.cgi?id=39674

[rjmccall]

Thanks, LGTM.

[romanovvlad]

cfe/trunk/lib/CodeGen/CGExpr.cpp

4268 ↗

(On Diff #174378)

Hi,

It seems this code doesn't work correctly(repro at the end). TBAA information is lost here because MakeNaturalAlignPointeeAddrLValue constructs LValue with alignment of poinee type but TBAA info is taken from pointer itself what is strange enough. As a result, for example, memcpy with wrong size is generated for copy constructors.

Repro:

class P {
public:
  P(const P &Rhs) = default;

  long a;
  long b;
};

__kernel void foo(__global P* GPtr) {
  P Val = GPtr[0];
}

As a solution the line could be replaced with the following:

return MakeAddrLValue(Address(V, LV.getAddress().getAlignment()),
                                       E->getType(), LV.getBaseInfo(),
                                       CGM.getTBAAInfoForSubobject(LV, E->getType()));

To take all the information from the original pointer.

What do you think about solution?

[rjmccall]

cfe/trunk/lib/CodeGen/CGExpr.cpp

4268 ↗

(On Diff #174378)

Oh, yes, this should absolutely not be using MakeNaturalAlignPointerAddrLValue; it should be preserving all of the extra information from the original l-value, as you say.

I think TBAA information is independent of address-space qualification and can just be taken from the original LV directly instead of using getTBAAInfoForSubobject.

lib/AST/Expr.cpp

case CK_CopyAndAutoreleaseBlockObject:
     goto CheckNoBasePath;
 
   case CK_FunctionToPointerDecay:
     assert(getType()->isPointerType());
     assert(getSubExpr()->getType()->isFunctionType());
     goto CheckNoBasePath;
 
   case CK_AddressSpaceConversion:
-    assert(getType()->isPointerType() || getType()->isBlockPointerType());
-    assert(getSubExpr()->getType()->isPointerType() ||
-           getSubExpr()->getType()->isBlockPointerType());
-    assert(getType()->getPointeeType().getAddressSpace() !=
-           getSubExpr()->getType()->getPointeeType().getAddressSpace());
-    LLVM_FALLTHROUGH;
+    assert(/*If pointer type then addr spaces for pointees must differ*/
+           (((getType()->isPointerType() &&
+              getSubExpr()->getType()->isPointerType()) ||
+             (getType()->isBlockPointerType() &&
+              getSubExpr()->getType()->isBlockPointerType())) &&
+            getType()->getPointeeType().getAddressSpace() !=
+                getSubExpr()->getType()->getPointeeType().getAddressSpace()) ||
+           /*For non pointer type, addr spaces of type must differ*/
+           (!getType()->isPointerType() &&
+            !getSubExpr()->getType()->isPointerType() &&
+            !getType()->isBlockPointerType() &&
+            !getSubExpr()->getType()->isBlockPointerType() &&
+            (getType().getAddressSpace() !=
+             getSubExpr()->getType().getAddressSpace())));
+    goto CheckNoBasePath;
   // These should not have an inheritance path.
   case CK_Dynamic:
   case CK_ToUnion:
   case CK_ArrayToPointerDecay:
   case CK_NullToMemberPointer:
   case CK_NullToPointer:
   case CK_ConstructorConversion:
   case CK_IntegralToPointer:

lib/CodeGen/CGExpr.cpp

LValue CodeGenFunction::EmitCastLValue(const CastExpr *E) {
   case CK_MemberPointerToBoolean:
   case CK_ReinterpretMemberPointer:
   case CK_AnyPointerToBlockPointerCast:
   case CK_ARCProduceObject:
   case CK_ARCConsumeObject:
   case CK_ARCReclaimReturnedObject:
   case CK_ARCExtendBlockObject:
   case CK_CopyAndAutoreleaseBlockObject:
-  case CK_AddressSpaceConversion:
   case CK_IntToOCLSampler:
     return EmitUnsupportedLValue(E, "unexpected cast lvalue");
 
   case CK_Dependent:
     llvm_unreachable("dependent cast kind in IR gen!");
 
   case CK_BuiltinFnToFnPtr:
     llvm_unreachable("builtin functions are handled elsewhere");
case CK_LValueBitCast: {
     if (SanOpts.has(SanitizerKind::CFIUnrelatedCast))
       EmitVTablePtrCheckForCast(E->getType(), V.getPointer(),
                                 /*MayBeNull=*/false, CFITCK_UnrelatedCast,
                                 E->getBeginLoc());
 
     return MakeAddrLValue(V, E->getType(), LV.getBaseInfo(),
                           CGM.getTBAAInfoForSubobject(LV, E->getType()));
   }
+  case CK_AddressSpaceConversion: {
+    LValue LV = EmitLValue(E->getSubExpr());
+    QualType DestTy = getContext().getPointerType(E->getType());
+    Address V =
+        Builder.CreateAddrSpaceCast(LV.getAddress(), ConvertType(DestTy));
+
+    return MakeAddrLValue(V, DestTy, LV.getBaseInfo(),
+                          CGM.getTBAAInfoForSubobject(LV, DestTy));
+  }
   case CK_ObjCObjectLValueCast: {
     LValue LV = EmitLValue(E->getSubExpr());
     Address V = Builder.CreateElementBitCast(LV.getAddress(),
                                              ConvertType(E->getType()));
     return MakeAddrLValue(V, E->getType(), LV.getBaseInfo(),
                           CGM.getTBAAInfoForSubobject(LV, E->getType()));
   }
   case CK_ZeroToOCLQueue:

lib/Sema/DeclSpec.cpp

bool DeclSpec::SetStorageClassSpec(Sema &S, SCS SC, SourceLocation Loc,
   // OpenCL v1.2 s6.8 changes this to "The auto and register storage-class
   // specifiers are not supported."
   if (S.getLangOpts().OpenCL &&
       !S.getOpenCLOptions().isEnabled("cl_clang_storage_class_specifiers")) {
     switch (SC) {
     case SCS_extern:
     case SCS_private_extern:
     case SCS_static:
-      if (S.getLangOpts().OpenCLVersion < 120) {
-        DiagID   = diag::err_opencl_unknown_type_specifier;
+      if (S.getLangOpts().OpenCLVersion < 120 &&
+          !S.getLangOpts().OpenCLCPlusPlus) {
+        DiagID = diag::err_opencl_unknown_type_specifier;
         PrevSpec = getSpecifierName(SC);
         return true;
       }
       break;
     case SCS_auto:
     case SCS_register:
       DiagID   = diag::err_opencl_unknown_type_specifier;
       PrevSpec = getSpecifierName(SC);

lib/Sema/SemaDecl.cpp

if (getLangOpts().OpenCL) {
     // OpenCL v2.0 s6.5.1 - Variables defined at program scope and static
     // variables inside a function can also be declared in the global
     // address space.
     if (NewVD->isFileVarDecl() || NewVD->isStaticLocal() ||
         NewVD->hasExternalStorage()) {
       if (!T->isSamplerT() &&
           !(T.getAddressSpace() == LangAS::opencl_constant ||
             (T.getAddressSpace() == LangAS::opencl_global &&
-             getLangOpts().OpenCLVersion == 200))) {
+             (getLangOpts().OpenCLVersion == 200 ||
+              getLangOpts().OpenCLCPlusPlus)))) {
         int Scope = NewVD->isStaticLocal() | NewVD->hasExternalStorage() << 1;
-        if (getLangOpts().OpenCLVersion == 200)
+        if (getLangOpts().OpenCLVersion == 200 || getLangOpts().OpenCLCPlusPlus)
           Diag(NewVD->getLocation(), diag::err_opencl_global_invalid_addr_space)
               << Scope << "global or constant";
         else
           Diag(NewVD->getLocation(), diag::err_opencl_global_invalid_addr_space)
               << Scope << "constant";
         NewVD->setInvalidDecl();
         return;
       }

lib/Sema/SemaExprCXX.cpp

case ICK_Function_Conversion:
 
     From = ImpCastExprToType(From, ToType, CK_NoOp,
                              VK_RValue, /*BasePath=*/nullptr, CCK).get();
     break;
 
   case ICK_Qualification: {
     // The qualification keeps the category of the inner expression, unless the
     // target type isn't a reference.
-    ExprValueKind VK = ToType->isReferenceType() ?
-                                  From->getValueKind() : VK_RValue;
-    From = ImpCastExprToType(From, ToType.getNonLValueExprType(Context),
-                             CK_NoOp, VK, /*BasePath=*/nullptr, CCK).get();
+    ExprValueKind VK =
+        ToType->isReferenceType() ? From->getValueKind() : VK_RValue;
+
+    CastKind CK = (ToType->getPointeeType().getAddressSpace() !=
+                   From->getType()->getPointeeType().getAddressSpace())
+                      ? CK_AddressSpaceConversion
+                      : CK_NoOp;
+
+    From = ImpCastExprToType(From, ToType.getNonLValueExprType(Context), CK, VK,
+                             /*BasePath=*/nullptr, CCK)
+               .get();
 
     if (SCS.DeprecatedStringLiteralToCharPtr &&
         !getLangOpts().WritableStrings) {
       Diag(From->getBeginLoc(),
            getLangOpts().CPlusPlus11
                ? diag::ext_deprecated_string_literal_conversion
                : diag::warn_deprecated_string_literal_conversion)
           << ToType.getNonReferenceType();

lib/Sema/SemaInit.cpp

case SK_UserConversion: {
       break;
     }
 
     case SK_QualificationConversionLValue:
     case SK_QualificationConversionXValue:
     case SK_QualificationConversionRValue: {
       // Perform a qualification conversion; these can never go wrong.
       ExprValueKind VK =
-          Step->Kind == SK_QualificationConversionLValue ?
-              VK_LValue :
-              (Step->Kind == SK_QualificationConversionXValue ?
-                   VK_XValue :
-                   VK_RValue);
-      CurInit = S.ImpCastExprToType(CurInit.get(), Step->Type, CK_NoOp, VK);
+          Step->Kind == SK_QualificationConversionLValue
+              ? VK_LValue
+              : (Step->Kind == SK_QualificationConversionXValue ? VK_XValue
+                                                                : VK_RValue);
+      CastKind CK = (Step->Type.getAddressSpace() !=
+                     CurInit.get()->getType().getAddressSpace())
+                        ? CK_AddressSpaceConversion
+                        : CK_NoOp;
+      CurInit = S.ImpCastExprToType(CurInit.get(), Step->Type, CK, VK);
       break;
     }
 
     case SK_AtomicConversion: {
       assert(CurInit.get()->isRValue() && "cannot convert glvalue to atomic");
       CurInit = S.ImpCastExprToType(CurInit.get(), Step->Type,
                                     CK_NonAtomicToAtomic, VK_RValue);
       break;

lib/Sema/SemaType.cpp

static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
   // space, except when it is used as a pointee type.
   //
   // Since LLVM assumes function type is in default address space, it should not
   // have address space.
   auto ChunkIndex = State.getCurrentChunkIndex();
   bool IsPointee =
       ChunkIndex > 0 &&
       (D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Pointer ||
-       D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer);
+       D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer ||
+       D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Reference);
   bool IsFuncReturnType =
       ChunkIndex > 0 &&
       D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function;
   bool IsFuncType =
       ChunkIndex < D.getNumTypeObjects() &&
       D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
   if ( // Do not deduce addr space for function return type and function type,
        // otherwise it will fail some sema check.

test/CodeGenOpenCLCXX/address-space-deduction.cl

+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -O0 -emit-llvm -o - | FileCheck %s -check-prefixes=COMMON,PTR
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -O0 -emit-llvm -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF
+
+#ifdef REF
+#define PTR &
+#define ADR(x) x
+#else
+#define PTR *
+#define ADR(x) &x
+#endif
+
+//COMMON: @glob = addrspace(1) global i32
+int glob;
+//PTR: @glob_p = addrspace(1) global i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @glob to i32 addrspace(4)*)
+//REF: @glob_p = addrspace(1) global i32 addrspace(4)* null
+int PTR glob_p = ADR(glob);
+
+//COMMON: @_ZZ3fooi{{P|R}}U3AS4iE6loc_st = internal addrspace(1) global i32
+//PTR: @_ZZ3fooiPU3AS4iE8loc_st_p = internal addrspace(1) global i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @_ZZ3fooiPU3AS4iE6loc_st to i32 addrspace(4)*)
+//REF: @_ZZ3fooiRU3AS4iE8loc_st_p = internal addrspace(1) global i32 addrspace(4)* null
+//COMMON: @loc_ext_p = external addrspace(1) {{global|constant}} i32 addrspace(4)*
+//COMMON: @loc_ext = external addrspace(1) global i32
+
+//REF: store i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @glob to i32 addrspace(4)*), i32 addrspace(4)* addrspace(1)* @glob_p
+
+//COMMON: define spir_func i32 @_Z3fooi{{P|R}}U3AS4i(i32 %par, i32 addrspace(4)*{{.*}} %par_p)
+int foo(int par, int PTR par_p){
+//COMMON: %loc = alloca i32
+  int loc;
+//COMMON: %loc_p = alloca i32 addrspace(4)*
+//COMMON: [[GAS:%[0-9]+]] = addrspacecast i32* %loc to i32 addrspace(4)*
+//COMMON: store i32 addrspace(4)* [[GAS]], i32 addrspace(4)** %loc_p
+  int PTR loc_p = ADR(loc);
+
+// CHECK directives for the following code are located above.
+  static int loc_st;
+  static int PTR loc_st_p = ADR(loc_st);
+  extern int loc_ext;
+  extern int PTR loc_ext_p;
+  (void)loc_ext_p;
+  return loc_ext;
+}