Merge branch 'dev' into bf-tests

Author: sleeepyjack

benchmarks/static_multimap/count_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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.
@@ -57,12 +57,12 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_count(
 
   state.add_element_count(num_keys);
 
-  cuco::static_multimap<Key, Value> map{
+  cuco::experimental::static_multimap<Key, Value> map{
     size, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
   map.insert(pairs.begin(), pairs.end());
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
-    auto count = map.count(keys.begin(), keys.end(), launch.get_stream());
+    auto count = map.count(keys.begin(), keys.end(), {launch.get_stream()});
   });
 }
 

benchmarks/static_multimap/insert_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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.
@@ -56,11 +56,18 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_insert(
 
   state.exec(nvbench::exec_tag::sync | nvbench::exec_tag::timer,
              [&](nvbench::launch& launch, auto& timer) {
-               cuco::static_multimap<Key, Value> map{
-                 size, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}, launch.get_stream()};
+               cuco::experimental::static_multimap<Key, Value> map{size,
+                                                                   cuco::empty_key<Key>{-1},
+                                                                   cuco::empty_value<Value>{-1},
+                                                                   {},
+                                                                   {},
+                                                                   {},
+                                                                   {},
+                                                                   {},
+                                                                   {launch.get_stream()}};
 
                timer.start();
-               map.insert(pairs.begin(), pairs.end(), launch.get_stream());
+               map.insert(pairs.begin(), pairs.end(), {launch.get_stream()});
                timer.stop();
              });
 }

include/cuco/detail/open_addressing/open_addressing_impl.cuh

