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/2] [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 094601a58d5085..deaac4178d178c 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 00000000000000..f588e047b4fe4f --- /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/2] 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 deaac4178d178c..f9cc8311144137 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 f588e047b4fe4f..de1ca6e68f5f5b 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}} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
