1 files changed, 69 insertions, 0 deletions
diff --git a/clang/test/SemaCUDA/host-device-constexpr.cu b/clang/test/SemaCUDA/host-device-constexpr.cu
new file mode 100644
index 00000000000..6625d722c19
--- /dev/null
+++ b/clang/test/SemaCUDA/host-device-constexpr.cu
@@ -0,0 +1,69 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s -fcuda-is-device
+
+#include "Inputs/cuda.h"
+
+// Declares one function and pulls it into namespace ns:
+//
+//   __device__ int OverloadMe();
+//   namespace ns { using ::OverloadMe; }
+//
+// Clang cares that this is done in a system header.
+#include <overload.h>
+
+// Opaque type used to determine which overload we're invoking.
+struct HostReturnTy {};
+
+// These shouldn't become host+device because they already have attributes.
+__host__ constexpr int HostOnly() { return 0; }
+// expected-note@-1 0+ {{not viable}}
+__device__ constexpr int DeviceOnly() { return 0; }
+// expected-note@-1 0+ {{not viable}}
+
+constexpr int HostDevice() { return 0; }
+
+// This should be a host-only function, because there's a previous __device__
+// overload in <overload.h>.
+constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
+
+namespace ns {
+// The "using" statement in overload.h should prevent OverloadMe from being
+// implicitly host+device.
+constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
+}  // namespace ns
+
+// This is an error, because NonSysHdrOverload was not defined in a system
+// header.
+__device__ int NonSysHdrOverload() { return 0; }
+// expected-note@-1 {{conflicting __device__ function declared here}}
+constexpr int NonSysHdrOverload() { return 0; }
+// expected-error@-1 {{constexpr function 'NonSysHdrOverload' without __host__ or __device__ attributes}}
+
+// Variadic device functions are not allowed, so this is just treated as
+// host-only.
+constexpr void Variadic(const char*, ...);
+// expected-note@-1 {{call to __host__ function from __device__ function}}
+
+__host__ void HostFn() {
+  HostOnly();
+  DeviceOnly(); // expected-error {{no matching function}}
+  HostReturnTy x = OverloadMe();
+  HostReturnTy y = ns::OverloadMe();
+  Variadic("abc", 42);
+}
+
+__device__ void DeviceFn() {
+  HostOnly(); // expected-error {{no matching function}}
+  DeviceOnly();
+  int x = OverloadMe();
+  int y = ns::OverloadMe();
+  Variadic("abc", 42); // expected-error {{no matching function}}
+}
+
+__host__ __device__ void HostDeviceFn() {
+#ifdef __CUDA_ARCH__
+  int y = OverloadMe();
+#else
+  constexpr HostReturnTy y = OverloadMe();
+#endif
+}