[OpenCL] Address space for default class members

[mikael]

Summary

• Implicity and explicity defaulted class members
• Resolved the FIXMEs in addrspace-of-this.cl

[mikael]

lib/Sema/SemaInit.cpp

4539

This is probably not the correct way to solve this. But it fixes the issues that i get.

The code that triggers this issue is found in test/CodeGenOpenCLCXX/addrspace-of-this.cl:47

C c3 = c1 + c2;

So the original (InitializedEntity) Entity.getType() looks like:

RValueReferenceType 0x717470 '__generic class C &&'
`-QualType 0x717458 '__generic class C' __generic
  `-RecordType 0x717160 'class C'
    `-CXXRecord 0x7170d0 'C'

It might be so that we actually want to qualify the RValueReferenceType rather than the RecordType here.

[rjmccall]

lib/AST/ASTContext.cpp

2781

You're trying to handle a method qualifier, not a type a functions that are themselves in some non-standard address space, right? The method qualifier should already be part of Proto->getExtProtoInfo(), so if there's an address space qualifier out here, something is very wrong.

[Anastasia]

lib/AST/ASTContext.cpp

2781

As far as I understand the new design, we have an address space qualifier on a method and as a part of the function prototype too. Are you saying that we need to make sure the prototype has an address space too?

[rjmccall]

lib/AST/ASTContext.cpp

2781

The address-space this qualifiers are a part of FunctionProtoType that's totally independent from the storage of top-level qualifiers that's part of QualType. Orig.getAddressSpace() is asking about the top-level qualifiers, not the this qualifiers. Generally, function types should not have top-level qualifiers at all (although of course a member pointer can have both top-level qualifiers and this qualifiers, with the former meaning something completely different from the latter).

[ebevhan]

lib/AST/ASTContext.cpp

2781

I suspect the reason there has to be qualifiers on the function type itself is because of the getType mistake I mentioned in the previous patch. If the this address space qualifiers aren't also on the function type, then GetThisType will not produce the correctly qualified type.

If the getType in GetThisType is changed to getTypeQualifiers, then the function type should not need any qualifiers.

[Anastasia]

lib/AST/ASTContext.cpp

2781

I fixed the issues reported here and in https://reviews.llvm.org/D54862 on the following review: https://reviews.llvm.org/D56066.

[Anastasia]

@mikael can you please close this review, because I opened another one to continue this: https://reviews.llvm.org/D56066.

Thanks for your work!

include/clang/Sema/Sema.h

if (New->isExternallyDeclarable()) {
       assert(Old->isExternallyDeclarable() &&
              "should not have found a non-externally-declarable previous decl");
       return true;
     }
     return false;
   }
   bool shouldLinkPossiblyHiddenDecl(LookupResult &Old, const NamedDecl *New);
 
+  void setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem,
+                                      QualType ResultTy,
+                                      ArrayRef<QualType> Args);
+
 public:
   typedef OpaquePtr<DeclGroupRef> DeclGroupPtrTy;
   typedef OpaquePtr<TemplateName> TemplateTy;
   typedef OpaquePtr<QualType> TypeTy;
 
   OpenCLOptions OpenCLFeatures;
   FPOptions FPFeatures;
 

lib/AST/ASTContext.cpp

if (const auto *AT = dyn_cast<AttributedType>(Orig))
     return getAttributedType(
         AT->getAttrKind(),
         getFunctionTypeWithExceptionSpec(AT->getModifiedType(), ESI),
         getFunctionTypeWithExceptionSpec(AT->getEquivalentType(), ESI));
 
   // Anything else must be a function type. Rebuild it with the new exception
   // specification.
   const auto *Proto = cast<FunctionProtoType>(Orig);
-  return getFunctionType(
-      Proto->getReturnType(), Proto->getParamTypes(),
+
+  QualType NewT =
+      getFunctionType(Proto->getReturnType(), Proto->getParamTypes(),
                       Proto->getExtProtoInfo().withExceptionSpec(ESI));
+
+  return getAddrSpaceQualType(NewT, Orig.getAddressSpace());
 }
 
 bool ASTContext::hasSameFunctionTypeIgnoringExceptionSpec(QualType T,
                                                           QualType U) {
   return hasSameType(T, U) ||
          (getLangOpts().CPlusPlus17 &&
           hasSameType(getFunctionTypeWithExceptionSpec(T, EST_None),
                       getFunctionTypeWithExceptionSpec(U, EST_None)));

lib/Sema/SemaDeclCXX.cpp

if (CSM == CXXCopyConstructor)
     CanHaveConstParam = RD->implicitCopyConstructorHasConstParam();
   else if (CSM == CXXCopyAssignment)
     CanHaveConstParam = RD->implicitCopyAssignmentHasConstParam();
 
   QualType ReturnType = Context.VoidTy;
   if (CSM == CXXCopyAssignment || CSM == CXXMoveAssignment) {
     // Check for return type matching.
     ReturnType = Type->getReturnType();
-    QualType ExpectedReturnType =
-        Context.getLValueReferenceType(Context.getTypeDeclType(RD));
+
+    QualType DeclType = Context.getTypeDeclType(RD);
+    DeclType = Context.getAddrSpaceQualType(DeclType, MD->getType().getAddressSpace());
+
+    QualType ExpectedReturnType = Context.getLValueReferenceType(DeclType);
+
     if (!Context.hasSameType(ReturnType, ExpectedReturnType)) {
       Diag(MD->getLocation(), diag::err_defaulted_special_member_return_type)
         << (CSM == CXXMoveAssignment) << ExpectedReturnType;
       HadError = true;
     }
 
     // A defaulted special member cannot have cv-qualifiers.
-    if (Type->getTypeQuals()) {
+    if (Type->getTypeQuals().hasConst() || Type->getTypeQuals().hasVolatile()) {
       if (DeleteOnTypeMismatch)
         ShouldDeleteForTypeMismatch = true;
       else {
         Diag(MD->getLocation(), diag::err_defaulted_special_member_quals)
           << (CSM == CXXMoveAssignment) << getLangOpts().CPlusPlus14;
         HadError = true;
       }
     }
for (auto *D : FD->getParent()->lookup(Name))
     if (auto *Acceptable = R.getAcceptableDecl(D))
       R.addDecl(Acceptable);
   R.resolveKind();
   R.suppressDiagnostics();
 
   CheckFunctionDeclaration(S, FD, R, /*IsMemberSpecialization*/false);
 }
 
+void Sema::setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem,
+                                          QualType ResultTy,
+                                          ArrayRef<QualType> Args) {
+  // Build an exception specification pointing back at this constructor.
+  FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, SpecialMem);
+
+  QualType QT;
+  if (getLangOpts().OpenCLCPlusPlus) {
+    // OpenCL: Implicitly defaulted special member are of the generic address
+    // space.
+    EPI.TypeQuals.addAddressSpace(LangAS::opencl_generic);
+
+    QT = Context.getFunctionType(ResultTy, Args, EPI);
+    QT = Context.getAddrSpaceQualType(QT, LangAS::opencl_generic);
+  } else {
+    QT = Context.getFunctionType(ResultTy, Args, EPI);
+  }
+
+  SpecialMem->setType(QT);
+}
+
 CXXConstructorDecl *Sema::DeclareImplicitDefaultConstructor(
                                                      CXXRecordDecl *ClassDecl) {
   // C++ [class.ctor]p5:
   //   A default constructor for a class X is a constructor of class X
   //   that can be called without an argument. If there is no
   //   user-declared constructor for class X, a default constructor is
   //   implicitly declared. An implicitly-declared default constructor
   //   is an inline public member of its class.
CXXConstructorDecl *Sema::DeclareImplicitDefaultConstructor(
 
   if (getLangOpts().CUDA) {
     inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXDefaultConstructor,
                                             DefaultCon,
                                             /* ConstRHS */ false,
                                             /* Diagnose */ false);
   }
 
-  // Build an exception specification pointing back at this constructor.
-  FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, DefaultCon);
-  DefaultCon->setType(Context.getFunctionType(Context.VoidTy, None, EPI));
+  setupImplicitSpecialMemberType(DefaultCon, Context.VoidTy, None);
 
   // We don't need to use SpecialMemberIsTrivial here; triviality for default
   // constructors is easy to compute.
   DefaultCon->setTrivial(ClassDecl->hasTrivialDefaultConstructor());
 
   // Note that we have declared this constructor.
   ++ASTContext::NumImplicitDefaultConstructorsDeclared;
 
CXXDestructorDecl *Sema::DeclareImplicitDestructor(CXXRecordDecl *ClassDecl) {
 
   if (getLangOpts().CUDA) {
     inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXDestructor,
                                             Destructor,
                                             /* ConstRHS */ false,
                                             /* Diagnose */ false);
   }
 
-  // Build an exception specification pointing back at this destructor.
-  FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, Destructor);
-  Destructor->setType(Context.getFunctionType(Context.VoidTy, None, EPI));
+  setupImplicitSpecialMemberType(Destructor, Context.VoidTy, None);
 
   // We don't need to use SpecialMemberIsTrivial here; triviality for
   // destructors is easy to compute.
   Destructor->setTrivial(ClassDecl->hasTrivialDestructor());
   Destructor->setTrivialForCall(ClassDecl->hasAttr<TrivialABIAttr>() ||
                                 ClassDecl->hasTrivialDestructorForCall());
 
   // Note that we have declared this destructor.
CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) {
   if (DSM.isAlreadyBeingDeclared())
     return nullptr;
 
   QualType ArgType = Context.getTypeDeclType(ClassDecl);
   QualType RetType = Context.getLValueReferenceType(ArgType);
   bool Const = ClassDecl->implicitCopyAssignmentHasConstParam();
   if (Const)
     ArgType = ArgType.withConst();
+
+  if (Context.getLangOpts().OpenCLCPlusPlus)
+    ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
+
   ArgType = Context.getLValueReferenceType(ArgType);
 
   bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
                                                      CXXCopyAssignment,
                                                      Const);
 
   //   An implicitly-declared copy assignment operator is an inline public
   //   member of its class.
CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) {
 
   if (getLangOpts().CUDA) {
     inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXCopyAssignment,
                                             CopyAssignment,
                                             /* ConstRHS */ Const,
                                             /* Diagnose */ false);
   }
 
-  // Build an exception specification pointing back at this member.
-  FunctionProtoType::ExtProtoInfo EPI =
-      getImplicitMethodEPI(*this, CopyAssignment);
-  CopyAssignment->setType(Context.getFunctionType(RetType, ArgType, EPI));
+  setupImplicitSpecialMemberType(CopyAssignment, RetType, ArgType);
 
   // Add the parameter to the operator.
   ParmVarDecl *FromParam = ParmVarDecl::Create(Context, CopyAssignment,
                                                ClassLoc, ClassLoc,
                                                /*Id=*/nullptr, ArgType,
                                                /*TInfo=*/nullptr, SC_None,
                                                nullptr);
   CopyAssignment->setParams(FromParam);
void Sema::DefineImplicitMoveAssignment(SourceLocation CurrentLocation,
 
   // The statements that form the synthesized function body.
   SmallVector<Stmt*, 8> Statements;
 
   // The parameter for the "other" object, which we are move from.
   ParmVarDecl *Other = MoveAssignOperator->getParamDecl(0);
   QualType OtherRefType = Other->getType()->
       getAs<RValueReferenceType>()->getPointeeType();
-  assert(!OtherRefType.getQualifiers() &&
-         "Bad argument type of defaulted move assignment");
 
   // Our location for everything implicitly-generated.
   SourceLocation Loc = MoveAssignOperator->getEndLoc().isValid()
                            ? MoveAssignOperator->getEndLoc()
                            : MoveAssignOperator->getLocation();
 
   // Builds a reference to the "other" object.
   RefBuilder OtherRef(Other, OtherRefType);
CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
   if (DSM.isAlreadyBeingDeclared())
     return nullptr;
 
   QualType ClassType = Context.getTypeDeclType(ClassDecl);
   QualType ArgType = ClassType;
   bool Const = ClassDecl->implicitCopyConstructorHasConstParam();
   if (Const)
     ArgType = ArgType.withConst();
+
+  if (Context.getLangOpts().OpenCLCPlusPlus)
+    ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
+
   ArgType = Context.getLValueReferenceType(ArgType);
 
   bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
                                                      CXXCopyConstructor,
                                                      Const);
 
   DeclarationName Name
     = Context.DeclarationNames.getCXXConstructorName(
CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
 
   if (getLangOpts().CUDA) {
     inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXCopyConstructor,
                                             CopyConstructor,
                                             /* ConstRHS */ Const,
                                             /* Diagnose */ false);
   }
 
-  // Build an exception specification pointing back at this member.
-  FunctionProtoType::ExtProtoInfo EPI =
-      getImplicitMethodEPI(*this, CopyConstructor);
-  CopyConstructor->setType(
-      Context.getFunctionType(Context.VoidTy, ArgType, EPI));
+  setupImplicitSpecialMemberType(CopyConstructor, Context.VoidTy, ArgType);
 
   // Add the parameter to the constructor.
   ParmVarDecl *FromParam = ParmVarDecl::Create(Context, CopyConstructor,
                                                ClassLoc, ClassLoc,
                                                /*IdentifierInfo=*/nullptr,
                                                ArgType, /*TInfo=*/nullptr,
                                                SC_None, nullptr);
   CopyConstructor->setParams(FromParam);
CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
                                                     CXXRecordDecl *ClassDecl) {
   assert(ClassDecl->needsImplicitMoveConstructor());
 
   DeclaringSpecialMember DSM(*this, ClassDecl, CXXMoveConstructor);
   if (DSM.isAlreadyBeingDeclared())
     return nullptr;
 
   QualType ClassType = Context.getTypeDeclType(ClassDecl);
-  QualType ArgType = Context.getRValueReferenceType(ClassType);
+
+  QualType ArgType;
+  if (Context.getLangOpts().OpenCLCPlusPlus)
+    ArgType = Context.getAddrSpaceQualType(ClassType, LangAS::opencl_generic);
+  ArgType = Context.getRValueReferenceType(ClassType);
 
   bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
                                                      CXXMoveConstructor,
                                                      false);
 
   DeclarationName Name
     = Context.DeclarationNames.getCXXConstructorName(
                                            Context.getCanonicalType(ClassType));
CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
 
   if (getLangOpts().CUDA) {
     inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXMoveConstructor,
                                             MoveConstructor,
                                             /* ConstRHS */ false,
                                             /* Diagnose */ false);
   }
 
-  // Build an exception specification pointing back at this member.
-  FunctionProtoType::ExtProtoInfo EPI =
-      getImplicitMethodEPI(*this, MoveConstructor);
-  MoveConstructor->setType(
-      Context.getFunctionType(Context.VoidTy, ArgType, EPI));
+  setupImplicitSpecialMemberType(MoveConstructor, Context.VoidTy, ArgType);
 
   // Add the parameter to the constructor.
   ParmVarDecl *FromParam = ParmVarDecl::Create(Context, MoveConstructor,
                                                ClassLoc, ClassLoc,
                                                /*IdentifierInfo=*/nullptr,
                                                ArgType, /*TInfo=*/nullptr,
                                                SC_None, nullptr);
   MoveConstructor->setParams(FromParam);

lib/Sema/SemaInit.cpp

static void TryReferenceInitializationCore(Sema &S,
   bool DerivedToBase = false;
   bool ObjCConversion = false;
   bool ObjCLifetimeConversion = false;
   Expr::Classification InitCategory = Initializer->Classify(S.Context);
   Sema::ReferenceCompareResult RefRelationship
     = S.CompareReferenceRelationship(DeclLoc, cv1T1, cv2T2, DerivedToBase,
                                      ObjCConversion, ObjCLifetimeConversion);
 
+  if (InitCategory.isPRValue() || InitCategory.isXValue())
+    T1Quals.removeAddressSpace();
+
   // C++0x [dcl.init.ref]p5:
   //   A reference to type "cv1 T1" is initialized by an expression of type
   //   "cv2 T2" as follows:
   //
   //     - If the reference is an lvalue reference and the initializer
   //       expression
   // Note the analogous bullet points for rvalue refs to functions. Because
   // there are no function rvalues in C++, rvalue refs to functions are treated

test/CodeGenOpenCLCXX/addrspace-of-this.cl

 // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s
 // expected-no-diagnostics
 
 // Test that the 'this' pointer is in the __generic address space.
 
 // FIXME: Add support for __constant address space.
 
 class C {
 public:
   int v;
   C() { v = 2; }
-  // FIXME: Does not work yet.
-  // C(C &&c) { v = c.v; }
+  C(C &&c) { v = c.v; }
   C(const C &c) { v = c.v; }
   C &operator=(const C &c) {
     v = c.v;
     return *this;
   }
-  // FIXME: Does not work yet.
-  //C &operator=(C&& c) & {
-  //  v = c.v;
-  //  return *this;
-  //}
+  C &operator=(C &&c) & {
+    v = c.v;
+    return *this;
+  }
+
+  C operator+(const C& c) {
+    v += c.v;
+    return *this;
+  }
 
   int get() { return v; }
 
   int outside();
 };
 
 int C::outside() {
   return v;
 }
 
 extern C&& foo();
 
 __global C c;
 
 __kernel void test__global() {
   int i = c.get();
   int i2 = c.outside();
   C c1(c);
   C c2;
   c2 = c1;
-  // FIXME: Does not work yet.
-  // C c3 = c1 + c2;
-  // C c4(foo());
-  // C c5 = foo();
-
+  C c3 = c1 + c2;
+  C c4(foo());
+  C c5 = foo();
 }
 
 // CHECK-LABEL: @__cxx_global_var_init()
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*)) #4
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test that the address space is __generic for the constructor
 // CHECK-LABEL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this)
 // CHECK: entry:
 // CHECK:   %this.addr = alloca %class.C addrspace(4)*, align 4
 // CHECK:   store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4
 // CHECK:   %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4
-// CHECK:   call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1) #4
+// CHECK:   call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1)
 // CHECK:   ret void
 
 // CHECK-LABEL: @_Z12test__globalv()
 
 // Test the address space of 'this' when invoking a method.
 // CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking a method that is declared in the file contex.
 // CHECK: %call1 = call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking copy-constructor.
 // CHECK: %0 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %0, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
 
 // Test the address space of 'this' when invoking a constructor.
 // CHECK:   %1 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK:   call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %1) #4
+// CHECK:   call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %1)
 
 // Test the address space of 'this' when invoking assignment operator.
 // CHECK:   %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // CHECK:   %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // CHECK:   %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2)
 
+// Test the address space of 'this' when invoking the operator+
+// CHECK:   %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK:   %5 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:   call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %c3, %class.C addrspace(4)* %4, %class.C addrspace(4)* dereferenceable(4) %5)
+
+// Test the address space of 'this' when invoking the move constructor
+// CHECK:   %6 = addrspacecast %class.C* %c4 to %class.C addrspace(4)*
+// CHECK:   %call3 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
+// CHECK:   call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %call3)
+
+// Test the address space of 'this' when invoking the move assignment
+// CHECK:   %7 = addrspacecast %class.C* %c5 to %class.C addrspace(4)*
+// CHECK:   %call4 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() #5
+// CHECK:   call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* %7, %class.C addrspace(4)* dereferenceable(4) %call4)
+
+
 #define TEST(AS)             \
   __kernel void test##AS() { \
     AS C c;                  \
     int i = c.get();         \
     C c1(c);                 \
     C c2;                    \
     c2 = c1;                 \
   }

 // CHECK:   %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // CHECK:   %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
 
 TEST()
 
 // CHECK-LABEL: @_Z4testv()
 // Test the address space of 'this' when invoking a method.
 // CHECK: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* %1) #4
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* %1)
 
 // Test the address space of 'this' when invoking a copy-constructor.
 // CHECK: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)*
 // CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3)
 
 // Test the address space of 'this' when invoking a constructor.
 // CHECK: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %4)
 
 // Test the address space of 'this' when invoking a copy-assignment.
 // CHECK: %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
 // CHECK: %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
 // CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)

test/CodeGenOpenCLCXX/explicitly-defaulted-methods.cl

+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s
+// expected-no-diagnostics
+
+// Test explicitly defaulted special member functions and their invocation.
+// The 'this' address space is expected to be generic.
+
+// FIXME: Other address spaces?
+
+// CHECK-LABEL:  _Z3foov
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %0)
+
+// CHECK-LABEL: @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this)
+// CHECK:   %this.addr = alloca %class.C addrspace(4)*, align 4
+// CHECK:   store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   ret void
+
+class C {
+public:
+  C() {}
+  C(const C &) = default;
+  C(C &&) = default;
+  C &operator=(const C &) = default;
+  C &operator=(C&&) & = default;
+};
+
+extern C&& bar();
+
+__kernel void foo() {
+  C c;
+  C c2(c);
+  C c3(bar());
+  C c5 = bar();
+  C c6 = c;
+}

test/CodeGenOpenCLCXX/implicitly-defaulted-methods.cl

+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0
+// expected-no-diagnostics
+
+// Test implicitly defaulted special member functions and their invocation.
+// The 'this' address space is expected to be generic.
+
+// CHECK-LABEL:  _Z3foov
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %0)
+
+// CHECK-LABEL: _ZNU3AS41CC1Ev(%class.C addrspace(4)* %this)
+// CHECK:   %this.addr = alloca %class.C addrspace(4)*, align 4
+// CHECK:   store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1) #4
+// CHECK:   ret void
+
+class C {
+public:
+  int a = 0;
+};
+
+extern C&& bar();
+
+__kernel void foo() {
+  C c;
+  C c2(c);
+  C c3(bar());
+  C c5 = bar();
+  C c6 = c;
+  // To force constructors to be CodeGen:ed
+  int x = c6.a + c.a;
+}
+