1/*===---- algorithm - CUDA wrapper for <algorithm> -------------------------=== 2 * 3 * Permission is hereby granted, free of charge, to any person obtaining a copy 4 * of this software and associated documentation files (the "Software"), to deal 5 * in the Software without restriction, including without limitation the rights 6 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 7 * copies of the Software, and to permit persons to whom the Software is 8 * furnished to do so, subject to the following conditions: 9 * 10 * The above copyright notice and this permission notice shall be included in 11 * all copies or substantial portions of the Software. 12 * 13 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 14 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 15 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 16 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 17 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 18 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 19 * THE SOFTWARE. 20 * 21 *===-----------------------------------------------------------------------=== 22 */ 23 24#ifndef __CLANG_CUDA_WRAPPERS_ALGORITHM 25#define __CLANG_CUDA_WRAPPERS_ALGORITHM 26 27// This header defines __device__ overloads of std::min/max. 28// 29// Ideally we'd declare these functions only if we're <= C++11. In C++14, 30// these functions are constexpr, and so are implicitly __host__ __device__. 31// 32// However, the compiler being in C++14 mode does not imply that the standard 33// library supports C++14. There is no macro we can test to check that the 34// stdlib has constexpr std::min/max. Thus we have to unconditionally define 35// our device overloads. 36// 37// A host+device function cannot be overloaded, and a constexpr function 38// implicitly become host device if there's no explicitly host or device 39// overload preceding it. So the simple thing to do would be to declare our 40// device min/max overloads, and then #include_next <algorithm>. This way our 41// device overloads would come first, and so if we have a C++14 stdlib, its 42// min/max won't become host+device and conflict with our device overloads. 43// 44// But that also doesn't work. libstdc++ is evil and declares std::min/max in 45// an internal header that is included *before* <algorithm>. Thus by the time 46// we're inside of this file, std::min/max may already have been declared, and 47// thus we can't prevent them from becoming host+device if they're constexpr. 48// 49// Therefore we perpetrate the following hack: We mark our __device__ overloads 50// with __attribute__((enable_if(true, ""))). This causes the signature of the 51// function to change without changing anything else about it. (Except that 52// overload resolution will prefer it over the __host__ __device__ version 53// rather than considering them equally good). 54 55#include_next <algorithm> 56 57// We need to define these overloads in exactly the namespace our standard 58// library uses (including the right inline namespace), otherwise they won't be 59// picked up by other functions in the standard library (e.g. functions in 60// <complex>). Thus the ugliness below. 61#ifdef _LIBCPP_BEGIN_NAMESPACE_STD 62_LIBCPP_BEGIN_NAMESPACE_STD 63#else 64namespace std { 65#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION 66_GLIBCXX_BEGIN_NAMESPACE_VERSION 67#endif 68#endif 69 70#pragma push_macro("_CPP14_CONSTEXPR") 71#if __cplusplus >= 201402L 72#define _CPP14_CONSTEXPR constexpr 73#else 74#define _CPP14_CONSTEXPR 75#endif 76 77template <class __T, class __Cmp> 78__attribute__((enable_if(true, ""))) 79inline _CPP14_CONSTEXPR __host__ __device__ const __T & 80max(const __T &__a, const __T &__b, __Cmp __cmp) { 81 return __cmp(__a, __b) ? __b : __a; 82} 83 84template <class __T> 85__attribute__((enable_if(true, ""))) 86inline _CPP14_CONSTEXPR __host__ __device__ const __T & 87max(const __T &__a, const __T &__b) { 88 return __a < __b ? __b : __a; 89} 90 91template <class __T, class __Cmp> 92__attribute__((enable_if(true, ""))) 93inline _CPP14_CONSTEXPR __host__ __device__ const __T & 94min(const __T &__a, const __T &__b, __Cmp __cmp) { 95 return __cmp(__b, __a) ? __b : __a; 96} 97 98template <class __T> 99__attribute__((enable_if(true, ""))) 100inline _CPP14_CONSTEXPR __host__ __device__ const __T & 101min(const __T &__a, const __T &__b) { 102 return __a < __b ? __a : __b; 103} 104 105#pragma pop_macro("_CPP14_CONSTEXPR") 106 107#ifdef _LIBCPP_END_NAMESPACE_STD 108_LIBCPP_END_NAMESPACE_STD 109#else 110#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION 111_GLIBCXX_END_NAMESPACE_VERSION 112#endif 113} // namespace std 114#endif 115 116#endif // __CLANG_CUDA_WRAPPERS_ALGORITHM 117