warpcore 0.0.1
Hashing at the Speed of Light on modern CUDA-accelerators
probing_schemes.cuh
Go to the documentation of this file.
1 #ifndef WARPCORE_PROBING_SCHEMES_CUH
2 #define WARPCORE_PROBING_SCHEMES_CUH
3 
4 namespace warpcore
5 {
6 
7 /*! \brief probing scheme iterators
8  */
9 namespace probing_schemes
10 {
11 
12 namespace cg = cooperative_groups;
13 namespace checks = warpcore::checks;
14 
15 //TODO add inner (warp-level) probing?
16 
17 /*! \brief double hashing scheme: \f$hash(k,i) = h_1(k)+i\cdot h_2(k)\f$
18  * \tparam Hasher1 first hash function
19  * \tparam Hasher1 second hash function
20  * \tparam CGSize cooperative group size
21  */
22 template <class Hasher1, class Hasher2, index_t CGSize = 1>
23 class DoubleHashing
24 {
25  static_assert(
26  checks::is_valid_cg_size(CGSize),
27  "invalid cooperative group size");
28 
29  static_assert(
30  std::is_same<
31  typename Hasher1::key_type,
32  typename Hasher2::key_type>::value,
33  "key types of both hashers must be the same");
34 
35 public:
36  using key_type = typename Hasher1::key_type;
37  using index_type = index_t;
38  using tag = typename std::conditional<
39  checks::is_true_permutation_hasher<Hasher1>() &&
40  checks::is_true_permutation_hasher<Hasher2>(),
41  tags::cycle_free_probing_scheme,
42  tags::probing_scheme>::type;
43 
44  /*! \brief get cooperative group size
45  * \return cooperative group size
46  */
47  HOSTDEVICEQUALIFIER INLINEQUALIFIER
48  static constexpr index_type cg_size() noexcept { return CGSize; }
49 
50  /*! \brief constructor
51  * \param[in] capacity capacity of the underlying hash table
52  * \param[in] probing_length number of probing attempts
53  * \param[in] group cooperative group
54  */
55  DEVICEQUALIFIER INLINEQUALIFIER
56  explicit DoubleHashing(
57  index_type capacity,
58  index_type probing_length,
59  const cg::thread_block_tile<CGSize>& group) :
60  capacity_(capacity),
61  probing_length_(SDIV(probing_length, group.size()) * group.size()),
62  group_(group)
63  {}
64 
65  template<class T>
66  DEVICEQUALIFIER INLINEQUALIFIER
67  index_type begin(T, T) = delete;
68 
69  template<class T>
70  DEVICEQUALIFIER INLINEQUALIFIER
71  index_type begin(T) = delete;
72 
73  /*! \brief begin probing sequence
74  * \param[in] key key to be probed
75  * \param[in] seed random seed
76  * \return initial probing index for \c key
77  */
78  DEVICEQUALIFIER INLINEQUALIFIER
79  index_type begin(key_type key, key_type seed = 0) noexcept
80  {
81  i_ = group_.thread_rank();
82  pos_ = Hasher1::hash(key + seed) + group_.thread_rank();
83  pos_ = pos_ % capacity_;
84  // step size in range [1, capacity-1] * group_size
85  base_ = (Hasher2::hash(key + seed + 1) % (capacity_ / group_.size() - 1) + 1) * group_.size();
86 
87  return pos_;
88  }
89 
90  /*! \brief next probing index for \c key
91  * \return next probing index for \c key
92  */
93  DEVICEQUALIFIER INLINEQUALIFIER
94  index_type next() noexcept
95  {
96  i_ += CGSize;
97  pos_ = (pos_ + base_) % capacity_;
98 
99  return (i_ < probing_length_) ? pos_ : end();
100  }
101 
102  /*! \brief end specifier of probing sequence
103  * \return end specifier
104  */
105  DEVICEQUALIFIER INLINEQUALIFIER
106  static constexpr index_type end() noexcept
107  {
108  return ~index_type(0);
109  }
110 
111 private:
112  const index_type capacity_; //< capacity of the underlying hash table
113  const index_type probing_length_; //< max number of probing attempts
114  const cg::thread_block_tile<CGSize>& group_; //< cooperative group
115 
116  index_type i_; //< current probe count
117  index_type pos_; //< current probing position
118  index_type base_; //< step size
119 
120 }; // class DoubleHashing
121 
122 /*! \brief linear probing scheme: \f$hash(k,i) = h(k)+i\f$
123  * \tparam Hasher hash function
124  * \tparam CGSize cooperative group size
125  */
126 template <class Hasher, index_t CGSize = 1>
127 class LinearProbing
128 {
129  static_assert(
130  checks::is_valid_cg_size(CGSize),
131  "invalid cooperative group size");
132 
133 public:
134  using key_type = typename Hasher::key_type;
135  using index_type = index_t;
136  using tag = tags::cycle_free_probing_scheme;
137 
138  /*! \brief get cooperative group size
139  * \return cooperative group size
140  */
141  HOSTDEVICEQUALIFIER INLINEQUALIFIER
142  static constexpr index_t cg_size() noexcept { return CGSize; }
143 
144  /*! \brief constructor
145  * \param[in] capacity capacity of the underlying hash table
146  * \param[in] probing_length number of probing attempts
147  * \param[in] group cooperative group
148  */
149  DEVICEQUALIFIER
150  explicit LinearProbing(
151  index_type capacity,
152  index_type probing_length,
153  const cg::thread_block_tile<CGSize>& group) :
154  capacity_(capacity),
155  probing_length_(SDIV(probing_length, group.size()) * group.size()),
156  group_(group)
157  {}
158 
159  template<class T>
160  DEVICEQUALIFIER INLINEQUALIFIER
161  index_type begin(T, T) = delete;
162 
163  template<class T>
164  DEVICEQUALIFIER INLINEQUALIFIER
165  index_type begin(T) = delete;
166 
167  /*! \brief begin probing sequence
168  * \param[in] key key to be probed
169  * \param[in] seed random seed
170  * \return initial probing index for \c key
171  */
172  DEVICEQUALIFIER INLINEQUALIFIER
173  index_type begin(key_type key, key_type seed = 0) noexcept
174  {
175  i_ = group_.thread_rank();
176  pos_ = Hasher::hash(key + seed) + group_.thread_rank();
177  pos_ = pos_ % capacity_;
178 
179  return pos_;
180  }
181 
182  /*! \brief next probing index for \c key
183  * \return next probing index
184  */
185  DEVICEQUALIFIER INLINEQUALIFIER
186  index_type next() noexcept
187  {
188  i_ += CGSize;
189  pos_ = (pos_ + CGSize) % capacity_;
190 
191  return (i_ < probing_length_) ? pos_ : end();
192  }
193 
194  /*! \brief end specifier of probing sequence
195  * \return end specifier
196  */
197  DEVICEQUALIFIER INLINEQUALIFIER
198  static constexpr index_type end() noexcept
199  {
200  return ~index_type(0);
201  }
202 
203 private:
204  const index_type capacity_; //< capacity of the underlying hash table
205  const index_type probing_length_; //< number of probing attempts
206  const cg::thread_block_tile<CGSize>& group_; //< cooperative group
207 
208  index_type i_; //< current probe count
209  index_type pos_; //< current probing position
210 
211 }; // class LinearProbing
212 
213 /*! \brief quadratic probing scheme: \f$hash(k,i) = h(k)+\frac{1}{2}\cdot i+\frac{1}{2}\cdot i^2\f$
214  * \tparam Hasher hash function
215  * \tparam CGSize cooperative group size
216  */
217 template <class Hasher, index_t CGSize = 1>
218 class QuadraticProbing
219 {
220  static_assert(
221  checks::is_valid_cg_size(CGSize),
222  "invalid cooperative group size");
223 
224 public:
225  using key_type = typename Hasher::key_type;
226  using index_type = index_t;
227  using tag = tags::probing_scheme;
228 
229  /*! \brief get cooperative group size
230  * \return cooperative group size
231  */
232  HOSTDEVICEQUALIFIER INLINEQUALIFIER
233  static constexpr index_type cg_size() noexcept { return CGSize; }
234 
235  /*! \brief constructor
236  * \param[in] capacity capacity of the underlying hash table
237  * \param[in] probing_length number of probing attempts
238  * \param[in] group cooperative group
239  */
240  DEVICEQUALIFIER
241  explicit QuadraticProbing(
242  index_type capacity,
243  index_type probing_length,
244  const cg::thread_block_tile<CGSize>& group) :
245  capacity_(capacity),
246  probing_length_(SDIV(probing_length, group.size()) * group.size()),
247  group_(group)
248  {}
249 
250  template<class T>
251  DEVICEQUALIFIER INLINEQUALIFIER
252  index_type begin(T, T) = delete;
253 
254  template<class T>
255  DEVICEQUALIFIER INLINEQUALIFIER
256  index_type begin(T) = delete;
257 
258  /*! \brief begin probing sequence
259  * \param[in] key key to be probed
260  * \param[in] seed random seed
261  * \return initial probing index for \c key
262  */
263  DEVICEQUALIFIER INLINEQUALIFIER
264  index_type begin(key_type key, key_type seed = 0) noexcept
265  {
266  i_ = group_.thread_rank();
267  step_ = 1;
268  pos_ = Hasher::hash(key + seed) + group_.thread_rank();
269  pos_ = pos_ % capacity_;
270 
271  return pos_;
272  }
273 
274  /*! \brief next probing index for \c key
275  * \return next probing index
276  */
277  DEVICEQUALIFIER INLINEQUALIFIER
278  index_type next() noexcept
279  {
280  i_ += CGSize;
281  pos_ = (pos_ + step_) % capacity_;
282  ++step_;
283 
284  return (i_ < probing_length_) ? pos_ : end();
285  }
286 
287  /*! \brief end specifier of probing sequence
288  * \return end specifier
289  */
290  DEVICEQUALIFIER INLINEQUALIFIER
291  static constexpr index_type end() noexcept
292  {
293  return ~index_type(0);
294  }
295 
296 private:
297  const index_type capacity_; //< capacity of the underlying hash table
298  const index_type probing_length_; //< number of probing attempts
299  const cg::thread_block_tile<CGSize>& group_; //< cooperative group
300 
301  index_type i_; //< current probe count
302  index_type pos_; //< current probing position
303  index_type step_; //< current step size
304 
305 }; // class QuadraticProbing
306 
307 } // namespace probing_schemes
308 
309 } // namespace warpcore
310 
311 #endif /* WARPCORE_PROBING_SCHEMES_CUH */