OmniSciDB  b24e664e58
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
TopKRuntime.cpp File Reference
+ Include dependency graph for TopKRuntime.cpp:
+ This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes

struct  KeyAccessor< KeyT, IndexT >
 
struct  KeyComparator< KeyT >
 

Macros

#define DEF_GET_BIN_FROM_K_HEAP(key_type)
 

Enumerations

enum  HeapOrdering { HeapOrdering::MIN, HeapOrdering::MAX }
 
enum  NullsOrdering { NullsOrdering::FIRST, NullsOrdering::LAST }
 

Functions

template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE void sift_down (NodeT *heap, const size_t heap_size, const NodeT curr_idx, const KeyComparator< KeyT > &compare, const KeyAccessor< KeyT, NodeT > &accessor)
 
template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE void sift_up (NodeT *heap, const NodeT curr_idx, const KeyComparator< KeyT > &compare, const KeyAccessor< KeyT, NodeT > &accessor)
 
template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE void push_heap (int64_t *heap_ptr, int64_t *rows_ptr, NodeT &node_count, const uint32_t row_size_quad, const uint32_t key_offset, const KeyComparator< KeyT > &comparator, const KeyAccessor< KeyT, NodeT > &accessor, const KeyT curr_key)
 
template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE bool pop_and_push_heap (int64_t *heap_ptr, int64_t *rows_ptr, const NodeT node_count, const uint32_t row_size_quad, const uint32_t key_offset, const KeyComparator< KeyT > &compare, const KeyAccessor< KeyT, NodeT > &accessor, const KeyT curr_key)
 
template<typename KeyT = int64_t>
ALWAYS_INLINE DEVICE int64_t * get_bin_from_k_heap_impl (int64_t *heaps, const uint32_t k, const uint32_t row_size_quad, const uint32_t key_offset, const bool min_heap, const bool has_null, const bool nulls_first, const KeyT null_key, const KeyT curr_key)
 

Macro Definition Documentation

#define DEF_GET_BIN_FROM_K_HEAP (   key_type)
Value:
extern "C" NEVER_INLINE DEVICE int64_t* get_bin_from_k_heap_##key_type( \
int64_t* heaps, \
const uint32_t k, \
const uint32_t row_size_quad, \
const uint32_t key_offset, \
const bool min_heap, \
const bool has_null, \
const bool nulls_first, \
const key_type null_key, \
const key_type curr_key) { \
return get_bin_from_k_heap_impl(heaps, \
k, \
row_size_quad, \
key_offset, \
min_heap, \
has_null, \
nulls_first, \
null_key, \
curr_key); \
}
#define DEVICE
ALWAYS_INLINE DEVICE int64_t * get_bin_from_k_heap_impl(int64_t *heaps, const uint32_t k, const uint32_t row_size_quad, const uint32_t key_offset, const bool min_heap, const bool has_null, const bool nulls_first, const KeyT null_key, const KeyT curr_key)
#define NEVER_INLINE

Definition at line 238 of file TopKRuntime.cpp.

Enumeration Type Documentation

enum HeapOrdering
strong
Enumerator
MIN 
MAX 

Definition at line 26 of file TopKRuntime.cpp.

enum NullsOrdering
strong
Enumerator
FIRST 
LAST 

Definition at line 28 of file TopKRuntime.cpp.

Function Documentation

template<typename KeyT = int64_t>
ALWAYS_INLINE DEVICE int64_t* get_bin_from_k_heap_impl ( int64_t *  heaps,
const uint32_t  k,
const uint32_t  row_size_quad,
const uint32_t  key_offset,
const bool  min_heap,
const bool  has_null,
const bool  nulls_first,
const KeyT  null_key,
const KeyT  curr_key 
)

Definition at line 184 of file TopKRuntime.cpp.

References FIRST, LAST, MAX, MIN, pop_and_push_heap(), pos_start_impl(), pos_step_impl(), and push_heap().

192  {
193  const int32_t thread_global_index = pos_start_impl(nullptr);
194  const int32_t thread_count = pos_step_impl();
195  int64_t& node_count = heaps[thread_global_index];
196  int64_t* heap_ptr = heaps + thread_count + thread_global_index * k;
197  int64_t* rows_ptr =
198  heaps + thread_count + thread_count * k + thread_global_index * row_size_quad * k;
200  has_null,
201  null_key,
202  nulls_first ? NullsOrdering::FIRST : NullsOrdering::LAST);
203  KeyAccessor<KeyT, int64_t> accessor(reinterpret_cast<int8_t*>(rows_ptr),
204  row_size_quad * sizeof(int64_t),
205  key_offset / sizeof(KeyT));
206  if (node_count < static_cast<int64_t>(k)) {
207  push_heap(heap_ptr,
208  rows_ptr,
209  node_count,
210  row_size_quad,
211  key_offset,
212  compare,
213  accessor,
214  curr_key);
215  const auto last_bin_index = node_count - 1;
216  auto row_ptr = rows_ptr + last_bin_index * row_size_quad;
217  row_ptr[0] = last_bin_index;
218  return row_ptr + 1;
219  } else {
220  const int64_t top_bin_idx = heap_ptr[0];
221  const bool rejected = !pop_and_push_heap(heap_ptr,
222  rows_ptr,
223  node_count,
224  row_size_quad,
225  key_offset,
226  compare,
227  accessor,
228  curr_key);
229  if (rejected) {
230  return nullptr;
231  }
232  auto row_ptr = rows_ptr + top_bin_idx * row_size_quad;
233  row_ptr[0] = top_bin_idx;
234  return row_ptr + 1;
235  }
236 }
ALWAYS_INLINE DEVICE void push_heap(int64_t *heap_ptr, int64_t *rows_ptr, NodeT &node_count, const uint32_t row_size_quad, const uint32_t key_offset, const KeyComparator< KeyT > &comparator, const KeyAccessor< KeyT, NodeT > &accessor, const KeyT curr_key)
__device__ int32_t pos_step_impl()
Definition: cuda_mapd_rt.cu:19
__device__ int32_t pos_start_impl(const int32_t *row_index_resume)
Definition: cuda_mapd_rt.cu:11
ALWAYS_INLINE DEVICE bool pop_and_push_heap(int64_t *heap_ptr, int64_t *rows_ptr, const NodeT node_count, const uint32_t row_size_quad, const uint32_t key_offset, const KeyComparator< KeyT > &compare, const KeyAccessor< KeyT, NodeT > &accessor, const KeyT curr_key)

