Author: tra Date: Mon May 2 15:30:03 2016 New Revision: 268299 URL: http://llvm.org/viewvc/llvm-project?rev=268299&view=rev Log: [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 Modified: cfe/trunk/lib/AST/ASTContext.cpp cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Modified: cfe/trunk/lib/AST/ASTContext.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=268299&r1=268298&r2=268299&view=diff ============================================================================== --- cfe/trunk/lib/AST/ASTContext.cpp (original) +++ cfe/trunk/lib/AST/ASTContext.cpp Mon May 2 15:30:03 2016 @@ -8418,22 +8418,29 @@ static GVALinkage basicGVALinkageForFunc 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 basicGVALinkageForVari } 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) { Modified: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu?rev=268299&r1=268298&r2=268299&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu (original) +++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Mon May 2 15:30:03 2016 @@ -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} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
