CUDAUtilities/interface/cudastdAlgorithm.h

0001 #ifndef HeterogeneousCore_CUDAUtilities_cudastdAlgorithm_h
0002 #define HeterogeneousCore_CUDAUtilities_cudastdAlgorithm_h
0003
0004 #include <utility>
0005
0006 #include <cuda_runtime.h>
0007
0008 // reimplementation of std algorithms able to compile with CUDA and run on GPUs,
0009 // mostly by declaringthem constexpr
0010
0011 namespace cuda_std {
0012
0013   template <typename T = void>
0014   struct less {
0015     __host__ __device__ constexpr bool operator()(const T &lhs, const T &rhs) const { return lhs < rhs; }
0016   };
0017
0018   template <>
0019   struct less<void> {
0020     template <typename T, typename U>
0021     __host__ __device__ constexpr bool operator()(const T &lhs, const U &rhs) const {
0022       return lhs < rhs;
0023     }
0024   };
0025
0026   template <typename RandomIt, typename T, typename Compare = less<T>>
0027   __host__ __device__ constexpr RandomIt lower_bound(RandomIt first, RandomIt last, const T &value, Compare comp = {}) {
0028     auto count = last - first;
0029
0030     while (count > 0) {
0031       auto it = first;
0032       auto step = count / 2;
0033       it += step;
0034       if (comp(*it, value)) {
0035         first = ++it;
0036         count -= step + 1;
0037       } else {
0038         count = step;
0039       }
0040     }
0041     return first;
0042   }
0043
0044   template <typename RandomIt, typename T, typename Compare = less<T>>
0045   __host__ __device__ constexpr RandomIt upper_bound(RandomIt first, RandomIt last, const T &value, Compare comp = {}) {
0046     auto count = last - first;
0047
0048     while (count > 0) {
0049       auto it = first;
0050       auto step = count / 2;
0051       it += step;
0052       if (!comp(value, *it)) {
0053         first = ++it;
0054         count -= step + 1;
0055       } else {
0056         count = step;
0057       }
0058     }
0059     return first;
0060   }
0061
0062   template <typename RandomIt, typename T, typename Compare = cuda_std::less<T>>
0063   __host__ __device__ constexpr RandomIt binary_find(RandomIt first, RandomIt last, const T &value, Compare comp = {}) {
0064     first = cuda_std::lower_bound(first, last, value, comp);
0065     return first != last && !comp(value, *first) ? first : last;
0066   }
0067
0068 }  // namespace cuda_std
0069
0070 #endif  // HeterogeneousCore_CUDAUtilities_cudastdAlgorithm_h