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 1/2] [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 >From 45309a263455eb20f596cb2a364572ab5879f9bc Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Thu, 15 Oct 2026 09:03:15 +0200 Subject: [PATCH 2/2] move to CUDA --- .../CIR/CodeGenCUDA/device-var-empty-ctor.cu | 40 +++++++++++++++++++ .../CIR/CodeGenHIP/device-var-empty-ctor.hip | 33 --------------- 2 files changed, 40 insertions(+), 33 deletions(-) create mode 100644 clang/test/CIR/CodeGenCUDA/device-var-empty-ctor.cu delete mode 100644 clang/test/CIR/CodeGenHIP/device-var-empty-ctor.hip 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 00000000000000..68481a1312420f --- /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 diff --git a/clang/test/CIR/CodeGenHIP/device-var-empty-ctor.hip b/clang/test/CIR/CodeGenHIP/device-var-empty-ctor.hip deleted file mode 100644 index 07133828408a96..00000000000000 --- a/clang/test/CIR/CodeGenHIP/device-var-empty-ctor.hip +++ /dev/null @@ -1,33 +0,0 @@ -// 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
