From c972ba4cd869b6340c447c9d1ebf5200907963df Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Mon, 7 Mar 2022 11:37:19 -0800
Subject: [PATCH] Revert "Remove template from hashing class, template on
 operator() and compute instead."

This reverts commit 466243b714cda4822eba5caecf5b284d4ddea474.
---
 .../cudf/detail/utilities/hash_functions.cuh  | 132 ++++++++++--------
 cpp/include/cudf/table/row_operators.cuh      |  10 +-
 cpp/src/hash/concurrent_unordered_map.cuh     |   2 +-
 .../hash/concurrent_unordered_multimap.cuh    |   4 +-
 cpp/src/hash/hashing.cu                       |   4 +-
 cpp/src/hash/unordered_multiset.cuh           |  10 +-
 cpp/src/io/json/json_gpu.cu                   |   4 +-
 cpp/src/io/parquet/chunk_dict.cu              |   4 +-
 cpp/src/partitioning/partitioning.cu          |   6 +-
 cpp/tests/hash_map/multimap_test.cu           |   4 +-
 10 files changed, 100 insertions(+), 80 deletions(-)

diff --git a/cpp/include/cudf/detail/utilities/hash_functions.cuh b/cpp/include/cudf/detail/utilities/hash_functions.cuh
index 3faff257cdf..1a56410e846 100644
--- a/cpp/include/cudf/detail/utilities/hash_functions.cuh
+++ b/cpp/include/cudf/detail/utilities/hash_functions.cuh
@@ -83,6 +83,7 @@ void __device__ inline uint32ToLowercaseHexString(uint32_t num, char* destinatio
 // algorithms are optimized for their respective platforms. You can still
 // compile and run any of them on any platform, but your performance with the
 // non-native version will be less than optimal.
+template <typename Key>
 struct MurmurHash3_32 {
   using result_type = hash_value_type;
 
@@ -114,8 +115,7 @@ struct MurmurHash3_32 {
   }
 
   // TODO Do we need this operator() and/or compute? Probably not both.
-  template <typename T>
-  [[nodiscard]] result_type __device__ inline operator()(T const& key) const
+  [[nodiscard]] result_type __device__ inline operator()(Key const& key) const
   {
     return compute(key);
   }
@@ -187,72 +187,82 @@ struct MurmurHash3_32 {
 };
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(bool const& key) const
+hash_value_type __device__ inline MurmurHash3_32<bool>::operator()(bool const& key) const
 {
   return this->compute(static_cast<uint8_t>(key));
 }
 
+/**
+ * @brief Specialization of MurmurHash3_32 operator for strings.
+ */
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(float const& key) const
+hash_value_type __device__ inline MurmurHash3_32<cudf::string_view>::operator()(
+  cudf::string_view const& key) const
 {
-  return this->compute_floating_point(key);
+  auto const data = reinterpret_cast<std::byte const*>(key.data());
+  auto const len  = key.size_bytes();
+  return this->compute_bytes(data, len);
 }
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(double const& key) const
+hash_value_type __device__ inline MurmurHash3_32<float>::operator()(float const& key) const
 {
   return this->compute_floating_point(key);
 }
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(cudf::string_view const& key) const
+hash_value_type __device__ inline MurmurHash3_32<double>::operator()(double const& key) const
 {
-  auto const data = reinterpret_cast<std::byte const*>(key.data());
-  auto const len  = key.size_bytes();
-  return this->compute_bytes(data, len);
+  return this->compute_floating_point(key);
 }
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(numeric::decimal32 const& key) const
+hash_value_type __device__ inline MurmurHash3_32<numeric::decimal32>::operator()(
+  numeric::decimal32 const& key) const
 {
   return this->compute(key.value());
 }
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(numeric::decimal64 const& key) const
+hash_value_type __device__ inline MurmurHash3_32<numeric::decimal64>::operator()(
+  numeric::decimal64 const& key) const
 {
   return this->compute(key.value());
 }
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(numeric::decimal128 const& key) const
+hash_value_type __device__ inline MurmurHash3_32<numeric::decimal128>::operator()(
+  numeric::decimal128 const& key) const
 {
   return this->compute(key.value());
 }
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(cudf::list_view const&) const
+hash_value_type __device__ inline MurmurHash3_32<cudf::list_view>::operator()(
+  cudf::list_view const& key) const
 {
   cudf_assert(false && "List column hashing is not supported");
   return 0;
 }
 
 template <>
-hash_value_type __device__ inline MurmurHash3_32::operator()(cudf::struct_view const&) const
+hash_value_type __device__ inline MurmurHash3_32<cudf::struct_view>::operator()(
+  cudf::struct_view const& key) const
 {
   cudf_assert(false && "Direct hashing of struct_view is not supported");
   return 0;
 }
 
+template <typename Key>
 struct SparkMurmurHash3_32 {
   using result_type = hash_value_type;
 
   SparkMurmurHash3_32() = default;
   constexpr SparkMurmurHash3_32(uint32_t seed) : m_seed(seed) {}
 
-  [[nodiscard]] __device__ inline uint32_t rotl32(uint32_t x, uint32_t r) const
+  __device__ inline uint32_t rotl32(uint32_t x, int8_t r) const
   {
-    return __funnelshift_l(x, x, r);  // Equivalent to (x << r) | (x >> (32 - r))
+    return (x << r) | (x >> (32 - r));
   }
 
   __device__ inline uint32_t fmix32(uint32_t h) const
@@ -265,11 +275,7 @@ struct SparkMurmurHash3_32 {
     return h;
   }
 
-  template <typename T>
-  result_type __device__ inline operator()(T const& key) const
-  {
-    return compute(key);
-  }
+  result_type __device__ inline operator()(Key const& key) const { return compute(key); }
 
   // compute wrapper for floating point types
   template <typename T, std::enable_if_t<std::is_floating_point_v<T>>* = nullptr>
@@ -346,72 +352,52 @@ struct SparkMurmurHash3_32 {
 };
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(bool const& key) const
+hash_value_type __device__ inline SparkMurmurHash3_32<bool>::operator()(bool const& key) const
 {
   return this->compute<uint32_t>(key);
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(int8_t const& key) const
+hash_value_type __device__ inline SparkMurmurHash3_32<int8_t>::operator()(int8_t const& key) const
 {
   return this->compute<uint32_t>(key);
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(int16_t const& key) const
+hash_value_type __device__ inline SparkMurmurHash3_32<uint8_t>::operator()(uint8_t const& key) const
 {
   return this->compute<uint32_t>(key);
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(uint8_t const& key) const
+hash_value_type __device__ inline SparkMurmurHash3_32<int16_t>::operator()(int16_t const& key) const
 {
   return this->compute<uint32_t>(key);
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(uint16_t const& key) const
+hash_value_type __device__ inline SparkMurmurHash3_32<uint16_t>::operator()(
+  uint16_t const& key) const
 {
   return this->compute<uint32_t>(key);
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(float const& key) const
-{
-  return this->compute_floating_point(key);
-}
-
-template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(double const& key) const
-{
-  return this->compute_floating_point(key);
-}
-
-template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(
-  cudf::string_view const& key) const
-{
-  auto const data = reinterpret_cast<std::byte const*>(key.data());
-  auto const len  = key.size_bytes();
-  return this->compute_bytes(data, len);
-}
-
-template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(
+hash_value_type __device__ inline SparkMurmurHash3_32<numeric::decimal32>::operator()(
   numeric::decimal32 const& key) const
 {
   return this->compute<uint64_t>(key.value());
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(
+hash_value_type __device__ inline SparkMurmurHash3_32<numeric::decimal64>::operator()(
   numeric::decimal64 const& key) const
 {
   return this->compute<uint64_t>(key.value());
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(
+hash_value_type __device__ inline SparkMurmurHash3_32<numeric::decimal128>::operator()(
   numeric::decimal128 const& key) const
 {
   // Generates the Spark MurmurHash3 hash value, mimicking the conversion:
@@ -453,37 +439,66 @@ hash_value_type __device__ inline SparkMurmurHash3_32::operator()(
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(cudf::list_view const&) const
+hash_value_type __device__ inline SparkMurmurHash3_32<cudf::list_view>::operator()(
+  cudf::list_view const& key) const
 {
   cudf_assert(false && "List column hashing is not supported");
   return 0;
 }
 
 template <>
-hash_value_type __device__ inline SparkMurmurHash3_32::operator()(cudf::struct_view const&) const
+hash_value_type __device__ inline SparkMurmurHash3_32<cudf::struct_view>::operator()(
+  cudf::struct_view const& key) const
 {
   cudf_assert(false && "Direct hashing of struct_view is not supported");
   return 0;
 }
 
+/**
+ * @brief Specialization of MurmurHash3_32 operator for strings.
+ */
+template <>
+hash_value_type __device__ inline SparkMurmurHash3_32<cudf::string_view>::operator()(
+  cudf::string_view const& key) const
+{
+  auto const data = reinterpret_cast<std::byte const*>(key.data());
+  auto const len  = key.size_bytes();
+  return this->compute_bytes(data, len);
+}
+
+template <>
+hash_value_type __device__ inline SparkMurmurHash3_32<float>::operator()(float const& key) const
+{
+  return this->compute_floating_point(key);
+}
+
+template <>
+hash_value_type __device__ inline SparkMurmurHash3_32<double>::operator()(double const& key) const
+{
+  return this->compute_floating_point(key);
+}
+
 /**
  * @brief  This hash function simply returns the value that is asked to be hash
  * reinterpreted as the result_type of the functor.
  */
+template <typename Key>
 struct IdentityHash {
   using result_type = hash_value_type;
   IdentityHash()    = default;
   constexpr IdentityHash(uint32_t seed) : m_seed(seed) {}
 
-  template <typename T, std::enable_if_t<!std::is_arithmetic_v<T>>* = nullptr>
-  constexpr result_type operator()(T const& key) const
+  template <typename return_type = result_type>
+  constexpr std::enable_if_t<!std::is_arithmetic_v<Key>, return_type> operator()(
+    Key const& key) const
   {
     cudf_assert(false && "IdentityHash does not support this data type");
     return 0;
   }
 
-  template <typename T, std::enable_if_t<std::is_arithmetic_v<T>>* = nullptr>
-  constexpr result_type operator()(T const& key) const
+  template <typename return_type = result_type>
+  constexpr std::enable_if_t<std::is_arithmetic_v<Key>, return_type> operator()(
+    Key const& key) const
   {
     return static_cast<result_type>(key);
   }
@@ -492,4 +507,5 @@ struct IdentityHash {
   uint32_t m_seed{cudf::DEFAULT_HASH_SEED};
 };
 
-using default_hash = MurmurHash3_32;
+template <typename Key>
+using default_hash = MurmurHash3_32<Key>;
diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh
index 730078a2675..6ff4833cb14 100644
--- a/cpp/include/cudf/table/row_operators.cuh
+++ b/cpp/include/cudf/table/row_operators.cuh
@@ -430,14 +430,14 @@ class row_lexicographic_comparator {
  * @tparam hash_function Hash functor to use for hashing elements.
  * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <typename hash_function, typename Nullate>
+template <template <typename> class hash_function, typename Nullate>
 class element_hasher {
  public:
   template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
     if (has_nulls && col.is_null(row_index)) { return std::numeric_limits<hash_value_type>::max(); }
-    return hash_function{}(col.element<T>(row_index));
+    return hash_function<T>{}(col.element<T>(row_index));
   }
 
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
@@ -450,7 +450,7 @@ class element_hasher {
   Nullate has_nulls;
 };
 
-template <typename hash_function, typename Nullate>
+template <template <typename> class hash_function, typename Nullate>
 class element_hasher_with_seed {
  public:
   __device__ element_hasher_with_seed(Nullate has_nulls, uint32_t seed)
@@ -467,7 +467,7 @@ class element_hasher_with_seed {
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
     if (_has_nulls && col.is_null(row_index)) { return _null_hash; }
-    return hash_function{_seed}(col.element<T>(row_index));
+    return hash_function<T>{_seed}(col.element<T>(row_index));
   }
 
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
@@ -489,7 +489,7 @@ class element_hasher_with_seed {
  * @tparam hash_function Hash functor to use for hashing elements.
  * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <typename hash_function, typename Nullate>
+template <template <typename> class hash_function, typename Nullate>
 class row_hasher {
  public:
   row_hasher() = delete;
diff --git a/cpp/src/hash/concurrent_unordered_map.cuh b/cpp/src/hash/concurrent_unordered_map.cuh
index 95e90087bdf..0ae0baa9908 100644
--- a/cpp/src/hash/concurrent_unordered_map.cuh
+++ b/cpp/src/hash/concurrent_unordered_map.cuh
@@ -113,7 +113,7 @@ union pair_packer<pair_type, std::enable_if_t<is_packable<pair_type>()>> {
  */
 template <typename Key,
           typename Element,
-          typename Hasher    = default_hash,
+          typename Hasher    = default_hash<Key>,
           typename Equality  = equal_to<Key>,
           typename Allocator = default_allocator<thrust::pair<Key, Element>>>
 class concurrent_unordered_map {
diff --git a/cpp/src/hash/concurrent_unordered_multimap.cuh b/cpp/src/hash/concurrent_unordered_multimap.cuh
index e1a2e4d0895..cdf5b6a8649 100644
--- a/cpp/src/hash/concurrent_unordered_multimap.cuh
+++ b/cpp/src/hash/concurrent_unordered_multimap.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2017-2020, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -52,7 +52,7 @@ template <typename Key,
           typename size_type,
           Key unused_key,
           Element unused_element,
-          typename Hasher       = default_hash,
+          typename Hasher       = default_hash<Key>,
           typename Equality     = equal_to<Key>,
           typename Allocator    = managed_allocator<thrust::pair<Key, Element>>,
           bool count_collisions = false>
diff --git a/cpp/src/hash/hashing.cu b/cpp/src/hash/hashing.cu
index 844c1bd20a6..ac2f06c0ea3 100644
--- a/cpp/src/hash/hashing.cu
+++ b/cpp/src/hash/hashing.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -49,7 +49,7 @@ std::vector<column_view> to_leaf_columns(IterType iter_begin, IterType iter_end)
 
 }  // namespace
 
-template <typename hash_function>
+template <template <typename> class hash_function>
 std::unique_ptr<column> serial_murmur_hash3_32(table_view const& input,
                                                uint32_t seed,
                                                rmm::cuda_stream_view stream,
diff --git a/cpp/src/hash/unordered_multiset.cuh b/cpp/src/hash/unordered_multiset.cuh
index 8764e9966d2..d28bf6f6fe5 100644
--- a/cpp/src/hash/unordered_multiset.cuh
+++ b/cpp/src/hash/unordered_multiset.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.  All rights reserved.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.  All rights reserved.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -32,7 +32,9 @@ namespace detail {
 /*
  *  Device view of the unordered multiset
  */
-template <typename Element, typename Hasher = default_hash, typename Equality = equal_to<Element>>
+template <typename Element,
+          typename Hasher   = default_hash<Element>,
+          typename Equality = equal_to<Element>>
 class unordered_multiset_device_view {
  public:
   unordered_multiset_device_view(size_type hash_size,
@@ -64,7 +66,9 @@ class unordered_multiset_device_view {
 /*
  * Fixed size set on a device.
  */
-template <typename Element, typename Hasher = default_hash, typename Equality = equal_to<Element>>
+template <typename Element,
+          typename Hasher   = default_hash<Element>,
+          typename Equality = equal_to<Element>>
 class unordered_multiset {
  public:
   /**
diff --git a/cpp/src/io/json/json_gpu.cu b/cpp/src/io/json/json_gpu.cu
index 2391e3f73dc..21455e3ab93 100644
--- a/cpp/src/io/json/json_gpu.cu
+++ b/cpp/src/io/json/json_gpu.cu
@@ -350,7 +350,7 @@ __device__ field_descriptor next_field_descriptor(const char* begin,
       ? field_descriptor{field_idx, begin, cudf::io::gpu::seek_field_end(begin, end, opts, true)}
       : [&]() {
           auto const key_range = get_next_key(begin, end, opts.quotechar);
-          auto const key_hash  = MurmurHash3_32{}(
+          auto const key_hash  = MurmurHash3_32<cudf::string_view>{}(
             cudf::string_view(key_range.first, key_range.second - key_range.first));
           auto const hash_col = col_map.find(key_hash);
           // Fall back to field index if not found (parsing error)
@@ -661,7 +661,7 @@ __global__ void collect_keys_info_kernel(parse_options_view const options,
       keys_info->column(0).element<uint64_t>(idx) = field_range.key_begin - data.begin();
       keys_info->column(1).element<uint16_t>(idx) = len;
       keys_info->column(2).element<uint32_t>(idx) =
-        MurmurHash3_32{}(cudf::string_view(field_range.key_begin, len));
+        MurmurHash3_32<cudf::string_view>{}(cudf::string_view(field_range.key_begin, len));
     }
   }
 }
diff --git a/cpp/src/io/parquet/chunk_dict.cu b/cpp/src/io/parquet/chunk_dict.cu
index 29e427f379f..5589f87e57e 100644
--- a/cpp/src/io/parquet/chunk_dict.cu
+++ b/cpp/src/io/parquet/chunk_dict.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -55,7 +55,7 @@ struct equality_functor {
 template <typename T>
 struct hash_functor {
   column_device_view const& col;
-  __device__ auto operator()(size_type idx) { return MurmurHash3_32{}(col.element<T>(idx)); }
+  __device__ auto operator()(size_type idx) { return MurmurHash3_32<T>{}(col.element<T>(idx)); }
 };
 
 struct map_insert_fn {
diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu
index 3f00b11c170..66b26148ede 100644
--- a/cpp/src/partitioning/partitioning.cu
+++ b/cpp/src/partitioning/partitioning.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -444,7 +444,7 @@ struct copy_block_partitions_dispatcher {
 };
 
 // NOTE hash_has_nulls must be true if table_to_hash has nulls
-template <typename hash_function, bool hash_has_nulls>
+template <template <typename> class hash_function, bool hash_has_nulls>
 std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
   table_view const& input,
   table_view const& table_to_hash,
@@ -711,7 +711,7 @@ struct dispatch_map_type {
 
 namespace detail {
 namespace local {
-template <typename hash_function>
+template <template <typename> class hash_function>
 std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition(
   table_view const& input,
   std::vector<size_type> const& columns_to_hash,
diff --git a/cpp/tests/hash_map/multimap_test.cu b/cpp/tests/hash_map/multimap_test.cu
index 34a901e3d20..456ba951a45 100644
--- a/cpp/tests/hash_map/multimap_test.cu
+++ b/cpp/tests/hash_map/multimap_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -49,7 +49,7 @@ class MultimapTest : public cudf::test::BaseFixture {
                                   size_type,
                                   std::numeric_limits<key_type>::max(),
                                   std::numeric_limits<value_type>::max(),
-                                  default_hash,
+                                  default_hash<key_type>,
                                   equal_to<key_type>,
                                   default_allocator<thrust::pair<key_type, value_type>>>;