llvm-project

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

20
void host_fn() {}
21
#ifdef __CUDA_ARCH__
22
  // expected-note@-2 {{declared here}}
23
#endif
24

25
__host__ __device__ void foo();
26
__device__ void bar() {
27
  foo();
28
#ifdef __CUDA_ARCH__
29
  // expected-note@-2 {{called by 'bar'}}
30
#endif
31
}
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
};
47
__device__ inline T::operator S() const { return S(); }
48

49
__device__ T t;
50
__device__ void boom() {
51
  S s = t;
52
}
53