117 lines · plain
1/*===---- algorithm - CUDA wrapper for <algorithm> -------------------------===2 *3 * Permission is hereby granted, free of charge, to any person obtaining a copy4 * of this software and associated documentation files (the "Software"), to deal5 * in the Software without restriction, including without limitation the rights6 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell7 * copies of the Software, and to permit persons to whom the Software is8 * furnished to do so, subject to the following conditions:9 *10 * The above copyright notice and this permission notice shall be included in11 * all copies or substantial portions of the Software.12 *13 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR14 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,15 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE16 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER17 * 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 IN19 * THE SOFTWARE.20 *21 *===-----------------------------------------------------------------------===22 */23 24#ifndef __CLANG_CUDA_WRAPPERS_ALGORITHM25#define __CLANG_CUDA_WRAPPERS_ALGORITHM26 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 standard33// library supports C++14. There is no macro we can test to check that the34// stdlib has constexpr std::min/max. Thus we have to unconditionally define35// our device overloads.36//37// A host+device function cannot be overloaded, and a constexpr function38// implicitly become host device if there's no explicitly host or device39// overload preceding it. So the simple thing to do would be to declare our40// device min/max overloads, and then #include_next <algorithm>. This way our41// device overloads would come first, and so if we have a C++14 stdlib, its42// 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 in45// an internal header that is included *before* <algorithm>. Thus by the time46// we're inside of this file, std::min/max may already have been declared, and47// 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__ overloads50// with __attribute__((enable_if(true, ""))). This causes the signature of the51// function to change without changing anything else about it. (Except that52// overload resolution will prefer it over the __host__ __device__ version53// rather than considering them equally good).54 55#include_next <algorithm>56 57// We need to define these overloads in exactly the namespace our standard58// library uses (including the right inline namespace), otherwise they won't be59// picked up by other functions in the standard library (e.g. functions in60// <complex>). Thus the ugliness below.61#ifdef _LIBCPP_BEGIN_NAMESPACE_STD62_LIBCPP_BEGIN_NAMESPACE_STD63#else64namespace std {65#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION66_GLIBCXX_BEGIN_NAMESPACE_VERSION67#endif68#endif69 70#pragma push_macro("_CPP14_CONSTEXPR")71#if __cplusplus >= 201402L72#define _CPP14_CONSTEXPR constexpr73#else74#define _CPP14_CONSTEXPR75#endif76 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 __b < __a ? __b : __a;103}104 105#pragma pop_macro("_CPP14_CONSTEXPR")106 107#ifdef _LIBCPP_END_NAMESPACE_STD108_LIBCPP_END_NAMESPACE_STD109#else110#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION111_GLIBCXX_END_NAMESPACE_VERSION112#endif113} // namespace std114#endif115 116#endif // __CLANG_CUDA_WRAPPERS_ALGORITHM117