diff options
| author | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2024-07-11 21:54:29 -0400 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2024-07-11 21:54:29 -0400 |
| commit | 3442d12b865a47fd6e604794f0b962a5956663e1 (patch) | |
| tree | 9c8ccde94215acc3eff7216cb9d6ab6980ceeb11 | |
| parent | fa11d1781b0826b9c64b65e6536fd9680a8f0496 (diff) | |
| download | llvm-users/yxsamliu/should-emit-var.zip llvm-users/yxsamliu/should-emit-var.tar.gz llvm-users/yxsamliu/should-emit-var.tar.bz2 | |
[CUDA][HIP] Fix template static member (#98544)users/yxsamliu/should-emit-var
Should check host/device attributes before emitting static member
of template instantiation.
Fixes: https://github.com/llvm/llvm-project/issues/98151
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 3 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/template-class-static-member.cu | 50 |
2 files changed, 52 insertions, 1 deletions
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 6c10b4a..599e206 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -5935,7 +5935,8 @@ static void ReplaceUsesOfNonProtoTypeWithRealFunction(llvm::GlobalValue *Old, void CodeGenModule::HandleCXXStaticMemberVarInstantiation(VarDecl *VD) { auto 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/CodeGenCUDA/template-class-static-member.cu b/clang/test/CodeGenCUDA/template-class-static-member.cu new file mode 100644 index 0000000..d790d2d --- /dev/null +++ b/clang/test/CodeGenCUDA/template-class-static-member.cu @@ -0,0 +1,50 @@ +// RUN: %clang_cc1 -triple amdgcn-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 amdgcn-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>; + +//DEV-DAG: @_ZN1AIiE8d_memberE = internal addrspace(1) global i32 0, comdat, align 4 +//DEV-DAG: @_ZN1AIiE8c_memberE = internal addrspace(4) global 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 + +//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 |