Home | History | Annotate | Line # | Download | only in cuda_wrappers 
      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
     64 namespace 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 
     77 template <class __T, class __Cmp>
     78 __attribute__((enable_if(true, "")))
     79 inline _CPP14_CONSTEXPR __host__ __device__ const __T &
     80 max(const __T &__a, const __T &__b, __Cmp __cmp) {
     81   return __cmp(__a, __b) ? __b : __a;
     82 }
     83 
     84 template <class __T>
     85 __attribute__((enable_if(true, "")))
     86 inline _CPP14_CONSTEXPR __host__ __device__ const __T &
     87 max(const __T &__a, const __T &__b) {
     88   return __a < __b ? __b : __a;
     89 }
     90 
     91 template <class __T, class __Cmp>
     92 __attribute__((enable_if(true, "")))
     93 inline _CPP14_CONSTEXPR __host__ __device__ const __T &
     94 min(const __T &__a, const __T &__b, __Cmp __cmp) {
     95   return __cmp(__b, __a) ? __b : __a;
     96 }
     97 
     98 template <class __T>
     99 __attribute__((enable_if(true, "")))
    100 inline _CPP14_CONSTEXPR __host__ __device__ const __T &
    101 min(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