OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
thread_search.cuh
Go to the documentation of this file.
1 /******************************************************************************
2  * Copyright (c) 2011, Duane Merrill. All rights reserved.
3  * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
4  *
5  * Redistribution and use in source and binary forms, with or without
6  * modification, are permitted provided that the following conditions are met:
7  * * Redistributions of source code must retain the above copyright
8  * notice, this list of conditions and the following disclaimer.
9  * * Redistributions in binary form must reproduce the above copyright
10  * notice, this list of conditions and the following disclaimer in the
11  * documentation and/or other materials provided with the distribution.
12  * * Neither the name of the NVIDIA CORPORATION nor the
13  * names of its contributors may be used to endorse or promote products
14  * derived from this software without specific prior written permission.
15  *
16  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
20  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26  *
27  ******************************************************************************/
28 
34 #pragma once
35 
36 #include "../util_namespace.cuh"
37 
39 CUB_NS_PREFIX
40 
42 namespace cub {
43 
44 
48 template <
49  typename AIteratorT,
50  typename BIteratorT,
51  typename OffsetT,
52  typename CoordinateT>
53 __host__ __device__ __forceinline__ void MergePathSearch(
54  OffsetT diagonal,
55  AIteratorT a,
56  BIteratorT b,
57  OffsetT a_len,
58  OffsetT b_len,
59  CoordinateT& path_coordinate)
60 {
62  typedef typename std::iterator_traits<AIteratorT>::value_type T;
63 
64  OffsetT split_min = CUB_MAX(diagonal - b_len, 0);
65  OffsetT split_max = CUB_MIN(diagonal, a_len);
66 
67  while (split_min < split_max)
68  {
69  OffsetT split_pivot = (split_min + split_max) >> 1;
70  if (a[split_pivot] <= b[diagonal - split_pivot - 1])
71  {
72  // Move candidate split range up A, down B
73  split_min = split_pivot + 1;
74  }
75  else
76  {
77  // Move candidate split range up B, down A
78  split_max = split_pivot;
79  }
80  }
81 
82  path_coordinate.x = CUB_MIN(split_min, a_len);
83  path_coordinate.y = diagonal - split_min;
84 }
85 
86 
87 
91 template <
92  typename InputIteratorT,
93  typename OffsetT,
94  typename T>
95 __device__ __forceinline__ OffsetT LowerBound(
96  InputIteratorT input,
97  OffsetT num_items,
98  T val)
99 {
100  OffsetT retval = 0;
101  while (num_items > 0)
102  {
103  OffsetT half = num_items >> 1;
104  if (input[retval + half] < val)
105  {
106  retval = retval + (half + 1);
107  num_items = num_items - (half + 1);
108  }
109  else
110  {
111  num_items = half;
112  }
113  }
114 
115  return retval;
116 }
117 
118 
122 template <
123  typename InputIteratorT,
124  typename OffsetT,
125  typename T>
126 __device__ __forceinline__ OffsetT UpperBound(
127  InputIteratorT input,
128  OffsetT num_items,
129  T val)
130 {
131  OffsetT retval = 0;
132  while (num_items > 0)
133  {
134  OffsetT half = num_items >> 1;
135  if (val < input[retval + half])
136  {
137  num_items = half;
138  }
139  else
140  {
141  retval = retval + (half + 1);
142  num_items = num_items - (half + 1);
143  }
144  }
145 
146  return retval;
147 }
148 
149 
150 
151 
152 
153 } // CUB namespace
154 CUB_NS_POSTFIX // Optional outer namespace(s)
cub::num_items
KeyT const ValueT ValueT OffsetT OffsetT num_items
[in] Total number of input data items
Definition: dispatch_radix_sort.cuh:164
cub
Optional outer namespace(s)
Definition: agent_histogram.cuh:48
cub::OffsetT
OffsetT OffsetT
[in] Total number of input data items
Definition: dispatch_radix_sort.cuh:74
cub::MergePathSearch
__host__ __device__ __forceinline__ void MergePathSearch(OffsetT diagonal, AIteratorT a, BIteratorT b, OffsetT a_len, OffsetT b_len, CoordinateT &path_coordinate)
Definition: thread_search.cuh:53
cub::UpperBound
__device__ __forceinline__ OffsetT UpperBound(InputIteratorT input, OffsetT num_items, T val)
Returns the offset of the first value within input which compares greater than val.
Definition: thread_search.cuh:126
CUB_MIN
#define CUB_MIN(a, b)
Select minimum(a, b)
Definition: util_macro.cuh:66
CUB_MAX
#define CUB_MAX(a, b)
Select maximum(a, b)
Definition: util_macro.cuh:61
cub::LowerBound
__device__ __forceinline__ OffsetT LowerBound(InputIteratorT input, OffsetT num_items, T val)
Returns the offset of the first value within input which does not compare less than val.
Definition: thread_search.cuh:95