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/ASTContext.cpp =================================================================== --- lib/AST/ASTContext.cpp +++ lib/AST/ASTContext.cpp @@ -2773,9 +2773,12 @@ // Anything else must be a function type. Rebuild it with the new exception // specification. const auto *Proto = cast(Orig); - return getFunctionType( - Proto->getReturnType(), Proto->getParamTypes(), - Proto->getExtProtoInfo().withExceptionSpec(ESI)); + + QualType NewT = + getFunctionType(Proto->getReturnType(), Proto->getParamTypes(), + Proto->getExtProtoInfo().withExceptionSpec(ESI)); + + return getAddrSpaceQualType(NewT, Orig.getAddressSpace()); } bool ASTContext::hasSameFunctionTypeIgnoringExceptionSpec(QualType T, Index: lib/Sema/SemaDeclCXX.cpp =================================================================== --- lib/Sema/SemaDeclCXX.cpp +++ lib/Sema/SemaDeclCXX.cpp @@ -6543,8 +6543,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->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; @@ -6552,7 +6556,7 @@ } // A defaulted special member cannot have cv-qualifiers. - if (Type->getTypeQuals()) { + if (Type->getTypeQuals().hasConst() || Type->getTypeQuals().hasVolatile()) { if (DeleteOnTypeMismatch) ShouldDeleteForTypeMismatch = true; else { @@ -10909,6 +10913,27 @@ 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); + + 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: @@ -10949,9 +10974,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. @@ -11222,9 +11245,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. @@ -11798,6 +11819,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, @@ -11824,10 +11849,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, @@ -12311,8 +12333,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() @@ -12494,6 +12514,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, @@ -12522,11 +12546,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, @@ -12623,7 +12643,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, @@ -12652,11 +12676,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: 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/CodeGenOpenCLCXX/explicitly-defaulted-methods.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCLCXX/explicitly-defaulted-methods.cl @@ -0,0 +1,35 @@ +// 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; +} Index: test/CodeGenOpenCLCXX/implicitly-defaulted-methods.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCLCXX/implicitly-defaulted-methods.cl @@ -0,0 +1,33 @@ +// 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; +} +