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>];
"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">;
CFT_Device,
CFT_Global,
CFT_Host,
- CFT_HostDevice
+ CFT_HostDevice,
+ CFT_InvalidTarget
};
CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
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
//@{
#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,
/// 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<CUDAInvalidTargetAttr>())
+ return CFT_InvalidTarget;
if (D->hasAttr<CUDAGlobalAttr>())
return CFT_Global;
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)
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<CUDAFunctionTarget> 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<const CXXBaseSpecifier *, 16> 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<RecordType>();
+ if (!BaseType) {
+ continue;
+ }
+
+ CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(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<RecordType>();
+ if (!FieldType) {
+ continue;
+ }
+
+ CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(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;
+}
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;
}
/*PrevDecl=*/nullptr);
getStdNamespace()->setImplicit(true);
}
-
+
return getStdNamespace();
}
// 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);
/*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);
/*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);
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);
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);
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);
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);
// (CUDA B.1): Check for invalid calls between targets.
if (getLangOpts().CUDA)
if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(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;
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<CXXMethodDecl>(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<ReferenceType>()) {
+ ConstRHS = RT->getPointeeType().isConstQualified();
+ }
+ }
+
+ S.inferCUDATargetForImplicitSpecialMember(ParentClass, CSM, Meth,
+ /* ConstRHS */ ConstRHS,
+ /* Diagnose */ true);
+ }
}
void DiagnoseFailedEnableIfAttr(Sema &S, OverloadCandidate *Cand) {
if (FunctionDecl *FunDecl = dyn_cast<FunctionDecl>(Fn)) {
if (S.getLangOpts().CUDA)
if (FunctionDecl *Caller = dyn_cast<FunctionDecl>(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
--- /dev/null
+// 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<C5_with_collision&&>(c)); // expected-error {{call to implicitly-deleted}}
+}
--- /dev/null
+// 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}}
+}
--- /dev/null
+// 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 <class T> struct remove_reference { typedef T type; };
+template <class T> struct remove_reference<T&> { typedef T type; };
+template <class T> struct remove_reference<T&&> { typedef T type; };
+
+template <class T> typename remove_reference<T>::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}}
+}
projects / clang / commitdiff