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. HostOnly()23__host__ constexpr int HostOnly() { return 0; } 24 // expected-note@-1 0+ {{not viable}} DeviceOnly()25__device__ constexpr int DeviceOnly() { return 0; } 26 // expected-note@-1 0+ {{not viable}} 27 HostDevice()28constexpr int HostDevice() { return 0; } 29 30 // This should be a host-only function, because there's a previous __device__ 31 // overload in <overload.h>. OverloadMe()32constexpr HostReturnTy OverloadMe() { return HostReturnTy(); } 33 34 namespace ns { 35 // The "using" statement in overload.h should prevent OverloadMe from being 36 // implicitly host+device. OverloadMe()37constexpr HostReturnTy OverloadMe() { return HostReturnTy(); } 38 } // namespace ns 39 40 // This is an error, because NonSysHdrOverload was not defined in a system 41 // header. NonSysHdrOverload()42__device__ int NonSysHdrOverload() { return 0; } 43 // expected-note@-1 {{conflicting __device__ function declared here}} NonSysHdrOverload()44constexpr 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 HostFn()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 DeviceFn()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 HostDeviceFn()68__host__ __device__ void HostDeviceFn() { 69 #ifdef __CUDA_ARCH__ 70 int y = OverloadMe(); 71 #else 72 constexpr HostReturnTy y = OverloadMe(); 73 #endif 74 } 75