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/3] [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/3] 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/3] 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'}} +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
