https://github.com/pjmc-oliveira updated https://github.com/llvm/llvm-project/pull/222985
>From 861dd9d74d154596d1eea6a517e91e3af58af391 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/lib/CodeGen/CGDecl.cpp | 11 +++- clang/lib/Sema/SemaCUDA.cpp | 13 ++++- .../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.cu | 58 +++++++++++++++++-- 6 files changed, 122 insertions(+), 7 deletions(-) diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp index 1ed8989d3f627..7ddddc9f0aba6 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,8 @@ CodeGenFunction::AddInitializerToStaticVarDecl(const VarDecl &D, // be constant. GV->setConstant(false); - EmitCXXGuardedInit(D, GV, /*PerformInit*/true); + if (!SkipCUDADeviceInit) + EmitCXXGuardedInit(D, GV, /*PerformInit*/ true); } return GV; } @@ -399,7 +406,7 @@ CodeGenFunction::AddInitializerToStaticVarDecl(const VarDecl &D, emitter.finalize(GV); - if (NeedsDtor && HaveInsertPoint()) { + if (NeedsDtor && HaveInsertPoint() && !SkipCUDADeviceInit) { // We have a constant initializer, but a nontrivial destructor. We still // need to perform a guarded "initialization" in order to register the // destructor. diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 090030ea82503..6dec73dcf1cf5 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -764,8 +764,19 @@ 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. Statics in + // implicit HD functions (such as lambdas) are host variables. + CUDAVariableTarget VT = IdentifyTarget(VD); + const auto *FD = dyn_cast_or_null<FunctionDecl>(VD->getDeclContext()); + bool IsDeviceLocalStatic = + !IsSharedVar && VD->isStaticLocal() && + (VT == CVT_Device || (VT == CVT_Both && getLangOpts().CUDAIsDevice && + FD && !isImplicitHostDeviceFunction(FD))); + const Expr *Init = VD->getInit(); - if (IsDeviceOrConstantVar || IsSharedVar) { + if (IsDeviceOrConstantVar || IsSharedVar || IsDeviceLocalStatic) { if (HasAllowedCUDADeviceStaticInitializer( *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant)) return; 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.cu b/clang/test/SemaCUDA/device-var-init.cu index a9e3557c20ebf..31283a0904169 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,6 +430,31 @@ __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 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static ECI l_eci; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static NEC l_nec; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static NED l_ned; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static VD l_vd; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static EC_I_EC1 l_ec_i_ec1; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static T_B_NEC l_t_b_nec; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static T_FA_NED l_t_fa_ned; + // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} } __host__ __device__ void hd_sema() { @@ -436,7 +462,7 @@ __host__ __device__ void hd_sema() { } inline __host__ __device__ void hd_emitted_host_only() { - static int x = 42; // no error on device because this is never codegen'ed there. + static int x = 42; // no error on device because this is constant initialized. } void call_hd_emitted_host_only() { hd_emitted_host_only(); } @@ -491,6 +517,30 @@ __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 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static HD_EC ec; + static int i = 42; +} + +__host__ __device__ void hd_local_static() { + static HD_NEC nec; + // dev-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + static int i = 42; +} + +inline __host__ __device__ void hd_local_static_host_only() { + static HD_NEC nec; + // dev-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} +} + +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; } + _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
