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