Bypass key equal check for multiset and multimap insertions (#769)

Closes #767 

This PR skips key equality checks during multiset and multimap
insertions. It also changes enum types from `int32_t` to `int8_t` to
reduce register usage, which may help address issue #761.

Author: PointKernel

include/cuco/detail/equal_wrapper.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,18 +17,17 @@
 
 #include <cuco/detail/bitwise_compare.cuh>
 
-#include <cstddef>
+#include <cuda/std/cstdint>
 
-namespace cuco {
-namespace detail {
+namespace cuco::detail {
 
 /**
  * @brief Enum of equality comparison results
  */
 // ENUM VALUE MATTERS, DO NOT CHANGE
-enum class equal_result : int32_t { UNEQUAL = 0, EQUAL = 1, EMPTY = 2, AVAILABLE = 3 };
+enum class equal_result : cuda::std::int8_t { UNEQUAL = 0, EQUAL = 1, EMPTY = 2, AVAILABLE = 3 };
 
-enum class is_insert : bool { YES, NO };
+enum class is_insert : cuda::std::int8_t { YES, NO };
 
 /**
  * @brief Key equality wrapper.
@@ -37,8 +36,9 @@ enum class is_insert : bool { YES, NO };
  *
  * @tparam T Right-hand side Element type
  * @tparam Equal Type of user-provided equality binary callable
+ * @tparam AllowsDuplicates Flag indicating whether duplicate keys are allowed
  */
-template <typename T, typename Equal>
+template <typename T, typename Equal, bool AllowsDuplicates>
 struct equal_wrapper {
   // TODO: Clean up the sentinel handling since it's duplicated in ref and equal wrapper
   T empty_sentinel_;   ///< Empty sentinel value
@@ -97,16 +97,22 @@ struct equal_wrapper {
   __device__ constexpr equal_result operator()(LHS const& lhs, RHS const& rhs) const noexcept
   {
     if constexpr (IsInsert == is_insert::YES) {
-      return (cuco::detail::bitwise_compare(rhs, empty_sentinel_) or
-              cuco::detail::bitwise_compare(rhs, erased_sentinel_))
-               ? equal_result::AVAILABLE
-               : this->equal_to(lhs, rhs);
+      if (cuco::detail::bitwise_compare(rhs, empty_sentinel_) or
+          cuco::detail::bitwise_compare(rhs, erased_sentinel_)) {
+        return equal_result::AVAILABLE;
+      }
+      // Optimization: For containers that allow duplicates, skip expensive key equality check
+      // during insertion since we always insert regardless of whether the key already exists
+      if constexpr (AllowsDuplicates) {
+        return equal_result::UNEQUAL;
+      } else {
+        return this->equal_to(lhs, rhs);
+      }
     } else {
       return cuco::detail::bitwise_compare(rhs, empty_sentinel_) ? equal_result::EMPTY
                                                                  : this->equal_to(lhs, rhs);
     }
   }
 };
 
-}  // namespace detail
-}  // namespace cuco
+}  // namespace cuco::detail

include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

@@ -42,7 +42,7 @@ namespace cuco {
 namespace detail {
 
 /// Three-way insert result enum
-enum class insert_result : cuda::std::int32_t { CONTINUE = 0, SUCCESS = 1, DUPLICATE = 2 };
+enum class insert_result : cuda::std::int8_t { CONTINUE = 0, SUCCESS = 1, DUPLICATE = 2 };
 
 /**
  * @brief Helper struct to store intermediate bucket probing results.
@@ -217,8 +217,9 @@ class open_addressing_ref_impl {
    *
    * @return The key equality predicate
    */
-  [[nodiscard]] __host__ __device__ constexpr detail::equal_wrapper<key_type, key_equal> predicate()
-    const noexcept
+  [[nodiscard]] __host__
+    __device__ constexpr detail::equal_wrapper<key_type, key_equal, allows_duplicates>
+    predicate() const noexcept
   {
     return this->predicate_;
   }
@@ -1927,9 +1928,10 @@ class open_addressing_ref_impl {
 
   // TODO: Clean up the sentinel handling since it's duplicated in ref and equal wrapper
   value_type empty_slot_sentinel_;  ///< Sentinel value indicating an empty slot
-  detail::equal_wrapper<key_type, key_equal> predicate_;  ///< Key equality binary callable
-  probing_scheme_type probing_scheme_;                    ///< Probing scheme
-  storage_ref_type storage_ref_;                          ///< Slot storage ref
+  detail::equal_wrapper<key_type, key_equal, allows_duplicates>
+    predicate_;                         ///< Key equality binary callable
+  probing_scheme_type probing_scheme_;  ///< Probing scheme
+  storage_ref_type storage_ref_;        ///< Slot storage ref
 };
 
 }  // namespace detail