This revision was automatically updated to reflect the committed changes.
Closed by commit rL274261: [CUDA] Give templated device functions internal
linkage, templated kernels… (authored by jlebar).
Changed prior to commit:
http://reviews.llvm.org/D21337?vs=60728&id=62391#toc
Repository:
rL LLVM
http://reviews.llvm.org/D21337
Files:
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
Index: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp
@@ -2671,9 +2671,18 @@
// 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.
Index: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
@@ -19,11 +19,11 @@
// 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() {
Index: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp
@@ -2671,9 +2671,18 @@
// 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.
Index: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
@@ -19,11 +19,11 @@
// 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
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
