~shimon/dismaltc

ref: 461e0617cb03f1d5027af4fbb3da0722791022e9 dismaltc/clang/disable-test/SemaCUDA/host-device-constexpr.cu -rw-r--r-- 2.2 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
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s -fcuda-is-device

#include "Inputs/cuda.h"

// Declares one function and pulls it into namespace ns:
//
//   __device__ int OverloadMe();
//   namespace ns { using ::OverloadMe; }
//
// Clang cares that this is done in a system header.
#include <overload.h>

// Opaque type used to determine which overload we're invoking.
struct HostReturnTy {};

// These shouldn't become host+device because they already have attributes.
__host__ constexpr int HostOnly() { return 0; }
// expected-note@-1 0+ {{not viable}}
__device__ constexpr int DeviceOnly() { return 0; }
// expected-note@-1 0+ {{not viable}}

constexpr int HostDevice() { return 0; }

// This should be a host-only function, because there's a previous __device__
// overload in <overload.h>.
constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }

namespace ns {
// The "using" statement in overload.h should prevent OverloadMe from being
// implicitly host+device.
constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
}  // namespace ns

// This is an error, because NonSysHdrOverload was not defined in a system
// header.
__device__ int NonSysHdrOverload() { return 0; }
// expected-note@-1 {{conflicting __device__ function declared here}}
constexpr int NonSysHdrOverload() { return 0; }
// expected-error@-1 {{constexpr function 'NonSysHdrOverload' without __host__ or __device__ attributes}}

// Variadic device functions are not allowed, so this is just treated as
// host-only.
constexpr void Variadic(const char*, ...);
// expected-note@-1 {{call to __host__ function from __device__ function}}

__host__ void HostFn() {
  HostOnly();
  DeviceOnly(); // expected-error {{no matching function}}
  HostReturnTy x = OverloadMe();
  HostReturnTy y = ns::OverloadMe();
  Variadic("abc", 42);
}

__device__ void DeviceFn() {
  HostOnly(); // expected-error {{no matching function}}
  DeviceOnly();
  int x = OverloadMe();
  int y = ns::OverloadMe();
  Variadic("abc", 42); // expected-error {{no matching function}}
}

__host__ __device__ void HostDeviceFn() {
#ifdef __CUDA_ARCH__
  int y = OverloadMe();
#else
  constexpr HostReturnTy y = OverloadMe();
#endif
}