Addresses Richard's review comments
http://reviews.llvm.org/D5199
Files:
include/clang/Basic/DiagnosticSemaKinds.td
include/clang/Sema/Sema.h
lib/Sema/SemaCUDA.cpp
lib/Sema/SemaDeclCXX.cpp
lib/Sema/SemaOverload.cpp
test/SemaCUDA/implicit-member-target-collision-cxx11.cu
test/SemaCUDA/implicit-member-target-collision.cu
test/SemaCUDA/implicit-member-target.cu
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -2997,6 +2997,16 @@
"constructor (inherited)}0 not viable: call to "
"%select{__device__|__global__|__host__|__host__ __device__}1 function from"
" %select{__device__|__global__|__host__|__host__ __device__}2 function">;
+def err_implicit_member_target_infer_collision : Error<
+ "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
@@ -8148,10 +8148,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,11 +38,6 @@
/// 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<CUDAGlobalAttr>())
return CFT_Global;
@@ -53,6 +50,12 @@
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) {
// CUDA B.1.1 "The __device__ qualifier declares a function that is...
@@ -74,3 +77,159 @@
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) {
+ assert((Target1 != Sema::CFT_Global && Target2 != Sema::CFT_Global) &&
+ "Special members cannot be marked global");
+
+ 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 (!ClassDecl->isAbstract() || !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::err_implicit_member_target_infer_collision)
+ << (unsigned)CSM << InferredTarget.getValue()
+ << BaseMethodTarget;
+ }
+ 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::err_implicit_member_target_infer_collision)
+ << (unsigned)CSM << InferredTarget.getValue()
+ << FieldMethodTarget;
+ }
+ 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
@@ -5562,6 +5562,20 @@
if (SMI.shouldDeleteForAllConstMembers())
return true;
+ if (getLangOpts().CUDA) {
+ // We should delete the special member in CUDA mode if target inference
+ // failed.
+ bool Const = false;
+ if ((CSM == CXXCopyConstructor &&
+ RD->implicitCopyConstructorHasConstParam()) ||
+ (CSM == CXXCopyAssignment &&
+ RD->implicitCopyAssignmentHasConstParam())) {
+ Const = true;
+ }
+ return inferCUDATargetForImplicitSpecialMember(RD, CSM, MD, Const,
+ Diagnose);
+ }
+
return false;
}
@@ -8500,7 +8514,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);
@@ -8526,6 +8540,16 @@
DefaultCon->setDefaulted();
DefaultCon->setImplicit();
+ if (getLangOpts().CUDA) {
+ // This has to happen before ShouldDeleteSpecialMember is called. In the
+ // absence of this inference, ShouldDeleteSpecialMember may wrongly decide
+ // to delete this ctor.
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXDefaultConstructor,
+ DefaultCon,
+ /* ConstRHS */ false,
+ /* 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));
@@ -8580,6 +8604,14 @@
}
DiagnoseUninitializedFields(*this, Constructor);
+
+ if (getLangOpts().CUDA) {
+ // Diagnose CUDA potential target inference errors.
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXDefaultConstructor,
+ Constructor,
+ /* ConstRHS */ false,
+ /* Diagnose */ true);
+ }
}
void Sema::ActOnFinishDelayedMemberInitializers(Decl *D) {
@@ -8981,6 +9013,13 @@
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);
Destructor->setType(Context.getFunctionType(Context.VoidTy, None, EPI));
@@ -9041,6 +9080,14 @@
if (ASTMutationListener *L = getASTMutationListener()) {
L->CompletedImplicitDefinition(Destructor);
}
+
+ if (getLangOpts().CUDA) {
+ // Diagnose CUDA potential target inference errors.
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXDestructor,
+ Destructor,
+ /* ConstRHS */ false,
+ /* Diagnose */ true);
+ }
}
/// \brief Perform any semantic analysis which needs to be delayed until all
@@ -9600,6 +9647,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);
@@ -9891,6 +9945,15 @@
if (ASTMutationListener *L = getASTMutationListener()) {
L->CompletedImplicitDefinition(CopyAssignOperator);
}
+
+ if (getLangOpts().CUDA) {
+ // Diagnose CUDA potential target inference errors.
+ bool Const = ClassDecl->implicitCopyAssignmentHasConstParam();
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXCopyAssignment,
+ CopyAssignOperator,
+ /* ConstRHS */ Const,
+ /* Diagnose */ true);
+ }
}
Sema::ImplicitExceptionSpecification
@@ -9977,6 +10040,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);
@@ -10313,6 +10383,14 @@
if (ASTMutationListener *L = getASTMutationListener()) {
L->CompletedImplicitDefinition(MoveAssignOperator);
}
+
+ if (getLangOpts().CUDA) {
+ // Diagnose CUDA potential target inference errors.
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXMoveAssignment,
+ MoveAssignOperator,
+ /* ConstRHS */ false,
+ /* Diagnose */ true);
+ }
}
Sema::ImplicitExceptionSpecification
@@ -10397,6 +10475,14 @@
Constexpr);
CopyConstructor->setAccess(AS_public);
CopyConstructor->setDefaulted();
+ CopyConstructor->setImplicit();
+
+ if (getLangOpts().CUDA) {
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXCopyConstructor,
+ CopyConstructor,
+ /* ConstRHS */ Const,
+ /* Diagnose */ false);
+ }
// Build an exception specification pointing back at this member.
FunctionProtoType::ExtProtoInfo EPI =
@@ -10471,6 +10557,15 @@
if (ASTMutationListener *L = getASTMutationListener()) {
L->CompletedImplicitDefinition(CopyConstructor);
}
+
+ if (getLangOpts().CUDA) {
+ // Diagnose CUDA potential target inference errors.
+ bool Const = ClassDecl->implicitCopyConstructorHasConstParam();
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXCopyConstructor,
+ CopyConstructor,
+ /* ConstRHS */ Const,
+ /* Diagnose */ true);
+ }
}
Sema::ImplicitExceptionSpecification
@@ -10562,6 +10657,14 @@
Constexpr);
MoveConstructor->setAccess(AS_public);
MoveConstructor->setDefaulted();
+ MoveConstructor->setImplicit();
+
+ if (getLangOpts().CUDA) {
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXMoveConstructor,
+ MoveConstructor,
+ /* ConstRHS */ false,
+ /* Diagnose */ false);
+ }
// Build an exception specification pointing back at this member.
FunctionProtoType::ExtProtoInfo EPI =
@@ -10631,6 +10734,14 @@
if (ASTMutationListener *L = getASTMutationListener()) {
L->CompletedImplicitDefinition(MoveConstructor);
}
+
+ if (getLangOpts().CUDA) {
+ // Diagnose CUDA potential target inference errors.
+ inferCUDATargetForImplicitSpecialMember(ClassDecl, CXXMoveConstructor,
+ MoveConstructor,
+ /* ConstRHS */ false,
+ /* Diagnose */ true);
+ }
}
bool Sema::isImplicitlyDeleted(FunctionDecl *FD) {
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<FunctionDecl>(CurContext))
- if (CheckCUDATarget(Caller, Function)) {
+ // Skip the check for callers that are implicit members, because in this
+ // case we still don't 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;
@@ -9868,7 +9872,7 @@
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
Index: test/SemaCUDA/implicit-member-target-collision-cxx11.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/implicit-member-target-collision-cxx11.cu
@@ -0,0 +1,50 @@
+// 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-error@-3 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+
+void hostfoo1() {
+ C1_with_collision c; // expected-error {{implicitly-deleted default constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: collision between two fields
+
+struct C2_with_collision {
+ A1_with_host_ctor aa;
+ B1_with_device_ctor bb;
+};
+
+// expected-error@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+
+void hostfoo2() {
+ C2_with_collision c; // expected-error {{implicitly-deleted default constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 3: collision between a field and a base
+
+struct C3_with_collision : A1_with_host_ctor {
+ B1_with_device_ctor bb;
+};
+
+// expected-error@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+
+void hostfoo4() {
+ C3_with_collision c; // expected-error {{implicitly-deleted default constructor}}
+}
Index: test/SemaCUDA/implicit-member-target-collision.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/implicit-member-target-collision.cu
@@ -0,0 +1,50 @@
+// 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-error@-3 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+
+void hostfoo1() {
+ C1_with_collision c;
+}
+
+//------------------------------------------------------------------------------
+// Test 2: collision between two fields
+
+struct C2_with_collision {
+ A1_with_host_ctor aa;
+ B1_with_device_ctor bb;
+};
+
+// expected-error@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+
+void hostfoo2() {
+ C2_with_collision c;
+}
+
+//------------------------------------------------------------------------------
+// Test 3: collision between a field and a base
+
+struct C3_with_collision : A1_with_host_ctor {
+ B1_with_device_ctor bb;
+};
+
+// expected-error@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+
+void hostfoo4() {
+ C3_with_collision c;
+}
Index: test/SemaCUDA/implicit-member-target.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/implicit-member-target.cu
@@ -0,0 +1,92 @@
+// RUN: %clang_cc1 -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 {{requires 1 argument}}
+
+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 {{requires 1 argument}}
+
+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 {{(the implicit copy constructor) not viable: call to __device__ function from __host__}}
+// expected-note@-4 {{requires 0 arguments}}
+
+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 {{no matching 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 {{requires 1 argument}}
+
+void hostfoo4() {
+ B4_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo4() {
+ B4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
