[clang] [CIR][CUDA][HIP] Skip initializer for static __shared__ variables (PR #231172)

[Arseniy Obolenskiy via cfe-commits]

https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/231172

>From 0cfaf59db3c5d9ae8a28878338d13864e9117da7 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 12 Oct 2026 15:51:39 +0200
Subject: [PATCH 1/2] [CIR][CUDA][HIP] Skip initializer for static __shared__
 variables

---
 clang/lib/CIR/CodeGen/CIRGenDecl.cpp         |  7 ++++---
 clang/test/CIR/CodeGenCUDA/address-spaces.cu | 10 ++++++++++
 2 files changed, 14 insertions(+), 3 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index f841cecbe0a85..883dadebe7e86 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -761,10 +761,11 @@ void CIRGenFunction::emitStaticVarDecl(const VarDecl &d,
 
   cir::GlobalOp var = globalOp;
 
-  assert(!cir::MissingFeatures::cudaSupport());
-
+  // Sema ensures any initializer of a CUDA __shared__ variable is a no-op.
+  bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+                         d.hasAttr<CUDASharedAttr>();
   // If this value has an initializer, emit it.
-  if (d.getInit())
+  if (d.getInit() && !isCudaSharedVar)
     var = addInitializerToStaticVarDecl(d, var, getAddrOp);
 
   var.setAlignment(alignment.getAsAlign().value());
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index 95bdcc0aa8e99..9a544dd7f627b 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -45,6 +45,8 @@
 
 // Verifies CIR emits correct address spaces for CUDA globals.
 
+// CIR-DEVICE: cir.global "private" internal dso_local target_address_space(3) @_ZZ8fnStructvE1s = #cir.undef : !rec_S
+// LLVM-DEVICE: @_ZZ8fnStructvE1s = internal addrspace(3) global %struct.S undef
 // CIR-DEVICE: cir.global "private" internal dso_local target_address_space(3) @_ZZ2fnvE1j = #cir.undef
 // LLVM-DEVICE: @_ZZ2fnvE1j = internal addrspace(3) global i32 undef
 
@@ -110,6 +112,7 @@ __device__ int ext_device_var_def = 1;
 // CIR-DEVICE: cir.global constant external target_address_space(4) @ext_constant_var_def = #cir.int<2>
 // LLVM-DEVICE: @ext_constant_var_def = addrspace(4) externally_initialized constant i32 2, align 4
 // OGCG-DEVICE: @ext_constant_var_def = addrspace(4) externally_initialized constant i32 2, align 4
+// OGCG-DEVICE: @_ZZ8fnStructvE1s = internal addrspace(3) global %struct.S undef
 // CIR-HOST: cir.global "private" internal {{.*}} @ext_constant_var_def = #cir.undef : {{.*}} {{.*}}{cu.var_registration = #cir.cu.var_registration<ext_constant_var_def, Variable, constant>}
 // LLVM-HOST: @ext_constant_var_def = internal global i32 undef, align 4
 // OGCG-HOST: @ext_constant_var_def = internal global i32
@@ -181,3 +184,10 @@ __global__ void fn() {
 // OGCG-DEVICE:   %[[VAL:.*]] = load i32, ptr %[[ALLOCA]], align 4
 // OGCG-DEVICE:   store i32 %[[VAL]], ptr addrspacecast (ptr addrspace(3) @_ZZ2fnvE1j to ptr), align 4
 // OGCG-DEVICE:   ret void
+
+struct S { int a; float b[4]; };
+
+__global__ void fnStruct() {
+  __shared__ S s;
+  s.a = 1;
+}

>From 367495ffbcf9fc8473325b302de16efe94ce57be Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 13 Oct 2026 06:06:04 +0200
Subject: [PATCH 2/2] Add CHECK-NOT

---
 clang/test/CIR/CodeGenCUDA/address-spaces.cu | 11 +++++++++--
 1 file changed, 9 insertions(+), 2 deletions(-)

diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
index 9a544dd7f627b..d182d76bc0d9d 100644
--- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -185,8 +185,15 @@ __global__ void fn() {
 // OGCG-DEVICE:   store i32 %[[VAL]], ptr addrspacecast (ptr addrspace(3) @_ZZ2fnvE1j to ptr), align 4
 // OGCG-DEVICE:   ret void
 
-struct S { int a; float b[4]; };
-
+struct S {
+  int a;
+  float b[4];
+  __device__ S() {}
+};
+
+// CIR-DEVICE-NOT: @_ZN1SC1Ev
+// LLVM-DEVICE-NOT: @_ZN1SC1Ev
+// OGCG-DEVICE-NOT: @_ZN1SC1Ev
 __global__ void fnStruct() {
   __shared__ S s;
   s.a = 1;