1 // Test the Sema analysis of caller-callee relationships of host device
2 // functions when compiling CUDA code. There are 4 permutations of this test as
3 // host and device compilation are separate compilation passes, and clang has
4 // an option to allow host calls from host device functions. __CUDA_ARCH__ is
5 // defined when compiling for the device and TEST_WARN_HD when host calls are
6 // allowed from host device functions. So for example, if __CUDA_ARCH__ is
7 // defined and TEST_WARN_HD is not then device compilation is happening but
8 // host device functions are not allowed to call device functions.
9
10 // RUN: %clang_cc1 -fsyntax-only -verify %s
11 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
12 // RUN: %clang_cc1 -fsyntax-only -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
13 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
14
15 #include "Inputs/cuda.h"
16
17 __host__ void hd1h(void);
18 #if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
19 // expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
20 #endif
21 __device__ void hd1d(void);
22 #ifndef __CUDA_ARCH__
23 // expected-note@-2 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
24 #endif
25 __host__ void hd1hg(void);
26 __device__ void hd1dg(void);
27 #ifdef __CUDA_ARCH__
28 __host__ void hd1hig(void);
29 #if !defined(TEST_WARN_HD)
30 // expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
31 #endif
32 #else
33 __device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
34 #endif
35 __host__ __device__ void hd1hd(void);
36 __global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
37
hd1(void)38 __host__ __device__ void hd1(void) {
39 #if defined(TEST_WARN_HD) && defined(__CUDA_ARCH__)
40 // expected-warning@-2 {{calling __host__ function hd1h from __host__ __device__ function hd1}}
41 // expected-warning@-3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}}
42 #endif
43 hd1d();
44 #ifndef __CUDA_ARCH__
45 // expected-error@-2 {{no matching function}}
46 #endif
47 hd1h();
48 #if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
49 // expected-error@-2 {{no matching function}}
50 #endif
51
52 // No errors as guarded
53 #ifdef __CUDA_ARCH__
54 hd1d();
55 #else
56 hd1h();
57 #endif
58
59 // Errors as incorrectly guarded
60 #ifndef __CUDA_ARCH__
61 hd1dig(); // expected-error {{no matching function}}
62 #else
63 hd1hig();
64 #ifndef TEST_WARN_HD
65 // expected-error@-2 {{no matching function}}
66 #endif
67 #endif
68
69 hd1hd();
70 hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
71 }
72