Hi pcc,
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.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,21 @@
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 defaulted (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.
+ /// The result of this call is implicit CUDA target attribute(s) attached to
+ /// the member declaration.
+ void inferCUDATargetForDefaultedSpecialMember(CXXRecordDecl *ClassDecl,
+ CXXSpecialMember CSM,
+ CXXMethodDecl *MemberDecl,
+ bool ConstRHS);
/// \name Code completion
//@{
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -36,11 +36,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 +48,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 +75,137 @@
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 true;
+}
+
+void Sema::inferCUDATargetForDefaultedSpecialMember(CXXRecordDecl *ClassDecl,
+ CXXSpecialMember CSM,
+ CXXMethodDecl *MemberDecl,
+ bool ConstRHS) {
+ CUDAFunctionTarget InferredTarget;
+ bool HasInferredTarget = false;
+
+ // 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.
+ for (const auto &B : ClassDecl->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 (!HasInferredTarget) {
+ HasInferredTarget = true;
+ InferredTarget = BaseMethodTarget;
+ } else {
+ bool ResolutionError = resolveCalleeCUDATargetConflict(
+ InferredTarget, BaseMethodTarget, &InferredTarget);
+ if (ResolutionError) {
+ Diag(ClassDecl->getLocation(),
+ diag::err_implicit_member_target_infer_collision)
+ << (unsigned)CSM << InferredTarget << BaseMethodTarget;
+ return;
+ }
+ }
+ }
+
+ // 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 (!HasInferredTarget) {
+ HasInferredTarget = true;
+ InferredTarget = FieldMethodTarget;
+ } else {
+ bool ResolutionError = resolveCalleeCUDATargetConflict(
+ InferredTarget, FieldMethodTarget, &InferredTarget);
+ if (ResolutionError) {
+ Diag(ClassDecl->getLocation(),
+ diag::err_implicit_member_target_infer_collision)
+ << (unsigned)CSM << InferredTarget << FieldMethodTarget;
+ return;
+ }
+ }
+ }
+
+ if (HasInferredTarget) {
+ if (InferredTarget == CFT_Device) {
+ MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ } else if (InferredTarget == 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));
+ }
+}
Index: lib/Sema/SemaDeclCXX.cpp
===================================================================
--- lib/Sema/SemaDeclCXX.cpp
+++ lib/Sema/SemaDeclCXX.cpp
@@ -8500,7 +8500,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 +8526,15 @@
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.
+ inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXDefaultConstructor,
+ DefaultCon,
+ false);
+ }
+
// Build an exception specification pointing back at this constructor.
FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, DefaultCon);
DefaultCon->setType(Context.getFunctionType(Context.VoidTy, None, EPI));
@@ -8981,6 +8990,12 @@
Destructor->setDefaulted();
Destructor->setImplicit();
+ if (getLangOpts().CUDA) {
+ inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXDestructor,
+ Destructor,
+ false);
+ }
+
// Build an exception specification pointing back at this destructor.
FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, Destructor);
Destructor->setType(Context.getFunctionType(Context.VoidTy, None, EPI));
@@ -9600,6 +9615,12 @@
CopyAssignment->setDefaulted();
CopyAssignment->setImplicit();
+ if (getLangOpts().CUDA) {
+ inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXCopyAssignment,
+ CopyAssignment,
+ Const);
+ }
+
// Build an exception specification pointing back at this member.
FunctionProtoType::ExtProtoInfo EPI =
getImplicitMethodEPI(*this, CopyAssignment);
@@ -9977,6 +9998,12 @@
MoveAssignment->setDefaulted();
MoveAssignment->setImplicit();
+ if (getLangOpts().CUDA) {
+ inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXMoveAssignment,
+ MoveAssignment,
+ false);
+ }
+
// Build an exception specification pointing back at this member.
FunctionProtoType::ExtProtoInfo EPI =
getImplicitMethodEPI(*this, MoveAssignment);
@@ -10397,6 +10424,13 @@
Constexpr);
CopyConstructor->setAccess(AS_public);
CopyConstructor->setDefaulted();
+ CopyConstructor->setImplicit();
+
+ if (getLangOpts().CUDA) {
+ inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXCopyConstructor,
+ CopyConstructor,
+ Const);
+ }
// Build an exception specification pointing back at this member.
FunctionProtoType::ExtProtoInfo EPI =
@@ -10562,6 +10596,13 @@
Constexpr);
MoveConstructor->setAccess(AS_public);
MoveConstructor->setDefaulted();
+ MoveConstructor->setImplicit();
+
+ if (getLangOpts().CUDA) {
+ inferCUDATargetForDefaultedSpecialMember(ClassDecl, CXXMoveConstructor,
+ MoveConstructor,
+ false);
+ }
// Build an exception specification pointing back at this member.
FunctionProtoType::ExtProtoInfo EPI =
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -5634,7 +5634,7 @@
// (CUDA B.1): Check for invalid calls between targets.
if (getLangOpts().CUDA)
if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
- if (CheckCUDATarget(Caller, Function)) {
+ if (!Caller->isImplicit() && CheckCUDATarget(Caller, Function)) {
Candidate.Viable = false;
Candidate.FailureKind = ovl_fail_bad_target;
return;
@@ -9868,7 +9868,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.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/implicit-member-target-collision.cu
@@ -0,0 +1,53 @@
+// 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}}
+// expected-error@-4 {{implicit copy constructor inferred target collision: call to both __host__ __device__ and __host__ __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}}
+// expected-error@-6 {{implicit copy constructor inferred target collision: call to both __host__ __device__ and __host__ __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}}
+// expected-error@-5 {{implicit copy constructor inferred target collision: call to both __host__ __device__ and __host__ __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
