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