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()) 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(); + 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 (!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(); + 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 (!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(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(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.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}} +}