diff options
author | Artem Belevich <tra@google.com> | 2016-05-02 20:30:03 +0000 |
---|---|---|
committer | Artem Belevich <tra@google.com> | 2016-05-02 20:30:03 +0000 |
commit | aba87a7ebcce24750d5d6845684c81345af0707b (patch) | |
tree | 0b9cc5ddbed8a75b1a4e00b52aea9d0b4f2a99c7 /lib | |
parent | c1e7ade641b996676baef60b3e46dc359589e191 (diff) | |
download | clang-aba87a7ebcce24750d5d6845684c81345af0707b.tar.gz |
[CUDA] Make sure device-side __global__ functions are always visible.
__global__ functions are a special case in CUDA.
Even when the symbol would normally not be externally
visible according to C++ rules, they still must be visible
in CUDA GPU object so host-side stub can launch them.
Differential Revision: http://reviews.llvm.org/D19748
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@268299 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'lib')
-rw-r--r-- | lib/AST/ASTContext.cpp | 19 |
1 files changed, 13 insertions, 6 deletions
diff --git a/lib/AST/ASTContext.cpp b/lib/AST/ASTContext.cpp index 4f7f8ee8b8..7a09b0b93b 100644 --- a/lib/AST/ASTContext.cpp +++ b/lib/AST/ASTContext.cpp @@ -8418,22 +8418,29 @@ static GVALinkage basicGVALinkageForFunction(const ASTContext &Context, return GVA_DiscardableODR; } -static GVALinkage adjustGVALinkageForAttributes(GVALinkage L, const Decl *D) { +static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context, + GVALinkage L, const Decl *D) { // See http://msdn.microsoft.com/en-us/library/xa0d9ste.aspx // dllexport/dllimport on inline functions. if (D->hasAttr<DLLImportAttr>()) { if (L == GVA_DiscardableODR || L == GVA_StrongODR) return GVA_AvailableExternally; - } else if (D->hasAttr<DLLExportAttr>() || D->hasAttr<CUDAGlobalAttr>()) { + } else if (D->hasAttr<DLLExportAttr>()) { if (L == GVA_DiscardableODR) return GVA_StrongODR; + } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice && + D->hasAttr<CUDAGlobalAttr>()) { + // Device-side functions with __global__ attribute must always be + // visible externally so they can be launched from host. + if (L == GVA_DiscardableODR || L == GVA_Internal) + return GVA_StrongODR; } return L; } GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const { - return adjustGVALinkageForAttributes(basicGVALinkageForFunction(*this, FD), - FD); + return adjustGVALinkageForAttributes( + *this, basicGVALinkageForFunction(*this, FD), FD); } static GVALinkage basicGVALinkageForVariable(const ASTContext &Context, @@ -8490,8 +8497,8 @@ static GVALinkage basicGVALinkageForVariable(const ASTContext &Context, } GVALinkage ASTContext::GetGVALinkageForVariable(const VarDecl *VD) { - return adjustGVALinkageForAttributes(basicGVALinkageForVariable(*this, VD), - VD); + return adjustGVALinkageForAttributes( + *this, basicGVALinkageForVariable(*this, VD), VD); } bool ASTContext::DeclMustBeEmitted(const Decl *D) { |