Add identity hasher (#514)

include/cuco/detail/hash_functions/identity_hash.cuh

@@ -0,0 +1,56 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <cuda/std/type_traits>
+#include <thrust/functional.h>
+
+namespace cuco::detail {
+
+/**
+ * @brief An Identity hash function to hash the given argument on host and device
+ *
+ * @note `identity_hash` is perfect if `hash_table_capacity >= |input set|`
+ *
+ * @note `identity_hash` is only intended to be used perfectly.
+ *
+ * @note Perfect hashes are deterministic, and thus do not need seeds.
+ *
+ * @tparam Key The type of the values to hash
+ */
+template <typename Key>
+struct identity_hash : private thrust::identity<Key> {
+  using argument_type = Key;  ///< The type of the values taken as argument
+  /// The type of the hash values produced
+  using result_type = cuda::std::conditional_t<sizeof(Key) <= 4, uint32_t, uint64_t>;
+
+  static_assert(cuda::std::is_convertible_v<Key, result_type>,
+                "Key type must be convertible to result_type");
+
+  /**
+   * @brief Returns a hash value for its argument, as a value of type `result_type`.
+   *
+   * @param x The input argument to hash
+   * @return A resulting hash value for `x`
+   */
+  __host__ __device__ result_type operator()(Key const& x) const
+  {
+    return static_cast<result_type>(thrust::identity<Key>::operator()(x));
+  }
+};  // identity_hash
+
+}  //  namespace cuco::detail

include/cuco/hash_functions.cuh

@@ -16,11 +16,24 @@
 
 #pragma once
 
+#include <cuco/detail/hash_functions/identity_hash.cuh>
 #include <cuco/detail/hash_functions/murmurhash3.cuh>
 #include <cuco/detail/hash_functions/xxhash.cuh>
 
+#include <thrust/functional.h>
+
 namespace cuco {
 
+/**
+ * @brief An Identity hash function to hash the given argument on host and device
+ *
+ * @throw A key must not be larger than uint64_t
+ *
+ * @tparam Key The type of the values to hash
+ */
+template <typename Key>
+using identity_hash = detail::identity_hash<Key>;
+
 /**
  * @brief The 32-bit integer finalizer function of `MurmurHash3` to hash the given argument on host
  * and device.

tests/utility/hash_test.cu

@@ -48,6 +48,59 @@ __host__ __device__ bool check_hash_result(typename Hash::argument_type const& k
   return (h(key) == expected);
 }
 
+// Overload for hash functions without a seed
+template <typename Hash>
+__host__ __device__ bool check_hash_result(typename Hash::argument_type const& key,
+                                           typename Hash::result_type expected) noexcept
+{
+  Hash h;
+  return (h(key) == expected);
+}
+
+template <typename OutputIter>
+__global__ void check_identity_hash_result_kernel(OutputIter result)
+{
+  int i = 0;
+
+  result[i++] = check_hash_result<cuco::identity_hash<signed char>>(0, 0);
+  result[i++] = check_hash_result<cuco::identity_hash<signed char>>(
+    std::numeric_limits<signed char>::max(), std::numeric_limits<signed char>::max());
+
+  result[i++] = check_hash_result<cuco::identity_hash<int32_t>>(0, 0);
+  result[i++] = check_hash_result<cuco::identity_hash<int32_t>>(
+    std::numeric_limits<int32_t>::max(), std::numeric_limits<int32_t>::max());
+
+  result[i++] = check_hash_result<cuco::identity_hash<int64_t>>(0, 0);
+  result[i++] = check_hash_result<cuco::identity_hash<int64_t>>(
+    std::numeric_limits<int64_t>::max(), std::numeric_limits<int64_t>::max());
+}
+
+TEST_CASE("Test cuco::identity_hash", "")
+{
+  SECTION("Check if host-generated hash values match the identity function.")
+  {
+    CHECK(check_hash_result<cuco::identity_hash<signed char>>(0, 0));
+    CHECK(check_hash_result<cuco::identity_hash<signed char>>(
+      std::numeric_limits<signed char>::max(), std::numeric_limits<signed char>::max()));
+
+    CHECK(check_hash_result<cuco::identity_hash<int32_t>>(0, 0));
+    CHECK(check_hash_result<cuco::identity_hash<int32_t>>(std::numeric_limits<int32_t>::max(),
+                                                          std::numeric_limits<int32_t>::max()));
+
+    CHECK(check_hash_result<cuco::identity_hash<int64_t>>(0, 0));
+    CHECK(check_hash_result<cuco::identity_hash<int64_t>>(std::numeric_limits<int64_t>::max(),
+                                                          std::numeric_limits<int64_t>::max()));
+  }
+  SECTION("Check if device-generated hash values match the identity function.")
+  {
+    thrust::device_vector<bool> result(7, true);
+
+    check_identity_hash_result_kernel<<<1, 1>>>(result.begin());
+
+    CHECK(cuco::test::all_of(result.begin(), result.end(), [] __device__(bool v) { return v; }));
+  }
+}
+
 template <typename OutputIter>
 __global__ void check_hash_result_kernel_64(OutputIter result)
 {
@@ -351,4 +404,4 @@ TEST_CASE("Test cuco::murmurhash3_x64_128", "")
 
     CHECK(cuco::test::all_of(result.begin(), result.end(), [] __device__(bool v) { return v; }));
   }
-}
+}