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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
