188 lines · plain
1// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s -Wno-defaulted-function-deleted2 3#include "Inputs/cuda.h"4 5//------------------------------------------------------------------------------6// Test 1: infer default ctor to be host.7 8struct A1_with_host_ctor {9 A1_with_host_ctor() {}10};11 12// The implicit default constructor is inferred to be host because it only needs13// to invoke a single host constructor (A1_with_host_ctor's). So we'll encounter14// an error when calling it from a __device__ function, but not from a __host__15// function.16struct 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 23void 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 34struct A2_with_device_ctor {35 __device__ A2_with_device_ctor() {}36};37 38struct 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 45void 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 ctor55 56struct A3_with_device_ctors {57 __host__ A3_with_device_ctors() {}58 __device__ A3_with_device_ctors(const A3_with_device_ctors&) {}59};60 61struct 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 67void hostfoo3() {68 B3_with_implicit_ctors b; // this is OK because the inferred default ctor69 // 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 base76 77struct A4_with_host_ctor {78 A4_with_host_ctor() {}79};80 81struct 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 89void 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 param99 100struct A5_copy_ctor_constness {101 __host__ A5_copy_ctor_constness() {}102 __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {}103};104 105struct 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 111void 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 have121// a host/device designation explicitly so no inference needs to be done.122 123struct A6_with_device_ctor {124 __device__ A6_with_device_ctor() {}125};126 127struct 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 operator141 142struct A7_with_copy_assign {143 A7_with_copy_assign() {}144 __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {}145};146 147struct 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 153void hostfoo7() {154 B7_with_copy_assign b1, b2;155 b1 = b2; // expected-error {{no viable overloaded '='}}156}157 158//------------------------------------------------------------------------------159// Test 8: move assignment operator160 161// definitions for std::move162namespace std {163inline namespace foo {164template <class T> struct remove_reference { typedef T type; };165template <class T> struct remove_reference<T&> { typedef T type; };166template <class T> struct remove_reference<T&&> { typedef T type; };167 168template <class T> typename remove_reference<T>::type&& move(T&& t);169}170}171 172struct 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 178struct 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 184void hostfoo8() {185 B8_with_move_assign b1, b2;186 b1 = std::move(b2); // expected-error {{no viable overloaded '='}}187}188