Make XXHash the default hash function (#318)

NVIDIA · Jun 29, 2023 · 6bc62c6 · 6bc62c6
1 parent 5e27a54
commit 6bc62c6
Show file tree

Hide file tree

Showing 19 changed files with 84 additions and 101 deletions.
diff --git a/include/cuco/dynamic_map.cuh b/include/cuco/dynamic_map.cuh
@@ -208,7 +208,7 @@ class dynamic_map {
    * @param stream Stream used for executing the kernels
    */
   template <typename InputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void insert(InputIt first,
               InputIt last,
@@ -247,7 +247,7 @@ class dynamic_map {
    * provided at construction
    */
   template <typename InputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void erase(InputIt first,
              InputIt last,
@@ -277,7 +277,7 @@ class dynamic_map {
    */
   template <typename InputIt,
             typename OutputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void find(InputIt first,
             InputIt last,
@@ -307,7 +307,7 @@ class dynamic_map {
    */
   template <typename InputIt,
             typename OutputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void contains(InputIt first,
                 InputIt last,

diff --git a/include/cuco/hash_functions.cuh b/include/cuco/hash_functions.cuh
@@ -44,27 +44,35 @@ template <typename Key>
 using murmurhash3_fmix_64 = detail::MurmurHash3_fmix64<Key>;
 
 /**
- * @brief A `murmurhash3_32` hash function to hash the given argument on host and device.
+ * @brief A 32-bit `MurmurHash3` hash function to hash the given argument on host and device.
  *
  * @tparam Key The type of the values to hash
  */
 template <typename Key>
 using murmurhash3_32 = detail::MurmurHash3_32<Key>;
 
 /**
- * @brief A `XXH32` hash function to hash the given argument on host and device.
+ * @brief A 32-bit `XXH32` hash function to hash the given argument on host and device.
  *
  * @tparam Key The type of the values to hash
  */
 template <typename Key>
 using xxhash_32 = detail::XXHash_32<Key>;
 
 /**
- * @brief A `XXH64` hash function to hash the given argument on host and device.
+ * @brief A 64-bit `XXH64` hash function to hash the given argument on host and device.
  *
  * @tparam Key The type of the values to hash
  */
 template <typename Key>
 using xxhash_64 = detail::XXHash_64<Key>;
 
+/**
+ * @brief Default hash function.
+ *
+ * @tparam Key The type of the values to hash
+ */
+template <typename Key>
+using default_hash_function = xxhash_32<Key>;
+
 }  // namespace cuco
diff --git a/include/cuco/probe_sequences.cuh b/include/cuco/probe_sequences.cuh
@@ -60,7 +60,7 @@ class linear_probing : public detail::probe_sequence_base<CGSize> {
  * @tparam Hash1 Unary callable type
  * @tparam Hash2 Unary callable type
  */
-template <uint32_t CGSize, typename Hash1, typename Hash2>
+template <uint32_t CGSize, typename Hash1, typename Hash2 = Hash1>
 class double_hashing : public detail::probe_sequence_base<CGSize> {
  public:
   using probe_sequence_base_type =

diff --git a/include/cuco/probing_scheme.cuh b/include/cuco/probing_scheme.cuh
@@ -96,7 +96,7 @@ class linear_probing : private detail::probing_scheme_base<CGSize> {
  * @tparam Hash1 Unary callable type
  * @tparam Hash2 Unary callable type
  */
-template <int32_t CGSize, typename Hash1, typename Hash2>
+template <int32_t CGSize, typename Hash1, typename Hash2 = Hash1>
 class double_hashing : private detail::probing_scheme_base<CGSize> {
  public:
   using probing_scheme_base_type =

diff --git a/include/cuco/static_map.cuh b/include/cuco/static_map.cuh
@@ -89,11 +89,11 @@ template <class Key,
           class Extent             = cuco::experimental::extent<std::size_t>,
           cuda::thread_scope Scope = cuda::thread_scope_device,
           class KeyEqual           = thrust::equal_to<Key>,
-          class ProbingScheme      = cuco::experimental::double_hashing<4,  // CG size
-                                                                   cuco::murmurhash3_32<Key>,
-                                                                   cuco::murmurhash3_32<Key>>,
-          class Allocator          = cuco::cuda_allocator<cuco::pair<Key, T>>,
-          class Storage            = cuco::experimental::aow_storage<1>>
+          class ProbingScheme =
+            cuco::experimental::double_hashing<4,  // CG size
+                                               cuco::default_hash_function<Key>>,
+          class Allocator = cuco::cuda_allocator<cuco::pair<Key, T>>,
+          class Storage   = cuco::experimental::aow_storage<1>>
 class static_map {
   static_assert(sizeof(Key) <= 4, "Container does not support key types larger than 4 bytes.");
 
@@ -691,7 +691,7 @@ class static_map {
    * @param stream Stream used for executing the kernels
    */
   template <typename InputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void insert(InputIt first,
               InputIt last,
@@ -725,7 +725,7 @@ class static_map {
   template <typename InputIt,
             typename StencilIt,
             typename Predicate,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void insert_if(InputIt first,
                  InputIt last,
@@ -763,7 +763,7 @@ class static_map {
    * provided at construction
    */
   template <typename InputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void erase(InputIt first,
              InputIt last,
@@ -792,7 +792,7 @@ class static_map {
    */
   template <typename InputIt,
             typename OutputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void find(InputIt first,
             InputIt last,
@@ -847,7 +847,7 @@ class static_map {
    */
   template <typename InputIt,
             typename OutputIt,
-            typename Hash     = cuco::murmurhash3_32<key_type>,
+            typename Hash     = cuco::default_hash_function<key_type>,
             typename KeyEqual = thrust::equal_to<key_type>>
   void contains(InputIt first,
                 InputIt last,
@@ -1369,7 +1369,7 @@ class static_map {
      * equality
      * @return `true` if the insert was successful, `false` otherwise.
      */
-    template <typename Hash     = cuco::murmurhash3_32<key_type>,
+    template <typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ bool insert(value_type const& insert_pair,
                            Hash hash          = Hash{},
@@ -1400,7 +1400,7 @@ class static_map {
      * @return a pair consisting of an iterator to the element and a bool,
      * either `true` if the insert was successful, `false` otherwise.
      */
-    template <typename Hash     = cuco::murmurhash3_32<key_type>,
+    template <typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ thrust::pair<iterator, bool> insert_and_find(
       value_type const& insert_pair, Hash hash = Hash{}, KeyEqual key_equal = KeyEqual{}) noexcept;
@@ -1425,7 +1425,7 @@ class static_map {
      * @return `true` if the insert was successful, `false` otherwise.
      */
     template <typename CG,
-              typename Hash     = cuco::murmurhash3_32<key_type>,
+              typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ bool insert(CG const& g,
                            value_type const& insert_pair,
@@ -1446,7 +1446,7 @@ class static_map {
      * equality
      * @return `true` if the erasure was successful, `false` otherwise.
      */
-    template <typename Hash     = cuco::murmurhash3_32<key_type>,
+    template <typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ bool erase(key_type const& k,
                           Hash hash          = Hash{},
@@ -1469,7 +1469,7 @@ class static_map {
      * @return `true` if the erasure was successful, `false` otherwise.
      */
     template <typename CG,
-              typename Hash     = cuco::murmurhash3_32<key_type>,
+              typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ bool erase(CG const& g,
                           key_type const& k,
@@ -1636,7 +1636,7 @@ class static_map {
      * @return An iterator to the position at which the key/value pair
      * containing `k` was inserted
      */
-    template <typename Hash     = cuco::murmurhash3_32<key_type>,
+    template <typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ iterator find(Key const& k,
                              Hash hash          = Hash{},
@@ -1656,7 +1656,7 @@ class static_map {
      * @return An iterator to the position at which the key/value pair
      * containing `k` was inserted
      */
-    template <typename Hash     = cuco::murmurhash3_32<key_type>,
+    template <typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ const_iterator find(Key const& k,
                                    Hash hash          = Hash{},
@@ -1683,7 +1683,7 @@ class static_map {
      * containing `k` was inserted
      */
     template <typename CG,
-              typename Hash     = cuco::murmurhash3_32<key_type>,
+              typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ iterator
     find(CG g, Key const& k, Hash hash = Hash{}, KeyEqual key_equal = KeyEqual{}) noexcept;
@@ -1709,7 +1709,7 @@ class static_map {
      * containing `k` was inserted
      */
     template <typename CG,
-              typename Hash     = cuco::murmurhash3_32<key_type>,
+              typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ const_iterator
     find(CG g, Key const& k, Hash hash = Hash{}, KeyEqual key_equal = KeyEqual{}) const noexcept;
@@ -1738,7 +1738,7 @@ class static_map {
      * containing `k` was inserted
      */
     template <typename ProbeKey,
-              typename Hash     = cuco::murmurhash3_32<key_type>,
+              typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ bool contains(ProbeKey const& k,
                              Hash hash          = Hash{},
@@ -1773,7 +1773,7 @@ class static_map {
      */
     template <typename CG,
               typename ProbeKey,
-              typename Hash     = cuco::murmurhash3_32<key_type>,
+              typename Hash     = cuco::default_hash_function<key_type>,
               typename KeyEqual = thrust::equal_to<key_type>>
     __device__ std::enable_if_t<std::is_invocable_v<KeyEqual, ProbeKey, Key>, bool> contains(
       CG const& g,

diff --git a/include/cuco/static_multimap.cuh b/include/cuco/static_multimap.cuh
@@ -130,8 +130,7 @@ template <typename Key,
           typename Value,
           cuda::thread_scope Scope = cuda::thread_scope_device,
           typename Allocator       = cuco::cuda_allocator<char>,
-          class ProbeSequence =
-            cuco::double_hashing<8, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>
+          class ProbeSequence      = cuco::double_hashing<8, cuco::default_hash_function<Key>>>
 class static_multimap {
   static_assert(
     cuco::is_bitwise_comparable_v<Key>,

diff --git a/include/cuco/static_set.cuh b/include/cuco/static_set.cuh
@@ -85,8 +85,7 @@ template <class Key,
           cuda::thread_scope Scope = cuda::thread_scope_device,
           class KeyEqual           = thrust::equal_to<Key>,
           class ProbingScheme      = experimental::double_hashing<4,  // CG size
-                                                             cuco::murmurhash3_32<Key>,
-                                                             cuco::murmurhash3_32<Key>>,
+                                                             cuco::default_hash_function<Key>>,
           class Allocator          = cuco::cuda_allocator<Key>,
           class Storage            = cuco::experimental::aow_storage<1>>
 class static_set {

diff --git a/tests/static_map/key_sentinel_test.cu b/tests/static_map/key_sentinel_test.cu
@@ -62,15 +62,17 @@ TEMPLATE_TEST_CASE_SIG(
                                pairs_begin + num_keys,
                                [m_view] __device__(cuco::pair<Key, Value> const& pair) mutable {
                                  return m_view.insert(
-                                   pair, cuco::murmurhash3_32<Key>{}, custom_equals<Key>{});
+                                   pair, cuco::default_hash_function<Key>{}, custom_equals<Key>{});
                                }));
   }
 
   SECTION(
     "Tests of CG insert: The custom `key_equal` can never be used to compare against sentinel")
   {
-    map.insert(
-      pairs_begin, pairs_begin + num_keys, cuco::murmurhash3_32<Key>{}, custom_equals<Key>{});
+    map.insert(pairs_begin,
+               pairs_begin + num_keys,
+               cuco::default_hash_function<Key>{},
+               custom_equals<Key>{});
     // All keys inserted via custom `key_equal` should be found
     REQUIRE(cuco::test::all_of(
       pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair<Key, Value> const& pair) {

diff --git a/tests/static_map/stream_test.cu b/tests/static_map/stream_test.cu
@@ -57,7 +57,7 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
     thrust::make_transform_iterator(thrust::make_counting_iterator<int>(0),
                                     [] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });
 
-  auto hash_fn  = cuco::murmurhash3_32<Key>{};
+  auto hash_fn  = cuco::default_hash_function<Key>{};
   auto equal_fn = thrust::equal_to<Value>{};
 
   // bulk function test cases

diff --git a/tests/static_multimap/custom_pair_retrieve_test.cu b/tests/static_multimap/custom_pair_retrieve_test.cu
@@ -196,10 +196,9 @@ TEMPLATE_TEST_CASE_SIG(
 {
   constexpr std::size_t num_pairs{200};
 
-  using probe = std::conditional_t<
-    Probe == cuco::test::probe_sequence::linear_probing,
-    cuco::linear_probing<1, cuco::murmurhash3_32<Key>>,
-    cuco::double_hashing<8, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>;
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<1, cuco::default_hash_function<Key>>,
+                                   cuco::double_hashing<8, cuco::default_hash_function<Key>>>;
 
   cuco::static_multimap<Key, Value, cuda::thread_scope_device, cuco::cuda_allocator<char>, probe>
     map{num_pairs * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};

diff --git a/tests/static_multimap/insert_if_test.cu b/tests/static_multimap/insert_if_test.cu
@@ -67,10 +67,9 @@ TEMPLATE_TEST_CASE_SIG(
                       return cuco::pair<Key, Value>{i, i};
                     });
 
-  using probe = std::conditional_t<
-    Probe == cuco::test::probe_sequence::linear_probing,
-    cuco::linear_probing<1, cuco::murmurhash3_32<Key>>,
-    cuco::double_hashing<8, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>;
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<1, cuco::default_hash_function<Key>>,
+                                   cuco::double_hashing<8, cuco::default_hash_function<Key>>>;
 
   cuco::static_multimap<Key, Value, cuda::thread_scope_device, cuco::cuda_allocator<char>, probe>
     map{num_keys * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};

diff --git a/tests/static_multimap/multiplicity_test.cu b/tests/static_multimap/multiplicity_test.cu
@@ -161,10 +161,9 @@ TEMPLATE_TEST_CASE_SIG(
 {
   constexpr std::size_t num_items{4};
 
-  using probe = std::conditional_t<
-    Probe == cuco::test::probe_sequence::linear_probing,
-    cuco::linear_probing<1, cuco::murmurhash3_32<Key>>,
-    cuco::double_hashing<8, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>;
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<1, cuco::default_hash_function<Key>>,
+                                   cuco::double_hashing<8, cuco::default_hash_function<Key>>>;
 
   cuco::static_multimap<Key, Value, cuda::thread_scope_device, cuco::cuda_allocator<char>, probe>
     map{5, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};

diff --git a/tests/static_multimap/non_match_test.cu b/tests/static_multimap/non_match_test.cu
@@ -139,16 +139,15 @@ TEMPLATE_TEST_CASE_SIG(
                       return cuco::pair<Key, Value>{i / 2, i};
                     });
 
-  using probe = std::conditional_t<
-    Probe == cuco::test::probe_sequence::linear_probing,
-    cuco::linear_probing<1, cuco::murmurhash3_32<Key>>,
-    cuco::double_hashing<8, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>;
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<1, cuco::default_hash_function<Key>>,
+                                   cuco::double_hashing<8, cuco::default_hash_function<Key>>>;
 
   cuco::static_multimap<Key,
                         Value,
                         cuda::thread_scope_device,
                         cuco::cuda_allocator<char>,
-                        cuco::linear_probing<1, cuco::murmurhash3_32<Key>>>
+                        cuco::linear_probing<1, cuco::default_hash_function<Key>>>
     map{num_keys * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
   test_non_matches<Key, Value>(map, d_pairs.begin(), d_keys.begin(), num_keys);
 }
diff --git a/tests/static_multimap/pair_function_test.cu b/tests/static_multimap/pair_function_test.cu
@@ -132,10 +132,9 @@ TEMPLATE_TEST_CASE_SIG(
                       return cuco::pair<Key, Value>{i / 2, i};
                     });
 
-  using probe = std::conditional_t<
-    Probe == cuco::test::probe_sequence::linear_probing,
-    cuco::linear_probing<1, cuco::murmurhash3_32<Key>>,
-    cuco::double_hashing<8, cuco::murmurhash3_32<Key>, cuco::murmurhash3_32<Key>>>;
+  using probe = std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
+                                   cuco::linear_probing<1, cuco::default_hash_function<Key>>,
+                                   cuco::double_hashing<8, cuco::default_hash_function<Key>>>;
 
   cuco::static_multimap<Key, Value, cuda::thread_scope_device, cuco::cuda_allocator<char>, probe>
     map{num_pairs * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};