Skip to content

[CUDA][HIP] Fix template static member #98544

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jul 12, 2024
Merged

[CUDA][HIP] Fix template static member #98544

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
Expand Down
50 changes: 50 additions & 0 deletions clang/test/CodeGenCUDA/template-class-static-member.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Loading