1*0a6a1f1dSLionel Sambuc // RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s 2*0a6a1f1dSLionel Sambuc 3*0a6a1f1dSLionel Sambuc #include "Inputs/cuda.h" 4*0a6a1f1dSLionel Sambuc 5*0a6a1f1dSLionel Sambuc //------------------------------------------------------------------------------ 6*0a6a1f1dSLionel Sambuc // Test 1: collision between two bases 7*0a6a1f1dSLionel Sambuc 8*0a6a1f1dSLionel Sambuc struct A1_with_host_ctor { A1_with_host_ctorA1_with_host_ctor9*0a6a1f1dSLionel Sambuc A1_with_host_ctor() {} 10*0a6a1f1dSLionel Sambuc }; 11*0a6a1f1dSLionel Sambuc 12*0a6a1f1dSLionel Sambuc struct B1_with_device_ctor { B1_with_device_ctorB1_with_device_ctor13*0a6a1f1dSLionel Sambuc __device__ B1_with_device_ctor() {} 14*0a6a1f1dSLionel Sambuc }; 15*0a6a1f1dSLionel Sambuc 16*0a6a1f1dSLionel Sambuc struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor { 17*0a6a1f1dSLionel Sambuc }; 18*0a6a1f1dSLionel Sambuc 19*0a6a1f1dSLionel Sambuc // expected-note@-3 {{candidate constructor (the implicit default constructor) not viable}} 20*0a6a1f1dSLionel Sambuc // expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} 21*0a6a1f1dSLionel Sambuc // expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}} 22*0a6a1f1dSLionel Sambuc // expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}} 23*0a6a1f1dSLionel Sambuc hostfoo1()24*0a6a1f1dSLionel Sambucvoid hostfoo1() { 25*0a6a1f1dSLionel Sambuc C1_with_collision c; // expected-error {{no matching constructor}} 26*0a6a1f1dSLionel Sambuc } 27*0a6a1f1dSLionel Sambuc 28*0a6a1f1dSLionel Sambuc //------------------------------------------------------------------------------ 29*0a6a1f1dSLionel Sambuc // Test 2: collision between two fields 30*0a6a1f1dSLionel Sambuc 31*0a6a1f1dSLionel Sambuc struct C2_with_collision { 32*0a6a1f1dSLionel Sambuc A1_with_host_ctor aa; 33*0a6a1f1dSLionel Sambuc B1_with_device_ctor bb; 34*0a6a1f1dSLionel Sambuc }; 35*0a6a1f1dSLionel Sambuc 36*0a6a1f1dSLionel Sambuc // expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable 37*0a6a1f1dSLionel Sambuc // expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} 38*0a6a1f1dSLionel Sambuc // expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable 39*0a6a1f1dSLionel Sambuc // expected-note@-8 {{candidate constructor (the implicit move constructor}} not viable 40*0a6a1f1dSLionel Sambuc hostfoo2()41*0a6a1f1dSLionel Sambucvoid hostfoo2() { 42*0a6a1f1dSLionel Sambuc C2_with_collision c; // expected-error {{no matching constructor}} 43*0a6a1f1dSLionel Sambuc } 44*0a6a1f1dSLionel Sambuc 45*0a6a1f1dSLionel Sambuc //------------------------------------------------------------------------------ 46*0a6a1f1dSLionel Sambuc // Test 3: collision between a field and a base 47*0a6a1f1dSLionel Sambuc 48*0a6a1f1dSLionel Sambuc struct C3_with_collision : A1_with_host_ctor { 49*0a6a1f1dSLionel Sambuc B1_with_device_ctor bb; 50*0a6a1f1dSLionel Sambuc }; 51*0a6a1f1dSLionel Sambuc 52*0a6a1f1dSLionel Sambuc // expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable 53*0a6a1f1dSLionel Sambuc // expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}} 54*0a6a1f1dSLionel Sambuc // expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable 55*0a6a1f1dSLionel Sambuc // expected-note@-7 {{candidate constructor (the implicit move constructor}} not viable 56*0a6a1f1dSLionel Sambuc hostfoo3()57*0a6a1f1dSLionel Sambucvoid hostfoo3() { 58*0a6a1f1dSLionel Sambuc C3_with_collision c; // expected-error {{no matching constructor}} 59*0a6a1f1dSLionel Sambuc } 60*0a6a1f1dSLionel Sambuc 61*0a6a1f1dSLionel Sambuc //------------------------------------------------------------------------------ 62*0a6a1f1dSLionel Sambuc // Test 4: collision on resolving a copy ctor 63*0a6a1f1dSLionel Sambuc 64*0a6a1f1dSLionel Sambuc struct A4_with_host_copy_ctor { A4_with_host_copy_ctorA4_with_host_copy_ctor65*0a6a1f1dSLionel Sambuc A4_with_host_copy_ctor() {} A4_with_host_copy_ctorA4_with_host_copy_ctor66*0a6a1f1dSLionel Sambuc A4_with_host_copy_ctor(const A4_with_host_copy_ctor&) {} 67*0a6a1f1dSLionel Sambuc }; 68*0a6a1f1dSLionel Sambuc 69*0a6a1f1dSLionel Sambuc struct B4_with_device_copy_ctor { B4_with_device_copy_ctorB4_with_device_copy_ctor70*0a6a1f1dSLionel Sambuc B4_with_device_copy_ctor() {} B4_with_device_copy_ctorB4_with_device_copy_ctor71*0a6a1f1dSLionel Sambuc __device__ B4_with_device_copy_ctor(const B4_with_device_copy_ctor&) {} 72*0a6a1f1dSLionel Sambuc }; 73*0a6a1f1dSLionel Sambuc 74*0a6a1f1dSLionel Sambuc struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor { 75*0a6a1f1dSLionel Sambuc }; 76*0a6a1f1dSLionel Sambuc 77*0a6a1f1dSLionel Sambuc // expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable 78*0a6a1f1dSLionel Sambuc // expected-note@-4 {{implicit copy constructor inferred target collision}} 79*0a6a1f1dSLionel Sambuc // expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable 80*0a6a1f1dSLionel Sambuc hostfoo4()81*0a6a1f1dSLionel Sambucvoid hostfoo4() { 82*0a6a1f1dSLionel Sambuc C4_with_collision c; 83*0a6a1f1dSLionel Sambuc C4_with_collision c2 = c; // expected-error {{no matching constructor}} 84*0a6a1f1dSLionel Sambuc } 85*0a6a1f1dSLionel Sambuc 86*0a6a1f1dSLionel Sambuc //------------------------------------------------------------------------------ 87*0a6a1f1dSLionel Sambuc // Test 5: collision on resolving a move ctor 88*0a6a1f1dSLionel Sambuc 89*0a6a1f1dSLionel Sambuc struct A5_with_host_move_ctor { A5_with_host_move_ctorA5_with_host_move_ctor90*0a6a1f1dSLionel Sambuc A5_with_host_move_ctor() {} A5_with_host_move_ctorA5_with_host_move_ctor91*0a6a1f1dSLionel Sambuc A5_with_host_move_ctor(A5_with_host_move_ctor&&) {} 92*0a6a1f1dSLionel Sambuc // expected-note@-1 {{copy constructor is implicitly deleted because 'A5_with_host_move_ctor' has a user-declared move constructor}} 93*0a6a1f1dSLionel Sambuc }; 94*0a6a1f1dSLionel Sambuc 95*0a6a1f1dSLionel Sambuc struct B5_with_device_move_ctor { B5_with_device_move_ctorB5_with_device_move_ctor96*0a6a1f1dSLionel Sambuc B5_with_device_move_ctor() {} B5_with_device_move_ctorB5_with_device_move_ctor97*0a6a1f1dSLionel Sambuc __device__ B5_with_device_move_ctor(B5_with_device_move_ctor&&) {} 98*0a6a1f1dSLionel Sambuc }; 99*0a6a1f1dSLionel Sambuc 100*0a6a1f1dSLionel Sambuc struct C5_with_collision : A5_with_host_move_ctor, B5_with_device_move_ctor { 101*0a6a1f1dSLionel Sambuc }; 102*0a6a1f1dSLionel Sambuc // expected-note@-2 {{deleted}} 103*0a6a1f1dSLionel Sambuc hostfoo5()104*0a6a1f1dSLionel Sambucvoid hostfoo5() { 105*0a6a1f1dSLionel Sambuc C5_with_collision c; 106*0a6a1f1dSLionel Sambuc // What happens here: 107*0a6a1f1dSLionel Sambuc // This tries to find the move ctor. Since the move ctor is deleted due to 108*0a6a1f1dSLionel Sambuc // collision, it then looks for a copy ctor. But copy ctors are implicitly 109*0a6a1f1dSLionel Sambuc // deleted when move ctors are declared explicitly. 110*0a6a1f1dSLionel Sambuc C5_with_collision c2(static_cast<C5_with_collision&&>(c)); // expected-error {{call to implicitly-deleted}} 111*0a6a1f1dSLionel Sambuc } 112