CMS 3D CMS Logo

cudastdAlgorithm.h
Go to the documentation of this file.
1 #ifndef HeterogeneousCore_CUDAUtilities_cudastdAlgorithm_h
2 #define HeterogeneousCore_CUDAUtilities_cudastdAlgorithm_h
3 
4 #include <utility>
5 
6 #include <cuda_runtime.h>
7 
8 // reimplementation of std algorithms able to compile with CUDA and run on GPUs,
9 // mostly by declaringthem constexpr
10 
11 namespace cuda_std {
12 
13  template <typename T = void>
14  struct less {
15  __host__ __device__ constexpr bool operator()(const T &lhs, const T &rhs) const { return lhs < rhs; }
16  };
17 
18  template <>
19  struct less<void> {
20  template <typename T, typename U>
21  __host__ __device__ constexpr bool operator()(const T &lhs, const U &rhs) const {
22  return lhs < rhs;
23  }
24  };
25 
26  template <typename RandomIt, typename T, typename Compare = less<T>>
27  __host__ __device__ constexpr RandomIt lower_bound(RandomIt first, RandomIt last, const T &value, Compare comp = {}) {
28  auto count = last - first;
29 
30  while (count > 0) {
31  auto it = first;
32  auto step = count / 2;
33  it += step;
34  if (comp(*it, value)) {
35  first = ++it;
36  count -= step + 1;
37  } else {
38  count = step;
39  }
40  }
41  return first;
42  }
43 
44  template <typename RandomIt, typename T, typename Compare = less<T>>
45  __host__ __device__ constexpr RandomIt upper_bound(RandomIt first, RandomIt last, const T &value, Compare comp = {}) {
46  auto count = last - first;
47 
48  while (count > 0) {
49  auto it = first;
50  auto step = count / 2;
51  it += step;
52  if (!comp(value, *it)) {
53  first = ++it;
54  count -= step + 1;
55  } else {
56  count = step;
57  }
58  }
59  return first;
60  }
61 
62  template <typename RandomIt, typename T, typename Compare = cuda_std::less<T>>
63  __host__ __device__ constexpr RandomIt binary_find(RandomIt first, RandomIt last, const T &value, Compare comp = {}) {
65  return first != last && !comp(value, *first) ? first : last;
66  }
67 
68 } // namespace cuda_std
69 
70 #endif // HeterogeneousCore_CUDAUtilities_cudastdAlgorithm_h
cuda_std::less
Definition: cudastdAlgorithm.h:14
step
step
Definition: StallMonitor.cc:94
cuda_std::less::operator()
__host__ constexpr __device__ bool operator()(const T &lhs, const T &rhs) const
Definition: cudastdAlgorithm.h:15
__device__
#define __device__
Definition: cudaCompat.h:85
cuda_std::less< void >::operator()
__host__ constexpr __device__ bool operator()(const T &lhs, const U &rhs) const
Definition: cudastdAlgorithm.h:21
dqmdumpme.first
first
Definition: dqmdumpme.py:55
AlCaHLTBitMon_QueryRunRegistry.comp
comp
Definition: AlCaHLTBitMon_QueryRunRegistry.py:249
dqmdumpme.last
last
Definition: dqmdumpme.py:56
cuda_std::upper_bound
__host__ constexpr __device__ RandomIt upper_bound(RandomIt first, RandomIt last, const T &value, Compare comp={})
Definition: cudastdAlgorithm.h:45
cuda_std::binary_find
__host__ constexpr __device__ RandomIt binary_find(RandomIt first, RandomIt last, const T &value, Compare comp={})
Definition: cudastdAlgorithm.h:63
cuda_std::lower_bound
__host__ constexpr __device__ RandomIt lower_bound(RandomIt first, RandomIt last, const T &value, Compare comp={})
Definition: cudastdAlgorithm.h:27
mitigatedMETSequence_cff.U
U
Definition: mitigatedMETSequence_cff.py:36
__host__
#define __host__
Definition: cudaCompat.h:84
KineDebug3::count
void count()
Definition: KinematicConstrainedVertexUpdatorT.h:21
cuda_std
Definition: cudastdAlgorithm.h:11
value
Definition: value.py:1
T
long double T
Definition: Basic3DVectorLD.h:48
funct::void
TEMPL(T2) struct Divides void
Definition: Factorize.h:29