[clang] [CIR][HIP] Skip global init for device variables (PR #231971)

[Arseniy Obolenskiy via cfe-commits]

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

>From ddc9b679e66f28abe37f7f8c45252bea8fb69a88 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Wed, 14 Oct 2026 17:53:26 +0200
Subject: [PATCH] [CIR][HIP] Skip global init for device variables

Without the skip, CIR registers a global ctor for __device__, __constant__ and __shared__ variables, which fails on SPIR-V
---
 clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp       |  7 +++-
 .../CIR/CodeGenHIP/device-var-empty-ctor.hip  | 33 +++++++++++++++++++
 2 files changed, 39 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/CIR/CodeGenHIP/device-var-empty-ctor.hip

diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclCXX.cpp
index 86c48c1373aa27..7ee18d20fc4103 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/CodeGenHIP/device-var-empty-ctor.hip b/clang/test/CIR/CodeGenHIP/device-var-empty-ctor.hip
new file mode 100644
index 00000000000000..07133828408a96
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/device-var-empty-ctor.hip
@@ -0,0 +1,33 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t-cir.ll
+// RUN: FileCheck --check-prefixes=LLVM,SPIRV %s --input-file=%t-cir.ll
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t.ll
+// RUN: FileCheck --check-prefixes=LLVM,SPIRV %s --input-file=%t.ll
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t-amdgcn-cir.ll
+// RUN: FileCheck --check-prefixes=LLVM,AMDGCN %s --input-file=%t-amdgcn-cir.ll
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip \
+// RUN:            -fcuda-is-device -emit-llvm -I%S/../CodeGenCUDA/Inputs %s -o %t-amdgcn.ll
+// RUN: FileCheck --check-prefixes=LLVM,AMDGCN %s --input-file=%t-amdgcn.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
+// AMDGCN: @c = addrspace(4) externally_initialized constant %struct.S zeroinitializer, align 4
+// LLVM-NOT: llvm.global_ctors
+// LLVM-NOT: llvm.global_dtors