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()) 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 InferredTarget; + + // We're going to invoke special member lookup; mark that these special + // members are called from this one, and not from its caller. + ContextRAII MethodContext(*this, MemberDecl); + + // Look for special members in base classes that should be invoked from here. + // Infer the target of this member base on the ones it should call. + // Skip direct and indirect virtual bases for abstract classes. + llvm::SmallVector Bases; + for (const auto &B : ClassDecl->bases()) { + if (!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(); + if (!BaseType) { + continue; + } + + CXXRecordDecl *BaseClassDecl = cast(BaseType->getDecl()); + Sema::SpecialMemberOverloadResult *SMOR = + LookupSpecialMember(BaseClassDecl, CSM, + /* ConstArg */ ConstRHS, + /* VolatileArg */ false, + /* RValueThis */ false, + /* ConstThis */ false, + /* VolatileThis */ false); + + if (!SMOR || !SMOR->getMethod()) { + continue; + } + + CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod()); + if (!InferredTarget.hasValue()) { + InferredTarget = BaseMethodTarget; + } else { + bool ResolutionError = resolveCalleeCUDATargetConflict( + InferredTarget.getValue(), BaseMethodTarget, + InferredTarget.getPointer()); + if (ResolutionError) { + if (Diagnose) { + Diag(ClassDecl->getLocation(), + diag::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(); + if (!FieldType) { + continue; + } + + CXXRecordDecl *FieldRecDecl = cast(FieldType->getDecl()); + Sema::SpecialMemberOverloadResult *SMOR = + LookupSpecialMember(FieldRecDecl, CSM, + /* ConstArg */ ConstRHS && !F->isMutable(), + /* VolatileArg */ false, + /* RValueThis */ false, + /* ConstThis */ false, + /* VolatileThis */ false); + + if (!SMOR || !SMOR->getMethod()) { + continue; + } + + CUDAFunctionTarget FieldMethodTarget = + IdentifyCUDATarget(SMOR->getMethod()); + if (!InferredTarget.hasValue()) { + InferredTarget = FieldMethodTarget; + } else { + bool ResolutionError = resolveCalleeCUDATargetConflict( + InferredTarget.getValue(), FieldMethodTarget, + InferredTarget.getPointer()); + if (ResolutionError) { + if (Diagnose) { + Diag(ClassDecl->getLocation(), + diag::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(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(Fn)) { if (S.getLangOpts().CUDA) if (FunctionDecl *Caller = dyn_cast(S.CurContext)) - if (S.CheckCUDATarget(Caller, FunDecl)) + if (!Caller->isImplicit() && S.CheckCUDATarget(Caller, FunDecl)) return false; // If any candidate has a placeholder return type, trigger its deduction Index: test/SemaCUDA/implicit-member-target-collision-cxx11.cu =================================================================== --- /dev/null +++ test/SemaCUDA/implicit-member-target-collision-cxx11.cu @@ -0,0 +1,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}} +}