1e8d8bef9SDimitry Andric/*===---- algorithm - CUDA wrapper for <algorithm> -------------------------=== 20b57cec5SDimitry Andric * 30b57cec5SDimitry Andric * Permission is hereby granted, free of charge, to any person obtaining a copy 40b57cec5SDimitry Andric * of this software and associated documentation files (the "Software"), to deal 50b57cec5SDimitry Andric * in the Software without restriction, including without limitation the rights 60b57cec5SDimitry Andric * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 70b57cec5SDimitry Andric * copies of the Software, and to permit persons to whom the Software is 80b57cec5SDimitry Andric * furnished to do so, subject to the following conditions: 90b57cec5SDimitry Andric * 100b57cec5SDimitry Andric * The above copyright notice and this permission notice shall be included in 110b57cec5SDimitry Andric * all copies or substantial portions of the Software. 120b57cec5SDimitry Andric * 130b57cec5SDimitry Andric * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 140b57cec5SDimitry Andric * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 150b57cec5SDimitry Andric * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 160b57cec5SDimitry Andric * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 170b57cec5SDimitry Andric * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 180b57cec5SDimitry Andric * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 190b57cec5SDimitry Andric * THE SOFTWARE. 200b57cec5SDimitry Andric * 210b57cec5SDimitry Andric *===-----------------------------------------------------------------------=== 220b57cec5SDimitry Andric */ 230b57cec5SDimitry Andric 240b57cec5SDimitry Andric#ifndef __CLANG_CUDA_WRAPPERS_ALGORITHM 250b57cec5SDimitry Andric#define __CLANG_CUDA_WRAPPERS_ALGORITHM 260b57cec5SDimitry Andric 270b57cec5SDimitry Andric// This header defines __device__ overloads of std::min/max. 280b57cec5SDimitry Andric// 290b57cec5SDimitry Andric// Ideally we'd declare these functions only if we're <= C++11. In C++14, 300b57cec5SDimitry Andric// these functions are constexpr, and so are implicitly __host__ __device__. 310b57cec5SDimitry Andric// 320b57cec5SDimitry Andric// However, the compiler being in C++14 mode does not imply that the standard 330b57cec5SDimitry Andric// library supports C++14. There is no macro we can test to check that the 340b57cec5SDimitry Andric// stdlib has constexpr std::min/max. Thus we have to unconditionally define 350b57cec5SDimitry Andric// our device overloads. 360b57cec5SDimitry Andric// 370b57cec5SDimitry Andric// A host+device function cannot be overloaded, and a constexpr function 380b57cec5SDimitry Andric// implicitly become host device if there's no explicitly host or device 390b57cec5SDimitry Andric// overload preceding it. So the simple thing to do would be to declare our 400b57cec5SDimitry Andric// device min/max overloads, and then #include_next <algorithm>. This way our 410b57cec5SDimitry Andric// device overloads would come first, and so if we have a C++14 stdlib, its 420b57cec5SDimitry Andric// min/max won't become host+device and conflict with our device overloads. 430b57cec5SDimitry Andric// 440b57cec5SDimitry Andric// But that also doesn't work. libstdc++ is evil and declares std::min/max in 450b57cec5SDimitry Andric// an internal header that is included *before* <algorithm>. Thus by the time 460b57cec5SDimitry Andric// we're inside of this file, std::min/max may already have been declared, and 470b57cec5SDimitry Andric// thus we can't prevent them from becoming host+device if they're constexpr. 480b57cec5SDimitry Andric// 490b57cec5SDimitry Andric// Therefore we perpetrate the following hack: We mark our __device__ overloads 500b57cec5SDimitry Andric// with __attribute__((enable_if(true, ""))). This causes the signature of the 510b57cec5SDimitry Andric// function to change without changing anything else about it. (Except that 520b57cec5SDimitry Andric// overload resolution will prefer it over the __host__ __device__ version 530b57cec5SDimitry Andric// rather than considering them equally good). 540b57cec5SDimitry Andric 550b57cec5SDimitry Andric#include_next <algorithm> 560b57cec5SDimitry Andric 570b57cec5SDimitry Andric// We need to define these overloads in exactly the namespace our standard 580b57cec5SDimitry Andric// library uses (including the right inline namespace), otherwise they won't be 590b57cec5SDimitry Andric// picked up by other functions in the standard library (e.g. functions in 600b57cec5SDimitry Andric// <complex>). Thus the ugliness below. 610b57cec5SDimitry Andric#ifdef _LIBCPP_BEGIN_NAMESPACE_STD 620b57cec5SDimitry Andric_LIBCPP_BEGIN_NAMESPACE_STD 630b57cec5SDimitry Andric#else 640b57cec5SDimitry Andricnamespace std { 650b57cec5SDimitry Andric#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION 660b57cec5SDimitry Andric_GLIBCXX_BEGIN_NAMESPACE_VERSION 670b57cec5SDimitry Andric#endif 680b57cec5SDimitry Andric#endif 690b57cec5SDimitry Andric 700b57cec5SDimitry Andric#pragma push_macro("_CPP14_CONSTEXPR") 710b57cec5SDimitry Andric#if __cplusplus >= 201402L 720b57cec5SDimitry Andric#define _CPP14_CONSTEXPR constexpr 730b57cec5SDimitry Andric#else 740b57cec5SDimitry Andric#define _CPP14_CONSTEXPR 750b57cec5SDimitry Andric#endif 760b57cec5SDimitry Andric 770b57cec5SDimitry Andrictemplate <class __T, class __Cmp> 780b57cec5SDimitry Andric__attribute__((enable_if(true, ""))) 790b57cec5SDimitry Andricinline _CPP14_CONSTEXPR __host__ __device__ const __T & 800b57cec5SDimitry Andricmax(const __T &__a, const __T &__b, __Cmp __cmp) { 810b57cec5SDimitry Andric return __cmp(__a, __b) ? __b : __a; 820b57cec5SDimitry Andric} 830b57cec5SDimitry Andric 840b57cec5SDimitry Andrictemplate <class __T> 850b57cec5SDimitry Andric__attribute__((enable_if(true, ""))) 860b57cec5SDimitry Andricinline _CPP14_CONSTEXPR __host__ __device__ const __T & 870b57cec5SDimitry Andricmax(const __T &__a, const __T &__b) { 880b57cec5SDimitry Andric return __a < __b ? __b : __a; 890b57cec5SDimitry Andric} 900b57cec5SDimitry Andric 910b57cec5SDimitry Andrictemplate <class __T, class __Cmp> 920b57cec5SDimitry Andric__attribute__((enable_if(true, ""))) 930b57cec5SDimitry Andricinline _CPP14_CONSTEXPR __host__ __device__ const __T & 940b57cec5SDimitry Andricmin(const __T &__a, const __T &__b, __Cmp __cmp) { 950b57cec5SDimitry Andric return __cmp(__b, __a) ? __b : __a; 960b57cec5SDimitry Andric} 970b57cec5SDimitry Andric 980b57cec5SDimitry Andrictemplate <class __T> 990b57cec5SDimitry Andric__attribute__((enable_if(true, ""))) 1000b57cec5SDimitry Andricinline _CPP14_CONSTEXPR __host__ __device__ const __T & 1010b57cec5SDimitry Andricmin(const __T &__a, const __T &__b) { 102*0fca6ea1SDimitry Andric return __b < __a ? __b : __a; 1030b57cec5SDimitry Andric} 1040b57cec5SDimitry Andric 1050b57cec5SDimitry Andric#pragma pop_macro("_CPP14_CONSTEXPR") 1060b57cec5SDimitry Andric 1070b57cec5SDimitry Andric#ifdef _LIBCPP_END_NAMESPACE_STD 1080b57cec5SDimitry Andric_LIBCPP_END_NAMESPACE_STD 1090b57cec5SDimitry Andric#else 1100b57cec5SDimitry Andric#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION 1110b57cec5SDimitry Andric_GLIBCXX_END_NAMESPACE_VERSION 1120b57cec5SDimitry Andric#endif 1130b57cec5SDimitry Andric} // namespace std 1140b57cec5SDimitry Andric#endif 1150b57cec5SDimitry Andric 1160b57cec5SDimitry Andric#endif // __CLANG_CUDA_WRAPPERS_ALGORITHM 117