maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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

llvm-svn: 268299
This commit is contained in:
Artem Belevich 2016-05-02 20:30:03 +00:00
parent 7bd8d99497
commit ca2b951cbc
2 changed files with 24 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
clang/lib/AST/ASTContext.cpp
View File

@ -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) {

13
clang/test/CodeGenCUDA/ptx-kernels.cu
View File

@ -19,8 +19,17 @@ __global__ void global_function() {
// Make sure host-instantiated kernels are preserved on device side.
template <typename T> __global__ void templated_kernel(T param) {}
// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_
void host_function() { templated_kernel<<<0,0>>>(0); }
// CHECK-DAG: define weak_odr void @_Z16templated_kernelIiEvT_(
namespace {
__global__ void anonymous_ns_kernel() {}
// CHECK-DAG: define weak_odr void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(
}
void host_function() {
templated_kernel<<<0, 0>>>(0);
anonymous_ns_kernel<<<0,0>>>();
}
// CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
// CHECK: !{{[0-9]+}} = !{void (i32)* @_Z16templated_kernelIiEvT_, !"kernel", i32 1}