hash_functions.h 3.3 KB
Newer Older
1
2
3
4
5
6
7
/*! @file hash_functions.h
 *  @brief Hash function code.
 */

#ifndef HASH_FUNCTIONS__H
#define HASH_FUNCTIONS__H

8
#include "definitions.h"
9
10
11
#include <tensorview/tensorview.h>
#include <vector_types.h>

traveller59's avatar
traveller59 committed
12
namespace cuhash {
13
14
15
16
17
18
19
20
21
22
23
24
25

//! 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.
26
27
    @param[in]  d_keys      Debug only: Device memory array containing the input
   keys.
28
29
    @param[in]  table_size  Debug only: Size of the hash table.
 */
30
31
32
void GenerateFunctions(const unsigned N, const unsigned num_keys,
                       const unsigned *d_keys, const unsigned table_size,
                       uint2 *constants);
33
34

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

  //! Generate new hash function constants.
41
42
  /*! The parameters are only used for debugging and examining the key
     distribution. \param[in] num_keys   Debug: Number of keys in the input.
43
44
45
      \param[in] d_keys     Debug: Device array of the input keys.
      \param[in] table_size Debug: Size of the hash table.
  */
46
47
  void Generate(const unsigned num_keys, const unsigned *d_keys,
                const unsigned table_size) {
48
49
50
51
52
53
54
55
56
    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.
 */
57
58
59
60
inline __device__ __host__ unsigned hash_function_inner(const uint2 constants,
                                                        const unsigned key) {
#if 1
  // Fast version.
61
62
63
64
65
  return ((constants.x ^ key) + constants.y) % kPrimeDivisor;
#else
  // Slow version.
  return ((unsigned long long)constants.x * key + constants.y) % kPrimeDivisor;
#endif
66
}
67
68
69
70
71
72
73
74

//! 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>
75
76
77
TV_HOST_DEVICE_INLINE unsigned
hash_function(const Functions<kNumHashFunctions> functions,
              const unsigned which_function, const unsigned key) {
78
79
80
81
82
  return hash_function_inner(functions.constants[which_function], key);
}

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

traveller59's avatar
traveller59 committed
87
unsigned generate_random_uint32();
88

89
}; // namespace cuhash
90
91

#endif