From: Justin Lebar <jlebar@google.com>
Date: Fri, 29 Jun 2018 22:27:56 +0000 (+0000)
Subject: [CUDA] Make min/max shims host+device.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=4a70ae36838d67e1bfce5b182ba96a1bd1496756;p=clang

[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 8ee84f7fb0..d8658bf1c7 100644
--- 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;
 }