https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/229008
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 >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] [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(); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
