[clang] [Sema][CUDA][HIP] Report deferred diagnostics of force-emitted functions (PR #229008)
Steffen Larsen via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 20 22:56:16 PDT 2026
https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/229008
>From 5fe90f33a2300af0dbbf2ccf2d137c2e7eb2e1a3 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Mon, 5 Oct 2026 03:19:41 -0500
Subject: [PATCH 1/7] [Sema][CUDA][HIP] Report deferred diagnostics of
force-emitted functions
Deferred diagnostics of a __host__ __device__ function are reported once
the function is known to be emitted, which Sema decided from its uses
and external linkage. CodeGen also emits a function regardless of its
uses if it has the used, constructor or destructor attribute, or with
-femit-all-decls, so such inline functions were emitted without their
diagnostics. Treat these functions as emitted, like what
ASTContext::DeclMustBeEmitted does.
Assisted-by: Claude Opus 5.5
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/Sema/SemaDecl.cpp | 8 ++
.../deferred-diags-forced-emission.cu | 82 +++++++++++++++++++
2 files changed, 90 insertions(+)
create mode 100644 clang/test/SemaCUDA/deferred-diags-forced-emission.cu
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 094601a58d508..deaac4178d178 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -21537,6 +21537,14 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
if (IsEmittedForExternalSymbol())
return FunctionEmissionStatus::Emitted;
+
+ // CodeGen also emits a function regardless of its uses if it is forced to,
+ // so its deferred diagnostics must not wait for a use.
+ const FunctionDecl *Def = FD->getDefinition();
+ if (Def && !Def->hasSkippedBody() &&
+ (LangOpts.EmitAllDecls || Def->hasAttr<UsedAttr>() ||
+ Def->hasAttr<ConstructorAttr>() || Def->hasAttr<DestructorAttr>()))
+ return FunctionEmissionStatus::Emitted;
}
// Otherwise, the function is known-emitted if it's in our set of
diff --git a/clang/test/SemaCUDA/deferred-diags-forced-emission.cu b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
new file mode 100644
index 0000000000000..f588e047b4fe4
--- /dev/null
+++ b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
@@ -0,0 +1,82 @@
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify=dev,dev-used %s
+// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify=host,host-used %s
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -femit-all-decls \
+// RUN: -fsyntax-only -verify=dev,dev-all %s
+// RUN: %clang_cc1 -fcxx-exceptions -femit-all-decls -fsyntax-only \
+// RUN: -verify=host,host-all %s
+// RUN: %clang_cc1 -x hip -fcxx-exceptions -fcuda-is-device -fsyntax-only \
+// RUN: -verify=dev,dev-used %s
+// RUN: %clang_cc1 -x hip -fcxx-exceptions -fsyntax-only -verify=host,host-used %s
+
+// The errors are reported in a compilation that emits code too, before the
+// functions are emitted.
+// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -emit-llvm -o /dev/null \
+// RUN: -verify=dev,dev-used %s
+
+// The deferred diagnostics of a __host__ __device__ function are reported if
+// the function is emitted. Besides the functions that are used, CodeGen emits
+// functions it is forced to emit, whether or not they are used.
+
+#include "Inputs/cuda.h"
+
+__device__ void device_only(); // #device_only
+// host-note@#device_only 4 {{'device_only' declared here}}
+// host-all-note@#device_only {{'device_only' declared here}}
+
+inline __host__ __device__ __attribute__((used)) void used_fn() {
+ throw NULL;
+ // dev-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+ device_only();
+ // host-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+}
+
+inline __host__ __device__ __attribute__((constructor)) void ctor_fn() {
+ // dev-error at -1 {{CUDA does not support global constructors for __device__ functions}}
+ device_only();
+ // host-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+}
+
+inline __host__ __device__ __attribute__((destructor)) void dtor_fn() {
+ // dev-error at -1 {{CUDA does not support global destructors for __device__ functions}}
+ device_only();
+ // host-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+}
+
+// Not used, so not emitted unless all declarations are.
+inline __host__ __device__ void unused_fn() {
+ throw NULL;
+ // dev-all-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+ device_only();
+ // host-all-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+}
+
+// A function used by a forced function is emitted as well.
+inline __host__ __device__ void callee() {
+ throw NULL;
+ // dev-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+ device_only();
+ // host-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+}
+inline __host__ __device__ __attribute__((used)) void caller() { callee(); }
+// Without -femit-all-decls, callee is only emitted as caller's callee.
+// dev-used-note at -2 {{called by 'caller'}}
+// host-used-note at -3 {{called by 'caller'}}
+
+// The definition inherits the attribute from an earlier declaration.
+__attribute__((used)) inline __host__ __device__ void inherited();
+inline __host__ __device__ void inherited() {
+ throw NULL;
+ // dev-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+}
+
+// Internal linkage instead of inline.
+static __host__ __device__ __attribute__((used)) void internal() {
+ throw NULL;
+ // dev-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+}
+
+// An uninstantiated template is not emitted, even with -femit-all-decls.
+template <class T> inline __host__ __device__ void uninstantiated() {
+ throw NULL;
+ device_only();
+}
>From f896b3914b3f62c1ae6af55799269e19ff28aaea Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Mon, 5 Oct 2026 10:13:26 -0500
Subject: [PATCH 2/7] Add more cases that do not emit
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/Sema/SemaDecl.cpp | 36 ++++-
.../deferred-diags-forced-emission.cu | 148 ++++++++++++++++--
2 files changed, 169 insertions(+), 15 deletions(-)
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index deaac4178d178..f9cc831114413 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -21540,10 +21540,38 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
// CodeGen also emits a function regardless of its uses if it is forced to,
// so its deferred diagnostics must not wait for a use.
- const FunctionDecl *Def = FD->getDefinition();
- if (Def && !Def->hasSkippedBody() &&
- (LangOpts.EmitAllDecls || Def->hasAttr<UsedAttr>() ||
- Def->hasAttr<ConstructorAttr>() || Def->hasAttr<DestructorAttr>()))
+ auto IsForcedToBeEmitted = [this, FD, Final]() {
+ const FunctionDecl *Def = FD->getDefinition();
+ if (!Def || Def->hasSkippedBody() ||
+ !(LangOpts.EmitAllDecls || getASTContext().DeclMustBeEmitted(Def)))
+ return false;
+ // Immediate functions are never emitted. Whether an immediate-escalating
+ // function is immediate is only known once its body is complete, so
+ // leave that to the check at the end of the translation unit.
+ if (Def->isImmediateFunction() ||
+ (!Final && LangOpts.CPlusPlus20 && Def->isImmediateEscalating()))
+ return false;
+ // Lambdas, implicit functions and functions defaulted on their first
+ // declaration are not handed to CodeGen, so they are only emitted when
+ // used.
+ if (isLambdaMethod(Def) || Def->isImplicit() ||
+ Def->getCanonicalDecl()->isDefaulted())
+ return false;
+ // An available externally definition is only emitted to be inlined into
+ // its callers.
+ if (getASTContext().GetGVALinkageForFunction(Def) ==
+ GVA_AvailableExternally)
+ return false;
+ // Implicit host device templates are only emitted for the device if they
+ // are used on the device.
+ if (LangOpts.CUDAIsDevice &&
+ LangOpts.OffloadImplicitHostDeviceTemplates &&
+ SemaCUDA::isImplicitHostDeviceFunction(Def) && !Def->isConstexpr())
+ return false;
+ return true;
+ };
+
+ if (IsForcedToBeEmitted())
return FunctionEmissionStatus::Emitted;
}
diff --git a/clang/test/SemaCUDA/deferred-diags-forced-emission.cu b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
index f588e047b4fe4..de1ca6e68f5f5 100644
--- a/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
+++ b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
@@ -1,17 +1,41 @@
-// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -fsyntax-only -verify=dev,dev-used %s
-// RUN: %clang_cc1 -fcxx-exceptions -fsyntax-only -verify=host,host-used %s
-// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -femit-all-decls \
+// RUN: %clang_cc1 -std=c++20 -triple nvptx64-nvidia-cuda \
+// RUN: -aux-triple x86_64-unknown-linux-gnu -fcxx-exceptions -fcuda-is-device \
+// RUN: -fsyntax-only -verify=dev,dev-used %s
+// RUN: %clang_cc1 -std=c++20 -triple x86_64-unknown-linux-gnu -fcxx-exceptions \
+// RUN: -fsyntax-only -verify=host,host-used,host-key %s
+// RUN: %clang_cc1 -std=c++20 -triple nvptx64-nvidia-cuda \
+// RUN: -aux-triple x86_64-unknown-linux-gnu -fcxx-exceptions -fcuda-is-device \
+// RUN: -femit-all-decls -fsyntax-only -verify=dev,dev-all %s
+// RUN: %clang_cc1 -std=c++20 -triple x86_64-unknown-linux-gnu -fcxx-exceptions \
+// RUN: -femit-all-decls -fsyntax-only -verify=host,host-all,host-key %s
+// RUN: %clang_cc1 -std=c++20 -x hip -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu -fcxx-exceptions -fcuda-is-device \
+// RUN: -fsyntax-only -verify=dev,dev-used %s
+// RUN: %clang_cc1 -std=c++20 -x hip -triple x86_64-unknown-linux-gnu \
+// RUN: -fcxx-exceptions -fsyntax-only -verify=host,host-used,host-key %s
+
+// Defaulted functions are immediate-escalating in C++20, which defers their
+// emission status, so check them in C++17 too.
+// RUN: %clang_cc1 -std=c++17 -triple nvptx64-nvidia-cuda \
+// RUN: -aux-triple x86_64-unknown-linux-gnu -fcxx-exceptions -fcuda-is-device \
+// RUN: -femit-all-decls -fsyntax-only -verify=dev,dev-all,dev-cxx17 %s
+
+// An inline function cannot be a key function in this ABI.
+// RUN: %clang_cc1 -std=c++20 -triple arm64-apple-macosx -fcxx-exceptions \
+// RUN: -fsyntax-only -verify=host,host-used %s
+
+// Implicit host device templates are only emitted for the device if they are
+// used there, even with -femit-all-decls.
+// RUN: %clang_cc1 -std=c++20 -x hip -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu -fcxx-exceptions -fcuda-is-device \
+// RUN: -foffload-implicit-host-device-templates -femit-all-decls \
// RUN: -fsyntax-only -verify=dev,dev-all %s
-// RUN: %clang_cc1 -fcxx-exceptions -femit-all-decls -fsyntax-only \
-// RUN: -verify=host,host-all %s
-// RUN: %clang_cc1 -x hip -fcxx-exceptions -fcuda-is-device -fsyntax-only \
-// RUN: -verify=dev,dev-used %s
-// RUN: %clang_cc1 -x hip -fcxx-exceptions -fsyntax-only -verify=host,host-used %s
// The errors are reported in a compilation that emits code too, before the
// functions are emitted.
-// RUN: %clang_cc1 -fcxx-exceptions -fcuda-is-device -emit-llvm -o /dev/null \
-// RUN: -verify=dev,dev-used %s
+// RUN: %clang_cc1 -std=c++20 -triple nvptx64-nvidia-cuda \
+// RUN: -aux-triple x86_64-unknown-linux-gnu -fcxx-exceptions -fcuda-is-device \
+// RUN: -emit-llvm -o /dev/null -verify=dev,dev-used %s
// The deferred diagnostics of a __host__ __device__ function are reported if
// the function is emitted. Besides the functions that are used, CodeGen emits
@@ -21,7 +45,9 @@
__device__ void device_only(); // #device_only
// host-note@#device_only 4 {{'device_only' declared here}}
-// host-all-note@#device_only {{'device_only' declared here}}
+// host-all-note@#device_only 2 {{'device_only' declared here}}
+// host-key-note@#device_only {{'device_only' declared here}}
+__host__ void host_only(); // #host_only
inline __host__ __device__ __attribute__((used)) void used_fn() {
throw NULL;
@@ -80,3 +106,103 @@ template <class T> inline __host__ __device__ void uninstantiated() {
throw NULL;
device_only();
}
+
+// An inline key function is emitted with the vtable.
+struct KeyFunction {
+ virtual __host__ __device__ void key();
+};
+inline __host__ __device__ void KeyFunction::key() {
+ throw NULL;
+ // dev-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+ device_only();
+ // host-key-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+}
+
+#if __cplusplus >= 202002L
+// dev-all-note@#host_only {{'host_only' declared here}}
+
+// Immediate functions are never emitted, even if forced.
+__attribute__((used)) consteval __host__ __device__ int immediate(int x) {
+ if (x) {
+ (void)&device_only;
+ throw NULL;
+ }
+ return 0;
+}
+
+// Neither is a function that becomes immediate by calling an immediate
+// function, which is only known once its body is complete.
+consteval int id(int x) { return x; }
+template <class T> constexpr __host__ __device__ int escalating(T x) {
+ if (x) {
+ device_only();
+ host_only();
+ }
+ return id(x);
+}
+int escalated = escalating(0);
+
+// Unlike a function that does not become immediate.
+template <class T> constexpr __host__ __device__ int not_escalating(T x) {
+ if (x) {
+ device_only();
+ // host-all-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+ host_only();
+ // dev-all-error at -1 {{reference to __host__ function 'host_only' in __host__ __device__ function}}
+ }
+ return x;
+}
+int not_escalated = not_escalating(0);
+#endif
+
+// Lambdas are only emitted if they are used.
+inline auto lambda = [] __attribute__((used)) __host__ __device__ {
+ throw NULL;
+ device_only();
+};
+
+// With -foffload-implicit-host-device-templates, this is an implicit host
+// device function only used on the host.
+template <class T> T implicit_hd(T x) {
+ host_only();
+ return x;
+}
+int host_user() { return implicit_hd(0); }
+
+// An available externally definition is only emitted to be inlined into its
+// callers.
+extern inline __attribute__((gnu_inline, used)) __host__ __device__ void
+available_externally() {
+ throw NULL;
+ device_only();
+}
+
+// Implicit functions and functions defaulted on their first declaration are
+// only emitted when used, here only by host functions.
+struct HostOnly {
+ __host__ HostOnly() {}
+ __host__ ~HostOnly() {} // #host_only_dtor
+ // dev-all-note@#host_only_dtor {{'~HostOnly' declared here}}
+};
+struct Defaulted {
+ HostOnly m;
+ __host__ __device__ ~Defaulted() = default;
+};
+void use_defaulted() { Defaulted d; }
+struct Base {
+ __host__ __device__ Base(int) {}
+};
+struct Inheriting : Base {
+ using Base::Base;
+ HostOnly m;
+};
+void use_inheriting() { Inheriting i(0); }
+
+// Unlike a function defaulted after its first declaration.
+struct DefaultedOutOfLine {
+ HostOnly m;
+ __host__ __device__ ~DefaultedOutOfLine();
+};
+inline __host__ __device__ DefaultedOutOfLine::~DefaultedOutOfLine() = default;
+// dev-all-error at -1 {{reference to __host__ function '~HostOnly' in __host__ __device__ function}}
+// dev-cxx17-note at -2 {{in defaulted destructor for 'DefaultedOutOfLine' first required here}}
>From c900851dfb7641ab7e80d68b8b51395619e82f0e Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 6 Oct 2026 03:31:37 -0500
Subject: [PATCH 3/7] Also consider ctors
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/Sema/Sema.cpp | 5 ++++
.../SemaCUDA/call-stack-for-deferred-err.cu | 25 +++++++++++++++++++
2 files changed, 30 insertions(+)
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 29158cfff6231..a3f577b693b46 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -2098,6 +2098,11 @@ class DeferredDiagnosticsEmitter
if (auto *S = FD->getBody()) {
this->Visit(S);
}
+ // A constructor initializes its bases and members, and so calls their
+ // constructors, in its initializer list rather than its body.
+ if (auto *Ctor = dyn_cast<CXXConstructorDecl>(FD))
+ for (CXXCtorInitializer *Init : Ctor->inits())
+ this->Visit(Init->getInit());
if (CXXDestructorDecl *Dtor = dyn_cast<CXXDestructorDecl>(FD))
asImpl().VisitCalledDestructors(Dtor);
UsePath.pop_back();
diff --git a/clang/test/SemaCUDA/call-stack-for-deferred-err.cu b/clang/test/SemaCUDA/call-stack-for-deferred-err.cu
index 76bb5116ea283..78bcd744422ed 100644
--- a/clang/test/SemaCUDA/call-stack-for-deferred-err.cu
+++ b/clang/test/SemaCUDA/call-stack-for-deferred-err.cu
@@ -16,3 +16,28 @@ __global__ void kernel() { device_fn2(); } // expected-note {{called by 'kernel'
inline __host__ __device__ void hd_fn(int n) {
int vla[n]; // expected-error {{variable-length array}}
}
+
+// A constructor calls the constructors of its bases and members from its
+// initializer list.
+struct Member {
+ __host__ __device__ Member() {
+ int n = 42;
+ int vla[n]; // expected-error {{variable-length array}}
+ }
+};
+struct HasMember { // expected-note {{called by 'HasMember'}}
+ Member m;
+};
+
+struct Base {
+ __host__ __device__ Base() {
+ int n = 42;
+ int vla[n]; // expected-error {{variable-length array}}
+ }
+};
+struct Derived : Base {}; // expected-note {{called by 'Derived'}}
+
+__global__ void ctor_kernel() {
+ HasMember h; // expected-note {{which is called by 'ctor_kernel'}}
+ Derived d; // expected-note {{which is called by 'ctor_kernel'}}
+}
>From 4b516656ebbf36601b458d4a9d08d2350cec4d14 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Fri, 16 Oct 2026 05:19:59 -0500
Subject: [PATCH 4/7] Refactor checks and use common helpers
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/Sema/SemaDecl.cpp | 73 ++++++++++++++++++-------------------
1 file changed, 36 insertions(+), 37 deletions(-)
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 0bb4425ae78d8..07cab8a90149e 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -21504,6 +21504,40 @@ void Sema::ActOnPragmaWeakAlias(IdentifierInfo* Name,
}
}
+/// Whether the definition \p FD is only handed to the ASTConsumer, and thus to
+/// CodeGen, once it is used, so that it is never emitted unless it is used.
+static bool isEmittedOnlyIfUsed(const FunctionDecl *FD) {
+ // Lambda call operators are defined by their lambda expressions, which are
+ // not top-level declarations. Implicit functions and functions defaulted on
+ // their first declaration are defined when they are used.
+ return isLambdaMethod(FD) || FD->isImplicit() ||
+ FD->getCanonicalDecl()->isDefaulted();
+}
+
+/// Whether CodeGen emits \p FD regardless of whether it is used.
+static bool isEmittedRegardlessOfUses(const Sema &S, const FunctionDecl *FD,
+ bool Final) {
+ const FunctionDecl *Def = FD->getDefinition();
+ if (!Def || Def->hasSkippedBody())
+ return false;
+ // Immediate functions are never emitted. Whether an immediate-escalating
+ // function is immediate is only known once its body is complete, so leave
+ // that to the check at the end of the translation unit.
+ const LangOptions &LangOpts = S.getLangOpts();
+ if (Def->isImmediateFunction() ||
+ (!Final && LangOpts.CPlusPlus20 && Def->isImmediateEscalating()))
+ return false;
+ if (isEmittedOnlyIfUsed(Def))
+ return false;
+ // An available externally definition is only emitted to be inlined into its
+ // callers.
+ ASTContext &Ctx = S.getASTContext();
+ if (Ctx.GetGVALinkageForFunction(Def) == GVA_AvailableExternally)
+ return false;
+ return Ctx.shouldEmitCUDADecl(Def) &&
+ (LangOpts.EmitAllDecls || Ctx.DeclMustBeEmitted(Def));
+}
+
Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
bool Final) {
assert(FD && "Expected non-null FunctionDecl");
@@ -21577,43 +21611,8 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
(T == CUDAFunctionTarget::Device || T == CUDAFunctionTarget::Global))
return FunctionEmissionStatus::CUDADiscarded;
- if (IsEmittedForExternalSymbol())
- return FunctionEmissionStatus::Emitted;
-
- // CodeGen also emits a function regardless of its uses if it is forced to,
- // so its deferred diagnostics must not wait for a use.
- auto IsForcedToBeEmitted = [this, FD, Final]() {
- const FunctionDecl *Def = FD->getDefinition();
- if (!Def || Def->hasSkippedBody() ||
- !(LangOpts.EmitAllDecls || getASTContext().DeclMustBeEmitted(Def)))
- return false;
- // Immediate functions are never emitted. Whether an immediate-escalating
- // function is immediate is only known once its body is complete, so
- // leave that to the check at the end of the translation unit.
- if (Def->isImmediateFunction() ||
- (!Final && LangOpts.CPlusPlus20 && Def->isImmediateEscalating()))
- return false;
- // Lambdas, implicit functions and functions defaulted on their first
- // declaration are not handed to CodeGen, so they are only emitted when
- // used.
- if (isLambdaMethod(Def) || Def->isImplicit() ||
- Def->getCanonicalDecl()->isDefaulted())
- return false;
- // An available externally definition is only emitted to be inlined into
- // its callers.
- if (getASTContext().GetGVALinkageForFunction(Def) ==
- GVA_AvailableExternally)
- return false;
- // Implicit host device templates are only emitted for the device if they
- // are used on the device.
- if (LangOpts.CUDAIsDevice &&
- LangOpts.OffloadImplicitHostDeviceTemplates &&
- SemaCUDA::isImplicitHostDeviceFunction(Def) && !Def->isConstexpr())
- return false;
- return true;
- };
-
- if (IsForcedToBeEmitted())
+ if (IsEmittedForExternalSymbol() ||
+ isEmittedRegardlessOfUses(*this, FD, Final))
return FunctionEmissionStatus::Emitted;
}
>From cae9a2d4ba552f4e327369dcff1292c1af798860 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 20 Oct 2026 03:08:16 -0500
Subject: [PATCH 5/7] Add RTTI test cases
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
.../deferred-diags-forced-emission.cu | 41 +++++++++++++++++++
clang/test/SemaCUDA/device-rtti.cu | 24 ++++++++++-
2 files changed, 64 insertions(+), 1 deletion(-)
diff --git a/clang/test/SemaCUDA/deferred-diags-forced-emission.cu b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
index de1ca6e68f5f5..662aac38db288 100644
--- a/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
+++ b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
@@ -206,3 +206,44 @@ struct DefaultedOutOfLine {
inline __host__ __device__ DefaultedOutOfLine::~DefaultedOutOfLine() = default;
// dev-all-error at -1 {{reference to __host__ function '~HostOnly' in __host__ __device__ function}}
// dev-cxx17-note at -2 {{in defaulted destructor for 'DefaultedOutOfLine' first required here}}
+
+// Uses of RTTI in device code are diagnosed the same way.
+namespace std {
+class type_info {};
+} // namespace std
+struct Polymorphic {
+ virtual __host__ __device__ ~Polymorphic() {}
+};
+struct DerivedPolymorphic : Polymorphic {};
+
+inline __host__ __device__ __attribute__((used)) void used_rtti(Polymorphic *p) {
+ (void)typeid(*p);
+ // dev-error at -1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}}
+ (void)dynamic_cast<DerivedPolymorphic *>(p);
+ // dev-error at -1 {{cannot use 'dynamic_cast' in __host__ __device__ function as RTTI is not available in device code}}
+}
+
+inline __host__ __device__ void unused_rtti(Polymorphic *p) {
+ (void)typeid(*p);
+ // dev-all-error at -1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}}
+}
+
+// A default argument is checked where it is used.
+inline __host__ __device__ void rtti_default_arg(const std::type_info *t = &typeid(int)) {}
+// dev-error at -1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}}
+inline __host__ __device__ __attribute__((used)) void use_rtti_default_arg() {
+ rtti_default_arg();
+ // dev-note at -1 {{default argument used here}}
+}
+
+// So is a default member initializer, here in the implicit constructor called
+// by a forced function.
+struct RTTIMemberInit { // #rtti_member_init
+ const std::type_info *t = &typeid(int);
+ // dev-error at -1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}}
+};
+// dev-note@#rtti_member_init {{default member initializer used here}}
+inline __host__ __device__ __attribute__((used)) void use_rtti_member_init() {
+ RTTIMemberInit m;
+ // dev-note at -1 {{called by 'use_rtti_member_init'}}
+}
diff --git a/clang/test/SemaCUDA/device-rtti.cu b/clang/test/SemaCUDA/device-rtti.cu
index c0375dd460ca9..6e508ab7634d0 100644
--- a/clang/test/SemaCUDA/device-rtti.cu
+++ b/clang/test/SemaCUDA/device-rtti.cu
@@ -90,6 +90,12 @@ inline __host__ __device__ void hd3(B *b) {
__device__ void call_hd3(B *b) { hd3(b); }
// dev-note at -1 {{called by 'call_hd3'}}
+// Error, emitted for device although it is not used there.
+inline __host__ __device__ __attribute__((used)) void hd_used(B *b) {
+ (void)typeid(*b);
+ // dev-error at -1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}}
+}
+
// Templates are checked when they are instantiated.
template <class T> __device__ T *tmpl_unused(B *b) {
return dynamic_cast<T *>(b);
@@ -233,13 +239,18 @@ __device__ void use_default_arg() {
// expected-note at -1 {{default argument used here}}
}
-inline __host__ __device__ void hd_default_arg(const std::type_info *t = &typeid(int)) {}
+inline __host__ __device__ void hd_default_arg(const std::type_info *t = &typeid(int)) {} // #hd_default_arg
// expected-error at -1 {{cannot use 'typeid' in __device__ function as RTTI is not available in device code}}
void use_hd_default_arg_host() { hd_default_arg(); }
__device__ void use_hd_default_arg_device() {
hd_default_arg();
// expected-note at -1 {{default argument used here}}
}
+// dev-error@#hd_default_arg {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}}
+inline __host__ __device__ __attribute__((used)) void use_hd_default_arg_used() {
+ hd_default_arg();
+ // dev-note at -1 {{default argument used here}}
+}
// Diagnosed once, at the use, although the default argument is parsed in a
// __device__ function.
@@ -316,6 +327,17 @@ __device__ void nested_default_arg() {
(void)n;
}
+// The constructor of a class calls the constructors of its members.
+struct Wrapper { // #wrapper_struct
+ NestedDefaultArg x;
+};
+// dev-note@#wrapper_struct {{called by 'Wrapper'}}
+__device__ void wrapper() {
+ Wrapper n;
+ // dev-note at -1 {{which is called by 'wrapper'}}
+ (void)n;
+}
+
// A dependent default argument used in a dependent default member initializer.
template <class T> struct DepNestedDefaultArg { // #dep_nested_struct
const std::type_info *t = (dep_default_arg<T>(), nullptr);
>From ce0e84e0975f93a795b214274310a3ce9241bb9d Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 20 Oct 2026 03:10:47 -0500
Subject: [PATCH 6/7] Remove now redundant change.
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/Sema/Sema.cpp | 6 +-----
1 file changed, 1 insertion(+), 5 deletions(-)
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 89a5549544629..bae41a6e6f5cf 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -2104,11 +2104,7 @@ class DeferredDiagnosticsEmitter
for (auto *Init : Ctor->inits())
this->Visit(Init->getInit());
}
- // A constructor initializes its bases and members, and so calls their
- // constructors, in its initializer list rather than its body.
- if (auto *Ctor = dyn_cast<CXXConstructorDecl>(FD))
- for (CXXCtorInitializer *Init : Ctor->inits())
- this->Visit(Init->getInit());
+
if (CXXDestructorDecl *Dtor = dyn_cast<CXXDestructorDecl>(FD))
asImpl().VisitCalledDestructors(Dtor);
UsePath.pop_back();
>From fddbf2dea87cc91881a4d332db0d48f4cfa2bfb1 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 20 Oct 2026 03:18:13 -0500
Subject: [PATCH 7/7] Add cases for called lambda and called external def
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
.../deferred-diags-forced-emission.cu | 30 ++++++++++++++++++-
1 file changed, 29 insertions(+), 1 deletion(-)
diff --git a/clang/test/SemaCUDA/deferred-diags-forced-emission.cu b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
index 662aac38db288..2176cd53ac0f6 100644
--- a/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
+++ b/clang/test/SemaCUDA/deferred-diags-forced-emission.cu
@@ -44,7 +44,7 @@
#include "Inputs/cuda.h"
__device__ void device_only(); // #device_only
-// host-note@#device_only 4 {{'device_only' declared here}}
+// host-note@#device_only 6 {{'device_only' declared here}}
// host-all-note@#device_only 2 {{'device_only' declared here}}
// host-key-note@#device_only {{'device_only' declared here}}
__host__ void host_only(); // #host_only
@@ -161,6 +161,19 @@ inline auto lambda = [] __attribute__((used)) __host__ __device__ {
device_only();
};
+// Such as by a forced function.
+inline __host__ __device__ __attribute__((used)) void lambda_caller() {
+ auto lambda = [] {
+ throw NULL;
+ // dev-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+ device_only();
+ // host-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+ };
+ lambda();
+ // dev-note at -1 {{called by 'lambda_caller'}}
+ // host-note at -2 {{called by 'lambda_caller'}}
+}
+
// With -foffload-implicit-host-device-templates, this is an implicit host
// device function only used on the host.
template <class T> T implicit_hd(T x) {
@@ -177,6 +190,21 @@ available_externally() {
device_only();
}
+// So it is checked when a forced function calls it.
+extern inline __attribute__((gnu_inline, used)) __host__ __device__ void
+called_available_externally() {
+ throw NULL;
+ // dev-error at -1 {{cannot use 'throw' in __host__ __device__ function}}
+ device_only();
+ // host-error at -1 {{reference to __device__ function 'device_only' in __host__ __device__ function}}
+}
+inline __host__ __device__ __attribute__((used)) void
+available_externally_caller() {
+ called_available_externally();
+ // dev-note at -1 {{called by 'available_externally_caller'}}
+ // host-note at -2 {{called by 'available_externally_caller'}}
+}
+
// Implicit functions and functions defaulted on their first declaration are
// only emitted when used, here only by host functions.
struct HostOnly {
More information about the cfe-commits
mailing list