1 // RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify %s
2 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
3
4 #include "Inputs/cuda.h"
5
6 #ifndef __CUDA_ARCH__
7 // expected-no-diagnostics
8 #endif
9
10 // When compiling for device, foo()'s call to host_fn() is an error, because
11 // foo() is known-emitted.
12 //
13 // The trickiness here comes from the fact that the FunctionDecl bar() sees
14 // foo() does not have the "inline" keyword, so we might incorrectly think that
15 // foo() is a priori known-emitted. This would prevent us from marking foo()
16 // as known-emitted when we see the call from bar() to foo(), which would
17 // prevent us from emitting an error for foo()'s call to host_fn() when we
18 // eventually see it.
19
host_fn()20 void host_fn() {}
21 #ifdef __CUDA_ARCH__
22 // expected-note@-2 {{declared here}}
23 #endif
24
25 __host__ __device__ void foo();
bar()26 __device__ void bar() {
27 foo();
28 #ifdef __CUDA_ARCH__
29 // expected-note@-2 {{called by 'bar'}}
30 #endif
31 }
foo()32 inline __host__ __device__ void foo() {
33 host_fn();
34 #ifdef __CUDA_ARCH__
35 // expected-error@-2 {{reference to __host__ function}}
36 #endif
37 }
38
39 // This is similar to the above, except there's no error here. This code used
40 // to trip an assertion due to us noticing, when emitting the definition of
41 // boom(), that T::operator S() was (incorrectly) considered a priori
42 // known-emitted.
43 struct S {};
44 struct T {
45 __device__ operator S() const;
46 };
operator S() const47 __device__ inline T::operator S() const { return S(); }
48
49 __device__ T t;
boom()50 __device__ void boom() {
51 S s = t;
52 }
53