236
|
1 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -emit-llvm -o - %s \
|
221
|
2 // RUN: | FileCheck -check-prefix=NORDC %s
|
236
|
3 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -emit-llvm -o - %s \
|
|
4 // RUN: | FileCheck -check-prefix=NORDC-NEG %s
|
|
5 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -fgpu-rdc -emit-llvm -o - %s \
|
221
|
6 // RUN: | FileCheck -check-prefix=RDC %s
|
236
|
7 // RUN: %clang_cc1 -triple nvptx -fcuda-is-device -fgpu-rdc -emit-llvm -o - %s \
|
|
8 // RUN: | FileCheck -check-prefix=RDC-NEG %s
|
221
|
9
|
|
10 #include "Inputs/cuda.h"
|
|
11
|
|
12 template <typename T> __device__ void func() {}
|
|
13 template <typename T> __global__ void kernel() {}
|
|
14
|
|
15 template __device__ void func<int>();
|
236
|
16 // NORDC: define internal void @_Z4funcIiEvv()
|
|
17 // RDC: define weak_odr void @_Z4funcIiEvv()
|
|
18
|
221
|
19 template __global__ void kernel<int>();
|
236
|
20 // NORDC: define void @_Z6kernelIiEvv()
|
|
21 // RDC: define weak_odr void @_Z6kernelIiEvv()
|
|
22
|
|
23 // Ensure that unused static device function is eliminated
|
|
24 static __device__ void static_func() {}
|
|
25 // NORDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv()
|
|
26 // RDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv[[FILEID:.*]]()
|
|
27
|
|
28 // Ensure that kernel function has external or weak_odr
|
|
29 // linkage regardless static specifier
|
|
30 static __global__ void static_kernel() {}
|
|
31 // NORDC: define void @_ZL13static_kernelv()
|
|
32 // RDC: define weak_odr void @_ZL13static_kernelv[[FILEID:.*]]()
|