+ Here is the call graph for this function:

template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE bool pop_and_push_heap ( int64_t *  heap_ptr,
int64_t *  rows_ptr,
const NodeT  node_count,
const uint32_t  row_size_quad,
const uint32_t  key_offset,
const KeyComparator< KeyT > &  compare,
const KeyAccessor< KeyT, NodeT > &  accessor,
const KeyT  curr_key 
)

Definition at line 161 of file TopKRuntime.cpp.

Referenced by get_bin_from_k_heap_impl().

168  {
169  const NodeT top_bin_idx = static_cast<NodeT>(heap_ptr[0]);
170  int8_t* top_row_ptr = reinterpret_cast<int8_t*>(rows_ptr + top_bin_idx * row_size_quad);
171  auto top_key = reinterpret_cast<KeyT*>(top_row_ptr + key_offset);
172  if (compare(curr_key, *top_key)) {
173  return false;
174  }
175  // kick out
176  *top_key = curr_key;
177  // sift down
178  sift_down<KeyT, NodeT>(heap_ptr, node_count, 0, compare, accessor);
179  return true;
180 }

+ Here is the caller graph for this function:

template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE void push_heap ( int64_t *  heap_ptr,
int64_t *  rows_ptr,
NodeT &  node_count,
const uint32_t  row_size_quad,
const uint32_t  key_offset,
const KeyComparator< KeyT > &  comparator,
const KeyAccessor< KeyT, NodeT > &  accessor,
const KeyT  curr_key 
)

Definition at line 143 of file TopKRuntime.cpp.

Referenced by get_bin_from_k_heap_impl().

150  {
151  const NodeT bin_index = node_count++;
152  heap_ptr[bin_index] = bin_index;
153  int8_t* row_ptr = reinterpret_cast<int8_t*>(rows_ptr + bin_index * row_size_quad);
154  auto key_ptr = reinterpret_cast<KeyT*>(row_ptr + key_offset);
155  *key_ptr = curr_key;
156  // sift up
157  sift_up<KeyT, NodeT>(heap_ptr, bin_index, comparator, accessor);
158 }

+ Here is the caller graph for this function:

template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE void sift_down ( NodeT *  heap,
const size_t  heap_size,
const NodeT  curr_idx,
const KeyComparator< KeyT > &  compare,
const KeyAccessor< KeyT, NodeT > &  accessor 
)

Definition at line 83 of file TopKRuntime.cpp.

References KeyAccessor< KeyT, IndexT >::get().

87  {
88  for (NodeT i = curr_idx, last = static_cast<NodeT>(heap_size); i < last;) {
89 #ifdef __CUDACC__
90  const auto left_child = min(2 * i + 1, last);
91  const auto right_child = min(2 * i + 2, last);
92 #else
93  const auto left_child = std::min(2 * i + 1, last);
94  const auto right_child = std::min(2 * i + 2, last);
95 #endif
96  auto candidate_idx = last;
97  if (left_child < last) {
98  if (right_child < last) {
99  const auto left_key = accessor.get(heap[left_child]);
100  const auto right_key = accessor.get(heap[right_child]);
101  candidate_idx = compare(left_key, right_key) ? left_child : right_child;
102  } else {
103  candidate_idx = left_child;
104  }
105  } else {
106  candidate_idx = right_child;
107  }
108  if (candidate_idx >= last) {
109  break;
110  }
111  const auto curr_key = accessor.get(heap[i]);
112  const auto candidate_key = accessor.get(heap[candidate_idx]);
113  if (compare(curr_key, candidate_key)) {
114  break;
115  }
116  auto temp_id = heap[i];
117  heap[i] = heap[candidate_idx];
118  heap[candidate_idx] = temp_id;
119  i = candidate_idx;
120  }
121 }
ALWAYS_INLINE DEVICE KeyT get(const IndexT rowid) const
Definition: TopKRuntime.cpp:36

+ Here is the call graph for this function:

template<typename KeyT = int64_t, typename NodeT = int64_t>
ALWAYS_INLINE DEVICE void sift_up ( NodeT *  heap,
const NodeT  curr_idx,
const KeyComparator< KeyT > &  compare,
const KeyAccessor< KeyT, NodeT > &  accessor 
)

Definition at line 124 of file TopKRuntime.cpp.

References KeyAccessor< KeyT, IndexT >::get().

127  {
128  for (NodeT i = curr_idx; i > 0 && (i - 1) < i;) {
129  const auto parent = (i - 1) / 2;
130  const auto curr_key = accessor.get(heap[i]);
131  const auto parent_key = accessor.get(heap[parent]);
132  if (compare(parent_key, curr_key)) {
133  break;
134  }
135  auto temp_id = heap[i];
136  heap[i] = heap[parent];
137  heap[parent] = temp_id;
138  i = parent;
139  }
140 }
ALWAYS_INLINE DEVICE KeyT get(const IndexT rowid) const
Definition: TopKRuntime.cpp:36

+ Here is the call graph for this function: