1 // RUN: %clang_cc1 -fsyntax-only -verify %s
2 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
3 
4 #include "Inputs/cuda.h"
5 
6 __host__ void h1h(void);
7 __device__ void h1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ function}}
8 __host__ __device__ void h1hd(void);
9 __global__ void h1g(void);
10 
11 struct h1ds { // expected-note {{requires 1 argument}}
12   __device__ h1ds(); // expected-note {{candidate constructor not viable: call to __device__ function from __host__ function}}
13 };
14 
h1(void)15 __host__ void h1(void) {
16   h1h();
17   h1d(); // expected-error {{no matching function}}
18   h1hd();
19   h1g<<<1, 1>>>();
20   h1ds x; // expected-error {{no matching constructor}}
21 }
22 
23 __host__ void d1h(void); // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
24 __device__ void d1d(void);
25 __host__ __device__ void d1hd(void);
26 __global__ void d1g(void); // expected-note {{'d1g' declared here}}
27 
d1(void)28 __device__ void d1(void) {
29   d1h(); // expected-error {{no matching function}}
30   d1d();
31   d1hd();
32   d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}}
33 }
34 
35 // Expected 0-1 as in one of host/device side compilation it is an error, while
36 // not in the other
37 __host__ void hd1h(void); // expected-note 0-1 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
38 __device__ void hd1d(void); // expected-note 0-1 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
39 __host__ void hd1hg(void);
40 __device__ void hd1dg(void);
41 #ifdef __CUDA_ARCH__
42 __host__ void hd1hig(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
43 #else
44 __device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
45 #endif
46 __host__ __device__ void hd1hd(void);
47 __global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
48 
hd1(void)49 __host__ __device__ void hd1(void) {
50   // Expected 0-1 as in one of host/device side compilation it is an error,
51   // while not in the other
52   hd1d(); // expected-error 0-1 {{no matching function}}
53   hd1h(); // expected-error 0-1 {{no matching function}}
54 
55   // No errors as guarded
56 #ifdef __CUDA_ARCH__
57   hd1d();
58 #else
59   hd1h();
60 #endif
61 
62   // Errors as incorrectly guarded
63 #ifndef __CUDA_ARCH__
64   hd1dig(); // expected-error {{no matching function}}
65 #else
66   hd1hig(); // expected-error {{no matching function}}
67 #endif
68 
69   hd1hd();
70   hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
71 }
72