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