OmniSciDB  72c90bc290
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
TopKRuntime.cpp File Reference

Structures and runtime functions of streaming top-k heap. More...

+ 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)
 

Detailed Description

Structures and runtime functions of streaming top-k heap.

Definition in file TopKRuntime.cpp.

Macro Definition Documentation

#define DEF_GET_BIN_FROM_K_HEAP (   key_type)
Value:
extern "C" RUNTIME_EXPORT 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
#define RUNTIME_EXPORT
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 237 of file TopKRuntime.cpp.

Enumeration Type Documentation

enum HeapOrdering
strong
Enumerator
MIN 
MAX 

Definition at line 25 of file TopKRuntime.cpp.

enum NullsOrdering
strong
Enumerator
FIRST 
LAST 

Definition at line 27 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 183 of file TopKRuntime.cpp.

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

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

Referenced by get_bin_from_k_heap_impl().

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

+ 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 142 of file TopKRuntime.cpp.

Referenced by get_bin_from_k_heap_impl().

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

+ 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 82 of file TopKRuntime.cpp.

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

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

+ 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 123 of file TopKRuntime.cpp.

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

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

+ Here is the call graph for this function: