1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
| // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -Wno-defaulted-function-deleted
#include "Inputs/cuda.h"
//------------------------------------------------------------------------------
// Test 1: infer inherited default ctor to be host.
struct A1_with_host_ctor {
A1_with_host_ctor() {}
};
// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
// The inherited default constructor is inferred to be host, so we'll encounter
// an error when calling it from a __device__ function, but not from a __host__
// function.
struct B1_with_implicit_default_ctor : A1_with_host_ctor {
using A1_with_host_ctor::A1_with_host_ctor;
};
// expected-note@-4 {{call to __host__ function from __device__}}
// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
// expected-note@-6 2{{constructor from base class 'A1_with_host_ctor' inherited here}}
void hostfoo() {
B1_with_implicit_default_ctor b;
}
__device__ void devicefoo() {
B1_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
}
//------------------------------------------------------------------------------
// Test 2: infer inherited default ctor to be device.
struct A2_with_device_ctor {
__device__ A2_with_device_ctor() {}
};
// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
struct B2_with_implicit_default_ctor : A2_with_device_ctor {
using A2_with_device_ctor::A2_with_device_ctor;
};
// expected-note@-4 {{call to __device__ function from __host__}}
// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
// expected-note@-6 2{{constructor from base class 'A2_with_device_ctor' inherited here}}
void hostfoo2() {
B2_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
}
__device__ void devicefoo2() {
B2_with_implicit_default_ctor b;
}
//------------------------------------------------------------------------------
// Test 3: infer inherited copy ctor
struct A3_with_device_ctors {
__host__ A3_with_device_ctors() {}
__device__ A3_with_device_ctors(const A3_with_device_ctors&) {}
};
struct B3_with_implicit_ctors : A3_with_device_ctors {
using A3_with_device_ctors::A3_with_device_ctors;
};
// expected-note@-3 2{{call to __device__ function from __host__ function}}
// expected-note@-4 {{default constructor}}
void hostfoo3() {
B3_with_implicit_ctors b; // this is OK because the inferred inherited default ctor
// here is __host__
B3_with_implicit_ctors b2 = b; // expected-error {{no matching constructor}}
}
//------------------------------------------------------------------------------
// Test 4: infer inherited default ctor from a field, not a base
struct A4_with_host_ctor {
A4_with_host_ctor() {}
};
struct B4_with_inherited_host_ctor : A4_with_host_ctor{
using A4_with_host_ctor::A4_with_host_ctor;
};
struct C4_with_implicit_default_ctor {
B4_with_inherited_host_ctor field;
};
// expected-note@-4 {{call to __host__ function from __device__}}
// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
void hostfoo4() {
C4_with_implicit_default_ctor b;
}
__device__ void devicefoo4() {
C4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
}
//------------------------------------------------------------------------------
// Test 5: inherited copy ctor with non-const param
struct A5_copy_ctor_constness {
__host__ A5_copy_ctor_constness() {}
__host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {}
};
struct B5_copy_ctor_constness : A5_copy_ctor_constness {
using A5_copy_ctor_constness::A5_copy_ctor_constness;
};
// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}}
// expected-note@-5 {{candidate constructor (the implicit default constructor) not viable}}
void hostfoo5(B5_copy_ctor_constness& b_arg) {
B5_copy_ctor_constness b = b_arg;
}
__device__ void devicefoo5(B5_copy_ctor_constness& b_arg) {
B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}}
}
//------------------------------------------------------------------------------
// Test 6: explicitly defaulted ctor
struct A6_with_device_ctor {
__device__ A6_with_device_ctor() {}
};
struct B6_with_defaulted_ctor : A6_with_device_ctor {
using A6_with_device_ctor::A6_with_device_ctor;
__host__ B6_with_defaulted_ctor() = default;
};
// expected-note@-3 {{explicitly defaulted function was implicitly deleted here}}
// expected-note@-6 {{default constructor of 'B6_with_defaulted_ctor' is implicitly deleted because base class 'A6_with_device_ctor' has no default constructor}}
void hostfoo6() {
B6_with_defaulted_ctor b; // expected-error {{call to implicitly-deleted default constructor}}
}
__device__ void devicefoo6() {
B6_with_defaulted_ctor b;
}
//------------------------------------------------------------------------------
// Test 7: inherited copy assignment operator
struct A7_with_copy_assign {
A7_with_copy_assign() {}
__device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {}
};
struct B7_with_copy_assign : A7_with_copy_assign {
using A7_with_copy_assign::A7_with_copy_assign;
};
// expected-note@-4 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
// expected-note@-5 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
void hostfoo7() {
B7_with_copy_assign b1, b2;
b1 = b2; // expected-error {{no viable overloaded '='}}
}
//------------------------------------------------------------------------------
// Test 8: inherited move assignment operator
// definitions for std::move
namespace std {
inline namespace foo {
template <class T> struct remove_reference { typedef T type; };
template <class T> struct remove_reference<T&> { typedef T type; };
template <class T> struct remove_reference<T&&> { typedef T type; };
template <class T> typename remove_reference<T>::type&& move(T&& t);
}
}
struct A8_with_move_assign {
A8_with_move_assign() {}
__device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {}
__device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {}
};
struct B8_with_move_assign : A8_with_move_assign {
using A8_with_move_assign::A8_with_move_assign;
};
// expected-note@-4 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
// expected-note@-5 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
void hostfoo8() {
B8_with_move_assign b1, b2;
b1 = std::move(b2); // expected-error {{no viable overloaded '='}}
}
|