r218624 - CUDA: Fix incorrect target inference for implicit members.
Eli Bendersky
eliben at google.com
Mon Sep 29 13:38:29 PDT 2014
Author: eliben
Date: Mon Sep 29 15:38:29 2014
New Revision: 218624
URL: http://llvm.org/viewvc/llvm-project?rev=218624&view=rev
Log:
CUDA: Fix incorrect target inference for implicit members.
As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device,
etc) for implicit members (constructors, etc.) incorrectly. This causes errors
and even assertions in Clang when compiling code (assertions in C++11 mode where
implicit move constructors are added into the mix).
Fix the problem by inferring the target from the methods the implicit member
should call (depending on its base classes and fields).
Added:
cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
cfe/trunk/test/SemaCUDA/implicit-member-target-collision.cu
cfe/trunk/test/SemaCUDA/implicit-member-target.cu
Modified:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaDeclCXX.cpp
cfe/trunk/lib/Sema/SemaOverload.cpp
Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=218624&r1=218623&r2=218624&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Mon Sep 29 15:38:29 2014
@@ -541,6 +541,13 @@ def CUDAHost : InheritableAttr {
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>];
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=218624&r1=218623&r2=218624&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Mon Sep 29 15:38:29 2014
@@ -3033,8 +3033,18 @@ def note_ovl_candidate_bad_target : Note
"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">;
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=218624&r1=218623&r2=218624&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Mon Sep 29 15:38:29 2014
@@ -8190,7 +8190,8 @@ public:
CFT_Device,
CFT_Global,
CFT_Host,
- CFT_HostDevice
+ CFT_HostDevice,
+ CFT_InvalidTarget
};
CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
@@ -8198,10 +8199,24 @@ public:
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
//@{
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=218624&r1=218623&r2=218624&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Mon Sep 29 15:38:29 2014
@@ -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 @@ ExprResult Sema::ActOnCUDAExecConfigExpr
/// 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;
@@ -53,8 +53,19 @@ Sema::CUDAFunctionTarget Sema::IdentifyC
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 @@ bool Sema::CheckCUDATarget(CUDAFunctionT
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;
+}
Modified: cfe/trunk/lib/Sema/SemaDeclCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclCXX.cpp?rev=218624&r1=218623&r2=218624&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclCXX.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclCXX.cpp Mon Sep 29 15:38:29 2014
@@ -5575,6 +5575,13 @@ bool Sema::ShouldDeleteSpecialMember(CXX
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 @@ NamespaceDecl *Sema::getOrCreateStdNames
/*PrevDecl=*/nullptr);
getStdNamespace()->setImplicit(true);
}
-
+
return getStdNamespace();
}
@@ -8516,7 +8523,7 @@ CXXConstructorDecl *Sema::DeclareImplici
// 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 @@ CXXConstructorDecl *Sema::DeclareImplici
/*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 @@ CXXDestructorDecl *Sema::DeclareImplicit
/*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 @@ CXXMethodDecl *Sema::DeclareImplicitCopy
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 @@ CXXMethodDecl *Sema::DeclareImplicitMove
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 @@ CXXConstructorDecl *Sema::DeclareImplici
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 @@ CXXConstructorDecl *Sema::DeclareImplici
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);
Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=218624&r1=218623&r2=218624&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Mon Sep 29 15:38:29 2014
@@ -5634,7 +5634,11 @@ Sema::AddOverloadCandidate(FunctionDecl
// (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;
@@ -9136,7 +9140,47 @@ void DiagnoseBadTarget(Sema &S, Overload
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) {
@@ -9877,7 +9921,7 @@ private:
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
Added: cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu?rev=218624&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu (added)
+++ cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu Mon Sep 29 15:38:29 2014
@@ -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 at -3 {{candidate constructor (the implicit default constructor) not viable}}
+// expected-note at -4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note at -5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note at -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 at -5 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note at -6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note at -7 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note at -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 at -4 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note at -5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note at -6 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note at -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 at -3 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note at -4 {{implicit copy constructor inferred target collision}}
+// expected-note at -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 at -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 at -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}}
+}
Added: cfe/trunk/test/SemaCUDA/implicit-member-target-collision.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-member-target-collision.cu?rev=218624&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target-collision.cu (added)
+++ cfe/trunk/test/SemaCUDA/implicit-member-target-collision.cu Mon Sep 29 15:38:29 2014
@@ -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 at -3 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note at -4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note at -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 at -5 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note at -6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note at -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 at -4 {{candidate constructor (the implicit default constructor}} not viable
+// expected-note at -5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note at -6 {{candidate constructor (the implicit copy constructor}} not viable
+
+void hostfoo3() {
+ C3_with_collision c; // expected-error {{no matching constructor}}
+}
Added: cfe/trunk/test/SemaCUDA/implicit-member-target.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-member-target.cu?rev=218624&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target.cu (added)
+++ cfe/trunk/test/SemaCUDA/implicit-member-target.cu Mon Sep 29 15:38:29 2014
@@ -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 at -3 {{call to __host__ function from __device__}}
+// expected-note at -4 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note at -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 at -3 {{call to __device__ function from __host__}}
+// expected-note at -4 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note at -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 at -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 at -4 {{call to __host__ function from __device__}}
+// expected-note at -5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note at -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 at -3 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}}
+// expected-note at -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 at -3 {{candidate constructor not viable: call to __host__ function from __device__ function}}
+// expected-note at -5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note at -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 at -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 at -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}}
+}
More information about the cfe-commits
mailing list