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