summaryrefslogtreecommitdiff
path: root/test/CodeGenCUDA
diff options
context:
space:
mode:
authorArtem Belevich <tra@google.com>2015-09-23 17:44:53 +0000
committerArtem Belevich <tra@google.com>2015-09-23 17:44:53 +0000
commitb7a9667a4192f47243f2fc4bb2665d421c3d52ad (patch)
treebb5582c16b76cf8b670ba2da72791696fe230953 /test/CodeGenCUDA
parent4bb5c31d2c26e4fe2cd6d0fede53293ea6dfff66 (diff)
downloadclang-b7a9667a4192f47243f2fc4bb2665d421c3d52ad.tar.gz
[CUDA] __global__ functions should always be visible externally.
Adjust __global__ functions with DiscardableODR linkage to use StrongODR linkage instead, so they are visible externally. Differential Revision: http://reviews.llvm.org/D13067 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@248400 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
-rw-r--r--test/CodeGenCUDA/ptx-kernels.cu7
1 files changed, 1 insertions, 6 deletions
diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu
index 3a8722a342..6280e604f2 100644
--- a/test/CodeGenCUDA/ptx-kernels.cu
+++ b/test/CodeGenCUDA/ptx-kernels.cu
@@ -6,11 +6,6 @@
#include "Inputs/cuda.h"
-// Make sure that all __global__ functions are added to @llvm.used
-// CHECK: @llvm.used = appending global
-// CHECK-SAME: @global_function
-// CHECK-SAME: @_Z16templated_kernelIiEvT_
-
// CHECK-LABEL: define void @device_function
extern "C"
__device__ void device_function() {}
@@ -24,7 +19,7 @@ __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 linkonce_odr void @_Z16templated_kernelIiEvT_
+// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_
void host_function() { templated_kernel<<<0,0>>>(0); }
// CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}