16
16
17
17
#include " Inputs/cuda.h"
18
18
19
- typedef int (*fp_t )(void );
20
- typedef void (*gp_t )(void );
19
+ typedef int (*fp_t )();
20
+ typedef void (*gp_t )();
21
21
22
- // Host and unattributed functions can't be overloaded
23
- __host__ int hh (void ) { return 1 ; } // expected-note {{previous definition is here}}
24
- int hh (void ) { return 1 ; } // expected-error {{redefinition of 'hh'}}
22
+ // Host and unattributed functions can't be overloaded.
23
+ __host__ void hh () {} // expected-note {{previous definition is here}}
24
+ void hh () {} // expected-error {{redefinition of 'hh'}}
25
25
26
- // H/D overloading is OK
27
- __host__ int dh (void ) { return 2 ; }
28
- __device__ int dh (void ) { return 2 ; }
26
+ // H/D overloading is OK.
27
+ __host__ int dh () { return 2 ; }
28
+ __device__ int dh () { return 2 ; }
29
29
30
- // H/HD and D/HD are not allowed
31
- __host__ __device__ int hdh (void ) { return 5 ; } // expected-note {{previous definition is here}}
32
- __host__ int hdh (void ) { return 4 ; } // expected-error {{redefinition of 'hdh'}}
30
+ // H/HD and D/HD are not allowed.
31
+ __host__ __device__ int hdh () { return 5 ; } // expected-note {{previous definition is here}}
32
+ __host__ int hdh () { return 4 ; } // expected-error {{redefinition of 'hdh'}}
33
33
34
- __host__ int hhd (void ) { return 4 ; } // expected-note {{previous definition is here}}
35
- __host__ __device__ int hhd (void ) { return 5 ; } // expected-error {{redefinition of 'hhd'}}
34
+ __host__ int hhd () { return 4 ; } // expected-note {{previous definition is here}}
35
+ __host__ __device__ int hhd () { return 5 ; } // expected-error {{redefinition of 'hhd'}}
36
36
// expected-warning@-1 {{attribute declaration must precede definition}}
37
37
// expected-note@-3 {{previous definition is here}}
38
38
39
- __host__ __device__ int hdd (void ) { return 7 ; } // expected-note {{previous definition is here}}
40
- __device__ int hdd (void ) { return 6 ; } // expected-error {{redefinition of 'hdd'}}
39
+ __host__ __device__ int hdd () { return 7 ; } // expected-note {{previous definition is here}}
40
+ __device__ int hdd () { return 6 ; } // expected-error {{redefinition of 'hdd'}}
41
41
42
- __device__ int dhd (void ) { return 6 ; } // expected-note {{previous definition is here}}
43
- __host__ __device__ int dhd (void ) { return 7 ; } // expected-error {{redefinition of 'dhd'}}
42
+ __device__ int dhd () { return 6 ; } // expected-note {{previous definition is here}}
43
+ __host__ __device__ int dhd () { return 7 ; } // expected-error {{redefinition of 'dhd'}}
44
44
// expected-warning@-1 {{attribute declaration must precede definition}}
45
45
// expected-note@-3 {{previous definition is here}}
46
46
47
- // Same tests for extern "C" functions
48
- extern " C" __host__ int chh (void ) {return 11 ;} // expected-note {{previous definition is here}}
49
- extern " C" int chh (void ) {return 11 ;} // expected-error {{redefinition of 'chh'}}
47
+ // Same tests for extern "C" functions.
48
+ extern " C" __host__ int chh () {return 11 ;} // expected-note {{previous definition is here}}
49
+ extern " C" int chh () {return 11 ;} // expected-error {{redefinition of 'chh'}}
50
50
51
- // H/D overloading is OK
52
- extern " C" __device__ int cdh (void ) {return 10 ;}
53
- extern " C" __host__ int cdh (void ) {return 11 ;}
51
+ // H/D overloading is OK.
52
+ extern " C" __device__ int cdh () {return 10 ;}
53
+ extern " C" __host__ int cdh () {return 11 ;}
54
54
55
55
// H/HD and D/HD overloading is not allowed.
56
- extern " C" __host__ __device__ int chhd1 (void ) {return 12 ;} // expected-note {{previous definition is here}}
57
- extern " C" __host__ int chhd1 (void ) {return 13 ;} // expected-error {{redefinition of 'chhd1'}}
56
+ extern " C" __host__ __device__ int chhd1 () {return 12 ;} // expected-note {{previous definition is here}}
57
+ extern " C" __host__ int chhd1 () {return 13 ;} // expected-error {{redefinition of 'chhd1'}}
58
58
59
- extern " C" __host__ int chhd2 (void ) {return 13 ;} // expected-note {{previous definition is here}}
60
- extern " C" __host__ __device__ int chhd2 (void ) {return 12 ;} // expected-error {{redefinition of 'chhd2'}}
59
+ extern " C" __host__ int chhd2 () {return 13 ;} // expected-note {{previous definition is here}}
60
+ extern " C" __host__ __device__ int chhd2 () {return 12 ;} // expected-error {{redefinition of 'chhd2'}}
61
61
// expected-warning@-1 {{attribute declaration must precede definition}}
62
62
// expected-note@-3 {{previous definition is here}}
63
63
64
64
// Helper functions to verify calling restrictions.
65
- __device__ int d (void ) { return 8 ; }
66
- __host__ int h (void ) { return 9 ; }
67
- __global__ void g (void ) {}
68
- extern " C" __device__ int cd (void ) {return 10 ;}
69
- extern " C" __host__ int ch (void ) {return 11 ;}
70
-
71
- __host__ void hostf (void ) {
65
+ __device__ int d () { return 8 ; }
66
+ // expected-note@-1 1+ {{'d' declared here}}
67
+ // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
68
+ // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
69
+
70
+ __host__ int h () { return 9 ; }
71
+ // expected-note@-1 1+ {{'h' declared here}}
72
+ // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
73
+ // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
74
+ // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
75
+
76
+ __global__ void g () {}
77
+ // expected-note@-1 1+ {{'g' declared here}}
78
+ // expected-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}}
79
+ // expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
80
+ // expected-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}}
81
+
82
+ extern " C" __device__ int cd () {return 10 ;}
83
+ // expected-note@-1 1+ {{'cd' declared here}}
84
+ // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
85
+ // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
86
+
87
+ extern " C" __host__ int ch () {return 11 ;}
88
+ // expected-note@-1 1+ {{'ch' declared here}}
89
+ // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
90
+ // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
91
+ // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
92
+
93
+ __host__ void hostf () {
72
94
fp_t dp = d;
73
95
// expected-error@-1 {{reference to __device__ function 'd' in __host__ function}}
74
- // expected-note@65 {{'d' declared here}}
75
96
fp_t cdp = cd;
76
97
// expected-error@-1 {{reference to __device__ function 'cd' in __host__ function}}
77
- // expected-note@68 {{'cd' declared here}}
78
98
fp_t hp = h;
79
99
fp_t chp = ch;
80
100
fp_t dhp = dh;
@@ -83,10 +103,8 @@ __host__ void hostf(void) {
83
103
84
104
d ();
85
105
// expected-error@-1 {{no matching function for call to 'd'}}
86
- // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}}
87
106
cd ();
88
107
// expected-error@-1 {{no matching function for call to 'cd'}}
89
- // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}}
90
108
h ();
91
109
ch ();
92
110
dh ();
@@ -95,108 +113,85 @@ __host__ void hostf(void) {
95
113
g<<<0 ,0 >>> ();
96
114
}
97
115
98
-
99
- __device__ void devicef (void ) {
116
+ __device__ void devicef () {
100
117
fp_t dp = d;
101
118
fp_t cdp = cd;
102
119
fp_t hp = h;
103
120
// expected-error@-1 {{reference to __host__ function 'h' in __device__ function}}
104
- // expected-note@66 {{'h' declared here}}
105
121
fp_t chp = ch;
106
122
// expected-error@-1 {{reference to __host__ function 'ch' in __device__ function}}
107
- // expected-note@69 {{'ch' declared here}}
108
123
fp_t dhp = dh;
109
124
fp_t cdhp = cdh;
110
125
gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
111
- // expected-note@67 {{'g' declared here}}
112
126
113
127
d ();
114
128
cd ();
115
129
h (); // expected-error {{no matching function for call to 'h'}}
116
- // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}}
117
130
ch (); // expected-error {{no matching function for call to 'ch'}}
118
- // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}}
119
131
dh ();
120
132
cdh ();
121
133
g (); // expected-error {{no matching function for call to 'g'}}
122
- // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}}
123
134
g<<<0 ,0 >>> (); // expected-error {{reference to __global__ function 'g' in __device__ function}}
124
- // expected-note@67 {{'g' declared here}}
125
135
}
126
136
127
- __global__ void globalf (void ) {
137
+ __global__ void globalf () {
128
138
fp_t dp = d;
129
139
fp_t cdp = cd;
130
140
fp_t hp = h;
131
141
// expected-error@-1 {{reference to __host__ function 'h' in __global__ function}}
132
- // expected-note@66 {{'h' declared here}}
133
142
fp_t chp = ch;
134
143
// expected-error@-1 {{reference to __host__ function 'ch' in __global__ function}}
135
- // expected-note@69 {{'ch' declared here}}
136
144
fp_t dhp = dh;
137
145
fp_t cdhp = cdh;
138
146
gp_t gp = g;
139
147
// expected-error@-1 {{reference to __global__ function 'g' in __global__ function}}
140
- // expected-note@67 {{'g' declared here}}
141
148
142
149
d ();
143
150
cd ();
144
151
h ();
145
152
// expected-error@-1 {{no matching function for call to 'h'}}
146
- // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}}
147
153
ch ();
148
154
// expected-error@-1 {{no matching function for call to 'ch'}}
149
- // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}}
150
155
dh ();
151
156
cdh ();
152
157
g (); // expected-error {{no matching function for call to 'g'}}
153
- // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}}
154
158
g<<<0 ,0 >>> (); // expected-error {{reference to __global__ function 'g' in __global__ function}}
155
- // expected-note@67 {{'g' declared here}}
156
159
}
157
160
158
- __host__ __device__ void hostdevicef (void ) {
161
+ __host__ __device__ void hostdevicef () {
159
162
fp_t dp = d;
160
163
fp_t cdp = cd;
164
+ #if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
165
+ // expected-error@-3 {{reference to __device__ function 'd' in __host__ __device__ function}}
166
+ // expected-error@-3 {{reference to __device__ function 'cd' in __host__ __device__ function}}
167
+ #endif
168
+
161
169
fp_t hp = h;
162
170
fp_t chp = ch;
163
- #if !defined(NOCHECKS)
164
- #if !defined(__CUDA_ARCH__)
165
- // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}}
166
- // expected-note@65 {{'d' declared here}}
167
- // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}}
168
- // expected-note@68 {{'cd' declared here}}
169
- #else
170
- // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}}
171
- // expected-note@66 {{'h' declared here}}
172
- // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}}
173
- // expected-note@69 {{'ch' declared here}}
174
- #endif
171
+ #if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
172
+ // expected-error@-3 {{reference to __host__ function 'h' in __host__ __device__ function}}
173
+ // expected-error@-3 {{reference to __host__ function 'ch' in __host__ __device__ function}}
175
174
#endif
175
+
176
176
fp_t dhp = dh;
177
177
fp_t cdhp = cdh;
178
178
gp_t gp = g;
179
179
#if defined(__CUDA_ARCH__)
180
180
// expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
181
- // expected-note@67 {{'g' declared here}}
182
181
#endif
183
182
184
183
d ();
185
184
cd ();
185
+ #if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
186
+ // expected-error@-3 {{no matching function for call to 'd'}}
187
+ // expected-error@-3 {{no matching function for call to 'cd'}}
188
+ #endif
189
+
186
190
h ();
187
191
ch ();
188
- #if !defined(NOCHECKS)
189
- #if !defined(__CUDA_ARCH__)
190
- // expected-error@-6 {{no matching function for call to 'd'}}
191
- // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
192
- // expected-error@-7 {{no matching function for call to 'cd'}}
193
- // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
194
- #else
195
- // expected-error@-9 {{no matching function for call to 'h'}}
196
- // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
197
- // expected-error@-10 {{no matching function for call to 'ch'}}
198
- // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
199
- #endif
192
+ #if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
193
+ // expected-error@-3 {{no matching function for call to 'h'}}
194
+ // expected-error@-3 {{no matching function for call to 'ch'}}
200
195
#endif
201
196
202
197
dh ();
@@ -207,9 +202,7 @@ __host__ __device__ void hostdevicef(void) {
207
202
// expected-error@-3 {{call to global function g not configured}}
208
203
#else
209
204
// expected-error@-5 {{no matching function for call to 'g'}}
210
- // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
211
- // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}}
212
- // expected-note@67 {{'g' declared here}}
205
+ // expected-error@-5 {{reference to __global__ function 'g' in __host__ __device__ function}}
213
206
#endif // __CUDA_ARCH__
214
207
}
215
208
0 commit comments