Skip to content

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

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.

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

Conversation

yxsamliu
Copy link
Collaborator

Should check host/device attributes before emitting static member of template instantiation.

Fixes: #98151

Should check host/device attributes before emitting static member
of template instantiation.

Fixes: llvm#98151
@llvmbot llvmbot added clang Clang issues not falling into any other category clang:codegen IR generation bugs: mangling, exceptions, etc. labels Jul 12, 2024
@llvmbot
Copy link
Member

llvmbot commented Jul 12, 2024

@llvm/pr-subscribers-clang-codegen

Author: Yaxun (Sam) Liu (yxsamliu)

Changes

Should check host/device attributes before emitting static member of template instantiation.

Fixes: #98151


Full diff: https://github.com/llvm/llvm-project/pull/98580.diff

2 Files Affected:

  • (modified) clang/lib/CodeGen/CodeGenModule.cpp (+2-1)
  • (added) clang/test/CodeGenCUDA/template-class-static-member.cu (+50)
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 6c10b4a2edef8..599e20634bf72 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 0000000000000..d790d2dea66ba
--- /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

@llvmbot
Copy link
Member

llvmbot commented Jul 12, 2024

@llvm/pr-subscribers-clang

Author: Yaxun (Sam) Liu (yxsamliu)

Changes

Should check host/device attributes before emitting static member of template instantiation.

Fixes: #98151


Full diff: https://github.com/llvm/llvm-project/pull/98580.diff

2 Files Affected:

  • (modified) clang/lib/CodeGen/CodeGenModule.cpp (+2-1)
  • (added) clang/test/CodeGenCUDA/template-class-static-member.cu (+50)
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 6c10b4a2edef8..599e20634bf72 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 0000000000000..d790d2dea66ba
--- /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

@yxsamliu yxsamliu requested a review from Artem-B July 12, 2024 02:09
@yxsamliu
Copy link
Collaborator Author

sorry for the trouble. It is the same change but rebased to main branch.

@yxsamliu yxsamliu merged commit 77fd30f into llvm:main Jul 12, 2024
10 checks passed
aaryanshukla pushed a commit to aaryanshukla/llvm-project that referenced this pull request Jul 14, 2024
Should check host/device attributes before emitting static member of
template instantiation.

Fixes: llvm#98151
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
clang:codegen IR generation bugs: mangling, exceptions, etc. clang Clang issues not falling into any other category
Projects
None yet
Development

Successfully merging this pull request may close these issues.

CUDA: Host-side static class member leaks into PTX as extern global
2 participants