|
150
|
1 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
|
|
|
2 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s \
|
|
|
3 // RUN: -fcuda-is-device
|
|
|
4 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
|
|
|
5 // RUN: -fopenmp %s
|
|
|
6 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs \
|
|
|
7 // RUN: -fopenmp %s -fcuda-is-device
|
|
|
8
|
|
|
9 #include "Inputs/cuda.h"
|
|
|
10
|
|
|
11 // Declares one function and pulls it into namespace ns:
|
|
|
12 //
|
|
|
13 // __device__ int OverloadMe();
|
|
|
14 // namespace ns { using ::OverloadMe; }
|
|
|
15 //
|
|
|
16 // Clang cares that this is done in a system header.
|
|
|
17 #include <overload.h>
|
|
|
18
|
|
|
19 // Opaque type used to determine which overload we're invoking.
|
|
|
20 struct HostReturnTy {};
|
|
|
21
|
|
|
22 // These shouldn't become host+device because they already have attributes.
|
|
|
23 __host__ constexpr int HostOnly() { return 0; }
|
|
|
24 // expected-note@-1 0+ {{not viable}}
|
|
|
25 __device__ constexpr int DeviceOnly() { return 0; }
|
|
|
26 // expected-note@-1 0+ {{not viable}}
|
|
|
27
|
|
|
28 constexpr int HostDevice() { return 0; }
|
|
|
29
|
|
|
30 // This should be a host-only function, because there's a previous __device__
|
|
|
31 // overload in <overload.h>.
|
|
|
32 constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
|
|
|
33
|
|
|
34 namespace ns {
|
|
|
35 // The "using" statement in overload.h should prevent OverloadMe from being
|
|
|
36 // implicitly host+device.
|
|
|
37 constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
|
|
|
38 } // namespace ns
|
|
|
39
|
|
|
40 // This is an error, because NonSysHdrOverload was not defined in a system
|
|
|
41 // header.
|
|
|
42 __device__ int NonSysHdrOverload() { return 0; }
|
|
|
43 // expected-note@-1 {{conflicting __device__ function declared here}}
|
|
|
44 constexpr int NonSysHdrOverload() { return 0; }
|
|
|
45 // expected-error@-1 {{constexpr function 'NonSysHdrOverload' without __host__ or __device__ attributes}}
|
|
|
46
|
|
|
47 // Variadic device functions are not allowed, so this is just treated as
|
|
|
48 // host-only.
|
|
|
49 constexpr void Variadic(const char*, ...);
|
|
|
50 // expected-note@-1 {{call to __host__ function from __device__ function}}
|
|
|
51
|
|
|
52 __host__ void HostFn() {
|
|
|
53 HostOnly();
|
|
|
54 DeviceOnly(); // expected-error {{no matching function}}
|
|
|
55 HostReturnTy x = OverloadMe();
|
|
|
56 HostReturnTy y = ns::OverloadMe();
|
|
|
57 Variadic("abc", 42);
|
|
|
58 }
|
|
|
59
|
|
|
60 __device__ void DeviceFn() {
|
|
|
61 HostOnly(); // expected-error {{no matching function}}
|
|
|
62 DeviceOnly();
|
|
|
63 int x = OverloadMe();
|
|
|
64 int y = ns::OverloadMe();
|
|
|
65 Variadic("abc", 42); // expected-error {{no matching function}}
|
|
|
66 }
|
|
|
67
|
|
|
68 __host__ __device__ void HostDeviceFn() {
|
|
|
69 #ifdef __CUDA_ARCH__
|
|
|
70 int y = OverloadMe();
|
|
|
71 #else
|
|
|
72 constexpr HostReturnTy y = OverloadMe();
|
|
|
73 #endif
|
|
|
74 }
|
