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 64 // expected-note@-3 {{copy constructor of 'B3_with_implicit_ctors' is implicitly deleted}} 65 66 void hostfoo3() { 67 B3_with_implicit_ctors b; // this is OK because the inferred default ctor 68 // here is __host__ 69 B3_with_implicit_ctors b2 = b; // expected-error {{call to implicitly-deleted copy constructor}} 70 71 } 72 73 //------------------------------------------------------------------------------ 74 // Test 4: infer default ctor from a field, not a base 75 76 struct A4_with_host_ctor { 77 A4_with_host_ctor() {} 78 }; 79 80 struct B4_with_implicit_default_ctor { 81 A4_with_host_ctor field; 82 }; 83 84 // expected-note@-4 {{call to __host__ function from __device__}} 85 // expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}} 86 // expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}} 87 88 void hostfoo4() { 89 B4_with_implicit_default_ctor b; 90 } 91 92 __device__ void devicefoo4() { 93 B4_with_implicit_default_ctor b; // expected-error {{no matching constructor}} 94 } 95 96 //------------------------------------------------------------------------------ 97 // Test 5: copy ctor with non-const param 98 99 struct A5_copy_ctor_constness { 100 __host__ A5_copy_ctor_constness() {} 101 __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {} 102 }; 103 104 struct B5_copy_ctor_constness : A5_copy_ctor_constness { 105 }; 106 107 // expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}} 108 // expected-note@-4 {{candidate constructor (the implicit default constructor) not viable}} 109 110 void hostfoo5(B5_copy_ctor_constness& b_arg) { 111 B5_copy_ctor_constness b = b_arg; 112 } 113 114 __device__ void devicefoo5(B5_copy_ctor_constness& b_arg) { 115 B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}} 116 } 117 118 //------------------------------------------------------------------------------ 119 // Test 6: explicitly defaulted ctor: since they are spelled out, they have 120 // a host/device designation explicitly so no inference needs to be done. 121 122 struct A6_with_device_ctor { 123 __device__ A6_with_device_ctor() {} 124 }; 125 126 struct B6_with_defaulted_ctor : A6_with_device_ctor { 127 __host__ B6_with_defaulted_ctor() = default; 128 }; 129 130 // expected-note@-3 {{candidate constructor not viable: call to __host__ function from __device__ function}} 131 // expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}} 132 // expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}} 133 134 __device__ void devicefoo6() { 135 B6_with_defaulted_ctor b; // expected-error {{no matching constructor}} 136 } 137 138 //------------------------------------------------------------------------------ 139 // Test 7: copy assignment operator 140 141 struct A7_with_copy_assign { 142 A7_with_copy_assign() {} 143 __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {} 144 }; 145 146 struct B7_with_copy_assign : A7_with_copy_assign { 147 }; 148 149 // expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}} 150 // expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}} 151 152 void hostfoo7() { 153 B7_with_copy_assign b1, b2; 154 b1 = b2; // expected-error {{no viable overloaded '='}} 155 } 156 157 //------------------------------------------------------------------------------ 158 // Test 8: move assignment operator 159 160 // definitions for std::move 161 namespace std { 162 inline namespace foo { 163 template <class T> struct remove_reference { typedef T type; }; 164 template <class T> struct remove_reference<T&> { typedef T type; }; 165 template <class T> struct remove_reference<T&&> { typedef T type; }; 166 167 template <class T> typename remove_reference<T>::type&& move(T&& t); 168 } 169 } 170 171 struct A8_with_move_assign { 172 A8_with_move_assign() {} 173 __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {} 174 __device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {} 175 }; 176 177 struct B8_with_move_assign : A8_with_move_assign { 178 }; 179 180 // expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}} 181 // expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}} 182 183 void hostfoo8() { 184 B8_with_move_assign b1, b2; 185 b1 = std::move(b2); // expected-error {{no viable overloaded '='}} 186 } 187
