diff options
Diffstat (limited to 'lib/Headers/cuda_wrappers/algorithm')
-rw-r--r-- | lib/Headers/cuda_wrappers/algorithm | 96 |
1 files changed, 96 insertions, 0 deletions
diff --git a/lib/Headers/cuda_wrappers/algorithm b/lib/Headers/cuda_wrappers/algorithm new file mode 100644 index 000000000000..95d9beb73c68 --- /dev/null +++ b/lib/Headers/cuda_wrappers/algorithm @@ -0,0 +1,96 @@ +/*===---- complex - CUDA wrapper for <algorithm> ----------------------------=== + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + *===-----------------------------------------------------------------------=== + */ + +#ifndef __CLANG_CUDA_WRAPPERS_ALGORITHM +#define __CLANG_CUDA_WRAPPERS_ALGORITHM + +// This header defines __device__ overloads of std::min/max, but only if we're +// <= C++11. In C++14, these functions are constexpr, and so are implicitly +// __host__ __device__. +// +// We don't support the initializer_list overloads because +// initializer_list::begin() and end() are not __host__ __device__ functions. +// +// When compiling in C++14 mode, we could force std::min/max to have different +// implementations for host and device, by declaring the device overloads +// before the constexpr overloads appear. We choose not to do this because + +// a) why write our own implementation when we can use one from the standard +// library? and +// b) libstdc++ is evil and declares min/max inside a header that is included +// *before* we include <algorithm>. So we'd have to unconditionally +// declare our __device__ overloads of min/max, but that would pollute +// things for people who choose not to include <algorithm>. + +#include_next <algorithm> + +#if __cplusplus <= 201103L + +// We need to define these overloads in exactly the namespace our standard +// library uses (including the right inline namespace), otherwise they won't be +// picked up by other functions in the standard library (e.g. functions in +// <complex>). Thus the ugliness below. +#ifdef _LIBCPP_BEGIN_NAMESPACE_STD +_LIBCPP_BEGIN_NAMESPACE_STD +#else +namespace std { +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_BEGIN_NAMESPACE_VERSION +#endif +#endif + +template <class __T, class __Cmp> +inline __device__ const __T & +max(const __T &__a, const __T &__b, __Cmp __cmp) { + return __cmp(__a, __b) ? __b : __a; +} + +template <class __T> +inline __device__ const __T & +max(const __T &__a, const __T &__b) { + return __a < __b ? __b : __a; +} + +template <class __T, class __Cmp> +inline __device__ const __T & +min(const __T &__a, const __T &__b, __Cmp __cmp) { + return __cmp(__b, __a) ? __b : __a; +} + +template <class __T> +inline __device__ const __T & +min(const __T &__a, const __T &__b) { + return __a < __b ? __b : __a; +} + +#ifdef _LIBCPP_END_NAMESPACE_STD +_LIBCPP_END_NAMESPACE_STD +#else +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_END_NAMESPACE_VERSION +#endif +} // namespace std +#endif + +#endif // __cplusplus <= 201103L +#endif // __CLANG_CUDA_WRAPPERS_ALGORITHM |