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;
}
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;
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
