Author: jlebar Date: Thu Jun 30 13:41:33 2016 New Revision: 274261 URL: http://llvm.org/viewvc/llvm-project?rev=274261&view=rev Log: [CUDA] Give templated device functions internal linkage, templated kernels external linkage.
Summary: This lets LLVM perform IPO over these functions. In particular, it allows LLVM to emit ld.global.nc for loads to __restrict pointers in kernels that are never written to. Reviewers: rsmith Subscribers: cfe-commits, tra Differential Revision: http://reviews.llvm.org/D21337 Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=274261&r1=274260&r2=274261&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original) +++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Thu Jun 30 13:41:33 2016 @@ -2671,9 +2671,18 @@ llvm::GlobalValue::LinkageTypes CodeGenM // explicit instantiations can occur in multiple translation units // and must all be equivalent. However, we are not allowed to // throw away these explicit instantiations. - if (Linkage == GVA_StrongODR) - return !Context.getLangOpts().AppleKext ? llvm::Function::WeakODRLinkage - : llvm::Function::ExternalLinkage; + // + // We don't currently support CUDA device code spread out across multiple TUs, + // so say that CUDA templates are either external (for kernels) or internal. + // This lets llvm perform aggressive inter-procedural optimizations. + if (Linkage == GVA_StrongODR) { + if (Context.getLangOpts().AppleKext) + return llvm::Function::ExternalLinkage; + if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice) + return D->hasAttr<CUDAGlobalAttr>() ? llvm::Function::ExternalLinkage + : llvm::Function::InternalLinkage; + return llvm::Function::WeakODRLinkage; + } // C++ doesn't have tentative definitions and thus cannot have common // linkage. Modified: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu?rev=274261&r1=274260&r2=274261&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu (original) +++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Thu Jun 30 13:41:33 2016 @@ -19,11 +19,11 @@ __global__ void global_function() { // Make sure host-instantiated kernels are preserved on device side. template <typename T> __global__ void templated_kernel(T param) {} -// CHECK-DAG: define weak_odr void @_Z16templated_kernelIiEvT_( +// CHECK-DAG: define void @_Z16templated_kernelIiEvT_( namespace { __global__ void anonymous_ns_kernel() {} -// CHECK-DAG: define weak_odr void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv( +// CHECK-DAG: define void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv( } void host_function() { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
