algorithm 3.5 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596
  1. /*===---- complex - 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. #ifndef __CLANG_CUDA_WRAPPERS_ALGORITHM
  24. #define __CLANG_CUDA_WRAPPERS_ALGORITHM
  25. // This header defines __device__ overloads of std::min/max, but only if we're
  26. // <= C++11. In C++14, these functions are constexpr, and so are implicitly
  27. // __host__ __device__.
  28. //
  29. // We don't support the initializer_list overloads because
  30. // initializer_list::begin() and end() are not __host__ __device__ functions.
  31. //
  32. // When compiling in C++14 mode, we could force std::min/max to have different
  33. // implementations for host and device, by declaring the device overloads
  34. // before the constexpr overloads appear. We choose not to do this because
  35. // a) why write our own implementation when we can use one from the standard
  36. // library? and
  37. // b) libstdc++ is evil and declares min/max inside a header that is included
  38. // *before* we include <algorithm>. So we'd have to unconditionally
  39. // declare our __device__ overloads of min/max, but that would pollute
  40. // things for people who choose not to include <algorithm>.
  41. #include_next <algorithm>
  42. #if __cplusplus <= 201103L
  43. // We need to define these overloads in exactly the namespace our standard
  44. // library uses (including the right inline namespace), otherwise they won't be
  45. // picked up by other functions in the standard library (e.g. functions in
  46. // <complex>). Thus the ugliness below.
  47. #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
  48. _LIBCPP_BEGIN_NAMESPACE_STD
  49. #else
  50. namespace std {
  51. #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
  52. _GLIBCXX_BEGIN_NAMESPACE_VERSION
  53. #endif
  54. #endif
  55. template <class __T, class __Cmp>
  56. inline __device__ const __T &
  57. max(const __T &__a, const __T &__b, __Cmp __cmp) {
  58. return __cmp(__a, __b) ? __b : __a;
  59. }
  60. template <class __T>
  61. inline __device__ const __T &
  62. max(const __T &__a, const __T &__b) {
  63. return __a < __b ? __b : __a;
  64. }
  65. template <class __T, class __Cmp>
  66. inline __device__ const __T &
  67. min(const __T &__a, const __T &__b, __Cmp __cmp) {
  68. return __cmp(__b, __a) ? __b : __a;
  69. }
  70. template <class __T>
  71. inline __device__ const __T &
  72. min(const __T &__a, const __T &__b) {
  73. return __a < __b ? __a : __b;
  74. }
  75. #ifdef _LIBCPP_END_NAMESPACE_STD
  76. _LIBCPP_END_NAMESPACE_STD
  77. #else
  78. #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
  79. _GLIBCXX_END_NAMESPACE_VERSION
  80. #endif
  81. } // namespace std
  82. #endif
  83. #endif // __cplusplus <= 201103L
  84. #endif // __CLANG_CUDA_WRAPPERS_ALGORITHM