Index: include/clang/Basic/Attr.td =================================================================== --- include/clang/Basic/Attr.td +++ include/clang/Basic/Attr.td @@ -541,6 +541,13 @@ let Documentation = [Undocumented]; } +def CUDAInvalidTarget : InheritableAttr { + let Spellings = []; + let Subjects = SubjectList<[Function]>; + let LangOpts = [CUDA]; + let Documentation = [Undocumented]; +} + def CUDALaunchBounds : InheritableAttr { let Spellings = [GNU<"launch_bounds">]; let Args = [IntArgument<"MaxThreads">, DefaultIntArgument<"MinBlocks", 0>]; Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -3024,8 +3024,18 @@ "function (the implicit copy assignment operator)|" "function (the implicit move assignment operator)|" "constructor (inherited)}0 not viable: call to " - "%select{__device__|__global__|__host__|__host__ __device__}1 function from" - " %select{__device__|__global__|__host__|__host__ __device__}2 function">; + "%select{__device__|__global__|__host__|__host__ __device__|invalid}1 function from" + " %select{__device__|__global__|__host__|__host__ __device__|invalid}2 function">; +def note_implicit_member_target_infer_collision : Note< + "implicit %select{" + "default constructor|" + "copy constructor|" + "move constructor|" + "copy assignment operator|" + "move assignment operator|" + "destructor}0 inferred target collision: call to both " + "%select{__device__|__global__|__host__|__host__ __device__}1 and " + "%select{__device__|__global__|__host__|__host__ __device__}2 members">; def note_ambiguous_type_conversion: Note< "because of ambiguity in conversion %diff{of $ to $|between types}0,1">; Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -8171,7 +8171,8 @@ CFT_Device, CFT_Global, CFT_Host, - CFT_HostDevice + CFT_HostDevice, + CFT_InvalidTarget }; CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D); @@ -8179,10 +8180,24 @@ bool CheckCUDATarget(CUDAFunctionTarget CallerTarget, CUDAFunctionTarget CalleeTarget); - bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee) { - return CheckCUDATarget(IdentifyCUDATarget(Caller), - IdentifyCUDATarget(Callee)); - } + bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee); + + /// Given a implicit special member, infer its CUDA target from the + /// calls it needs to make to underlying base/field special members. + /// \param ClassDecl the class for which the member is being created. + /// \param CSM the kind of special member. + /// \param MemberDecl the special member itself. + /// \param ConstRHS true if this is a copy operation with a const object on + /// its RHS. + /// \param Diagnose true if this call should emit diagnostics. + /// \return true if there was an error inferring. + /// The result of this call is implicit CUDA target attribute(s) attached to + /// the member declaration. + bool inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, + CXXSpecialMember CSM, + CXXMethodDecl *MemberDecl, + bool ConstRHS, + bool Diagnose); /// \name Code completion //@{ Index: lib/Sema/SemaCUDA.cpp =================================================================== --- lib/Sema/SemaCUDA.cpp +++ lib/Sema/SemaCUDA.cpp @@ -15,6 +15,8 @@ #include "clang/AST/ASTContext.h" #include "clang/AST/Decl.h" #include "clang/Sema/SemaDiagnostic.h" +#include "llvm/ADT/Optional.h" +#include "llvm/ADT/SmallVector.h" using namespace clang; ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, @@ -36,10 +38,8 @@ /// IdentifyCUDATarget - Determine the CUDA compilation target for this function Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { - // Implicitly declared functions (e.g. copy constructors) are - // __host__ __device__ - if (D->isImplicit()) - return CFT_HostDevice; + if (D->hasAttr()) + return CFT_InvalidTarget; if (D->hasAttr()) return CFT_Global; @@ -53,8 +53,19 @@ return CFT_Host; } +bool Sema::CheckCUDATarget(const FunctionDecl *Caller, + const FunctionDecl *Callee) { + return CheckCUDATarget(IdentifyCUDATarget(Caller), + IdentifyCUDATarget(Callee)); +} + bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget, CUDAFunctionTarget CalleeTarget) { + // If one of the targets is invalid, the check always fails, no matter what + // the other target is. + if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget) + return true; + // CUDA B.1.1 "The __device__ qualifier declares a function that is... // Callable from the device only." if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device) @@ -74,3 +85,164 @@ return false; } +/// When an implicitly-declared special member has to invoke more than one +/// base/field special member, conflicts may occur in the targets of these +/// members. For example, if one base's member __host__ and another's is +/// __device__, it's a conflict. +/// This function figures out if the given targets \param Target1 and +/// \param Target2 conflict, and if they do not it fills in +/// \param ResolvedTarget with a target that resolves for both calls. +/// \return true if there's a conflict, false otherwise. +static bool +resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1, + Sema::CUDAFunctionTarget Target2, + Sema::CUDAFunctionTarget *ResolvedTarget) { + if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) { + // TODO: this shouldn't happen, really. Methods cannot be marked __global__. + // Clang should detect this earlier and produce an error. Then this + // condition can be changed to an assertion. + return true; + } + + if (Target1 == Sema::CFT_HostDevice) { + *ResolvedTarget = Target2; + } else if (Target2 == Sema::CFT_HostDevice) { + *ResolvedTarget = Target1; + } else if (Target1 != Target2) { + return true; + } else { + *ResolvedTarget = Target1; + } + + return false; +} + +bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, + CXXSpecialMember CSM, + CXXMethodDecl *MemberDecl, + bool ConstRHS, + bool Diagnose) { + llvm::Optional InferredTarget; + + // We're going to invoke special member lookup; mark that these special + // members are called from this one, and not from its caller. + ContextRAII MethodContext(*this, MemberDecl); + + // Look for special members in base classes that should be invoked from here. + // Infer the target of this member base on the ones it should call. + // Skip direct and indirect virtual bases for abstract classes. + llvm::SmallVector Bases; + for (const auto &B : ClassDecl->bases()) { + if (!B.isVirtual()) { + Bases.push_back(&B); + } + } + + if (!ClassDecl->isAbstract()) { + for (const auto &VB : ClassDecl->vbases()) { + Bases.push_back(&VB); + } + } + + for (const auto *B : Bases) { + const RecordType *BaseType = B->getType()->getAs(); + if (!BaseType) { + continue; + } + + CXXRecordDecl *BaseClassDecl = cast(BaseType->getDecl()); + Sema::SpecialMemberOverloadResult *SMOR = + LookupSpecialMember(BaseClassDecl, CSM, + /* ConstArg */ ConstRHS, + /* VolatileArg */ false, + /* RValueThis */ false, + /* ConstThis */ false, + /* VolatileThis */ false); + + if (!SMOR || !SMOR->getMethod()) { + continue; + } + + CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod()); + if (!InferredTarget.hasValue()) { + InferredTarget = BaseMethodTarget; + } else { + bool ResolutionError = resolveCalleeCUDATargetConflict( + InferredTarget.getValue(), BaseMethodTarget, + InferredTarget.getPointer()); + if (ResolutionError) { + if (Diagnose) { + Diag(ClassDecl->getLocation(), + diag::note_implicit_member_target_infer_collision) + << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget; + } + MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context)); + return true; + } + } + } + + // Same as for bases, but now for special members of fields. + for (const auto *F : ClassDecl->fields()) { + if (F->isInvalidDecl()) { + continue; + } + + const RecordType *FieldType = + Context.getBaseElementType(F->getType())->getAs(); + if (!FieldType) { + continue; + } + + CXXRecordDecl *FieldRecDecl = cast(FieldType->getDecl()); + Sema::SpecialMemberOverloadResult *SMOR = + LookupSpecialMember(FieldRecDecl, CSM, + /* ConstArg */ ConstRHS && !F->isMutable(), + /* VolatileArg */ false, + /* RValueThis */ false, + /* ConstThis */ false, + /* VolatileThis */ false); + + if (!SMOR || !SMOR->getMethod()) { + continue; + } + + CUDAFunctionTarget FieldMethodTarget = + IdentifyCUDATarget(SMOR->getMethod()); + if (!InferredTarget.hasValue()) { + InferredTarget = FieldMethodTarget; + } else { + bool ResolutionError = resolveCalleeCUDATargetConflict( + InferredTarget.getValue(), FieldMethodTarget, + InferredTarget.getPointer()); + if (ResolutionError) { + if (Diagnose) { + Diag(ClassDecl->getLocation(), + diag::note_implicit_member_target_infer_collision) + << (unsigned)CSM << InferredTarget.getValue() + << FieldMethodTarget; + } + MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context)); + return true; + } + } + } + + if (InferredTarget.hasValue()) { + if (InferredTarget.getValue() == CFT_Device) { + MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); + } else if (InferredTarget.getValue() == CFT_Host) { + MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); + } else { + MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); + MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); + } + } else { + // If no target was inferred, mark this member as __host__ __device__; + // it's the least restrictive option that can be invoked from any target. + MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); + MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); + } + + return false; +} Index: lib/Sema/SemaDeclCXX.cpp =================================================================== --- lib/Sema/SemaDeclCXX.cpp +++ lib/Sema/SemaDeclCXX.cpp @@ -5575,6 +5575,13 @@ if (SMI.shouldDeleteForAllConstMembers()) return true; + if (getLangOpts().CUDA) { + // We should delete the special member in CUDA mode if target inference + // failed. + return inferCUDATargetForImplicitSpecialMember(RD, CSM, MD, SMI.ConstArg, + Diagnose); + } + return false; } @@ -6984,7 +6991,7 @@ /*PrevDecl=*/nullptr); getStdNamespace()->setImplicit(true); } - + return getStdNamespace(); } @@ -8516,7 +8523,7 @@ // 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. - assert(ClassDecl->needsImplicitDefaultConstructor() && + assert(ClassDecl->needsImplicitDefaultConstructor() && "Should not build implicit default constructor!"); DeclaringSpecialMember DSM(*this, ClassDecl, CXXDefaultConstructor); @@ -8540,7 +8547,13 @@ /*isImplicitlyDeclared=*/true, Constexpr); DefaultCon->setAccess(AS_public); DefaultCon->setDefaulted(); - DefaultCon->setImplicit(); + + 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); @@ -9000,7 +9013,13 @@ /*isImplicitlyDeclared=*/true); Destructor->setAccess(AS_public); Destructor->setDefaulted(); - Destructor->setImplicit(); + + 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); @@ -9626,6 +9645,13 @@ CopyAssignment->setDefaulted(); CopyAssignment->setImplicit(); + 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); @@ -10008,6 +10034,13 @@ MoveAssignment->setDefaulted(); MoveAssignment->setImplicit(); + if (getLangOpts().CUDA) { + inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXMoveAssignment, + MoveAssignment, + /* ConstRHS */ false, + /* Diagnose */ false); + } + // Build an exception specification pointing back at this member. FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, MoveAssignment); @@ -10434,6 +10467,13 @@ CopyConstructor->setAccess(AS_public); CopyConstructor->setDefaulted(); + 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); @@ -10604,6 +10644,13 @@ MoveConstructor->setAccess(AS_public); MoveConstructor->setDefaulted(); + 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); Index: lib/Sema/SemaOverload.cpp =================================================================== --- lib/Sema/SemaOverload.cpp +++ lib/Sema/SemaOverload.cpp @@ -5634,7 +5634,11 @@ // (CUDA B.1): Check for invalid calls between targets. if (getLangOpts().CUDA) if (const FunctionDecl *Caller = dyn_cast(CurContext)) - if (CheckCUDATarget(Caller, Function)) { + // Skip the check for callers that are implicit members, because in this + // case we may not yet know what the member's target is; the target is + // inferred for the member automatically, based on the bases and fields of + // the class. + if (!Caller->isImplicit() && CheckCUDATarget(Caller, Function)) { Candidate.Viable = false; Candidate.FailureKind = ovl_fail_bad_target; return; @@ -9136,7 +9140,47 @@ OverloadCandidateKind FnKind = ClassifyOverloadCandidate(S, Callee, FnDesc); S.Diag(Callee->getLocation(), diag::note_ovl_candidate_bad_target) - << (unsigned) FnKind << CalleeTarget << CallerTarget; + << (unsigned)FnKind << CalleeTarget << CallerTarget; + + // This could be an implicit constructor for which we could not infer the + // target due to a collsion. Diagnose that case. + CXXMethodDecl *Meth = dyn_cast(Callee); + if (Meth != nullptr && Meth->isImplicit()) { + CXXRecordDecl *ParentClass = Meth->getParent(); + Sema::CXXSpecialMember CSM; + + switch (FnKind) { + default: + return; + case oc_implicit_default_constructor: + CSM = Sema::CXXDefaultConstructor; + break; + case oc_implicit_copy_constructor: + CSM = Sema::CXXCopyConstructor; + break; + case oc_implicit_move_constructor: + CSM = Sema::CXXMoveConstructor; + break; + case oc_implicit_copy_assignment: + CSM = Sema::CXXCopyAssignment; + break; + case oc_implicit_move_assignment: + CSM = Sema::CXXMoveAssignment; + break; + }; + + bool ConstRHS = false; + if (Meth->getNumParams()) { + if (const ReferenceType *RT = + Meth->getParamDecl(0)->getType()->getAs()) { + ConstRHS = RT->getPointeeType().isConstQualified(); + } + } + + S.inferCUDATargetForImplicitSpecialMember(ParentClass, CSM, Meth, + /* ConstRHS */ ConstRHS, + /* Diagnose */ true); + } } void DiagnoseFailedEnableIfAttr(Sema &S, OverloadCandidate *Cand) { @@ -9877,7 +9921,7 @@ if (FunctionDecl *FunDecl = dyn_cast(Fn)) { if (S.getLangOpts().CUDA) if (FunctionDecl *Caller = dyn_cast(S.CurContext)) - if (S.CheckCUDATarget(Caller, FunDecl)) + if (!Caller->isImplicit() && S.CheckCUDATarget(Caller, FunDecl)) return false; // If any candidate has a placeholder return type, trigger its deduction Index: test/SemaCUDA/implicit-member-target-collision-cxx11.cu =================================================================== --- /dev/null +++ test/SemaCUDA/implicit-member-target-collision-cxx11.cu @@ -0,0 +1,111 @@ +// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s + +#include "Inputs/cuda.h" + +//------------------------------------------------------------------------------ +// Test 1: collision between two bases + +struct A1_with_host_ctor { + A1_with_host_ctor() {} +}; + +struct B1_with_device_ctor { + __device__ B1_with_device_ctor() {} +}; + +struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor { +}; + +// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable +// expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} +// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable +// expected-note@-6 {{candidate constructor (the implicit move constructor}} not viable + +void hostfoo1() { + C1_with_collision c; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 2: collision between two fields + +struct C2_with_collision { + A1_with_host_ctor aa; + B1_with_device_ctor bb; +}; + +// expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable +// expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} +// expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable +// expected-note@-8 {{candidate constructor (the implicit move constructor}} not viable + +void hostfoo2() { + C2_with_collision c; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 3: collision between a field and a base + +struct C3_with_collision : A1_with_host_ctor { + B1_with_device_ctor bb; +}; + +// expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable +// expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} +// expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable +// expected-note@-7 {{candidate constructor (the implicit move constructor}} not viable + +void hostfoo3() { + C3_with_collision c; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 4: collision on resolving a copy ctor + +struct A4_with_host_copy_ctor { + A4_with_host_copy_ctor() {} + A4_with_host_copy_ctor(const A4_with_host_copy_ctor&) {} +}; + +struct B4_with_device_copy_ctor { + B4_with_device_copy_ctor() {} + __device__ B4_with_device_copy_ctor(const B4_with_device_copy_ctor&) {} +}; + +struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor { +}; + +// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable +// expected-note@-4 {{implicit copy constructor inferred target collision}} +// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable + +void hostfoo4() { + C4_with_collision c; + C4_with_collision c2 = c; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 5: collision on resolving a move ctor + +struct A5_with_host_move_ctor { + A5_with_host_move_ctor() {} + A5_with_host_move_ctor(A5_with_host_move_ctor&&) {} +// expected-note@-1 {{copy constructor is implicitly deleted because 'A5_with_host_move_ctor' has a user-declared move constructor}} +}; + +struct B5_with_device_move_ctor { + B5_with_device_move_ctor() {} + __device__ B5_with_device_move_ctor(B5_with_device_move_ctor&&) {} +}; + +struct C5_with_collision : A5_with_host_move_ctor, B5_with_device_move_ctor { +}; +// expected-note@-2 {{deleted}} + +void hostfoo5() { + C5_with_collision c; + // What happens here: + // This tries to find the move ctor. Since the move ctor is deleted due to + // collision, it then looks for a copy ctor. But copy ctors are implicitly + // deleted when move ctors are declared explicitly. + C5_with_collision c2(static_cast(c)); // expected-error {{call to implicitly-deleted}} +} Index: test/SemaCUDA/implicit-member-target-collision.cu =================================================================== --- /dev/null +++ test/SemaCUDA/implicit-member-target-collision.cu @@ -0,0 +1,57 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s + +#include "Inputs/cuda.h" + +//------------------------------------------------------------------------------ +// Test 1: collision between two bases + +struct A1_with_host_ctor { + A1_with_host_ctor() {} +}; + +struct B1_with_device_ctor { + __device__ B1_with_device_ctor() {} +}; + +struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor { +}; + +// expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable +// expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} +// expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable + +void hostfoo1() { + C1_with_collision c; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 2: collision between two fields + +struct C2_with_collision { + A1_with_host_ctor aa; + B1_with_device_ctor bb; +}; + +// expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable +// expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} +// expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable + +void hostfoo2() { + C2_with_collision c; // expected-error {{no matching constructor}} + +} + +//------------------------------------------------------------------------------ +// Test 3: collision between a field and a base + +struct C3_with_collision : A1_with_host_ctor { + B1_with_device_ctor bb; +}; + +// expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable +// expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} +// expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable + +void hostfoo3() { + C3_with_collision c; // expected-error {{no matching constructor}} +} Index: test/SemaCUDA/implicit-member-target.cu =================================================================== --- /dev/null +++ test/SemaCUDA/implicit-member-target.cu @@ -0,0 +1,184 @@ +// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s + +#include "Inputs/cuda.h" + +//------------------------------------------------------------------------------ +// Test 1: infer default ctor to be host. + +struct A1_with_host_ctor { + A1_with_host_ctor() {} +}; + +// The implicit default constructor is inferred to be host because it only needs +// to invoke a single host constructor (A1_with_host_ctor's). So we'll encounter +// an error when calling it from a __device__ function, but not from a __host__ +// function. +struct B1_with_implicit_default_ctor : A1_with_host_ctor { +}; + +// expected-note@-3 {{call to __host__ function from __device__}} +// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}} +// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}} + +void hostfoo() { + B1_with_implicit_default_ctor b; +} + +__device__ void devicefoo() { + B1_with_implicit_default_ctor b; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 2: infer default ctor to be device. + +struct A2_with_device_ctor { + __device__ A2_with_device_ctor() {} +}; + +struct B2_with_implicit_default_ctor : A2_with_device_ctor { +}; + +// expected-note@-3 {{call to __device__ function from __host__}} +// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}} +// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}} + +void hostfoo2() { + B2_with_implicit_default_ctor b; // expected-error {{no matching constructor}} +} + +__device__ void devicefoo2() { + B2_with_implicit_default_ctor b; +} + +//------------------------------------------------------------------------------ +// Test 3: infer copy ctor + +struct A3_with_device_ctors { + __host__ A3_with_device_ctors() {} + __device__ A3_with_device_ctors(const A3_with_device_ctors&) {} +}; + +struct B3_with_implicit_ctors : A3_with_device_ctors { +}; + +// expected-note@-3 {{copy constructor of 'B3_with_implicit_ctors' is implicitly deleted}} + +void hostfoo3() { + B3_with_implicit_ctors b; // this is OK because the inferred default ctor + // here is __host__ + B3_with_implicit_ctors b2 = b; // expected-error {{call to implicitly-deleted copy constructor}} + +} + +//------------------------------------------------------------------------------ +// Test 4: infer default ctor from a field, not a base + +struct A4_with_host_ctor { + A4_with_host_ctor() {} +}; + +struct B4_with_implicit_default_ctor { + A4_with_host_ctor field; +}; + +// expected-note@-4 {{call to __host__ function from __device__}} +// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}} +// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}} + +void hostfoo4() { + B4_with_implicit_default_ctor b; +} + +__device__ void devicefoo4() { + B4_with_implicit_default_ctor b; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 5: copy ctor with non-const param + +struct A5_copy_ctor_constness { + __host__ A5_copy_ctor_constness() {} + __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {} +}; + +struct B5_copy_ctor_constness : A5_copy_ctor_constness { +}; + +// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}} +// expected-note@-4 {{candidate constructor (the implicit default constructor) not viable}} + +void hostfoo5(B5_copy_ctor_constness& b_arg) { + B5_copy_ctor_constness b = b_arg; +} + +__device__ void devicefoo5(B5_copy_ctor_constness& b_arg) { + B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 6: explicitly defaulted ctor: since they are spelled out, they have +// a host/device designation explicitly so no inference needs to be done. + +struct A6_with_device_ctor { + __device__ A6_with_device_ctor() {} +}; + +struct B6_with_defaulted_ctor : A6_with_device_ctor { + __host__ B6_with_defaulted_ctor() = default; +}; + +// expected-note@-3 {{candidate constructor not viable: call to __host__ function from __device__ function}} +// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}} +// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}} + +__device__ void devicefoo6() { + B6_with_defaulted_ctor b; // expected-error {{no matching constructor}} +} + +//------------------------------------------------------------------------------ +// Test 7: copy assignment operator + +struct A7_with_copy_assign { + A7_with_copy_assign() {} + __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {} +}; + +struct B7_with_copy_assign : A7_with_copy_assign { +}; + +// expected-note@-3 {{copy assignment operator of 'B7_with_copy_assign' is implicitly deleted}} + +void hostfoo7() { + B7_with_copy_assign b1, b2; + b1 = b2; // expected-error {{object of type 'B7_with_copy_assign' cannot be assigned because its copy assignment operator is implicitly deleted}} +} + +//------------------------------------------------------------------------------ +// Test 8: move assignment operator + +// definitions for std::move +namespace std { +inline namespace foo { +template struct remove_reference { typedef T type; }; +template struct remove_reference { typedef T type; }; +template struct remove_reference { typedef T type; }; + +template typename remove_reference::type&& move(T&& t); +} +} + +struct A8_with_move_assign { + A8_with_move_assign() {} + __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {} + __device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {} +}; + +struct B8_with_move_assign : A8_with_move_assign { +}; + +// expected-note@-3 {{copy assignment operator of 'B8_with_move_assign' is implicitly deleted because base class 'A8_with_move_assign' has no copy assignment operator}} + +void hostfoo8() { + B8_with_move_assign b1, b2; + b1 = std::move(b2); // expected-error {{object of type 'B8_with_move_assign' cannot be assigned because its copy assignment operator is implicitly deleted}} +}