diff options
author | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2022-04-07 22:57:56 -0400 |
---|---|---|
committer | Tom Stellard <tstellar@redhat.com> | 2022-05-24 15:02:58 -0700 |
commit | e6de9ed37308e46560243229dd78e84542f37ead (patch) | |
tree | 3150a97301bd76708dc4e767001e707e0af3ed42 | |
parent | fecfc8394484be0ff686e2c936eb494ce6a19645 (diff) | |
download | llvm-e6de9ed37308e46560243229dd78e84542f37ead.tar.gz |
[CUDA][HIP] Externalize kernels in anonymous name space
kernels in anonymous name space needs to have unique name
to avoid duplicate symbols.
Fixes: https://github.com/llvm/llvm-project/issues/54560
Reviewed by: Artem Belevich
Differential Revision: https://reviews.llvm.org/D123353
(cherry picked from commit 4ea1d435099f992cc16127619b0feb64e070630d)
-rw-r--r-- | clang/include/clang/AST/ASTContext.h | 4 | ||||
-rw-r--r-- | clang/lib/AST/ASTContext.cpp | 10 | ||||
-rw-r--r-- | clang/lib/CodeGen/CGCUDANV.cpp | 2 | ||||
-rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 11 | ||||
-rw-r--r-- | clang/lib/CodeGen/CodeGenModule.h | 7 | ||||
-rw-r--r-- | clang/test/CodeGenCUDA/kernel-in-anon-ns.cu | 24 |
6 files changed, 43 insertions, 15 deletions
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 63c11e237d6c..1bd5d7a6c1d7 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -3279,10 +3279,10 @@ public: /// Return a new OMPTraitInfo object owned by this context. OMPTraitInfo &getNewOMPTraitInfo(); - /// Whether a C++ static variable may be externalized. + /// Whether a C++ static variable or CUDA/HIP kernel may be externalized. bool mayExternalizeStaticVar(const Decl *D) const; - /// Whether a C++ static variable should be externalized. + /// Whether a C++ static variable or CUDA/HIP kernel should be externalized. bool shouldExternalizeStaticVar(const Decl *D) const; StringRef getCUIDHash() const; diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index c873ff0515e1..b554cf833b44 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12263,14 +12263,16 @@ bool ASTContext::mayExternalizeStaticVar(const Decl *D) const { (D->hasAttr<CUDAConstantAttr>() && !D->getAttr<CUDAConstantAttr>()->isImplicit()); // CUDA/HIP: static managed variables need to be externalized since it is - // a declaration in IR, therefore cannot have internal linkage. - return IsStaticVar && - (D->hasAttr<HIPManagedAttr>() || IsExplicitDeviceVar); + // a declaration in IR, therefore cannot have internal linkage. Kernels in + // anonymous name space needs to be externalized to avoid duplicate symbols. + return (IsStaticVar && + (D->hasAttr<HIPManagedAttr>() || IsExplicitDeviceVar)) || + (D->hasAttr<CUDAGlobalAttr>() && D->isInAnonymousNamespace()); } bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const { return mayExternalizeStaticVar(D) && - (D->hasAttr<HIPManagedAttr>() || + (D->hasAttr<HIPManagedAttr>() || D->hasAttr<CUDAGlobalAttr>() || CUDADeviceVarODRUsedByHost.count(cast<VarDecl>(D))); } diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index c4e3f7f54f4f..414e61f25fb3 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -287,7 +287,7 @@ std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) { SmallString<256> Buffer; llvm::raw_svector_ostream Out(Buffer); Out << DeviceSideName; - CGM.printPostfixForExternalizedStaticVar(Out); + CGM.printPostfixForExternalizedDecl(Out, ND); DeviceSideName = std::string(Out.str()); } return DeviceSideName; diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 29806b65e984..65b9f4e40dc1 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1367,7 +1367,7 @@ static std::string getMangledNameImpl(CodeGenModule &CGM, GlobalDecl GD, if (CGM.getContext().shouldExternalizeStaticVar(ND) && CGM.getLangOpts().GPURelocatableDeviceCode && CGM.getLangOpts().CUDAIsDevice && !CGM.getLangOpts().CUID.empty()) - CGM.printPostfixForExternalizedStaticVar(Out); + CGM.printPostfixForExternalizedDecl(Out, ND); return std::string(Out.str()); } @@ -1455,7 +1455,7 @@ StringRef CodeGenModule::getMangledName(GlobalDecl GD) { // directly between host- and device-compilations, the host- and // device-mangling in host compilation could help catching certain ones. assert(!isa<FunctionDecl>(ND) || !ND->hasAttr<CUDAGlobalAttr>() || - getLangOpts().CUDAIsDevice || + getContext().shouldExternalizeStaticVar(ND) || getLangOpts().CUDAIsDevice || (getContext().getAuxTargetInfo() && (getContext().getAuxTargetInfo()->getCXXABI() != getContext().getTargetInfo().getCXXABI())) || @@ -6645,7 +6645,8 @@ bool CodeGenModule::stopAutoInit() { return false; } -void CodeGenModule::printPostfixForExternalizedStaticVar( - llvm::raw_ostream &OS) const { - OS << "__static__" << getContext().getCUIDHash(); +void CodeGenModule::printPostfixForExternalizedDecl(llvm::raw_ostream &OS, + const Decl *D) const { + OS << (isa<VarDecl>(D) ? "__static__" : ".anon.") + << getContext().getCUIDHash(); } diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index 1fcd5d4d808a..a8a63c8da57f 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1447,9 +1447,10 @@ public: TBAAAccessInfo *TBAAInfo = nullptr); bool stopAutoInit(); - /// Print the postfix for externalized static variable for single source - /// offloading languages CUDA and HIP. - void printPostfixForExternalizedStaticVar(llvm::raw_ostream &OS) const; + /// Print the postfix for externalized static variable or kernels for single + /// source offloading languages CUDA and HIP. + void printPostfixForExternalizedDecl(llvm::raw_ostream &OS, + const Decl *D) const; private: llvm::Constant *GetOrCreateLLVMFunction( diff --git a/clang/test/CodeGenCUDA/kernel-in-anon-ns.cu b/clang/test/CodeGenCUDA/kernel-in-anon-ns.cu new file mode 100644 index 000000000000..4243cec796a8 --- /dev/null +++ b/clang/test/CodeGenCUDA/kernel-in-anon-ns.cu @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \ +// RUN: -aux-triple x86_64-unknown-linux-gnu -std=c++11 -fgpu-rdc \ +// RUN: -emit-llvm -o - -x hip %s > %t.dev + +// RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \ +// RUN: -aux-triple amdgcn-amd-amdhsa -std=c++11 -fgpu-rdc \ +// RUN: -emit-llvm -o - -x hip %s > %t.host + +// RUN: cat %t.dev %t.host | FileCheck %s + +#include "Inputs/cuda.h" + +// CHECK: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv\.anon\.b04fd23c98500190]]( +// CHECK: @[[STR:.*]] = {{.*}} c"[[KERN]]\00" +// CHECK: call i32 @__hipRegisterFunction({{.*}}@[[STR]] + +namespace { +__global__ void kernel() { +} +} + +void test() { + kernel<<<1, 1>>>(); +} |