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

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Arseniy Obolenskiy (aobolensk)

<details>
<summary>Changes</summary>

Align behavior with classic codegen

---
Full diff: https://github.com/llvm/llvm-project/pull/231172.diff


2 Files Affected:

- (modified) clang/lib/CIR/CodeGen/CIRGenDecl.cpp (+4-3) 
- (modified) clang/test/CIR/CodeGenCUDA/address-spaces.cu (+10) 


``````````diff
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;
+}

``````````

</details>


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