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 <[email protected]> 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 <[email protected]> --- 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@-1 {{cannot use 'throw' in __host__ __device__ function}} + device_only(); + // host-error@-1 {{reference to __device__ function 'device_only' in __host__ __device__ function}} +} + +inline __host__ __device__ __attribute__((constructor)) void ctor_fn() { + // dev-error@-1 {{CUDA does not support global constructors for __device__ functions}} + device_only(); + // host-error@-1 {{reference to __device__ function 'device_only' in __host__ __device__ function}} +} + +inline __host__ __device__ __attribute__((destructor)) void dtor_fn() { + // dev-error@-1 {{CUDA does not support global destructors for __device__ functions}} + device_only(); + // host-error@-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@-1 {{cannot use 'throw' in __host__ __device__ function}} + device_only(); + // host-all-error@-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@-1 {{cannot use 'throw' in __host__ __device__ function}} + device_only(); + // host-error@-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@-2 {{called by 'caller'}} +// host-used-note@-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@-1 {{cannot use 'throw' in __host__ __device__ function}} +} + +// Internal linkage instead of inline. +static __host__ __device__ __attribute__((used)) void internal() { + throw NULL; + // dev-error@-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 <[email protected]> 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 <[email protected]> --- 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@-1 {{cannot use 'throw' in __host__ __device__ function}} + device_only(); + // host-key-error@-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@-1 {{reference to __device__ function 'device_only' in __host__ __device__ function}} + host_only(); + // dev-all-error@-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@-1 {{reference to __host__ function '~HostOnly' in __host__ __device__ function}} +// dev-cxx17-note@-2 {{in defaulted destructor for 'DefaultedOutOfLine' first required here}} >From c900851dfb7641ab7e80d68b8b51395619e82f0e Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 6 Oct 2026 03:31:37 -0500 Subject: [PATCH 3/7] Also consider ctors Signed-off-by: Steffen Holst Larsen <[email protected]> --- 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 <[email protected]> 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 <[email protected]> --- 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 <[email protected]> Date: Tue, 20 Oct 2026 03:08:16 -0500 Subject: [PATCH 5/7] Add RTTI test cases Signed-off-by: Steffen Holst Larsen <[email protected]> --- .../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@-1 {{reference to __host__ function '~HostOnly' in __host__ __device__ function}} // dev-cxx17-note@-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@-1 {{cannot use 'typeid' in __host__ __device__ function as RTTI is not available in device code}} + (void)dynamic_cast<DerivedPolymorphic *>(p); + // dev-error@-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@-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@-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@-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@-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@-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@-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@-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@-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@-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@-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@-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@-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 <[email protected]> Date: Tue, 20 Oct 2026 03:10:47 -0500 Subject: [PATCH 6/7] Remove now redundant change. Signed-off-by: Steffen Holst Larsen <[email protected]> --- 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 <[email protected]> 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 <[email protected]> --- .../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@-1 {{cannot use 'throw' in __host__ __device__ function}} + device_only(); + // host-error@-1 {{reference to __device__ function 'device_only' in __host__ __device__ function}} + }; + lambda(); + // dev-note@-1 {{called by 'lambda_caller'}} + // host-note@-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@-1 {{cannot use 'throw' in __host__ __device__ function}} + device_only(); + // host-error@-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@-1 {{called by 'available_externally_caller'}} + // host-note@-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 { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
