1*67e74705SXin Li // RUN: %clang_cc1 -fsyntax-only -verify %s 2*67e74705SXin Li 3*67e74705SXin Li #include "Inputs/cuda.h" 4*67e74705SXin Li 5*67e74705SXin Li //------------------------------------------------------------------------------ 6*67e74705SXin Li // Test 1: host method called from device function 7*67e74705SXin Li 8*67e74705SXin Li struct S1 { methodS19*67e74705SXin Li void method() {} 10*67e74705SXin Li }; 11*67e74705SXin Li foo1(S1 & s)12*67e74705SXin Li__device__ void foo1(S1& s) { 13*67e74705SXin Li s.method(); // expected-error {{reference to __host__ function 'method' in __device__ function}} 14*67e74705SXin Li } 15*67e74705SXin Li 16*67e74705SXin Li //------------------------------------------------------------------------------ 17*67e74705SXin Li // Test 2: host method called from device function, for overloaded method 18*67e74705SXin Li 19*67e74705SXin Li struct S2 { methodS220*67e74705SXin Li void method(int) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}} methodS221*67e74705SXin Li void method(float) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}} 22*67e74705SXin Li }; 23*67e74705SXin Li foo2(S2 & s,int i,float f)24*67e74705SXin Li__device__ void foo2(S2& s, int i, float f) { 25*67e74705SXin Li s.method(f); // expected-error {{no matching member function}} 26*67e74705SXin Li } 27*67e74705SXin Li 28*67e74705SXin Li //------------------------------------------------------------------------------ 29*67e74705SXin Li // Test 3: device method called from host function 30*67e74705SXin Li 31*67e74705SXin Li struct S3 { methodS332*67e74705SXin Li __device__ void method() {} 33*67e74705SXin Li }; 34*67e74705SXin Li foo3(S3 & s)35*67e74705SXin Livoid foo3(S3& s) { 36*67e74705SXin Li s.method(); // expected-error {{reference to __device__ function 'method' in __host__ function}} 37*67e74705SXin Li } 38*67e74705SXin Li 39*67e74705SXin Li //------------------------------------------------------------------------------ 40*67e74705SXin Li // Test 4: device method called from host&device function 41*67e74705SXin Li 42*67e74705SXin Li struct S4 { methodS443*67e74705SXin Li __device__ void method() {} 44*67e74705SXin Li }; 45*67e74705SXin Li foo4(S4 & s)46*67e74705SXin Li__host__ __device__ void foo4(S4& s) { 47*67e74705SXin Li s.method(); 48*67e74705SXin Li } 49*67e74705SXin Li 50*67e74705SXin Li //------------------------------------------------------------------------------ 51*67e74705SXin Li // Test 5: overloaded operators 52*67e74705SXin Li 53*67e74705SXin Li struct S5 { S5S554*67e74705SXin Li S5() {} operator =S555*67e74705SXin Li S5& operator=(const S5&) {return *this;} // expected-note {{candidate function not viable}} 56*67e74705SXin Li }; 57*67e74705SXin Li foo5(S5 & s,S5 & t)58*67e74705SXin Li__device__ void foo5(S5& s, S5& t) { 59*67e74705SXin Li s = t; // expected-error {{no viable overloaded '='}} 60*67e74705SXin Li } 61*67e74705SXin Li 62*67e74705SXin Li //------------------------------------------------------------------------------ 63*67e74705SXin Li // Test 6: call method through pointer 64*67e74705SXin Li 65*67e74705SXin Li struct S6 { methodS666*67e74705SXin Li void method() {} 67*67e74705SXin Li }; 68*67e74705SXin Li foo6(S6 * s)69*67e74705SXin Li__device__ void foo6(S6* s) { 70*67e74705SXin Li s->method(); // expected-error {{reference to __host__ function 'method' in __device__ function}} 71*67e74705SXin Li } 72