1 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -fsyntax-only -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s
2 // RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s
3 
4 #include "Inputs/cuda.h"
5 
6 __device__ void device_fn() {
7   auto f1 = [&] {};
8   f1(); // implicitly __device__
9 
10   auto f2 = [&] __device__ {};
11   f2();
12 
13   auto f3 = [&] __host__ {};
14   f3();  // expected-error {{no matching function}}
15 
16   auto f4 = [&] __host__ __device__ {};
17   f4();
18 
19   // Now do it all again with '()'s in the lambda declarations: This is a
20   // different parse path.
21   auto g1 = [&]() {};
22   g1(); // implicitly __device__
23 
24   auto g2 = [&]() __device__ {};
25   g2();
26 
27   auto g3 = [&]() __host__ {};
28   g3();  // expected-error {{no matching function}}
29 
30   auto g4 = [&]() __host__ __device__ {};
31   g4();
32 
33   // Once more, with the '()'s in a different place.
34   auto h1 = [&]() {};
35   h1(); // implicitly __device__
36 
37   auto h2 = [&] __device__ () {};
38   h2();
39 
40   auto h3 = [&] __host__ () {};
41   h3();  // expected-error {{no matching function}}
42 
43   auto h4 = [&] __host__ __device__ () {};
44   h4();
45 }
46 
47 // Behaves identically to device_fn.
48 __global__ void kernel_fn() {
49   auto f1 = [&] {};
50   f1(); // implicitly __device__
51 
52   auto f2 = [&] __device__ {};
53   f2();
54 
55   auto f3 = [&] __host__ {};
56   f3();  // expected-error {{no matching function}}
57 
58   auto f4 = [&] __host__ __device__ {};
59   f4();
60 
61   // No need to re-test all the parser contortions we test in the device
62   // function.
63 }
64 
65 __host__ void host_fn() {
66   auto f1 = [&] {};
67   f1(); // implicitly __host__ (i.e., no magic)
68 
69   auto f2 = [&] __device__ {};
70   f2();  // expected-error {{no matching function}}
71 
72   auto f3 = [&] __host__ {};
73   f3();
74 
75   auto f4 = [&] __host__ __device__ {};
76   f4();
77 }
78 
79 __host__ __device__ void hd_fn() {
80   auto f1 = [&] {};
81   f1(); // implicitly __host__ __device__
82 
83   auto f2 = [&] __device__ {};
84   f2();
85 #ifndef __CUDA_ARCH__
86   // expected-error@-2 {{reference to __device__ function}}
87 #endif
88 
89   auto f3 = [&] __host__ {};
90   f3();
91 #ifdef __CUDA_ARCH__
92   // expected-error@-2 {{reference to __host__ function}}
93 #endif
94 
95   auto f4 = [&] __host__ __device__ {};
96   f4();
97 }
98 
99 // The special treatment above only applies to lambdas.
100 __device__ void foo() {
101   struct X {
102     void foo() {}
103   };
104   X x;
105   x.foo(); // expected-error {{reference to __host__ function 'foo' in __device__ function}}
106 }
107