Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -307,6 +307,10 @@ } bool shouldLinkPossiblyHiddenDecl(LookupResult &Old, const NamedDecl *New); + void setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem, + QualType ResultTy, + ArrayRef Args); + public: typedef OpaquePtr DeclGroupPtrTy; typedef OpaquePtr TemplateTy; Index: lib/CodeGen/CGCall.cpp =================================================================== --- lib/CodeGen/CGCall.cpp +++ lib/CodeGen/CGCall.cpp @@ -74,7 +74,8 @@ 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)); } Index: lib/Sema/SemaDecl.cpp =================================================================== --- lib/Sema/SemaDecl.cpp +++ lib/Sema/SemaDecl.cpp @@ -3192,12 +3192,7 @@ if (RequiresAdjustment) { const FunctionType *AdjustedType = New->getType()->getAs(); 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(NewQType); } Index: lib/Sema/SemaDeclCXX.cpp =================================================================== --- lib/Sema/SemaDeclCXX.cpp +++ lib/Sema/SemaDeclCXX.cpp @@ -6546,8 +6546,12 @@ 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; @@ -6555,7 +6559,7 @@ } // A defaulted special member cannot have cv-qualifiers. - if (Type->getTypeQuals()) { + if (Type->getTypeQuals().hasConst() || Type->getTypeQuals().hasVolatile()) { if (DeleteOnTypeMismatch) ShouldDeleteForTypeMismatch = true; else { @@ -10913,6 +10917,22 @@ CheckFunctionDeclaration(S, FD, R, /*IsMemberSpecialization*/false); } +void Sema::setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem, + QualType ResultTy, + ArrayRef 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: @@ -10953,9 +10973,7 @@ /* 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. @@ -11226,9 +11244,7 @@ /* 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. @@ -11802,6 +11818,10 @@ 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, @@ -11828,10 +11848,7 @@ /* 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, @@ -12315,8 +12332,6 @@ ParmVarDecl *Other = MoveAssignOperator->getParamDecl(0); QualType OtherRefType = Other->getType()-> getAs()->getPointeeType(); - assert(!OtherRefType.getQualifiers() && - "Bad argument type of defaulted move assignment"); // Our location for everything implicitly-generated. SourceLocation Loc = MoveAssignOperator->getEndLoc().isValid() @@ -12498,6 +12513,10 @@ 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, @@ -12526,11 +12545,7 @@ /* 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, @@ -12627,7 +12642,11 @@ 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, @@ -12656,11 +12675,7 @@ /* 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, Index: lib/Sema/SemaInit.cpp =================================================================== --- lib/Sema/SemaInit.cpp +++ lib/Sema/SemaInit.cpp @@ -4535,6 +4535,9 @@ = 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: Index: lib/Sema/SemaType.cpp =================================================================== --- lib/Sema/SemaType.cpp +++ lib/Sema/SemaType.cpp @@ -4839,11 +4839,8 @@ 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; } Index: test/CodeGenOpenCLCXX/addrspace-of-this.cl =================================================================== --- test/CodeGenOpenCLCXX/addrspace-of-this.cl +++ test/CodeGenOpenCLCXX/addrspace-of-this.cl @@ -9,18 +9,21 @@ 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; } @@ -41,15 +44,13 @@ 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) @@ -57,7 +58,7 @@ // 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() @@ -74,13 +75,28 @@ // 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; \ @@ -137,7 +153,7 @@ // 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)* Index: test/SemaOpenCLCXX/address-space-templates.cl =================================================================== --- test/SemaOpenCLCXX/address-space-templates.cl +++ test/SemaOpenCLCXX/address-space-templates.cl @@ -2,11 +2,10 @@ template 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}} + 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