1 // RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s
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