[CUDA][HIP] check dtor in deferred diag

[yxsamliu]

Currently the deferred diag fails to diagnose calling of host function in host device function in device compilation triggered by destructors.

This can be further divided into two issuse:

1. the deferred diag visitor does not visit dtor of member and parent class when visiting dtor, which it should

2. the deferred diag visitor does not visit virtual dtor of explicit template class instantiation, which it should

Due to these issues, some constexpr functions which call host functions are emitted on device side, which causes undefind symbols in linking stage, as revealed by
#108548

By fixing these issue, clang will diag the issues early during compilation instead of linking.

[llvmbot]

@llvm/pr-subscribers-clang

Author: Yaxun (Sam) Liu (yxsamliu)

Changes

Currently the deferred diag fails to diagnose calling of host function in host device function in device compilation triggered by destructors.

This can be further divided into two issuse:

1. the deferred diag visitor does not visit dtor of member and parent class when visiting dtor, which it should

2. the deferred diag visitor does not visit virtual dtor of explicit template class instantiation, which it should

Due to these issues, some constexpr functions which call host functions are emitted on device side, which causes undefind symbols in linking stage, as revealed by
#108548

By fixing these issue, clang will diag the issues early during compilation instead of linking.

---

Full diff: https://github.com/llvm/llvm-project/pull/129117.diff

5 Files Affected:

• (modified) clang/include/clang/Sema/Sema.h (+1-1)
• (modified) clang/lib/Sema/Sema.cpp (+58)
• (modified) clang/lib/Sema/SemaCUDA.cpp (+15)
• (modified) clang/lib/Sema/SemaDecl.cpp (+15)
• (added) clang/test/SemaCUDA/dtor.cu (+102)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index ebdbc69384efb..3b2be86a88e82 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -4392,11 +4392,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;
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index c699e92985156..fa9e6db62a3a0 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1798,6 +1798,62 @@ 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");
@@ -1839,6 +1895,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);
   }
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 0e1bf727d72d2..473956c37bb51 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<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.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 285bd27a35a76..ab86d2a2d920a 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -20469,6 +20469,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
diff --git a/clang/test/SemaCUDA/dtor.cu b/clang/test/SemaCUDA/dtor.cu
new file mode 100644
index 0000000000000..c266e51f5c29e
--- /dev/null
+++ b/clang/test/SemaCUDA/dtor.cu
@@ -0,0 +1,102 @@
+// 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<float>::~B() is inferred to
+// have implicit device attr since dtors of its members and parent classes can
+// be executed on device. This causes a diagnostic since B<float>::~B() must
+// be emitted, and it eventually causes host_fun() called on device side.
+namespace ExplicitInstantiationDtorNoAttr {
+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 ~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;
+}
+}

[Artem-B]

LGTM functionally, some style nits.