1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3
4 // Make sure we handle target overloads correctly.
5 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
6 // RUN: -fcuda-target-overloads -emit-llvm -o - %s \
7 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
8 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
9 // RUN: -fcuda-target-overloads -emit-llvm -o - %s \
10 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
11
12 // Check target overloads handling with disabled call target checks.
13 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
14 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
15 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
16 // RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
17 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
18 // RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \
19 // RUN: -fcuda-is-device -o - %s \
20 // RUN: | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
21 // RUN: -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s
22
23 #include "Inputs/cuda.h"
24
25 typedef int (*fp_t)(void);
26 typedef void (*gp_t)(void);
27
28 // CHECK-HOST: @hp = global i32 ()* @_Z1hv
29 // CHECK-HOST: @chp = global i32 ()* @ch
30 // CHECK-HOST: @dhp = global i32 ()* @_Z2dhv
31 // CHECK-HOST: @cdhp = global i32 ()* @cdh
32 // CHECK-HOST: @gp = global void ()* @_Z1gv
33
34 // CHECK-BOTH-LABEL: define i32 @_Z2dhv()
dh(void)35 __device__ int dh(void) { return 1; }
36 // CHECK-DEVICE: ret i32 1
dh(void)37 __host__ int dh(void) { return 2; }
38 // CHECK-HOST: ret i32 2
39
40 // CHECK-BOTH-LABEL: define i32 @_Z2hdv()
hd(void)41 __host__ __device__ int hd(void) { return 3; }
42 // CHECK-BOTH: ret i32 3
43
44 // CHECK-DEVICE-LABEL: define i32 @_Z1dv()
d(void)45 __device__ int d(void) { return 8; }
46 // CHECK-DEVICE: ret i32 8
47
48 // CHECK-HOST-LABEL: define i32 @_Z1hv()
h(void)49 __host__ int h(void) { return 9; }
50 // CHECK-HOST: ret i32 9
51
52 // CHECK-BOTH-LABEL: define void @_Z1gv()
g(void)53 __global__ void g(void) {}
54 // CHECK-BOTH: ret void
55
56 // mangled names of extern "C" __host__ __device__ functions clash
57 // with those of their __host__/__device__ counterparts, so
58 // overloading of extern "C" functions can only happen for __host__
59 // and __device__ functions -- we never codegen them in the same
60 // compilation and therefore mangled name conflict is not a problem.
61
62 // CHECK-BOTH-LABEL: define i32 @cdh()
cdh(void)63 extern "C" __device__ int cdh(void) {return 10;}
64 // CHECK-DEVICE: ret i32 10
cdh(void)65 extern "C" __host__ int cdh(void) {return 11;}
66 // CHECK-HOST: ret i32 11
67
68 // CHECK-DEVICE-LABEL: define i32 @cd()
cd(void)69 extern "C" __device__ int cd(void) {return 12;}
70 // CHECK-DEVICE: ret i32 12
71
72 // CHECK-HOST-LABEL: define i32 @ch()
ch(void)73 extern "C" __host__ int ch(void) {return 13;}
74 // CHECK-HOST: ret i32 13
75
76 // CHECK-BOTH-LABEL: define i32 @chd()
chd(void)77 extern "C" __host__ __device__ int chd(void) {return 14;}
78 // CHECK-BOTH: ret i32 14
79
80 // CHECK-HOST-LABEL: define void @_Z5hostfv()
hostf(void)81 __host__ void hostf(void) {
82 #if defined (NOCHECKS)
83 fp_t dp = d; // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
84 fp_t cdp = cd; // CHECK-HOST-NC: store {{.*}} @cd, {{.*}} %cdp,
85 #endif
86 fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
87 fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp,
88 fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp,
89 fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp,
90 fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp,
91 fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
92 gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp,
93
94 #if defined (NOCHECKS)
95 d(); // CHECK-HOST-NC: call i32 @_Z1dv()
96 cd(); // CHECK-HOST-NC: call i32 @cd()
97 #endif
98 h(); // CHECK-HOST: call i32 @_Z1hv()
99 ch(); // CHECK-HOST: call i32 @ch()
100 dh(); // CHECK-HOST: call i32 @_Z2dhv()
101 cdh(); // CHECK-HOST: call i32 @cdh()
102 g<<<0,0>>>(); // CHECK-HOST: call void @_Z1gv()
103 }
104
105 // CHECK-DEVICE-LABEL: define void @_Z7devicefv()
devicef(void)106 __device__ void devicef(void) {
107 fp_t dp = d; // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
108 fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp,
109 #if defined (NOCHECKS)
110 fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
111 fp_t chp = ch; // CHECK-DEVICE-NC: store {{.*}} @ch, {{.*}} %chp,
112 #endif
113 fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp,
114 fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp,
115 fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp,
116 fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
117
118 d(); // CHECK-DEVICE: call i32 @_Z1dv()
119 cd(); // CHECK-DEVICE: call i32 @cd()
120 #if defined (NOCHECKS)
121 h(); // CHECK-DEVICE-NC: call i32 @_Z1hv()
122 ch(); // CHECK-DEVICE-NC: call i32 @ch()
123 #endif
124 dh(); // CHECK-DEVICE: call i32 @_Z2dhv()
125 cdh(); // CHECK-DEVICE: call i32 @cdh()
126 }
127
128 // CHECK-BOTH-LABEL: define void @_Z11hostdevicefv()
hostdevicef(void)129 __host__ __device__ void hostdevicef(void) {
130 #if defined (NOCHECKS)
131 fp_t dp = d; // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
132 fp_t cdp = cd; // CHECK-BOTH-NC: store {{.*}} @cd, {{.*}} %cdp,
133 fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
134 fp_t chp = ch; // CHECK-BOTH-NC: store {{.*}} @ch, {{.*}} %chp,
135 #endif
136 fp_t dhp = dh; // CHECK-BOTH: store {{.*}} @_Z2dhv, {{.*}} %dhp,
137 fp_t cdhp = cdh; // CHECK-BOTH: store {{.*}} @cdh, {{.*}} %cdhp,
138 fp_t hdp = hd; // CHECK-BOTH: store {{.*}} @_Z2hdv, {{.*}} %hdp,
139 fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp,
140 #if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
141 gp_t gp = g; // CHECK-HOST-NC: store {{.*}} @_Z1gv, {{.*}} %gp,
142 #endif
143
144 #if defined (NOCHECKS)
145 d(); // CHECK-BOTH-NC: call i32 @_Z1dv()
146 cd(); // CHECK-BOTH-NC: call i32 @cd()
147 h(); // CHECK-BOTH-NC: call i32 @_Z1hv()
148 ch(); // CHECK-BOTH-NC: call i32 @ch()
149 #endif
150 dh(); // CHECK-BOTH: call i32 @_Z2dhv()
151 cdh(); // CHECK-BOTH: call i32 @cdh()
152 #if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
153 g<<<0,0>>>(); // CHECK-HOST-NC: call void @_Z1gv()
154 #endif
155 }
156
157 // Test for address of overloaded function resolution in the global context.
158 fp_t hp = h;
159 fp_t chp = ch;
160 fp_t dhp = dh;
161 fp_t cdhp = cdh;
162 gp_t gp = g;
163
164 int x;
165 // Check constructors/destructors for D/H functions
166 struct s_cd_dh {
s_cd_dhs_cd_dh167 __host__ s_cd_dh() { x = 11; }
s_cd_dhs_cd_dh168 __device__ s_cd_dh() { x = 12; }
~s_cd_dhs_cd_dh169 __host__ ~s_cd_dh() { x = 21; }
~s_cd_dhs_cd_dh170 __device__ ~s_cd_dh() { x = 22; }
171 };
172
173 struct s_cd_hd {
s_cd_hds_cd_hd174 __host__ __device__ s_cd_hd() { x = 31; }
~s_cd_hds_cd_hd175 __host__ __device__ ~s_cd_hd() { x = 32; }
176 };
177
178 // CHECK-BOTH: define void @_Z7wrapperv
179 #if defined(__CUDA_ARCH__)
180 __device__
181 #else
182 __host__
183 #endif
wrapper()184 void wrapper() {
185 s_cd_dh scddh;
186 // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
187 s_cd_hd scdhd;
188 // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
189
190 // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
191 // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev(
192 }
193 // CHECK-BOTH: ret void
194
195 // Now it's time to check what's been generated for the methods we used.
196
197 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
198 // CHECK-HOST: store i32 11,
199 // CHECK-DEVICE: store i32 12,
200 // CHECK-BOTH: ret void
201
202 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
203 // CHECK-BOTH: store i32 31,
204 // CHECK-BOTH: ret void
205
206 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
207 // CHECK-BOTH: store i32 32,
208 // CHECK-BOTH: ret void
209
210 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
211 // CHECK-HOST: store i32 21,
212 // CHECK-DEVICE: store i32 22,
213 // CHECK-BOTH: ret void
214
215