@@ -578,8 +578,8 @@ class open_addressing_impl {
    * a match with its key equivalent to the query key.
    *
    * @note If `pred( *(stencil + i) )` is true, stores the payload of the
-   * matched key or the `empty_value_sentienl` to `(output_begin + i)`. If `pred( *(stencil + i) )`
-   * is false, stores `empty_value_sentienl` to `(output_begin + i)`.
+   * matched key or the `empty_value_sentinel` to `(output_begin + i)`. If `pred( *(stencil + i) )`
+   * is false, stores `empty_value_sentinel` to `(output_begin + i)`.
    *
    * @tparam InputIt Device accessible input iterator
    * @tparam StencilIt Device accessible random access iterator whose value_type is
@@ -953,7 +953,7 @@ class open_addressing_impl {
    * @note Behavior is undefined if the desired `extent` is insufficient to store all of the
    * contained elements.
    *
-   * @note This function is not available if the conatiner's `extent_type` is static.
+   * @note This function is not available if the container's `extent_type` is static.
    *
    * @tparam Container The container type this function operates on
    *
@@ -994,7 +994,7 @@ class open_addressing_impl {
    * @note Behavior is undefined if the desired `extent` is insufficient to store all of the
    * contained elements.
    *
-   * @note This function is not available if the conatiner's `extent_type` is static.
+   * @note This function is not available if the container's `extent_type` is static.
    *
    * @tparam Container The container type this function operates on
    *

include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

@@ -309,7 +309,7 @@ class open_addressing_ref_impl {
    *
    * @tparam CG The type of the cooperative thread group
    *
-   * @param g The ooperative thread group used to copy the data structure
+   * @param g The cooperative thread group used to copy the data structure
    * @param memory_to_use Array large enough to support `capacity` elements. Object does not take
    * the ownership of the memory
    */
@@ -426,7 +426,7 @@ class open_addressing_ref_impl {
    *
    * @return True if the given element is successfully inserted
    */
-  template <typename Value>
+  template <bool SupportsErase, typename Value>
   __device__ bool insert(cooperative_groups::thread_block_tile<cg_size> const& group,
                          Value const& value) noexcept
   {
@@ -466,12 +466,20 @@ class open_addressing_ref_impl {
       auto const group_contains_available = group.ballot(state == detail::equal_result::AVAILABLE);
       if (group_contains_available) {
         auto const src_lane = __ffs(group_contains_available) - 1;
-        auto const status =
-          (group.thread_rank() == src_lane)
-            ? attempt_insert((storage_ref_.data() + *probing_iter)->data() + intra_bucket_index,
+        auto status         = insert_result::CONTINUE;
+        if (group.thread_rank() == src_lane) {
+          if constexpr (SupportsErase) {
+            status =
+              attempt_insert((storage_ref_.data() + *probing_iter)->data() + intra_bucket_index,
                              bucket_slots[intra_bucket_index],
-                             val)
-            : insert_result::CONTINUE;
+                             val);
+          } else {
+            status =
+              attempt_insert((storage_ref_.data() + *probing_iter)->data() + intra_bucket_index,
+                             this->empty_slot_sentinel(),
+                             val);
+          }
+        }
 
         switch (group.shfl(status, src_lane)) {
           case insert_result::SUCCESS: return true;

include/cuco/detail/operator.inl

@@ -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.
@@ -48,7 +48,7 @@ class operator_impl {
  * @return `true` if `Operator` is contained in `Operators`, `false` otherwise.
  */
 template <typename Operator, typename... Operators>
-static constexpr bool has_operator()
+__host__ __device__ static constexpr bool has_operator()
 {
   return ((std::is_same_v<Operators, Operator>) || ...);
 }

include/cuco/detail/static_map/static_map_ref.inl

@@ -449,7 +449,11 @@ class operator_impl<
                          Value const& value) noexcept
   {
     auto& ref_ = static_cast<ref_type&>(*this);
-    return ref_.impl_.insert(group, value);
+    if (ref_.erased_key_sentinel() != ref_.empty_key_sentinel()) {
+      return ref_.impl_.insert<true>(group, value);
+    } else {
+      return ref_.impl_.insert<false>(group, value);
+    }
   }
 };
 

include/cuco/detail/static_multimap/static_multimap_ref.inl

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-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.
@@ -455,7 +455,11 @@ class operator_impl<
                          Value const& value) noexcept
   {
     auto& ref_ = static_cast<ref_type&>(*this);
-    return ref_.impl_.insert(group, value);
+    if (ref_.erased_key_sentinel() != ref_.empty_key_sentinel()) {
+      return ref_.impl_.insert<true>(group, value);
+    } else {
+      return ref_.impl_.insert<false>(group, value);
+    }
   }
 };
 

include/cuco/detail/static_multiset/static_multiset_ref.inl

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2024, NVIDIA CORPORATION.
+ * Copyright (c) 2024-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.
@@ -358,7 +358,11 @@ class operator_impl<
                          Value const& value) noexcept
   {
     auto& ref_ = static_cast<ref_type&>(*this);
-    return ref_.impl_.insert(group, value);
+    if (ref_.erased_key_sentinel() != ref_.empty_key_sentinel()) {
+      return ref_.impl_.insert<true>(group, value);
+    } else {
+      return ref_.impl_.insert<false>(group, value);
+    }
   }
 };
 

include/cuco/detail/static_set/static_set_ref.inl

@@ -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.
@@ -391,7 +391,11 @@ class operator_impl<op::insert_tag,
                          Value const& value) noexcept
   {
     auto& ref_ = static_cast<ref_type&>(*this);
-    return ref_.impl_.insert(group, value);
+    if (ref_.erased_key_sentinel() != ref_.empty_key_sentinel()) {
+      return ref_.impl_.insert<true>(group, value);
+    } else {
+      return ref_.impl_.insert<false>(group, value);
+    }
   }
 };
 

include/cuco/static_map.cuh

@@ -50,12 +50,12 @@ namespace cuco {
  *
  * The host-side bulk operations include `insert`, `contains`, etc. These APIs should be used when
  * there are a large number of keys to modify or lookup. For example, given a range of keys
- * specified by device-accessible iterators, the bulk `insert` function will insert all keys into
+ * specified by device-accessible iterators, the bulk `insert` function inserts all keys into
  * the map.
  *
  * The singular device-side operations allow individual threads (or cooperative groups) to perform
  * independent modify or lookup operations from device code. These operations are accessed through
- * non-owning, trivially copyable reference types (or "ref"). User can combine any arbitrary
+ * non-owning, trivially copyable reference types (or "ref"). Users can combine any arbitrary
  * operators (see options in `include/cuco/operator.hpp`) when creating the ref. Concurrent modify
  * and lookup will be supported if both kinds of operators are specified during the ref
  * construction.
@@ -799,8 +799,8 @@ class static_map {
    * query key.
    *
    * @note If `pred( *(stencil + i) )` is true, stores the payload of the
-   * matched key or the `empty_value_sentienl` to `(output_begin + i)`. If `pred( *(stencil + i) )`
-   * is false, always stores the `empty_value_sentienl` to `(output_begin + i)`.
+   * matched key or the `empty_value_sentinel` to `(output_begin + i)`. If `pred( *(stencil + i) )`
+   * is false, always stores the `empty_value_sentinel` to `(output_begin + i)`.
    * @note This function synchronizes the given stream. For asynchronous execution use
    * `find_if_async`.
    *
@@ -832,8 +832,8 @@ class static_map {
    * a match with its key equivalent to the query key.
    *
    * @note If `pred( *(stencil + i) )` is true, stores the payload of the
-   * matched key or the `empty_value_sentienl` to `(output_begin + i)`. If `pred( *(stencil + i) )`
-   * is false, always stores the `empty_value_sentienl` to `(output_begin + i)`.
+   * matched key or the `empty_value_sentinel` to `(output_begin + i)`. If `pred( *(stencil + i) )`
+   * is false, always stores the `empty_value_sentinel` to `(output_begin + i)`.
    *
    * @tparam InputIt Device accessible input iterator
    * @tparam StencilIt Device accessible random access iterator whose `value_type` is convertible to
@@ -863,8 +863,8 @@ class static_map {
    * a match with its key equivalent to the query key.
    *
    * @note If `pred( *(stencil + i) )` is true, stores the payload of the
-   * matched key or the `empty_value_sentienl` to `(output_begin + i)`. If `pred( *(stencil + i) )`
-   * is false, always stores the `empty_value_sentienl` to `(output_begin + i)`.
+   * matched key or the `empty_value_sentinel` to `(output_begin + i)`. If `pred( *(stencil + i) )`
+   * is false, always stores the `empty_value_sentinel` to `(output_begin + i)`.
    *
    * @tparam InputIt Device accessible input iterator
    * @tparam StencilIt Device accessible random access iterator whose `value_type` is convertible to
@@ -1029,7 +1029,7 @@ class static_map {
    *
    * @tparam KeyOut Device accessible random access output iterator whose `value_type` is
    * convertible from `key_type`.
-   * @tparam ValueOut Device accesible random access output iterator whose `value_type` is
+   * @tparam ValueOut Device accessible random access output iterator whose `value_type` is
    * convertible from `mapped_type`.
    *
    * @param keys_out Beginning output iterator for keys
@@ -1066,7 +1066,7 @@ class static_map {
    * @note Behavior is undefined if the desired `capacity` is insufficient to store all of the
    * contained elements.
    *
-   * @note This function is not available if the conatiner's `extent_type` is static.
+   * @note This function is not available if the container's `extent_type` is static.
    *
    * @param capacity New capacity of the container
    * @param stream CUDA stream used for this operation
@@ -1091,7 +1091,7 @@ class static_map {
    * @note Behavior is undefined if the desired `capacity` is insufficient to store all of the
    * contained elements.
    *
-   * @note This function is not available if the conatiner's `extent_type` is static.
+   * @note This function is not available if the container's `extent_type` is static.
    *
    * @param capacity New capacity of the container
    * @param stream CUDA stream used for this operation
@@ -1512,7 +1512,7 @@ class static_map {
    *
    * @tparam KeyOut Device accessible random access output iterator whose `value_type` is
    * convertible from `key_type`.
-   * @tparam ValueOut Device accesible random access output iterator whose `value_type` is
+   * @tparam ValueOut Device accessible random access output iterator whose `value_type` is
    * convertible from `mapped_type`.
    * @param keys_out Beginning output iterator for keys
    * @param values_out Beginning output iterator for values
@@ -2280,7 +2280,7 @@ class static_map {
      * @endcode
      *
      * @tparam CG The type of the cooperative thread group
-     * @param g The ooperative thread group used to copy the slots
+     * @param g The cooperative thread group used to copy the slots
      * @param source_device_view `device_view` to copy from
      * @param memory_to_use Array large enough to support `capacity` elements. Object does not take
      * the ownership of the memory

include/cuco/static_map_ref.cuh

@@ -276,7 +276,7 @@ class static_map_ref
    * @tparam CG The type of the cooperative thread group
    * @tparam NewScope The thread scope of the newly created device ref
    *
-   * @param tile The ooperative thread group used to copy the data structure
+   * @param tile The cooperative thread group used to copy the data structure
    * @param memory_to_use Array large enough to support `capacity` elements. Object does not take
    * the ownership of the memory
    * @param scope The thread scope of the newly created device ref