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