~shimon/dismaltc

ref: 461e0617cb03f1d5027af4fbb3da0722791022e9 dismaltc/clang/disable-test/SemaCUDA/implicit-device-lambda.cu -rw-r--r-- 2.4 KiB
461e0617Itoh Shimon re-enable tests. 1 year, 6 months ago
                                                                                
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
// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -verify -fsyntax-only -verify-ignore-unexpected=warning -verify-ignore-unexpected=note %s
// RUN: %clang_cc1 -std=c++11 -verify -fsyntax-only -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(); // expected-error {{reference to __host__ function 'foo' in __device__ function}}
}