[CUDA][HIP] fix virtual dtor host/device attr

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
 ==============================================================
 

clang/include/clang/Sema/Sema.h

@@ -4336,11 +4336,11 @@ class Sema final : public SemaBase {
   // 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<Decl *, 4> 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<const NamedDecl *, const NamedDecl *> ShadowingDecls;

clang/lib/Sema/Sema.cpp

@@ -1789,6 +1789,47 @@ class DeferredDiagnosticsEmitter
       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<RecordType>())
+        if (const auto *ClassDecl = dyn_cast<CXXRecordDecl>(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<RecordType>())
+        if (const auto *BaseDecl = dyn_cast<CXXRecordDecl>(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<VarDecl>(D))
+        if (VD->isThisDeclarationADefinition() &&
+            VD->needsDestruction(S.Context)) {
+          QualType VT = VD->getType();
+          if (const auto *RT = VT->getAs<RecordType>())
+            if (const auto *ClassDecl = dyn_cast<CXXRecordDecl>(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 @@ class DeferredDiagnosticsEmitter
     if (auto *S = FD->getBody()) {
       this->Visit(S);
     }
+    if (CXXDestructorDecl *Dtor = dyn_cast<CXXDestructorDecl>(FD))
+      asImpl().VisitCalledDestructors(Dtor);
     UsePath.pop_back();
     InUsePath.erase(FD);
   }

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<CXXDestructorDecl>(MemberDecl) && MemberDecl->isVirtual()) {
+    if (auto *Spec = dyn_cast<ClassTemplateSpecializationDecl>(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 {

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<CXXDestructorDecl>(FD)) {
+      if (Destructor->isVirtual()) {
+        if (auto *Spec = dyn_cast<ClassTemplateSpecializationDecl>(
+                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

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<float> must
+// be emitted, which causes host_fun() called.
+namespace ExplicitInstantiationExplicitDevDtor {
+void host_fun() // dev-note {{'host_fun' declared here}}
+{}
+
+template <unsigned>
+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 <typename T>
+struct B {
+public:
+  virtual __device__ ~B() = default;
+  A _a;
+};
+
+template class B<float>;
+}
+
+// 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<float>::~B() should not have implicit device attr.
+// However C<float>::~C() should have implicit device attr since
+// it is trivial.
+namespace ExplicitInstantiationDtorNoAttr {
+void host_fun()
+{}
+
+template <unsigned>
+constexpr void hd_fun() {
+  host_fun();
+}
+
+struct A {
+  constexpr ~A() {
+     hd_fun<8>();
+  }
+};
+
+template <typename T>
+struct B {
+public:
+  virtual ~B() = default;
+  A _a;
+};
+
+template <typename T>
+struct C {
+public:
+  virtual ~C() = default;
+};
+
+template class B<float>;
+template class C<float>;
+__device__ void foo() {
+  C<float> 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 <unsigned>
+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 <typename T>
+struct B {
+public:
+  ~B() = default; // dev-note {{called by 'foo'}}
+  A _a;
+};
+
+__device__ void foo() {
+  B<float> x;
+}
+}