[OpenCL] Address space for default class members

[Anastasia]

Summary

Fixed several issues reported in https://reviews.llvm.org/D55656 and https://reviews.llvm.org/D54862 - mainly removed addr space from the QualType of the method itself.

This should make all implicitly and explicity defaulted special class members work now!

Since Mikael is on holiday and I have no permission to upload to his review, I have to create a new one.

[rjmccall]

lib/CodeGen/CGCall.cpp

78 ↗

(On Diff #179477)

Would you mind making a patch to rename this (and the method on FunctionProtoType) to something like getMethodQualifiers? It can be a follow-up instead of blocking this.

lib/Sema/SemaDeclCXX.cpp

12649 ↗

(On Diff #179477)

You're implicitly dropping the address space qualifier here by going back to ClassType.

lib/Sema/SemaInit.cpp

4539 ↗

(On Diff #179477)

I can understand why a pr-value wouldn't have an address space, but an x-value's address space is surely important.

[Anastasia]

Addressed review comments.

[Anastasia]

lib/CodeGen/CGCall.cpp

78 ↗

(On Diff #179477)

Sure. I will prepare a separate patch for this.

[rjmccall]

lib/Sema/SemaInit.cpp

4539 ↗

(On Diff #179477)

No, wait, I think this is more deeply wrong. We're initializing a reference to type cv1 T1, and if we're initializing it with a temporary, we're dropping the address space from cv1? That's definitely wrong.

If we're starting with a pr-value, reference-initialization normally has to start by initializing a temporary. (The exception is it's a class type with a conversion operator to a reference type; C++ is abysmally complicated. Let's ignore this for a moment.) Temporaries are always in the default address space (the address space of local variables) because the language (neither OpenCL nor Embedded C) does not give us any facility for allocating temporary memory anywhere else. So reference-initialization from a pr-value creates a temporary in the default address space and then attempts to initialize the destination reference type with that temporary. That initialization should fail if there's no conversion from the default address space to the destination address space. For example, consider initializing a const int __global & from a pr-value: this should clearly not succeed, because we have no way of putting a temporary in the global address space. That reference can only be initialized with a gl-value that's already in the __global address space.

On the other hand, we can successfully initialize a const int & with a gl-value of type const int __global not by direct reference initialization but by loading from the operand and then materializing the result into a new temporary.

I think what this means is:

• We need to be doing this checking as if pr-values were in __private. This includes both (1) expressions that were originally pr-values and (2) expressions that have been converted to pr-values due to a failure to perform direct reference initialization.

• We need to make sure that reference compatibility correctly prevents references from being bound to gl-values in incompatible address spaces.

[ebevhan]

lib/Sema/SemaInit.cpp

4539 ↗

(On Diff #179477)

On the other hand, we can successfully initialize a const int & with a gl-value of type const int __global not by direct reference initialization but by loading from the operand and then materializing the result into a new temporary.

Is this the right thing to do? It means that initializing such a reference would incur a cost under the hood that might be unnoticed/unwanted by the user.

[rjmccall]

lib/Sema/SemaInit.cpp

4539 ↗

(On Diff #179477)

Hmm. I was going to say that it's consistent with the normal behavior of C++, but looking at the actual rule, it's more complicated than that: C++ will only do this if the types are not reference-related or if the gl-value is a bit-field. So this would only be allowed if the reference type was e.g. const long &. It's a strange rule, but it's straightforward to stay consistent with it.

I think we will eventually find ourselves needing a special-case rule for copy-initializing and copy-assigning trivially-copyable types, but we're not there yet.

[Anastasia]

lib/Sema/SemaInit.cpp

4539 ↗

(On Diff #179477)

I need some help to understand what's the right way to fix this problem.

As far as I gathered from this particular example. The issues is caused by the following statement of addrspace-of-this.cl line 47:

C c3 = c1 + c2;

If I remove this new if statement in SemaInit.cpp and disable some asserts about the incorrect address space cast the AST that Clang is trying to create is as follows:

| | |-DeclStmt 0x9d6e70 <line:47:3, col:17>
| | | `-VarDecl 0x9d6b28 <col:3, col:15> col:5 c3 'C' cinit
| | |   `-ExprWithCleanups 0x9d6e58 <col:10, col:15> 'C'
| | |     `-CXXConstructExpr 0x9d6e20 <col:10, col:15> 'C' 'void (__generic C &&) __generic' elidable
| | |       `-MaterializeTemporaryExpr 0x9d6e08 <col:10, col:15> '__generic C' xvalue
| | |         `-ImplicitCastExpr 0x9d6df0 <col:10, col:15> '__generic C' <AddressSpaceConversion>
| | |           `-CXXOperatorCallExpr 0x9d6c90 <col:10, col:15> 'C'
| | |             |-ImplicitCastExpr 0x9d6c78 <col:13> 'C (*)(const __generic C &) __generic' <FunctionToPointerDecay>
| | |             | `-DeclRefExpr 0x9d6bf8 <col:13> 'C (const __generic C &) __generic' lvalue CXXMethod 0x9d4da0 'operator+' 'C (const __generic C &) __generic'
| | |             |-ImplicitCastExpr 0x9d6be0 <col:10> '__generic C' lvalue <AddressSpaceConversion>
| | |             | `-DeclRefExpr 0x9d6b88 <col:10> 'C' lvalue Var 0x9d6830 'c1' 'C'
| | |             `-ImplicitCastExpr 0x9d6bc8 <col:15> 'const __generic C' lvalue <AddressSpaceConversion>
| | |               `-DeclRefExpr 0x9d6ba8 <col:15> 'C' lvalue Var 0x9d6928 'c2' 'C'

I feel that this AST (kind of) makes sense? The result of operator+ returns class C object but as it's directly used to copy construct c3 it seems logical to cast it to __generic C.

Before references with address spaces existed we only had to cast pointers to different address spaces... but now in this particular case of initializing a reference, casting an address space of non-pointer type seem to make sense?

Do you think the AST is correct and I should just change the address space related asserts in casting and teach CodeGen how to generate addrspacecast for such reference initialization? Or do you think the AST has to be created differently?

As a side note: I wish we would have some 'special' AST node to obtain a reference to an object just like we are taking an address of an object to become a pointer, that we could construct implicitly. It would make our life easier for such cases.

[rjmccall]

lib/Sema/SemaInit.cpp

4539 ↗

(On Diff #179477)

This mostly looks right, but I think the MaterializeTemporaryExpr ought to create the temporary in the private address space, and then there should be an ICE to convert that to the generic address space, rather than the other way around.

[Anastasia]

Modified conversion sequence to perform address space conversion after temporary materialization conversion.

[Anastasia]

Removed FIXME that this patch fixes.

[rjmccall]

lib/AST/Expr.cpp

1618 ↗

(On Diff #181114)

What's this about? I would expect that the VK of this and its sub-expression are always the same.

lib/CodeGen/CGExprAgg.cpp

813 ↗

(On Diff #181114)

If there's a Dest, you might need to reverse-convert its address space, right?

lib/Sema/SemaInit.cpp

4686 ↗

(On Diff #181114)

Okay. The existing behavior seems a little strange, but, well, it's existing behavior.

[Anastasia]

• Set correctly value kind of address space conversion after materializing the temporary
• Removed changes in various places including CodeGen for address space conversion that became unnecessary after the AST is generated correctly.
• Updated tests to check for the right IR and also not to use hard-coded registers.

[Anastasia]

lib/CodeGen/CGExprAgg.cpp

813 ↗

(On Diff #181114)

After I fixed the AST (by setting VK correctly on ICE) we are no longer hitting this path. So I removed the unnecessary changes.

[rjmccall]

Great, LGTM.

[Anastasia]

lib/CodeGen/CGCall.cpp

78 ↗

(On Diff #179477)

FYI, covered by: https://reviews.llvm.org/D56792

cfe/trunk/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;
 

cfe/trunk/lib/AST/Expr.cpp

case CK_FunctionToPointerDecay:
     assert(getType()->isPointerType());
     assert(getSubExpr()->getType()->isFunctionType());
     goto CheckNoBasePath;
 
   case CK_AddressSpaceConversion: {
     auto Ty = getType();
     auto SETy = getSubExpr()->getType();
     assert(getValueKindForType(Ty) == Expr::getValueKindForType(SETy));
-    if (!isGLValue())
+    if (isRValue()) {
       Ty = Ty->getPointeeType();
-    if (!isGLValue())
       SETy = SETy->getPointeeType();
+    }
     assert(!Ty.isNull() && !SETy.isNull() &&
            Ty.getAddressSpace() != SETy.getAddressSpace());
     goto CheckNoBasePath;
   }
   // These should not have an inheritance path.
   case CK_Dynamic:
   case CK_ToUnion:
   case CK_ArrayToPointerDecay:

cfe/trunk/lib/CodeGen/CGCall.cpp

 }
 
 /// Derives the 'this' type for codegen purposes, i.e. ignoring method CVR
 /// qualification.
 static CanQualType GetThisType(ASTContext &Context, const CXXRecordDecl *RD,
                                const CXXMethodDecl *MD) {
   QualType RecTy = Context.getTagDeclType(RD)->getCanonicalTypeInternal();
   if (MD)
-    RecTy = Context.getAddrSpaceQualType(RecTy, MD->getType().getAddressSpace());
+    RecTy = Context.getAddrSpaceQualType(RecTy, MD->getTypeQualifiers().getAddressSpace());
   return Context.getPointerType(CanQualType::CreateUnsafe(RecTy));
 }
 
 /// Returns the canonical formal type of the given C++ method.
 static CanQual<FunctionProtoType> GetFormalType(const CXXMethodDecl *MD) {
   return MD->getType()->getCanonicalTypeUnqualified()
            .getAs<FunctionProtoType>();
 }

cfe/trunk/lib/Sema/SemaDecl.cpp

if (OldTypeInfo.getNoCallerSavedRegs() !=
 
     NewTypeInfo = NewTypeInfo.withNoCallerSavedRegs(true);
     RequiresAdjustment = true;
   }
 
   if (RequiresAdjustment) {
     const FunctionType *AdjustedType = New->getType()->getAs<FunctionType>();
     AdjustedType = Context.adjustFunctionType(AdjustedType, NewTypeInfo);
-
-    QualType AdjustedQT = QualType(AdjustedType, 0);
-    LangAS AS = Old->getType().getAddressSpace();
-    AdjustedQT = Context.getAddrSpaceQualType(AdjustedQT, AS);
-
-    New->setType(AdjustedQT);
+    New->setType(QualType(AdjustedType, 0));
     NewQType = Context.getCanonicalType(New->getType());
     NewType = cast<FunctionType>(NewQType);
   }
 
   // If this redeclaration makes the function inline, we may need to add it to
   // UndefinedButUsed.
   if (!Old->isInlined() && New->isInlined() &&
       !New->hasAttr<GNUInlineAttr>() &&

cfe/trunk/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->getTypeQualifiers().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);
+
+  if (getLangOpts().OpenCLCPlusPlus) {
+    // OpenCL: Implicitly defaulted special member are of the generic address
+    // space.
+    EPI.TypeQuals.addAddressSpace(LangAS::opencl_generic);
+  }
+
+  auto 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 = ClassType;
+  if (Context.getLangOpts().OpenCLCPlusPlus)
+    ArgType = Context.getAddrSpaceQualType(ClassType, LangAS::opencl_generic);
+  ArgType = Context.getRValueReferenceType(ArgType);
 
   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);

cfe/trunk/lib/Sema/SemaInit.cpp

if (InitCategory.isPRValue() && T2->isRecordType()) {
       else if (S.getLangOpts().CPlusPlus11)
         CheckCXX98CompatAccessibleCopy(S, Entity, Initializer);
     }
 
     // C++1z [dcl.init.ref]/5.2.1.2:
     //   If the converted initializer is a prvalue, its type T4 is adjusted
     //   to type "cv1 T4" and the temporary materialization conversion is
     //   applied.
+    // Postpone address space conversions to after the temporary materialization
+    // conversion to allow creating temporaries in the alloca address space.
+    auto AS1 = T1Quals.getAddressSpace();
+    auto AS2 = T2Quals.getAddressSpace();
+    T1Quals.removeAddressSpace();
+    T2Quals.removeAddressSpace();
     QualType cv1T4 = S.Context.getQualifiedType(cv2T2, T1Quals);
     if (T1Quals != T2Quals)
       Sequence.AddQualificationConversionStep(cv1T4, ValueKind);
     Sequence.AddReferenceBindingStep(cv1T4, ValueKind == VK_RValue);
     ValueKind = isLValueRef ? VK_LValue : VK_XValue;
+    if (AS1 != AS2) {
+      T1Quals.addAddressSpace(AS1);
+      QualType cv1AST4 = S.Context.getQualifiedType(cv2T2, T1Quals);
+      Sequence.AddQualificationConversionStep(cv1AST4, ValueKind);
+    }
 
     //   In any case, the reference is bound to the resulting glvalue (or to
     //   an appropriate base class subobject).
     if (DerivedToBase)
       Sequence.AddDerivedToBaseCastStep(cv1T1, ValueKind);
     else if (ObjCConversion)
       Sequence.AddObjCObjectConversionStep(cv1T1);
     return;

cfe/trunk/lib/Sema/SemaType.cpp

case DeclaratorChunk::Function: {
              state.getDeclarator().getContext() ==
                  DeclaratorContext::MemberContext)) {
           LangAS CurAS = EPI.TypeQuals.getAddressSpace();
           // If a class member function's address space is not set, set it to
           // __generic.
           LangAS AS =
               (CurAS == LangAS::Default ? LangAS::opencl_generic : CurAS);
           EPI.TypeQuals.addAddressSpace(AS);
-          T = Context.getFunctionType(T, ParamTys, EPI);
-          T = state.getSema().Context.getAddrSpaceQualType(T, AS);
-        } else {
-          T = Context.getFunctionType(T, ParamTys, EPI);
         }
+        T = Context.getFunctionType(T, ParamTys, EPI);
       }
       break;
     }
     case DeclaratorChunk::MemberPointer: {
       // The scope spec must refer to a class, or be dependent.
       CXXScopeSpec &SS = DeclType.Mem.Scope();
       QualType ClsType;
 

cfe/trunk/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)*))
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %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:   [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:   call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // 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)
+// CHECK:  [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK:  [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:  %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
+
+// Test the address space of 'this' when invoking the operator+
+// CHECK: [[C3GEN:%[0-9]+]] = addrspacecast %class.C* %c3 to %class.C addrspace(4)*
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %ref.tmp, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]])
+// CHECK: [[REFGEN:%[0-9]+]] = addrspacecast %class.C* %ref.tmp to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C3GEN]], %class.C addrspace(4)* dereferenceable(4) [[REFGEN]])
+
+// Test the address space of 'this' when invoking the move constructor
+// CHECK: [[C4GEN:%[0-9]+]] = 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)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) %call3)
+
+// Test the address space of 'this' when invoking the move assignment
+// CHECK: [[C5GEN:%[0-9]+]] = 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)* [[C5GEN:%[0-9]+]], %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;                 \
   }
 
 TEST(__local)
 
 // CHECK-LABEL: _Z11test__localv
 // CHECK: @__cxa_guard_acquire
 
-// Test the address space of 'this' when invoking a method.
+// Test the address space of 'this' when invoking a constructor for an object in non-default address space
 // CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
 
-// Test the address space of 'this' when invoking copy-constructor.
+// Test the address space of 'this' when invoking a method.
 // CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
 
+
+// Test the address space of 'this' when invoking copy-constructor.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+
 // Test the address space of 'this' when invoking a constructor.
-// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %3)
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // Test the address space of 'this' when invoking assignment operator.
-// CHECK:  %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK:  %5 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK:  %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %4)
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
 
 TEST(__private)
 
 // CHECK-LABEL: @_Z13test__privatev
 
+// Test the address space of 'this' when invoking a constructor for an object in non-default address space
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
+
 // 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)
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
 
 // 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)
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]])
 
 // 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)
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // 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)
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
 
 TEST()
 
 // CHECK-LABEL: @_Z4testv()
+
+// Test the address space of 'this' when invoking a constructor for an object in non-default address space
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
+
 // 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: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
 
 // 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)
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]])
 
 // 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)
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
 
 // 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)
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])

cfe/trunk/test/SemaOpenCLCXX/address-space-templates.cl

 //RUN: %clang_cc1 %s -cl-std=c++ -pedantic -verify -fsyntax-only
 
 template <typename T>
 struct S {
   T a;        // expected-error{{field may not be qualified with an address space}}
   T f1();     // expected-error{{function type may not be qualified with an address space}}
-  // FIXME: Should only get the error message once.
-  void f2(T); // expected-error{{parameter may not be qualified with an address space}} expected-error{{parameter may not be qualified with an address space}}
-
+  void f2(T); // expected-error{{parameter may not be qualified with an address space}}
 };
 
 template <typename T>
 T foo1(__global T *i) { // expected-note{{candidate template ignored: substitution failure [with T = __local int]: conflicting address space qualifiers are provided between types '__global T' and '__local int'}}
   return *i;
 }
 
 template <typename T>