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: 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