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

[Justin Lebar via Phabricator via cfe-commits]

jlebar created this revision.
jlebar added a reviewer: rsmith.
Herald added a subscriber: sanjoy.

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

Testcase in a separate test-suite commit.


https://reviews.llvm.org/D48036

Files:
  clang/lib/Headers/cuda_wrappers/algorithm


Index: clang/lib/Headers/cuda_wrappers/algorithm
===================================================================
--- clang/lib/Headers/cuda_wrappers/algorithm
+++ clang/lib/Headers/cuda_wrappers/algorithm
@@ -69,28 +69,28 @@
 
 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;
 }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D48036.150790.patch
Type: text/x-patch
Size: 1125 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20180611/b2fa04da/attachment-0001.bin>