1*67e74705SXin Li // RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
2*67e74705SXin Li // RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fcuda-is-device -fsyntax-only -verify %s
3*67e74705SXin Li
4*67e74705SXin Li struct CopyableH {
operator =CopyableH5*67e74705SXin Li const CopyableH& operator=(const CopyableH& x) { return *this; }
6*67e74705SXin Li };
7*67e74705SXin Li struct CopyableD {
operator =CopyableD8*67e74705SXin Li __attribute__((device)) const CopyableD& operator=(const CopyableD x) { return *this; }
9*67e74705SXin Li };
10*67e74705SXin Li
11*67e74705SXin Li struct SimpleH {
12*67e74705SXin Li CopyableH b;
13*67e74705SXin Li };
14*67e74705SXin Li // expected-note@-3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __host__ function from __device__ function}}
15*67e74705SXin Li // expected-note@-4 2 {{candidate function (the implicit move assignment operator) not viable: call to __host__ function from __device__ function}}
16*67e74705SXin Li
17*67e74705SXin Li struct SimpleD {
18*67e74705SXin Li CopyableD b;
19*67e74705SXin Li };
20*67e74705SXin Li // expected-note@-3 2 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
21*67e74705SXin Li // expected-note@-4 2 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
22*67e74705SXin Li
foo1hh()23*67e74705SXin Li void foo1hh() {
24*67e74705SXin Li SimpleH a, b;
25*67e74705SXin Li a = b;
26*67e74705SXin Li }
foo1hd()27*67e74705SXin Li __attribute__((device)) void foo1hd() {
28*67e74705SXin Li SimpleH a, b;
29*67e74705SXin Li a = b; // expected-error {{no viable overloaded}}
30*67e74705SXin Li }
foo1dh()31*67e74705SXin Li void foo1dh() {
32*67e74705SXin Li SimpleD a, b;
33*67e74705SXin Li a = b; // expected-error {{no viable overloaded}}
34*67e74705SXin Li }
foo1dd()35*67e74705SXin Li __attribute__((device)) void foo1dd() {
36*67e74705SXin Li SimpleD a, b;
37*67e74705SXin Li a = b;
38*67e74705SXin Li }
39*67e74705SXin Li
foo2hh(SimpleH & a,SimpleH & b)40*67e74705SXin Li void foo2hh(SimpleH &a, SimpleH &b) {
41*67e74705SXin Li a = b;
42*67e74705SXin Li }
foo2hd(SimpleH & a,SimpleH & b)43*67e74705SXin Li __attribute__((device)) void foo2hd(SimpleH &a, SimpleH &b) {
44*67e74705SXin Li a = b; // expected-error {{no viable overloaded}}
45*67e74705SXin Li }
foo2dh(SimpleD & a,SimpleD & b)46*67e74705SXin Li void foo2dh(SimpleD &a, SimpleD &b) {
47*67e74705SXin Li a = b; // expected-error {{no viable overloaded}}
48*67e74705SXin Li }
foo2dd(SimpleD & a,SimpleD & b)49*67e74705SXin Li __attribute__((device)) void foo2dd(SimpleD &a, SimpleD &b) {
50*67e74705SXin Li a = b;
51*67e74705SXin Li }
52