hash_functions.h 3.5 KB
Newer Older
1
2
3
4
5
6
7
8
9
10
11
/*! @file hash_functions.h
 *  @brief Hash function code.
 */

#ifndef HASH_FUNCTIONS__H
#define HASH_FUNCTIONS__H

#include <tensorview/tensorview.h>
#include <vector_types.h>
#include "definitions.h"

traveller59's avatar
traveller59 committed
12
namespace cuhash {
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91

//! Prime number larger than the largest practical hash table size.
const unsigned kPrimeDivisor = 4294967291u;
// https://www.alpertron.com.ar/ECM.HTM
// const unsigned long kPrimeDivisor = 18446744073709551557lu
// const long kPrimeDivisor = 9223372036854775783l
// const Entry kPrimeDivisor = 4300000013lu;
// const unsigned kPrimeDivisor = 334214459;

//! Generates a set of linear hash function constants.
/*! @param[in]  N           Number of hash functions.
    @param[out] constants   CPU pointer to the constants.
    @param[in]  num_keys    Debug only: How many keys are in the input.
    @param[in]  d_keys      Debug only: Device memory array containing the input keys.
    @param[in]  table_size  Debug only: Size of the hash table.
 */
void GenerateFunctions(const unsigned  N,
                       const unsigned  num_keys,
                       const unsigned *d_keys,
                       const unsigned  table_size,
                             uint2    *constants);

//! Container for all of the hash functions.
template <unsigned N>
struct Functions {
  //! The constants required for all of the hash functions, including the stash.  Each function requires 2.
  uint2 constants[N];

  //! Generate new hash function constants.
  /*! The parameters are only used for debugging and examining the key distribution.
      \param[in] num_keys   Debug: Number of keys in the input.
      \param[in] d_keys     Debug: Device array of the input keys.
      \param[in] table_size Debug: Size of the hash table.
  */
  void Generate(const unsigned  num_keys,
                const unsigned *d_keys,
                const unsigned  table_size) {
    GenerateFunctions(N, num_keys, d_keys, table_size, constants);
  }
};

//! Computes the value of a hash function for a given key.
/*! \param[in] constants  Constants used by the hash function.
  ! \param[in] key        Key being hashed.
  ! \returns              The value of the hash function for the key.
 */
inline __device__ __host__
unsigned hash_function_inner(const uint2    constants,
                             const unsigned key) {
#if 1                             
  // Fast version.                             
  return ((constants.x ^ key) + constants.y) % kPrimeDivisor;
#else
  // Slow version.
  return ((unsigned long long)constants.x * key + constants.y) % kPrimeDivisor;
#endif
}                             

//! Computes the value of a hash function for a given key.
/*! \param[in] functions        All of the constants used by the hash functions.
  ! \param[in] which_function   Which hash function is being used.
  ! \param[in] key              Key being hashed.
  ! \returns                    The value of a hash function with a given key.
 */
template <unsigned kNumHashFunctions>
TV_HOST_DEVICE_INLINE
unsigned hash_function(const Functions<kNumHashFunctions> functions,
                       const unsigned which_function,
                       const unsigned key) {
  return hash_function_inner(functions.constants[which_function], key);
}

//! Simple hash function used by the stash.
TV_HOST_DEVICE_INLINE
unsigned stash_hash_function(const uint2 stash_constants,
                             const unsigned key) {
  return (stash_constants.x ^ key + stash_constants.y) % kStashSize;
}

traveller59's avatar
traveller59 committed
92
unsigned generate_random_uint32();
93
94
95
96

};  // namespace CuckooHashing

#endif