[llvm-branch-commits] [clang] [CUDA][HIP] fix virtual dtor host/device attr (PR #130126)
Tom Stellard via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Mar 18 14:15:25 PDT 2025
https://github.com/tstellar updated https://github.com/llvm/llvm-project/pull/130126
>From 0ceb4efefeaa650166254ec8f1836a2ee76be207 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 28 Feb 2025 09:58:19 -0500
Subject: [PATCH] [CUDA][HIP] fix virtual dtor host/device attr (#128926)
When inferring host device attr of virtual dtor of explicit
template class instantiation, clang should be conservative.
This guarantees dtors that may call host functions not to
have implicit device attr, therefore will not be emitted
on device side.
Backports: 0f0665db067f d37a39207bc1
Fixes: #108548
---
clang/docs/HIPSupport.rst | 20 ++++++
clang/include/clang/Sema/Sema.h | 2 +-
clang/lib/Sema/Sema.cpp | 43 +++++++++++++
clang/lib/Sema/SemaCUDA.cpp | 23 ++++++-
clang/lib/Sema/SemaDecl.cpp | 15 +++++
clang/test/SemaCUDA/dtor.cu | 104 ++++++++++++++++++++++++++++++++
6 files changed, 204 insertions(+), 3 deletions(-)
create mode 100644 clang/test/SemaCUDA/dtor.cu
diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst
index 481ed39230813..8f473c21e1918 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 a30a7076ea5d4..af648d7f9c63f 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/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;
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 9507d7602aa40..e0eac690e6e65 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/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);
}
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 0e1bf727d72d2..0e5fc5e1a40b4 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.
@@ -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 01f09aba8c2ad..f70401ea33b4a 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<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..cc37837e70791
--- /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<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;
+}
+}
More information about the llvm-branch-commits
mailing list