[clang] [CUDA] Treat function-scope statics in device code as device variables (PR #222985)
Pedro Oliveira via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 11 09:46:11 PDT 2026
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 <poliveira at nvidia.com>
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 at -1 {{initialization is not supported for __shared__ variables}}
static __constant__ T_FA_NED c_t_fa_ned;
// expected-error at -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 at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ static ECI l_eci;
+ // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ static NEC l_nec;
+ // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ static NED l_ned;
+ // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ static VD l_vd;
+ // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ static EC_I_EC1 l_ec_i_ec1;
+ // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ static T_B_NEC l_t_b_nec;
+ // expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+ static T_FA_NED l_t_fa_ned;
+ // expected-error at -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 at -1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
__device__ [[gnu::constructor(101)]] void ctor() {}
-// expected-error at -1 {{CUDA does not support global constructors for __device__ functions}}
+// dev-error at -1 {{CUDA does not support global constructors for __device__ functions}}
__device__ [[gnu::destructor(101)]] void dtor() {}
-// expected-error at -1 {{CUDA does not support global destructors for __device__ functions}}
+// dev-error at -1 {{CUDA does not support global destructors for __device__ functions}}
+
+__global__ void gf_local_static() {
+ static HD_NEC nec;
+ // expected-error at -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 at -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 at -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; }
+
More information about the cfe-commits
mailing list