| Artem Belevich | c3fa25d | 2015-09-22 17:22:51 +0000 | [diff] [blame^] | 1 | // Make sure that __global__ functions are emitted along with correct |
| 2 | // annotations and are added to @llvm.used to prevent their elimination. |
| 3 | // REQUIRES: nvptx-registered-target |
| 4 | // |
| Justin Holewinski | 83e9668 | 2012-05-24 17:43:12 +0000 | [diff] [blame] | 5 | // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s |
| Peter Collingbourne | 5bad4af | 2011-10-06 16:49:54 +0000 | [diff] [blame] | 6 | |
| Eli Bendersky | 3468d9d | 2014-04-28 22:21:28 +0000 | [diff] [blame] | 7 | #include "Inputs/cuda.h" |
| Peter Collingbourne | 5bad4af | 2011-10-06 16:49:54 +0000 | [diff] [blame] | 8 | |
| Artem Belevich | c3fa25d | 2015-09-22 17:22:51 +0000 | [diff] [blame^] | 9 | // Make sure that all __global__ functions are added to @llvm.used |
| 10 | // CHECK: @llvm.used = appending global |
| 11 | // CHECK-SAME: @global_function |
| 12 | // CHECK-SAME: @_Z16templated_kernelIiEvT_ |
| 13 | |
| Stephen Lin | 4362261 | 2013-08-15 06:47:53 +0000 | [diff] [blame] | 14 | // CHECK-LABEL: define void @device_function |
| Justin Holewinski | 3683743 | 2013-03-30 14:38:24 +0000 | [diff] [blame] | 15 | extern "C" |
| Peter Collingbourne | 5bad4af | 2011-10-06 16:49:54 +0000 | [diff] [blame] | 16 | __device__ void device_function() {} |
| 17 | |
| Stephen Lin | 4362261 | 2013-08-15 06:47:53 +0000 | [diff] [blame] | 18 | // CHECK-LABEL: define void @global_function |
| Justin Holewinski | 3683743 | 2013-03-30 14:38:24 +0000 | [diff] [blame] | 19 | extern "C" |
| Peter Collingbourne | 5bad4af | 2011-10-06 16:49:54 +0000 | [diff] [blame] | 20 | __global__ void global_function() { |
| Justin Holewinski | 3683743 | 2013-03-30 14:38:24 +0000 | [diff] [blame] | 21 | // CHECK: call void @device_function |
| Peter Collingbourne | 5bad4af | 2011-10-06 16:49:54 +0000 | [diff] [blame] | 22 | device_function(); |
| 23 | } |
| Justin Holewinski | 3683743 | 2013-03-30 14:38:24 +0000 | [diff] [blame] | 24 | |
| Artem Belevich | c3fa25d | 2015-09-22 17:22:51 +0000 | [diff] [blame^] | 25 | // Make sure host-instantiated kernels are preserved on device side. |
| 26 | template <typename T> __global__ void templated_kernel(T param) {} |
| 27 | // CHECK-LABEL: define linkonce_odr void @_Z16templated_kernelIiEvT_ |
| 28 | void host_function() { templated_kernel<<<0,0>>>(0); } |
| 29 | |
| Duncan P. N. Exon Smith | b3a6669 | 2014-12-15 19:10:08 +0000 | [diff] [blame] | 30 | // CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1} |
| Artem Belevich | c3fa25d | 2015-09-22 17:22:51 +0000 | [diff] [blame^] | 31 | // CHECK: !{{[0-9]+}} = !{void (i32)* @_Z16templated_kernelIiEvT_, !"kernel", i32 1} |