diff options
| author | Artem Belevich <tra@google.com> | 2016-05-02 20:30:03 +0000 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2016-05-02 20:30:03 +0000 |
| commit | aba87a7ebcce24750d5d6845684c81345af0707b (patch) | |
| tree | 0b9cc5ddbed8a75b1a4e00b52aea9d0b4f2a99c7 /test/CodeGenCUDA | |
| parent | c1e7ade641b996676baef60b3e46dc359589e191 (diff) | |
[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
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@268299 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
| -rw-r--r-- | test/CodeGenCUDA/ptx-kernels.cu | 13 |
1 files changed, 11 insertions, 2 deletions
diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu index 6280e604f2..034cef6b60 100644 --- a/test/CodeGenCUDA/ptx-kernels.cu +++ b/test/CodeGenCUDA/ptx-kernels.cu @@ -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} |