| 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 | |
|---|