[CUDA] Make min/max shims host+device.

Summary:
Fixes PR37753: min/max can't be called from __host__ __device__
functions in C++14 mode.

Testcase in a separate test-suite commit.

Reviewers: rsmith

Subscribers: sanjoy, lahwaacz, cfe-commits

Differential Revision: https://reviews.llvm.org/D48036

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@336025 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Headers/cuda_wrappers/algorithm b/lib/Headers/cuda_wrappers/algorithm
index 8ee84f7fb02c822ef6035e8192a7d5b3d34b52a0..d8658bf1c7d8fe5fb52d0537836c07e50727c00e 100644 (file)
--- a/lib/Headers/cuda_wrappers/algorithm
+++ b/lib/Headers/cuda_wrappers/algorithm
@@ -69,28 +69,28 @@ _GLIBCXX_BEGIN_NAMESPACE_VERSION
 
 template <class __T, class __Cmp>
 __attribute__((enable_if(true, "")))
-inline __device__ const __T &
+inline __host__ __device__ const __T &
 max(const __T &__a, const __T &__b, __Cmp __cmp) {
   return __cmp(__a, __b) ? __b : __a;
 }
 
 template <class __T>
 __attribute__((enable_if(true, "")))
-inline __device__ const __T &
+inline __host__ __device__ const __T &
 max(const __T &__a, const __T &__b) {
   return __a < __b ? __b : __a;
 }
 
 template <class __T, class __Cmp>
 __attribute__((enable_if(true, "")))
-inline __device__ const __T &
+inline __host__ __device__ const __T &
 min(const __T &__a, const __T &__b, __Cmp __cmp) {
   return __cmp(__b, __a) ? __b : __a;
 }
 
 template <class __T>
 __attribute__((enable_if(true, "")))
-inline __device__ const __T &
+inline __host__ __device__ const __T &
 min(const __T &__a, const __T &__b) {
   return __a < __b ? __a : __b;
 }