https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/231172
Align behavior with classic codegen >From 0cfaf59db3c5d9ae8a28878338d13864e9117da7 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 12 Oct 2026 15:51:39 +0200 Subject: [PATCH] [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; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
