https://github.com/pjmc-oliveira updated https://github.com/llvm/llvm-project/pull/222985
>From c6907ad4d45d138f9e0055a7aefefc3046d35e8b Mon Sep 17 00:00:00 2001 From: Pedro Oliveira <[email protected]> Date: Fri, 4 Sep 2026 16:17:51 +0000 Subject: [PATCH] [CUDA] Treat function-scope statics in device code as device variables For function-scope statics without an explicit host/device annotation (so execution space is implied by the enclosing function) there were two issues: - Allowed cases (empty constructor, or constant initializer with an empty destructor) got a guard variable. - Disallowed cases (non-empty/non-constant initializer) were not diagnosed. Fixes https://github.com/llvm/llvm-project/issues/117023 Assisted-by: Claude Opus 5 --- .../clang/Basic/DiagnosticSemaKinds.td | 4 + clang/lib/CodeGen/CGDecl.cpp | 43 ++++++++-- clang/lib/Sema/SemaCUDA.cpp | 26 +++++- .../CodeGenCUDA/Inputs/cuda-initializers.h | 12 +++ clang/test/CodeGenCUDA/device-var-init.cu | 23 ++++++ .../test/SemaCUDA/Inputs/cuda-initializers.h | 12 +++ clang/test/SemaCUDA/device-var-init-cxx20.cu | 26 ++++++ clang/test/SemaCUDA/device-var-init.cu | 81 +++++++++++++++++-- 8 files changed, 215 insertions(+), 12 deletions(-) create mode 100644 clang/test/SemaCUDA/device-var-init-cxx20.cu diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 8e54962a3d383..74c7f08c06482 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9827,6 +9827,10 @@ def err_cuda_device_exceptions : Error< def err_dynamic_var_init : Error< "dynamic initialization is not supported for " "__device__, __constant__, __shared__, and __managed__ variables">; +def err_cuda_static_local_var : Error< + "cannot use 'static' local variable requiring runtime initialization or " + "destruction in " + "%select{__device__|__global__|__host__|__host__ __device__}0 function">; def err_cuda_ctor_dtor_attrs : Error<"CUDA does not support global %0 for __device__ functions">; def err_shared_var_init : Error< diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp index 019b7ac22272e..4b4835c732905 100644 --- a/clang/lib/CodeGen/CGDecl.cpp +++ b/clang/lib/CodeGen/CGDecl.cpp @@ -363,6 +363,12 @@ CodeGenFunction::AddInitializerToStaticVarDecl(const VarDecl &D, ConstantEmitter emitter(*this); llvm::Constant *Init = emitter.tryEmitForInitializer(D); + // CUDA device compilation only. Sema has verified that a function-scope + // static in device code has an empty/constant initializer and an empty + // destructor, so neither needs to be emitted here. + const bool SkipCUDADeviceInit = + getLangOpts().CUDAIsDevice && !getLangOpts().GPUAllowDeviceInit; + // If constant emission failed, then this should be a C++ static // initializer. if (!Init) { @@ -375,7 +381,20 @@ CodeGenFunction::AddInitializerToStaticVarDecl(const VarDecl &D, // be constant. GV->setConstant(false); - EmitCXXGuardedInit(D, GV, /*PerformInit*/true); +#ifndef NDEBUG + // Constant emission failed, so in device code Sema can only have accepted + // this via its "empty constructor" rule. + if (SkipCUDADeviceInit) { + const auto *CE = dyn_cast<CXXConstructExpr>(D.getInit()); + const CXXConstructorDecl *Ctor = CE ? CE->getConstructor() : nullptr; + assert(Ctor && Ctor->hasTrivialBody() && Ctor->getNumParams() == 0 && + "non-empty initializer for a device-side static should have " + "been diagnosed by Sema"); + } +#endif + + if (!SkipCUDADeviceInit) + EmitCXXGuardedInit(D, GV, /*PerformInit*/ true); } return GV; } @@ -399,10 +418,24 @@ CodeGenFunction::AddInitializerToStaticVarDecl(const VarDecl &D, emitter.finalize(GV); - if (NeedsDtor && HaveInsertPoint()) { - // We have a constant initializer, but a nontrivial destructor. We still - // need to perform a guarded "initialization" in order to register the - // destructor. +#ifndef NDEBUG + if (SkipCUDADeviceInit && NeedsDtor) { + const auto *RD = + D.getType()->getBaseElementTypeUnsafe()->getAsCXXRecordDecl(); + const CXXDestructorDecl *Dtor = RD ? RD->getDestructor() : nullptr; + assert(Dtor && Dtor->hasTrivialBody() && + "non-empty destructor for a device-side static should have been " + "diagnosed by Sema"); + } +#endif + + // We have a constant initializer, but a nontrivial destructor. We still need + // to perform a guarded "initialization" in order to register the destructor. + // + // CUDA allows a device-side static whose destructor is non-trivial, but + // empty. A user-provided destructor with an empty body is non-trivial but + // does nothing. + if (NeedsDtor && HaveInsertPoint() && !SkipCUDADeviceInit) { EmitCXXGuardedInit(D, GV, /*PerformInit*/false); } diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 090030ea82503..d3d4dca8d129b 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -764,11 +764,35 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage() || IsDependentVar(VD)) return; + + // A function-scope static is a device variable when it is emitted on the + // device side, and has the same initialization restrictions. + CUDAVariableTarget VT = IdentifyTarget(VD); + + // CVT_Both means the enclosing function is __host__ __device__, so the + // variable is emitted on both sides and only the device-side copy is + // restricted. + const bool IsDeviceCopyOfHDStatic = + VT == CVT_Both && getLangOpts().CUDAIsDevice; + + // constexpr implies constant initialization and constant destruction. + bool IsDeviceLocalStatic = !IsSharedVar && !IsDeviceOrConstantVar && + VD->isStaticLocal() && !VD->isConstexpr() && + (VT == CVT_Device || IsDeviceCopyOfHDStatic); + const Expr *Init = VD->getInit(); - if (IsDeviceOrConstantVar || IsSharedVar) { + if (IsDeviceOrConstantVar || IsSharedVar || IsDeviceLocalStatic) { if (HasAllowedCUDADeviceStaticInitializer( *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant)) return; + // Defer the diagnostic until we know whether a __host__ __device__ + // function is emitted on the device side. + if (IsDeviceLocalStatic) { + if (DiagIfDeviceCode(VD->getLocation(), diag::err_cuda_static_local_var) + << CurrentTarget() << Init->getSourceRange()) + VD->setInvalidDecl(); + return; + } Diag(VD->getLocation(), IsSharedVar ? diag::err_shared_var_init : diag::err_dynamic_var_init) << Init->getSourceRange(); diff --git a/clang/test/CodeGenCUDA/Inputs/cuda-initializers.h b/clang/test/CodeGenCUDA/Inputs/cuda-initializers.h index 186b160276512..bcab7f852e912 100644 --- a/clang/test/CodeGenCUDA/Inputs/cuda-initializers.h +++ b/clang/test/CodeGenCUDA/Inputs/cuda-initializers.h @@ -15,6 +15,12 @@ struct EC { __device__ EC(int) {} // -- not allowed }; +// host/device, empty constructor +struct HD_EC { + int hd_ec; + __host__ __device__ HD_EC() {} // -- allowed +}; + // empty destructor struct ED { __device__ ~ED() {} // -- allowed @@ -54,6 +60,12 @@ struct NEC { __device__ NEC() { nec = 1; } }; +// host/device, non-empty constructor -- not allowed +struct HD_NEC { + int hd_nec; + __host__ __device__ HD_NEC() { hd_nec = 1; } +}; + // non-empty destructor -- not allowed struct NED { int ned; diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu index 8c7a2884ad328..b204ea39b250b 100644 --- a/clang/test/CodeGenCUDA/device-var-init.cu +++ b/clang/test/CodeGenCUDA/device-var-init.cu @@ -12,6 +12,9 @@ // RUN: %clang_cc1 -triple amdgpu -fcuda-is-device -std=c++11 \ // RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,AMDGCN %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \ +// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefix=DEVICE-NEG %s + #ifdef __clang__ #include "Inputs/cuda.h" #endif @@ -162,6 +165,9 @@ __constant__ EC_I_EC c_ec_i_ec; // DEVICE: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5] // DEVICE: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123 +// DEVICE: @_ZZ15hd_local_staticvE2ec = internal addrspace(1) global %struct.HD_EC zeroinitializer +// DEVICE: @_ZZ20df_local_static_dtorvE4s_ed = internal addrspace(1) global %struct.ED zeroinitializer + // We should not emit global initializers for device-side variables. // DEVICE-NOT: @__cxx_global_var_init @@ -305,3 +311,20 @@ __device__ void df() { // We should not emit global init function. // DEVICE-NOT: @_GLOBAL__sub_I + +// host/device, empty constructor -- allowed, but needs no guard on the device +__host__ __device__ void hd_local_static() { + static HD_EC ec; + // HOST: @_ZGVZ15hd_local_staticvE2ec = internal global i8 0 +} + +// trivial constructor, empty destructor -- allowed, but the destructor must +// not be registered +__device__ void df_local_static_dtor() { + static ED s_ed; +} + +// We should not emit guard variables or destructor registration for +// device-side statics. +// DEVICE-NEG-NOT: _ZGV +// DEVICE-NEG-NOT: __cxa_atexit diff --git a/clang/test/SemaCUDA/Inputs/cuda-initializers.h b/clang/test/SemaCUDA/Inputs/cuda-initializers.h index b1e7a1bd48fb5..5eb9aaf20849b 100644 --- a/clang/test/SemaCUDA/Inputs/cuda-initializers.h +++ b/clang/test/SemaCUDA/Inputs/cuda-initializers.h @@ -15,6 +15,12 @@ struct EC { __device__ EC(int) {} // -- not allowed }; +// host/device, empty constructor +struct HD_EC { + int hd_ec; + __host__ __device__ HD_EC() {} // -- allowed +}; + // empty destructor struct ED { __device__ ~ED() {} // -- allowed @@ -54,6 +60,12 @@ struct NEC { __device__ NEC() { nec = 1; } }; +// host/device, non-empty constructor -- not allowed +struct HD_NEC { + int hd_nec; + __host__ __device__ HD_NEC() { hd_nec = 1; } +}; + // non-empty destructor -- not allowed struct NED { int ned; diff --git a/clang/test/SemaCUDA/device-var-init-cxx20.cu b/clang/test/SemaCUDA/device-var-init-cxx20.cu new file mode 100644 index 0000000000000..403c1452d79ac --- /dev/null +++ b/clang/test/SemaCUDA/device-var-init-cxx20.cu @@ -0,0 +1,26 @@ +// REQUIRES: nvptx-registered-target + +// C++20 cases split out of device-var-init.cu. + +// RUN: %clang_cc1 -verify %s -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++20 +// RUN: %clang_cc1 -verify %s -std=c++20 + +#include "Inputs/cuda.h" + +struct CE_NED { + int x; + constexpr CE_NED() { x = 43; } + constexpr ~CE_NED() { x = 0; } +}; + +__device__ void df_local_static_constexpr() { + static constexpr CE_NED ce; + static constinit CE_NED ci; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static CE_NED ned; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} +} + +__host__ __device__ void hd_local_static_constexpr() { + static constexpr CE_NED ce; +} diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu index a9e3557c20ebf..63f1e366c4066 100644 --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -3,7 +3,8 @@ // Make sure we don't allow dynamic initialization for device // variables, but accept empty constructors allowed by CUDA. -// RUN: %clang_cc1 -verify %s -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 %s +// RUN: %clang_cc1 -verify=expected,dev %s -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 +// RUN: %clang_cc1 -verify=expected %s -std=c++11 #ifdef __clang__ #include "Inputs/cuda.h" @@ -429,16 +430,42 @@ __device__ void df_sema() { // expected-error@-1 {{initialization is not supported for __shared__ variables}} static __constant__ T_FA_NED c_t_fa_ned; // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + + static T l_t; + static EC l_ec; + static ECD l_ecd; + static EC_I_EC l_ec_i_ec; + static CEEC l_ceec; + static CGTC l_cgtc; + static NCFS l_ncfs; + + static EC l_ec_i(3); + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static ECI l_eci; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static NEC l_nec; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static NED l_ned; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static VD l_vd; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static EC_I_EC1 l_ec_i_ec1; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static T_B_NEC l_t_b_nec; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} + static T_FA_NED l_t_fa_ned; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __device__ function}} } __host__ __device__ void hd_sema() { static int x = 42; } -inline __host__ __device__ void hd_emitted_host_only() { - static int x = 42; // no error on device because this is never codegen'ed there. +inline __host__ __device__ void hd_const_init() { + static int x = 42; // no error on device because this is constant initialized. } -void call_hd_emitted_host_only() { hd_emitted_host_only(); } +void call_hd_const_init_from_host() { hd_const_init(); } +__device__ void call_hd_const_init_from_device() { hd_const_init(); } // Verify that we also check field initializers in instantiated structs. struct NontrivialInitializer { @@ -491,6 +518,48 @@ __device__ void *ptr2 = ptr1; // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} __device__ [[gnu::constructor(101)]] void ctor() {} -// expected-error@-1 {{CUDA does not support global constructors for __device__ functions}} +// dev-error@-1 {{CUDA does not support global constructors for __device__ functions}} __device__ [[gnu::destructor(101)]] void dtor() {} -// expected-error@-1 {{CUDA does not support global destructors for __device__ functions}} +// dev-error@-1 {{CUDA does not support global destructors for __device__ functions}} + +__global__ void gf_local_static() { + static HD_NEC nec; + // expected-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __global__ function}} + static HD_EC ec; + static int i = 42; +} + +__host__ __device__ void hd_local_static() { + static HD_NEC nec; + // dev-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __host__ __device__ function}} + static int i = 42; +} + +inline __host__ __device__ void hd_local_static_host_only() { + static HD_NEC nec; +} + +void call_hd_local_static_host_only() { hd_local_static_host_only(); } + +__host__ void h_local_static() { static HD_NEC nec; } +void plain_local_static() { static HD_NEC nec; } + +__device__ int df_lambda_called_from_device() { + auto l = []() { + static HD_NEC nec; + // dev-error@-1 {{cannot use 'static' local variable requiring runtime initialization or destruction in __host__ __device__ function}} + return nec.hd_nec; + }; + return l(); + // dev-note@-1 {{called by 'df_lambda_called_from_device'}} +} + +int h_lambda_called_from_host() { + auto l = []() { + static HD_NEC nec; + return nec.hd_nec; + }; + return l(); +} + + _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
