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 
device_fn()11 __device__ void device_fn() {}
12 // expected-note@-1 5 {{'device_fn' declared here}}
13 
14 struct S {
SS15   __device__ S() {}
16   // expected-note@-1 2 {{'S' declared here}}
~SS17   __device__ ~S() { device_fn(); }
18   // expected-note@-1 {{'~S' declared here}}
19   int x;
20 };
21 
22 struct T {
hdT23   __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.
hd2T28   __host__ __device__ void hd2() { device_fn(); }
29 
30   __host__ __device__ void hd3();
31 
dT32   __device__ void d() {}
33   // expected-note@-1 {{'d' declared here}}
34 };
35 
hd3()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 
hd2()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}}
host_fn()43 void host_fn() { hd2<int>(); }
44 
hd()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.
hd3()49 template <typename T> __host__ __device__ void hd3() { device_fn(); }
50 
local_var()51 __host__ __device__ void local_var() {
52   S s;
53   // expected-error@-1 {{reference to __device__ function 'S' in __host__ __device__ function}}
54 }
55 
placement_new(char * ptr)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 
explicit_destructor(S * s)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 
hd_member_fn()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 
h_member_fn()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 
fn_ptr()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>
fn_ptr_template()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.
hd_func()92 static __host__ __device__ void hd_func() { device_fn(); }
kernel()93 __global__ void kernel() { hd_func(); }
host_func(void)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();
t()98 template<void(*F)()> __global__ void t() { F(); }
g()99 __host__ void g() { t<f><<<1,1>>>(); }
100