~shimon/dismaltc

ref: 461e0617cb03f1d5027af4fbb3da0722791022e9 dismaltc/clang/disable-test/SemaCUDA/add-inline-in-definition.cu -rw-r--r-- 1.5 KiB
461e0617Itoh Shimon re-enable tests. 1 year, 6 months ago
                                                                                
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -fsyntax-only -verify %s
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s

#include "Inputs/cuda.h"

#ifndef __CUDA_ARCH__
// expected-no-diagnostics
#endif

// When compiling for device, foo()'s call to host_fn() is an error, because
// foo() is known-emitted.
//
// The trickiness here comes from the fact that the FunctionDecl bar() sees
// foo() does not have the "inline" keyword, so we might incorrectly think that
// foo() is a priori known-emitted.  This would prevent us from marking foo()
// as known-emitted when we see the call from bar() to foo(), which would
// prevent us from emitting an error for foo()'s call to host_fn() when we
// eventually see it.

void host_fn() {}
#ifdef __CUDA_ARCH__
  // expected-note@-2 {{declared here}}
#endif

__host__ __device__ void foo();
__device__ void bar() {
  foo();
#ifdef __CUDA_ARCH__
  // expected-note@-2 {{called by 'bar'}}
#endif
}
inline __host__ __device__ void foo() {
  host_fn();
#ifdef __CUDA_ARCH__
  // expected-error@-2 {{reference to __host__ function}}
#endif
}

// This is similar to the above, except there's no error here.  This code used
// to trip an assertion due to us noticing, when emitting the definition of
// boom(), that T::operator S() was (incorrectly) considered a priori
// known-emitted.
struct S {};
struct T {
  __device__ operator S() const;
};
__device__ inline T::operator S() const { return S(); }

__device__ T t;
__device__ void boom() {
  S s = t;
}