|
150
|
1 // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
|
|
|
2 // RUN: -verify -verify-ignore-unexpected=note
|
|
|
3 // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
|
|
|
4 // RUN: -verify -verify-ignore-unexpected=note -fopenmp
|
|
|
5
|
|
|
6 // Note: This test won't work with -fsyntax-only, because some of these errors
|
|
|
7 // are emitted during codegen.
|
|
|
8
|
|
|
9 #include "Inputs/cuda.h"
|
|
|
10
|
|
|
11 __device__ void device_fn() {}
|
|
|
12 // expected-note@-1 5 {{'device_fn' declared here}}
|
|
|
13
|
|
|
14 struct S {
|
|
|
15 __device__ S() {}
|
|
|
16 // expected-note@-1 2 {{'S' declared here}}
|
|
|
17 __device__ ~S() { device_fn(); }
|
|
|
18 // expected-note@-1 {{'~S' declared here}}
|
|
|
19 int x;
|
|
|
20 };
|
|
|
21
|
|
|
22 struct T {
|
|
|
23 __host__ __device__ void hd() { device_fn(); }
|
|
|
24 // expected-error@-1 {{reference to __device__ function 'device_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() { device_fn(); }
|
|
|
29
|
|
|
30 __host__ __device__ void hd3();
|
|
|
31
|
|
|
32 __device__ void d() {}
|
|
|
33 // expected-note@-1 {{'d' declared here}}
|
|
|
34 };
|
|
|
35
|
|
|
36 __host__ __device__ void T::hd3() {
|
|
|
37 device_fn();
|
|
|
38 // expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
|
|
|
39 }
|
|
|
40
|
|
|
41 template <typename T> __host__ __device__ void hd2() { device_fn(); }
|
|
|
42 // expected-error@-1 2 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
|
|
|
43 void host_fn() { hd2<int>(); }
|
|
|
44
|
|
|
45 __host__ __device__ void hd() { device_fn(); }
|
|
|
46 // expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
|
|
|
47
|
|
|
48 // No error because this is never instantiated.
|
|
|
49 template <typename T> __host__ __device__ void hd3() { device_fn(); }
|
|
|
50
|
|
|
51 __host__ __device__ void local_var() {
|
|
|
52 S s;
|
|
|
53 // expected-error@-1 {{reference to __device__ function 'S' in __host__ __device__ function}}
|
|
|
54 }
|
|
|
55
|
|
|
56 __host__ __device__ void placement_new(char *ptr) {
|
|
|
57 ::new(ptr) S();
|
|
|
58 // expected-error@-1 {{reference to __device__ function 'S' in __host__ __device__ function}}
|
|
|
59 }
|
|
|
60
|
|
|
61 __host__ __device__ void explicit_destructor(S *s) {
|
|
|
62 s->~S();
|
|
|
63 // expected-error@-1 {{reference to __device__ function '~S' in __host__ __device__ function}}
|
|
|
64 }
|
|
|
65
|
|
|
66 __host__ __device__ void hd_member_fn() {
|
|
|
67 T t;
|
|
|
68 // Necessary to trigger an error on T::hd. It's (implicitly) inline, so
|
|
|
69 // isn't codegen'ed until we call it.
|
|
|
70 t.hd();
|
|
|
71 }
|
|
|
72
|
|
|
73 __host__ __device__ void h_member_fn() {
|
|
|
74 T t;
|
|
|
75 t.d();
|
|
|
76 // expected-error@-1 {{reference to __device__ function 'd' in __host__ __device__ function}}
|
|
|
77 }
|
|
|
78
|
|
|
79 __host__ __device__ void fn_ptr() {
|
|
|
80 auto* ptr = &device_fn;
|
|
|
81 // expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
|
|
|
82 }
|
|
|
83
|
|
|
84 template <typename T>
|
|
|
85 __host__ __device__ void fn_ptr_template() {
|
|
|
86 auto* ptr = &device_fn; // Not an error because the template isn't instantiated.
|
|
|
87 }
|
|
|
88
|
|
|
89 // Launching a kernel from a host function does not result in code generation
|
|
|
90 // for it, so calling HD function which calls a D function should not trigger
|
|
|
91 // errors.
|
|
|
92 static __host__ __device__ void hd_func() { device_fn(); }
|
|
|
93 __global__ void kernel() { hd_func(); }
|
|
|
94 void host_func(void) { kernel<<<1, 1>>>(); }
|
|
|
95
|
|
|
96 // Should allow host function call kernel template with device function argument.
|
|
|
97 __device__ void f();
|
|
|
98 template<void(*F)()> __global__ void t() { F(); }
|
|
|
99 __host__ void g() { t<f><<<1,1>>>(); }
|
