implicit-member-target.cu | Explore in Territory

// RUN: %clang_cc1 -std=gnu++11 -fsyntax-only -verify %s -Wno-defaulted-function-deleted

#include "Inputs/cuda.h"

//------------------------------------------------------------------------------
// Test 1: infer default ctor to be host.

struct A1_with_host_ctor {
  A1_with_host_ctor() {}
};

// The implicit default constructor is inferred to be host because it only needs
// to invoke a single host constructor (A1_with_host_ctor's). So we'll encounter
// an error when calling it from a __device__ function, but not from a __host__
// function.
struct B1_with_implicit_default_ctor : A1_with_host_ctor {
};

// expected-note@-3 {{call to __host__ function from __device__}}
// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}}

void hostfoo() {
  B1_with_implicit_default_ctor b;
}

__device__ void devicefoo() {
  B1_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
}

//------------------------------------------------------------------------------
// Test 2: infer default ctor to be device.

struct A2_with_device_ctor {
  __device__ A2_with_device_ctor() {}
};

struct B2_with_implicit_default_ctor : A2_with_device_ctor {
};

// expected-note@-3 {{call to __device__ function from __host__}}
// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-5 {{candidate constructor (the implicit move constructor) not viable}}

void hostfoo2() {
  B2_with_implicit_default_ctor b;  // expected-error {{no matching constructor}}
}

__device__ void devicefoo2() {
  B2_with_implicit_default_ctor b;
}

//------------------------------------------------------------------------------
// Test 3: infer copy ctor

struct A3_with_device_ctors {
  __host__ A3_with_device_ctors() {}
  __device__ A3_with_device_ctors(const A3_with_device_ctors&) {}
};

struct B3_with_implicit_ctors : A3_with_device_ctors {
};
// expected-note@-2 2{{call to __device__ function from __host__ function}}
// expected-note@-3 {{default constructor}}


void hostfoo3() {
  B3_with_implicit_ctors b;  // this is OK because the inferred default ctor
                             // here is __host__
  B3_with_implicit_ctors b2 = b; // expected-error {{no matching constructor}}

}

//------------------------------------------------------------------------------
// Test 4: infer default ctor from a field, not a base

struct A4_with_host_ctor {
  A4_with_host_ctor() {}
};

struct B4_with_implicit_default_ctor {
  A4_with_host_ctor field;
};

// expected-note@-4 {{call to __host__ function from __device__}}
// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}

void hostfoo4() {
  B4_with_implicit_default_ctor b;
}

__device__ void devicefoo4() {
  B4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
}

//------------------------------------------------------------------------------
// Test 5: copy ctor with non-const param

struct A5_copy_ctor_constness {
  __host__ A5_copy_ctor_constness() {}
  __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {}
};

struct B5_copy_ctor_constness : A5_copy_ctor_constness {
};

// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}}
// expected-note@-4 {{candidate constructor (the implicit default constructor) not viable}}

void hostfoo5(B5_copy_ctor_constness& b_arg) {
  B5_copy_ctor_constness b = b_arg;
}

__device__ void devicefoo5(B5_copy_ctor_constness& b_arg) {
  B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}}
}

//------------------------------------------------------------------------------
// Test 6: explicitly defaulted ctor: since they are spelled out, they have
// a host/device designation explicitly so no inference needs to be done.

struct A6_with_device_ctor {
  __device__ A6_with_device_ctor() {}
};

struct B6_with_defaulted_ctor : A6_with_device_ctor {
  __host__ B6_with_defaulted_ctor() = default;
};

// expected-note@-3 {{candidate constructor not viable: call to __host__ function from __device__ function}}
// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}

__device__ void devicefoo6() {
  B6_with_defaulted_ctor b; // expected-error {{no matching constructor}}
}

//------------------------------------------------------------------------------
// Test 7: copy assignment operator

struct A7_with_copy_assign {
  A7_with_copy_assign() {}
  __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {}
};

struct B7_with_copy_assign : A7_with_copy_assign {
};

// expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
// expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}

void hostfoo7() {
  B7_with_copy_assign b1, b2;
  b1 = b2; // expected-error {{no viable overloaded '='}}
}

//------------------------------------------------------------------------------
// Test 8: move assignment operator

// definitions for std::move
namespace std {
inline namespace foo {
template <class T> struct remove_reference { typedef T type; };
template <class T> struct remove_reference<T&> { typedef T type; };
template <class T> struct remove_reference<T&&> { typedef T type; };

template <class T> typename remove_reference<T>::type&& move(T&& t);
}
}

struct A8_with_move_assign {
  A8_with_move_assign() {}
  __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {}
  __device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {}
};

struct B8_with_move_assign : A8_with_move_assign {
};

// expected-note@-3 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
// expected-note@-4 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}

void hostfoo8() {
  B8_with_move_assign b1, b2;
  b1 = std::move(b2); // expected-error {{no viable overloaded '='}}
}
llvm/clang/test/SemaCUDA/implicit-member-target.cu
