diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst index 481ed3923081..8f473c21e191 100644 --- a/clang/docs/HIPSupport.rst +++ b/clang/docs/HIPSupport.rst @@ -286,6 +286,26 @@ Example Usage basePtr->virtualFunction(); // Allowed since obj is constructed in device code } +Host and Device Attributes of Default Destructors +=================================================== + +If a default destructor does not have explicit host or device attributes, +clang infers these attributes based on the destructors of its data members +and base classes. If any conflicts are detected among these destructors, +clang diagnoses the issue. Otherwise, clang adds an implicit host or device +attribute according to whether the data members's and base classes's +destructors can execute on the host or device side. + +For explicit template classes with virtual destructors, which must be emitted, +the inference adopts a conservative approach. In this case, implicit host or +device attributes from member and base class destructors are ignored. This +precaution is necessary because, although a constexpr destructor carries +implicit host or device attributes, a constexpr function may call a +non-constexpr function, which is by default a host function. + +Users can override the inferred host and device attributes of default +destructors by adding explicit host and device attributes to them. + C++ Standard Parallelism Offload Support: Compiler And Runtime ============================================================== diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index a30a7076ea5d..af648d7f9c63 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -4336,11 +4336,11 @@ public: // Whether the callee should be ignored in CUDA/HIP/OpenMP host/device check. bool shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee); -private: /// Function or variable declarations to be checked for whether the deferred /// diagnostics should be emitted. llvm::SmallSetVector DeclsToCheckForDeferredDiags; +private: /// Map of current shadowing declarations to shadowed declarations. Warn if /// it looks like the user is trying to modify the shadowing declaration. llvm::DenseMap ShadowingDecls; diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 9507d7602aa4..e0eac690e6e6 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -1789,6 +1789,47 @@ public: Inherited::visitUsedDecl(Loc, D); } + // Visitor member and parent dtors called by this dtor. + void VisitCalledDestructors(CXXDestructorDecl *DD) { + const CXXRecordDecl *RD = DD->getParent(); + + // Visit the dtors of all members + for (const FieldDecl *FD : RD->fields()) { + QualType FT = FD->getType(); + if (const auto *RT = FT->getAs()) + if (const auto *ClassDecl = dyn_cast(RT->getDecl())) + if (ClassDecl->hasDefinition()) + if (CXXDestructorDecl *MemberDtor = ClassDecl->getDestructor()) + asImpl().visitUsedDecl(MemberDtor->getLocation(), MemberDtor); + } + + // Also visit base class dtors + for (const auto &Base : RD->bases()) { + QualType BaseType = Base.getType(); + if (const auto *RT = BaseType->getAs()) + if (const auto *BaseDecl = dyn_cast(RT->getDecl())) + if (BaseDecl->hasDefinition()) + if (CXXDestructorDecl *BaseDtor = BaseDecl->getDestructor()) + asImpl().visitUsedDecl(BaseDtor->getLocation(), BaseDtor); + } + } + + void VisitDeclStmt(DeclStmt *DS) { + // Visit dtors called by variables that need destruction + for (auto *D : DS->decls()) + if (auto *VD = dyn_cast(D)) + if (VD->isThisDeclarationADefinition() && + VD->needsDestruction(S.Context)) { + QualType VT = VD->getType(); + if (const auto *RT = VT->getAs()) + if (const auto *ClassDecl = dyn_cast(RT->getDecl())) + if (ClassDecl->hasDefinition()) + if (CXXDestructorDecl *Dtor = ClassDecl->getDestructor()) + asImpl().visitUsedDecl(Dtor->getLocation(), Dtor); + } + + Inherited::VisitDeclStmt(DS); + } void checkVar(VarDecl *VD) { assert(VD->isFileVarDecl() && "Should only check file-scope variables"); @@ -1830,6 +1871,8 @@ public: if (auto *S = FD->getBody()) { this->Visit(S); } + if (CXXDestructorDecl *Dtor = dyn_cast(FD)) + asImpl().VisitCalledDestructors(Dtor); UsePath.pop_back(); InUsePath.erase(FD); } diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 0e1bf727d72d..0e5fc5e1a40b 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -372,6 +372,21 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, CXXMethodDecl *MemberDecl, bool ConstRHS, bool Diagnose) { + // If MemberDecl is virtual destructor of an explicit template class + // instantiation, it must be emitted, therefore it needs to be inferred + // conservatively by ignoring implicit host/device attrs of member and parent + // dtors called by it. Also, it needs to be checed by deferred diag visitor. + bool IsExpVDtor = false; + if (isa(MemberDecl) && MemberDecl->isVirtual()) { + if (auto *Spec = dyn_cast(ClassDecl)) { + TemplateSpecializationKind TSK = Spec->getTemplateSpecializationKind(); + IsExpVDtor = TSK == TSK_ExplicitInstantiationDeclaration || + TSK == TSK_ExplicitInstantiationDefinition; + } + } + if (IsExpVDtor) + SemaRef.DeclsToCheckForDeferredDiags.insert(MemberDecl); + // If the defaulted special member is defined lexically outside of its // owning class, or the special member already has explicit device or host // attributes, do not infer. @@ -422,7 +437,9 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, if (!SMOR.getMethod()) continue; - CUDAFunctionTarget BaseMethodTarget = IdentifyTarget(SMOR.getMethod()); + CUDAFunctionTarget BaseMethodTarget = + IdentifyTarget(SMOR.getMethod(), IsExpVDtor); + if (!InferredTarget) { InferredTarget = BaseMethodTarget; } else { @@ -466,7 +483,9 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, if (!SMOR.getMethod()) continue; - CUDAFunctionTarget FieldMethodTarget = IdentifyTarget(SMOR.getMethod()); + CUDAFunctionTarget FieldMethodTarget = + IdentifyTarget(SMOR.getMethod(), IsExpVDtor); + if (!InferredTarget) { InferredTarget = FieldMethodTarget; } else { diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 01f09aba8c2a..f70401ea33b4 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -20388,6 +20388,21 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD, if (IsEmittedForExternalSymbol()) return FunctionEmissionStatus::Emitted; + + // If FD is a virtual destructor of an explicit instantiation + // of a template class, return Emitted. + if (auto *Destructor = dyn_cast(FD)) { + if (Destructor->isVirtual()) { + if (auto *Spec = dyn_cast( + Destructor->getParent())) { + TemplateSpecializationKind TSK = + Spec->getTemplateSpecializationKind(); + if (TSK == TSK_ExplicitInstantiationDeclaration || + TSK == TSK_ExplicitInstantiationDefinition) + return FunctionEmissionStatus::Emitted; + } + } + } } // Otherwise, the function is known-emitted if it's in our set of diff --git a/clang/test/SemaCUDA/dtor.cu b/clang/test/SemaCUDA/dtor.cu new file mode 100644 index 000000000000..cc37837e7079 --- /dev/null +++ b/clang/test/SemaCUDA/dtor.cu @@ -0,0 +1,104 @@ +// RUN: %clang_cc1 %s -std=c++20 -fsyntax-only -verify=host +// RUN: %clang_cc1 %s -std=c++20 -fcuda-is-device -fsyntax-only -verify=dev + +// host-no-diagnostics + +#include "Inputs/cuda.h" + +// Virtual dtor ~B() of explicit instantiation B must +// be emitted, which causes host_fun() called. +namespace ExplicitInstantiationExplicitDevDtor { +void host_fun() // dev-note {{'host_fun' declared here}} +{} + +template +constexpr void hd_fun() { + host_fun(); // dev-error {{reference to __host__ function 'host_fun' in __host__ __device__ function}} +} + +struct A { + constexpr ~A() { // dev-note {{called by '~B'}} + hd_fun<8>(); // dev-note {{called by '~A'}} + } +}; + +template +struct B { +public: + virtual __device__ ~B() = default; + A _a; +}; + +template class B; +} + +// The implicit host/device attrs of virtual dtor ~B() should be +// conservatively inferred, where constexpr member dtor's should +// not be considered device since they may call host functions. +// Therefore B::~B() should not have implicit device attr. +// However C::~C() should have implicit device attr since +// it is trivial. +namespace ExplicitInstantiationDtorNoAttr { +void host_fun() +{} + +template +constexpr void hd_fun() { + host_fun(); +} + +struct A { + constexpr ~A() { + hd_fun<8>(); + } +}; + +template +struct B { +public: + virtual ~B() = default; + A _a; +}; + +template +struct C { +public: + virtual ~C() = default; +}; + +template class B; +template class C; +__device__ void foo() { + C x; +} +} + +// Dtors of implicit template class instantiation are not +// conservatively inferred because the invalid usage can +// be diagnosed. +namespace ImplicitInstantiation { +void host_fun() // dev-note {{'host_fun' declared here}} +{} + +template +constexpr void hd_fun() { + host_fun(); // dev-error {{reference to __host__ function 'host_fun' in __host__ __device__ function}} +} + +struct A { + constexpr ~A() { // dev-note {{called by '~B'}} + hd_fun<8>(); // dev-note {{called by '~A'}} + } +}; + +template +struct B { +public: + ~B() = default; // dev-note {{called by 'foo'}} + A _a; +}; + +__device__ void foo() { + B x; +} +}