diff options
| author | Yaxun (Sam) Liu <yaxun.liu@amd.com> | 2022-04-21 11:44:11 -0400 |
|---|---|---|
| committer | Tom Stellard <tstellar@redhat.com> | 2022-05-24 15:02:58 -0700 |
| commit | 29f1039a7285a5c3a9c353d054140bf2556d4c4d (patch) | |
| tree | 2aceb3755416852906643b4734fa9b9c0940203e /clang/lib/AST/ASTContext.cpp | |
| parent | e6de9ed37308e46560243229dd78e84542f37ead (diff) | |
| download | llvmorg-14.0.4.tar.gz | |
[CUDA][HIP] Externalize kernels with internal linkagellvmorg-14.0.4
This patch is a continuation of https://reviews.llvm.org/D123353.
Not only kernels in anonymous namespace, but also template
kernels with template arguments in anonymous namespace
need to be externalized.
To be more generic, this patch checks the linkage of a kernel
assuming the kernel does not have __global__ attribute. If
the linkage is internal then clang will externalize it.
This patch also fixes the postfix for externalized symbol
since nvptx does not allow '.' in symbol name.
Reviewed by: Artem Belevich
Differential Revision: https://reviews.llvm.org/D124189
Fixes: https://github.com/llvm/llvm-project/issues/54560
(cherry picked from commit 04fb81674ed7981397ffe70fe6a07b7168f6fe2f)
Diffstat (limited to 'clang/lib/AST/ASTContext.cpp')
| -rw-r--r-- | clang/lib/AST/ASTContext.cpp | 4 |
1 files changed, 3 insertions, 1 deletions
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index b554cf833b44..e4b3827b8714 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12267,7 +12267,9 @@ bool ASTContext::mayExternalizeStaticVar(const Decl *D) const { // anonymous name space needs to be externalized to avoid duplicate symbols. return (IsStaticVar && (D->hasAttr<HIPManagedAttr>() || IsExplicitDeviceVar)) || - (D->hasAttr<CUDAGlobalAttr>() && D->isInAnonymousNamespace()); + (D->hasAttr<CUDAGlobalAttr>() && + basicGVALinkageForFunction(*this, cast<FunctionDecl>(D)) == + GVA_Internal); } bool ASTContext::shouldExternalizeStaticVar(const Decl *D) const { |