[CUDA] Make unattributed constexpr functions implicitly host+device.

With this patch, by a constexpr function is implicitly host+device
unless:

 a) it's a variadic function (variadic functions are not allowed on the
    device side), or
 b) it's preceeded by a __device__ overload in a system header.

The restriction on overloading __host__ __device__ functions on the
basis of their CUDA attributes remains in place, but we use (b) to allow
us to define __device__ overloads for constexpr functions in cmath,
which would otherwise be __host__ __device__ and thus not overloadable.

You can disable this behavior with -fno-cuda-host-device-constexpr.

Reviewers: tra, rnk, rsmith

Subscribers: cfe-commits

Differential Revision: http://reviews.llvm.org/D18380

llvm-svn: 264964

3 files changed, 97 insertions, 0 deletions

diff --git a/clang/test/SemaCUDA/Inputs/overload.h b/clang/test/SemaCUDA/Inputs/overload.h
new file mode 100644
index 00000000000..1c021f1ec57
--- /dev/null
+++ b/clang/test/SemaCUDA/Inputs/overload.h
@@ -0,0 +1,8 @@
+// This header is used by tests which are interested in __device__ functions
+// which appear in a system header.
+
+__device__ int OverloadMe();
+
+namespace ns {
+using ::OverloadMe;
+}
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
+}
diff --git a/clang/test/SemaCUDA/no-host-device-constexpr.cu b/clang/test/SemaCUDA/no-host-device-constexpr.cu
new file mode 100644
index 00000000000..c70d97d61e4
--- /dev/null
+++ b/clang/test/SemaCUDA/no-host-device-constexpr.cu
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fno-cuda-host-device-constexpr -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fno-cuda-host-device-constexpr -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+// Check that, with -fno-cuda-host-device-constexpr, constexpr functions are
+// host-only, and __device__ constexpr functions are still device-only.
+
+constexpr int f() { return 0; } // expected-note {{not viable}}
+__device__ constexpr int g() { return 0; } // expected-note {{not viable}}
+
+void __device__ foo() {
+  f(); // expected-error {{no matching function}}
+  g();
+}
+
+void __host__ foo() {
+  f();
+  g(); // expected-error {{no matching function}}
+}