[clang] 187d6dd - [CIR][CUDA][HIP] Skip global init for device variables (#231971)

[via cfe-commits]

Author: Arseniy Obolenskiy
Date: 2026-10-15T14:02:56+02:00
New Revision: 187d6ddf58430f6e78d65907af22791727ae30e3

URL: https://github.com/llvm/llvm-project/commit/187d6ddf58430f6e78d65907af22791727ae30e3
DIFF: https://github.com/llvm/llvm-project/commit/187d6ddf58430f6e78d65907af22791727ae30e3.diff

LOG: [CIR][CUDA][HIP] Skip global init for device variables (#231971)

Sema allows only empty or constant initializers and empty dtors on
__device__, __constant__ and __shared__ variables, so they need no
dynamic init

CIR emitted a dtor registration that failed to verify

Added: 
    clang/test/CIR/CodeGenCUDA/device-var-empty-ctor.cu

Modified: 
    clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp
index 86c48c1373aa2..7ee18d20fc410 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp
@@ -97,7 +97,12 @@ void CIRGenModule::setGlobalTlsReferences(const VarDecl &vd,
 void CIRGenModule::emitCXXGlobalVarDeclInitFunc(const VarDecl *vd,
                                                 cir::GlobalOp addr,
                                                 bool performInit) {
-  assert(!cir::MissingFeatures::cudaSupport());
+  // Sema has already checked that the initializers of device-side variables
+  // are empty, so they can be ignored.
+  if (getLangOpts().CUDAIsDevice && !getLangOpts().GPUAllowDeviceInit &&
+      (vd->hasAttr<CUDADeviceAttr>() || vd->hasAttr<CUDAConstantAttr>() ||
+       vd->hasAttr<CUDASharedAttr>()))
+    return;
 
   // Classic CodeGen dispatches guarded initialization through
   // CGCXXABI::EmitGuardedInit, which the Microsoft ABI overrides with a

diff  --git a/clang/test/CIR/CodeGenCUDA/device-var-empty-ctor.cu b/clang/test/CIR/CodeGenCUDA/device-var-empty-ctor.cu
new file mode 100644
index 0000000000000..68481a1312420
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/device-var-empty-ctor.cu
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/Inputs %s -o %t-spirv-cir.ll
+// RUN: FileCheck --check-prefixes=LLVM,SPIRV %s --input-file=%t-spirv-cir.ll
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/Inputs %s -o %t-spirv.ll
+// RUN: FileCheck --check-prefixes=LLVM,SPIRV %s --input-file=%t-spirv.ll
+
+// RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/Inputs %s -o %t-amdgpu-cir.ll
+// RUN: FileCheck --check-prefixes=LLVM,AS4 %s --input-file=%t-amdgpu-cir.ll
+
+// RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -x hip \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/Inputs %s -o %t-amdgpu.ll
+// RUN: FileCheck --check-prefixes=LLVM,AS4 %s --input-file=%t-amdgpu.ll
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/Inputs %s -o %t-nvptx-cir.ll
+// RUN: FileCheck --check-prefixes=LLVM,AS4 %s --input-file=%t-nvptx-cir.ll
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/Inputs %s -o %t-nvptx.ll
+// RUN: FileCheck --check-prefixes=LLVM,AS4 %s --input-file=%t-nvptx.ll
+
+#include "cuda.h"
+
+struct S {
+  int x;
+  __device__ S() {}
+  __device__ ~S() {}
+};
+
+__device__ S d;
+__constant__ S c;
+
+// LLVM: @d = addrspace(1) externally_initialized global %struct.S zeroinitializer, align 4
+// SPIRV: @c = addrspace(1) externally_initialized constant %struct.S zeroinitializer, align 4
+// AS4: @c = addrspace(4) externally_initialized constant %struct.S zeroinitializer, align 4
+// LLVM-NOT: llvm.global_ctors
+// LLVM-NOT: llvm.global_dtors