warpcore 0.0.1
Hashing at the Speed of Light on modern CUDA-accelerators
counting_hash_table.cuh
Go to the documentation of this file.
1 #ifndef WARPCORE_COUNTING_HASH_TABLE_CUH
2 #define WARPCORE_COUNTING_HASH_TABLE_CUH
3 
4 #include "single_value_hash_table.cuh"
5 
6 namespace warpcore
7 {
8 
9 /*! \brief counting hash table
10  * \tparam Key key type (\c std::uint32_t or \c std::uint64_t)
11  * \tparam Value value type
12  * \tparam EmptyKey key which represents an empty slot
13  * \tparam TombstoneKey key which represents an erased slot
14  * \tparam ProbingScheme probing scheme from \c warpcore::probing_schemes
15  * \tparam TableStorage internal storage class \c warpcore::storage::key_value
16  * \tparam TempMemoryBytes size of temporary storage (typically a few kB)
17  */
18 template<
19  class Key,
20  class Value = index_t,
21  Key EmptyKey = defaults::empty_key<Key>(),
22  Key TombstoneKey = defaults::tombstone_key<Key>(),
23  class ProbingScheme = defaults::probing_scheme_t<Key, 4>,
24  class TableStorage = defaults::table_storage_t<Key, Value>,
25  index_t TempMemoryBytes = defaults::temp_memory_bytes()>
26 class CountingHashTable
27 {
28  using base_type = SingleValueHashTable<
29  Key, Value, EmptyKey, TombstoneKey, ProbingScheme, TableStorage>;
30 
31  static_assert(
32  checks::is_valid_counter_type<typename base_type::value_type>(),
33  "counter type must be either int, std::uint32_t or std::uint64_t");
34 
35 public:
36  using key_type = typename base_type::key_type;
37  using value_type = typename base_type::value_type;
38  using index_type = typename base_type::index_type;
39  using status_type = typename base_type::status_type;
40 
41  /*! \brief get empty key
42  * \return empty key
43  */
44  HOSTDEVICEQUALIFIER INLINEQUALIFIER
45  static constexpr key_type empty_key() noexcept
46  {
47  return base_type::empty_key();
48  }
49 
50  /*! \brief get tombstone key
51  * \return tombstone key
52  */
53  HOSTDEVICEQUALIFIER INLINEQUALIFIER
54  static constexpr key_type tombstone_key() noexcept
55  {
56  return base_type::tombstone_key();
57  }
58 
59  /*! \brief get cooperative group size
60  * \return cooperative group size
61  */
62  HOSTDEVICEQUALIFIER INLINEQUALIFIER
63  static constexpr index_type cg_size() noexcept
64  {
65  return base_type::cg_size();
66  }
67 
68  /*! \brief constructor
69  * \param[in] min_capacity minimum number of slots in the hash table
70  * \param[in] seed random seed
71  * \param[in] max_count count after which to stop counting new occurrences
72  * \param[in] no_init whether to initialize the table at construction or not
73  */
74  HOSTQUALIFIER INLINEQUALIFIER
75  explicit CountingHashTable(
76  index_type min_capacity,
77  key_type seed = defaults::seed<key_type>(),
78  value_type max_count = std::numeric_limits<value_type>::max(),
79  bool no_init = false) noexcept :
80  base_table_(min_capacity, seed, true),
81  max_count_(max_count),
82  is_copy_(false)
83  {
84  if(!no_init) init(seed);
85  }
86 
87  /*! \brief copy-constructor (shallow)
88  * \param[in] object to be copied
89  */
90  HOSTDEVICEQUALIFIER INLINEQUALIFIER
91  CountingHashTable(const CountingHashTable& o) noexcept :
92  base_table_(o.base_table_),
93  max_count_(o.max_count_),
94  is_copy_(true)
95  {}
96 
97  /*! \brief move-constructor
98  * \param[in] object to be moved
99  */
100  HOSTQUALIFIER INLINEQUALIFIER
101  CountingHashTable(CountingHashTable&& o) noexcept :
102  base_table_(std::move(o.base_table_)),
103  max_count_(std::move(o.max_count_)),
104  is_copy_(std::move(o.is_copy_))
105  {
106  o.is_copy_ = true;
107  }
108 
109  /*! \brief (re)initialize the hash table
110  * \param[in] seed random seed
111  * \param[in] stream CUDA stream in which this operation is executed in
112  */
113  HOSTQUALIFIER INLINEQUALIFIER
114  void init(
115  const key_type seed,
116  const cudaStream_t stream = 0) noexcept
117  {
118  base_table_.init(seed, stream);
119 
120  if(base_table_.is_initialized())
121  {
122  base_table_.table_.init_values(0, stream);
123  }
124  }
125 
126  /*! \brief (re)initialize the hash table
127  * \param[in] stream CUDA stream in which this operation is executed in
128  */
129  HOSTQUALIFIER INLINEQUALIFIER
130  void init(const cudaStream_t stream = 0) noexcept
131  {
132  init(base_table_.seed(), stream);
133  }
134 
135  /*! \brief inserts a key into the hash table
136  * \param[in] key_in key to insert into the hash table
137  * \param[in] group cooperative group
138  * \param[in] probing_length maximum number of probing attempts
139  * \return status (per thread)
140  */
141  DEVICEQUALIFIER INLINEQUALIFIER
142  status_type insert(
143  key_type key_in,
144  const cg::thread_block_tile<cg_size()>& group,
145  index_type probing_length = defaults::probing_length()) noexcept
146  {
147  status_type status = status_type::unknown_error();
148 
149  value_type * value_ptr =
150  base_table_.insert_impl(key_in, status, group, probing_length);
151 
152  if(group.thread_rank() == 0 && value_ptr != nullptr)
153  {
154  // this may not exactly stop counting when max_count is reached
155  // however, this is resolved during retireval
156  if(*value_ptr < max_count_)
157  {
158  const value_type old = atomicAdd(value_ptr, 1);
159 
160  // guard wrap-around
161  if(old > *value_ptr)
162  {
163  status += status_type::index_overflow();
164  base_table_.status_->atomic_join(
165  status_type::index_overflow());
166  atomicExch(value_ptr, max_count_);
167  }
168  }
169  }
170 
171  return status;
172  }
173 
174  /*! \brief insert a set of keys into the hash table
175  * \tparam StatusHandler handles status per key (see \c status_handlers)
176  * \param[in] keys_in pointer to keys to insert into the hash table
177  * \param[in] num_in number of keys to insert
178  * \param[in] stream CUDA stream in which this operation is executed in
179  * \param[in] probing_length maximum number of probing attempts
180  * \param[out] status_out status information per key
181  */
182  template<class StatusHandler = defaults::status_handler_t>
183  HOSTQUALIFIER INLINEQUALIFIER
184  void insert(
185  const key_type * keys_in,
186  index_type num_in,
187  cudaStream_t stream = 0,
188  index_type probing_length = defaults::probing_length(),
189  typename StatusHandler::base_type * status_out = nullptr) noexcept
190  {
191  kernels::insert<CountingHashTable, StatusHandler>
192  <<<SDIV(num_in * cg_size(), WARPCORE_BLOCKSIZE), WARPCORE_BLOCKSIZE, 0, stream>>>
193  (keys_in, num_in, *this, probing_length, status_out);
194  }
195 
196  /*! \brief retrieves a key from the hash table
197  * \param[in] key_in key to retrieve from the hash table
198  * \param[out] value_out count of \c key_in
199  * \param[in] group cooperative group
200  * \param[in] probing_length maximum number of probing attempts
201  * \return status (per thread)
202  */
203  DEVICEQUALIFIER INLINEQUALIFIER
204  status_type retrieve(
205  key_type key_in,
206  value_type& value_out,
207  const cg::thread_block_tile<cg_size()>& group,
208  index_type probing_length = defaults::probing_length()) const noexcept
209  {
210  value_type out;
211 
212  const status_type status =
213  base_table_.retrieve(key_in, out, group, probing_length);
214 
215  if(status.has_any())
216  {
217  value_out = 0;
218  }
219  else
220  {
221  value_out = (out > max_count_) ? max_count_ : out;
222  }
223 
224  return status;
225  }
226 
227  /*! \brief retrieve a set of keys from the hash table
228  * \tparam StatusHandler handles status per key (see \c status_handlers)
229  * \param[in] keys_in pointer to keys to retrieve from the hash table
230  * \param[in] num_in number of keys to retrieve
231  * \param[out] values_out corresponding counts of keys in \c key_in
232  * \param[in] stream CUDA stream in which this operation is executed in
233  * \param[in] probing_length maximum number of probing attempts
234  * \param[out] status_out status information (per key)
235  */
236  template<class StatusHandler = defaults::status_handler_t>
237  HOSTQUALIFIER INLINEQUALIFIER
238  void retrieve(
239  const key_type * keys_in,
240  index_type num_in,
241  value_type * values_out,
242  cudaStream_t stream = 0,
243  index_type probing_length = defaults::probing_length(),
244  typename StatusHandler::base_type * status_out = nullptr) const noexcept
245  {
246  kernels::retrieve<CountingHashTable, StatusHandler>
247  <<<SDIV(num_in * cg_size(), WARPCORE_BLOCKSIZE), WARPCORE_BLOCKSIZE, 0, stream>>>
248  (keys_in, num_in, values_out, *this, probing_length, status_out);
249  }
250 
251  /*! \brief retrieves all elements from the hash table
252  * \param[out] keys_out location to store all retrieved keys
253  * \param[out] values_out location to store all retrieved counts
254  * \param[out] num_out number of of key/value pairs retrieved
255  * \param[in] stream CUDA stream in which this operation is executed in
256  */
257  HOSTQUALIFIER INLINEQUALIFIER
258  void retrieve_all(
259  key_type * keys_out,
260  value_type * values_out,
261  index_t& num_out,
262  cudaStream_t stream = 0) const noexcept
263  {
264  base_table_.retrieve_all(keys_out, values_out, num_out, stream);
265  }
266 
267  /*! \brief erases a key from the hash table
268  * \param[in] key_in key to erase from the hash table
269  * \param[in] group cooperative group
270  * \param[in] probing_length maximum number of probing attempts
271  * \return status (per thread)
272  */
273  DEVICEQUALIFIER INLINEQUALIFIER
274  status_type erase(
275  key_type key_in,
276  const cg::thread_block_tile<cg_size()>& group,
277  index_type probing_length = defaults::probing_length()) noexcept
278  {
279  return base_table_.erase(key_in, group, probing_length);
280  }
281 
282  /*! \brief erases a set of keys from the hash table
283  * \tparam StatusHandler handles status per key (see \c status_handlers)
284  * \param[in] keys_in pointer to keys to erase from the hash table
285  * \param[in] num_in number of keys to erase
286  * \param[in] stream CUDA stream in which this operation is executed in
287  * \param[in] probing_length maximum number of probing attempts
288  * \param[out] status_out status information (per key)
289  */
290  template<class StatusHandler = defaults::status_handler_t>
291  HOSTQUALIFIER INLINEQUALIFIER
292  void erase(
293  key_type * keys_in,
294  index_type num_in,
295  cudaStream_t stream = 0,
296  index_type probing_length = defaults::probing_length(),
297  typename StatusHandler::base_type * status_out = nullptr) noexcept
298  {
299  base_table_.template erase<StatusHandler>(
300  keys_in, num_in, probing_length, status_out);
301  }
302 
303  /*! \brief applies a funtion on all key value pairs inside the table
304  * \tparam Func type of map i.e. CUDA device lambda
305  * \param[in] f map to apply
306  * \param[in] stream CUDA stream in which this operation is executed in
307  * \param[in] size of shared memory to reserve for this execution
308  */
309  template<class Func>
310  HOSTQUALIFIER INLINEQUALIFIER
311  void for_each(
312  Func f,
313  cudaStream_t stream = 0,
314  index_type smem_bytes = 0) const noexcept
315  {
316  base_table_.template for_each<Func>(f, stream, smem_bytes);
317  }
318 
319  /*! \brief number of key/value pairs stored inside the hash table
320  * \param stream CUDA stream in which this operation is executed in
321  * \return the number of key/value pairs inside the hash table
322  */
323  HOSTQUALIFIER INLINEQUALIFIER
324  index_type size(cudaStream_t stream = 0) const noexcept
325  {
326  return base_table_.size(stream);
327  }
328 
329  /*! \brief current load factor of the hash table
330  * \param stream CUDA stream in which this operation is executed in
331  * \return load factor
332  */
333  HOSTQUALIFIER INLINEQUALIFIER
334  float load_factor(cudaStream_t stream = 0) const noexcept
335  {
336  return base_table_.load_factor(stream);
337  }
338 
339  /*! \brief current storage density of the hash table
340  * \param stream CUDA stream in which this operation is executed in
341  * \return storage density
342  */
343  HOSTQUALIFIER INLINEQUALIFIER
344  float storage_density(cudaStream_t stream = 0) const noexcept
345  {
346  return base_table_.storage_density(stream);
347  }
348 
349  /*! \brief get the capacity of the hash table
350  * \return number of slots in the hash table
351  */
352  HOSTDEVICEQUALIFIER INLINEQUALIFIER
353  index_type capacity() const noexcept
354  {
355  return base_table_.capacity();
356  }
357 
358  /*! \brief get the status of the hash table
359  * \param stream CUDA stream in which this operation is executed in
360  * \return the status
361  */
362  HOSTQUALIFIER INLINEQUALIFIER
363  status_type peek_status(cudaStream_t stream = 0) const noexcept
364  {
365  return base_table_.peek_status(stream);
366  }
367 
368  /*! \brief get and reset the status of the hash table
369  * \param stream CUDA stream in which this operation is executed in
370  * \return the status
371  */
372  HOSTQUALIFIER INLINEQUALIFIER
373  status_type pop_status(cudaStream_t stream = 0) noexcept
374  {
375  return base_table_.pop_status(stream);
376  }
377 
378  /*! \brief checks if \c key is equal to \c EmptyKey
379  * \return \c bool
380  */
381  HOSTDEVICEQUALIFIER INLINEQUALIFIER
382  static constexpr bool is_empty_key(key_type key) noexcept
383  {
384  return base_type::is_empty_key(key);
385  }
386 
387  /*! \brief checks if \c key is equal to \c TombstoneKey
388  * \return \c bool
389  */
390  HOSTDEVICEQUALIFIER INLINEQUALIFIER
391  static constexpr bool is_tombstone_key(key_type key) noexcept
392  {
393  return base_type::is_tombstone_key(key);
394  }
395 
396  /*! \brief checks if \c key is equal to \c (EmptyKey||TombstoneKey)
397  * \return \c bool
398  */
399  HOSTDEVICEQUALIFIER INLINEQUALIFIER
400  static constexpr bool is_valid_key(key_type key) noexcept
401  {
402  return base_type::is_valid_key(key);
403  }
404 
405  /*! \brief get random seed
406  * \return seed
407  */
408  HOSTDEVICEQUALIFIER INLINEQUALIFIER
409  key_type seed() const noexcept
410  {
411  return base_table_.seed();
412  }
413 
414  /*! \brief indicates if this object is a shallow copy
415  * \return \c bool
416  */
417  HOSTDEVICEQUALIFIER INLINEQUALIFIER
418  bool is_copy() const noexcept
419  {
420  return is_copy_;
421  }
422 
423 private:
424  base_type base_table_; //< base type aka SingleValueHashTable
425  const value_type max_count_; //< count after which new occurrences are ignored
426  bool is_copy_; //< indicates if this object is a shallow copy
427 
428 }; // class CountingHashTable
429 
430 } // namespace warpcore
431 
432 #endif /* WARPCORE_COUNTING_HASH_TABLE_CUH */