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
