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/AST/Expr.cpp =================================================================== --- lib/AST/Expr.cpp +++ lib/AST/Expr.cpp @@ -1615,10 +1615,10 @@ auto Ty = getType(); auto SETy = getSubExpr()->getType(); assert(getValueKindForType(Ty) == Expr::getValueKindForType(SETy)); - if (!isGLValue()) + if (!isGLValue() && !getSubExpr()->isXValue()) { Ty = Ty->getPointeeType(); - if (!isGLValue()) SETy = SETy->getPointeeType(); + } assert(!Ty.isNull() && !SETy.isNull() && Ty.getAddressSpace() != SETy.getAddressSpace()); goto CheckNoBasePath; Index: lib/CodeGen/CGCall.cpp =================================================================== --- lib/CodeGen/CGCall.cpp +++ lib/CodeGen/CGCall.cpp @@ -74,7 +74,7 @@ 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/CodeGen/CGExprAgg.cpp =================================================================== --- lib/CodeGen/CGExprAgg.cpp +++ lib/CodeGen/CGExprAgg.cpp @@ -809,6 +809,11 @@ case CK_LValueBitCast: llvm_unreachable("should not be emitting lvalue bitcast as rvalue"); + case CK_AddressSpaceConversion: { + Visit(E->getSubExpr()); + break; + } + case CK_Dependent: case CK_BitCast: case CK_ArrayToPointerDecay: @@ -852,7 +857,6 @@ case CK_CopyAndAutoreleaseBlockObject: case CK_BuiltinFnToFnPtr: case CK_ZeroToOCLOpaqueType: - case CK_AddressSpaceConversion: case CK_IntToOCLSampler: case CK_FixedPointCast: case CK_FixedPointToBoolean: Index: lib/Sema/Sema.cpp =================================================================== --- lib/Sema/Sema.cpp +++ lib/Sema/Sema.cpp @@ -481,6 +481,7 @@ default: llvm_unreachable("can't implicitly cast lvalue to rvalue with this cast " "kind"); + case CK_AddressSpaceConversion: case CK_LValueToRValue: case CK_ArrayToPointerDecay: case CK_FunctionToPointerDecay: 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,11 @@ 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 +6558,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 +10916,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 +10972,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 +11243,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 +11817,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 +11847,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 +12331,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 +12512,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 +12544,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 +12641,11 @@ 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, @@ -12656,11 +12674,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 @@ -4669,10 +4669,21 @@ // 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); + if (AS1 != AS2) { + T1Quals.addAddressSpace(AS1); + QualType cv1AST4 = S.Context.getQualifiedType(cv2T2, T1Quals); + Sequence.AddQualificationConversionStep(cv1AST4, ValueKind); + } ValueKind = isLValueRef ? VK_LValue : VK_XValue; // In any case, the reference is bound to the resulting glvalue (or to 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,29 @@ // 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 +154,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 @@ -4,9 +4,7 @@ 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