1 | // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \ |
2 | // RUN: -emit-llvm -o /dev/null -verify -verify-ignore-unexpected=note |
3 | |
4 | // Note: This test won't work with -fsyntax-only, because some of these errors |
5 | // are emitted during codegen. |
6 | |
7 | #include "Inputs/cuda.h" |
8 | |
9 | extern "C" void host_fn() {} |
10 | // expected-note@-1 7 {{'host_fn' declared here}} |
11 | |
12 | struct Dummy {}; |
13 | |
14 | struct S { |
15 | S() {} |
16 | // expected-note@-1 2 {{'S' declared here}} |
17 | ~S() { host_fn(); } |
18 | // expected-note@-1 {{'~S' declared here}} |
19 | int x; |
20 | }; |
21 | |
22 | struct T { |
23 | __host__ __device__ void hd() { host_fn(); } |
24 | // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} |
25 | |
26 | // No error; this is (implicitly) inline and is never called, so isn't |
27 | // codegen'ed. |
28 | __host__ __device__ void hd2() { host_fn(); } |
29 | |
30 | __host__ __device__ void hd3(); |
31 | |
32 | void h() {} |
33 | // expected-note@-1 {{'h' declared here}} |
34 | |
35 | void operator+(); |
36 | // expected-note@-1 {{'operator+' declared here}} |
37 | |
38 | void operator-(const T&) {} |
39 | // expected-note@-1 {{'operator-' declared here}} |
40 | |
41 | operator Dummy() { return Dummy(); } |
42 | // expected-note@-1 {{'operator Dummy' declared here}} |
43 | |
44 | __host__ void operator delete(void *) { host_fn(); }; |
45 | __device__ void operator delete(void*, __SIZE_TYPE__); |
46 | }; |
47 | |
48 | struct U { |
49 | __device__ void operator delete(void*, __SIZE_TYPE__) = delete; |
50 | __host__ __device__ void operator delete(void*); |
51 | }; |
52 | |
53 | __host__ __device__ void T::hd3() { |
54 | host_fn(); |
55 | // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} |
56 | } |
57 | |
58 | template <typename T> __host__ __device__ void hd2() { host_fn(); } |
59 | // expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} |
60 | __global__ void kernel() { hd2<int>(); } |
61 | |
62 | __host__ __device__ void hd() { host_fn(); } |
63 | // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} |
64 | |
65 | template <typename T> __host__ __device__ void hd3() { host_fn(); } |
66 | // expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} |
67 | __device__ void device_fn() { hd3<int>(); } |
68 | |
69 | // No error because this is never instantiated. |
70 | template <typename T> __host__ __device__ void hd4() { host_fn(); } |
71 | |
72 | __host__ __device__ void local_var() { |
73 | S s; |
74 | // expected-error@-1 {{reference to __host__ function 'S' in __host__ __device__ function}} |
75 | } |
76 | |
77 | __host__ __device__ void placement_new(char *ptr) { |
78 | ::new(ptr) S(); |
79 | // expected-error@-1 {{reference to __host__ function 'S' in __host__ __device__ function}} |
80 | } |
81 | |
82 | __host__ __device__ void explicit_destructor(S *s) { |
83 | s->~S(); |
84 | // expected-error@-1 {{reference to __host__ function '~S' in __host__ __device__ function}} |
85 | } |
86 | |
87 | __host__ __device__ void class_specific_delete(T *t, U *u) { |
88 | delete t; // ok, call sized device delete even though host has preferable non-sized version |
89 | delete u; // ok, call non-sized HD delete rather than sized D delete |
90 | } |
91 | |
92 | __host__ __device__ void hd_member_fn() { |
93 | T t; |
94 | // Necessary to trigger an error on T::hd. It's (implicitly) inline, so |
95 | // isn't codegen'ed until we call it. |
96 | t.hd(); |
97 | } |
98 | |
99 | __host__ __device__ void h_member_fn() { |
100 | T t; |
101 | t.h(); |
102 | // expected-error@-1 {{reference to __host__ function 'h' in __host__ __device__ function}} |
103 | } |
104 | |
105 | __host__ __device__ void fn_ptr() { |
106 | auto* ptr = &host_fn; |
107 | // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} |
108 | } |
109 | |
110 | template <typename T> |
111 | __host__ __device__ void fn_ptr_template() { |
112 | auto* ptr = &host_fn; // Not an error because the template isn't instantiated. |
113 | } |
114 | |
115 | __host__ __device__ void unaryOp() { |
116 | T t; |
117 | (void) +t; // expected-error {{reference to __host__ function 'operator+' in __host__ __device__ function}} |
118 | } |
119 | |
120 | __host__ __device__ void binaryOp() { |
121 | T t; |
122 | (void) (t - t); // expected-error {{reference to __host__ function 'operator-' in __host__ __device__ function}} |
123 | } |
124 | |
125 | __host__ __device__ void implicitConversion() { |
126 | T t; |
127 | Dummy d = t; // expected-error {{reference to __host__ function 'operator Dummy' in __host__ __device__ function}} |
128 | } |
129 | |
130 | template <typename T> |
131 | struct TmplStruct { |
132 | template <typename U> __host__ __device__ void fn() {} |
133 | }; |
134 | |
135 | template <> |
136 | template <> |
137 | __host__ __device__ void TmplStruct<int>::fn<int>() { host_fn(); } |
138 | // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} |
139 | |
140 | __device__ void double_specialization() { TmplStruct<int>().fn<int>(); } |
141 | |