https://github.com/Fznamznon updated https://github.com/llvm/llvm-project/pull/219536
>From 962189f5472df0a952332e9cc8fc899e967fef01 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Fri, 28 Aug 2026 12:10:08 -0500 Subject: [PATCH 1/2] [CIR][CUDA/HIP] Fix template static member filtering Check host/device attributes before emitting static member of template instantiation. This is basically a CIR version of 77fd30f7ce0795b4bbc22e65b3ff42856839d708 . --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 3 +- .../template-class-static-member.cu | 79 +++++++++++++++++++ 2 files changed, 81 insertions(+), 1 deletion(-) create mode 100644 clang/test/CIR/CodeGenCUDA/template-class-static-member.cu diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index cc4ce9a786327..c5bb68e03e353 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -733,7 +733,8 @@ void CIRGenModule::addGlobalDtor(cir::FuncOp dtor, void CIRGenModule::handleCXXStaticMemberVarInstantiation(VarDecl *vd) { VarDecl::DefinitionKind dk = vd->isThisDeclarationADefinition(); - if (dk == VarDecl::Definition && vd->hasAttr<DLLImportAttr>()) + if ((dk == VarDecl::Definition && vd->hasAttr<DLLImportAttr>()) || + (langOpts.CUDA && !shouldEmitCUDAGlobalVar(vd))) return; TemplateSpecializationKind tsk = vd->getTemplateSpecializationKind(); diff --git a/clang/test/CIR/CodeGenCUDA/template-class-static-member.cu b/clang/test/CIR/CodeGenCUDA/template-class-static-member.cu new file mode 100644 index 0000000000000..cd639ae34debd --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/template-class-static-member.cu @@ -0,0 +1,79 @@ +// Based on clang/test/CodeGenCUDA/template-class-static-member.cu. + +// CIR tests + +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \ +// RUN: -emit-cir -o - -x hip %s | FileCheck -check-prefix=CIR-DEV %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ +// RUN: -emit-cir -o - -x hip %s | FileCheck -check-prefix=CIR-HOST %s + +// Negative tests. + +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=CIR-DEV-NEG %s + +// LLVM-IR tests + +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV %s + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=HOST %s + +// Negative tests. + +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device \ +// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefix=DEV-NEG %s + +#include "Inputs/cuda.h" + +template <class T> +class A { + static int h_member; + __device__ static int d_member; + __constant__ static int c_member; + __managed__ static int m_member; + const static int const_member = 0; +}; + +template <class T> +int A<T>::h_member; + +template <class T> +__device__ int A<T>::d_member; + +template <class T> +__constant__ int A<T>::c_member; + +template <class T> +__managed__ int A<T>::m_member; + +template <class T> +const int A<T>::const_member; + +template class A<int>; + +//CIR-DEV-DAG: cir.global "private" internal comdat dso_local target_address_space(1) @_ZN1AIiE8d_memberE +//CIR-DEV-DAG: cir.global "private" constant internal comdat dso_local target_address_space(4) @_ZN1AIiE8c_memberE +//CIR-DEV-DAG: cir.global "private" internal comdat dso_local target_address_space(1) @_ZN1AIiE8m_memberE +//CIR-DEV-DAG: cir.global "private" constant internal comdat dso_local target_address_space(4) @_ZN1AIiE12const_memberE +//CIR-DEV-NEG-NOT: @_ZN1AIiE8h_memberE + +//DEV-DAG: @_ZN1AIiE8d_memberE = internal addrspace(1) global i32 0, comdat, align 4 +//DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) constant i32 0, comdat, align 4 +//DEV-DAG: @_ZN1AIiE8m_memberE = internal addrspace(1) externally_initialized global ptr addrspace(1) null +//DEV-DAG: @_ZN1AIiE12const_memberE = internal addrspace(4) constant i32 0, comdat, align 4 +//DEV-NEG-NOT: @_ZN1AIiE8h_memberE + +// CIR-HOST-DAG: cir.global weak_odr comdat @_ZN1AIiE8h_memberE = #cir.int<0> +// CIR-HOST-DAG: cir.global "private" internal comdat dso_local @_ZN1AIiE8d_memberE = #cir.undef +// CIR-HOST-DAG: cir.global "private" internal comdat dso_local @_ZN1AIiE8c_memberE = #cir.undef +// CIR-HOST-DAG: cir.global "private" internal comdat dso_local @_ZN1AIiE8m_memberE = #cir.int<0> +// CIR-HOST-DAG: cir.global constant weak_odr comdat @_ZN1AIiE12const_memberE = #cir.int<0> + +//HOST-DAG: @_ZN1AIiE8h_memberE = weak_odr global i32 0, comdat, align 4 +//HOST-DAG: @_ZN1AIiE8d_memberE = internal global i32 undef, comdat, align 4 +//HOST-DAG: @_ZN1AIiE8c_memberE = internal global i32 undef, comdat, align 4 +//HOST-DAG: @_ZN1AIiE8m_memberE = internal externally_initialized global ptr null +//HOST-DAG: @_ZN1AIiE12const_memberE = weak_odr constant i32 0, comdat, align 4 >From 9a2d88a687b5c92f5f656dee0c2e8ae70224d779 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Mon, 31 Aug 2026 05:18:40 -0500 Subject: [PATCH 2/2] Check initializers of device-side variables, make the test a bit nicer --- .../template-class-static-member.cu | 32 +++++++++---------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/clang/test/CIR/CodeGenCUDA/template-class-static-member.cu b/clang/test/CIR/CodeGenCUDA/template-class-static-member.cu index cd639ae34debd..fccfddc2b8018 100644 --- a/clang/test/CIR/CodeGenCUDA/template-class-static-member.cu +++ b/clang/test/CIR/CodeGenCUDA/template-class-static-member.cu @@ -54,17 +54,17 @@ const int A<T>::const_member; template class A<int>; -//CIR-DEV-DAG: cir.global "private" internal comdat dso_local target_address_space(1) @_ZN1AIiE8d_memberE -//CIR-DEV-DAG: cir.global "private" constant internal comdat dso_local target_address_space(4) @_ZN1AIiE8c_memberE -//CIR-DEV-DAG: cir.global "private" internal comdat dso_local target_address_space(1) @_ZN1AIiE8m_memberE -//CIR-DEV-DAG: cir.global "private" constant internal comdat dso_local target_address_space(4) @_ZN1AIiE12const_memberE -//CIR-DEV-NEG-NOT: @_ZN1AIiE8h_memberE - -//DEV-DAG: @_ZN1AIiE8d_memberE = internal addrspace(1) global i32 0, comdat, align 4 -//DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) constant i32 0, comdat, align 4 -//DEV-DAG: @_ZN1AIiE8m_memberE = internal addrspace(1) externally_initialized global ptr addrspace(1) null -//DEV-DAG: @_ZN1AIiE12const_memberE = internal addrspace(4) constant i32 0, comdat, align 4 -//DEV-NEG-NOT: @_ZN1AIiE8h_memberE +// CIR-DEV-DAG: cir.global "private" internal comdat dso_local target_address_space(1) @_ZN1AIiE8d_memberE = #cir.int<0> +// CIR-DEV-DAG: cir.global "private" constant internal comdat dso_local target_address_space(4) @_ZN1AIiE8c_memberE = #cir.int<0> +// CIR-DEV-DAG: cir.global "private" internal comdat dso_local target_address_space(1) @_ZN1AIiE8m_memberE = #cir.int<0> +// CIR-DEV-DAG: cir.global "private" constant internal comdat dso_local target_address_space(4) @_ZN1AIiE12const_memberE = #cir.int<0> +// CIR-DEV-NEG-NOT: @_ZN1AIiE8h_memberE + +// DEV-DAG: @_ZN1AIiE8d_memberE = internal addrspace(1) global i32 0, comdat, align 4 +// DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) constant i32 0, comdat, align 4 +// DEV-DAG: @_ZN1AIiE8m_memberE = internal addrspace(1) externally_initialized global ptr addrspace(1) null +// DEV-DAG: @_ZN1AIiE12const_memberE = internal addrspace(4) constant i32 0, comdat, align 4 +// DEV-NEG-NOT: @_ZN1AIiE8h_memberE // CIR-HOST-DAG: cir.global weak_odr comdat @_ZN1AIiE8h_memberE = #cir.int<0> // CIR-HOST-DAG: cir.global "private" internal comdat dso_local @_ZN1AIiE8d_memberE = #cir.undef @@ -72,8 +72,8 @@ template class A<int>; // CIR-HOST-DAG: cir.global "private" internal comdat dso_local @_ZN1AIiE8m_memberE = #cir.int<0> // CIR-HOST-DAG: cir.global constant weak_odr comdat @_ZN1AIiE12const_memberE = #cir.int<0> -//HOST-DAG: @_ZN1AIiE8h_memberE = weak_odr global i32 0, comdat, align 4 -//HOST-DAG: @_ZN1AIiE8d_memberE = internal global i32 undef, comdat, align 4 -//HOST-DAG: @_ZN1AIiE8c_memberE = internal global i32 undef, comdat, align 4 -//HOST-DAG: @_ZN1AIiE8m_memberE = internal externally_initialized global ptr null -//HOST-DAG: @_ZN1AIiE12const_memberE = weak_odr constant i32 0, comdat, align 4 +// HOST-DAG: @_ZN1AIiE8h_memberE = weak_odr global i32 0, comdat, align 4 +// HOST-DAG: @_ZN1AIiE8d_memberE = internal global i32 undef, comdat, align 4 +// HOST-DAG: @_ZN1AIiE8c_memberE = internal global i32 undef, comdat, align 4 +// HOST-DAG: @_ZN1AIiE8m_memberE = internal externally_initialized global ptr null +// HOST-DAG: @_ZN1AIiE12const_memberE = weak_odr constant i32 0, comdat, align 4 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
