Add make_copy and initialize for multiset ref (#693)

This PR adds `make_copy` and `initialize` member functions for multiset
ref.

Author: PointKernel

include/cuco/detail/static_multiset/static_multiset_ref.inl

@@ -305,6 +305,45 @@ static_multiset_ref<Key, Scope, KeyEqual, ProbingScheme, StorageRef, Operators..
                                            {},
                                            this->storage_ref()};
 }
+template <typename Key,
+          cuda::thread_scope Scope,
+          typename KeyEqual,
+          typename ProbingScheme,
+          typename StorageRef,
+          typename... Operators>
+template <typename CG, cuda::thread_scope NewScope>
+__device__ constexpr auto
+static_multiset_ref<Key, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::make_copy(
+  CG const& tile,
+  bucket_type* const memory_to_use,
+  cuda_thread_scope<NewScope> scope) const noexcept
+{
+  auto const storage_ref = this->storage_ref().make_copy(tile, memory_to_use);
+  return static_multiset_ref<Key,
+                             NewScope,
+                             KeyEqual,
+                             ProbingScheme,
+                             decltype(storage_ref),
+                             Operators...>{cuco::empty_key<Key>{this->empty_key_sentinel()},
+                                           this->key_eq(),
+                                           this->probing_scheme(),
+                                           scope,
+                                           storage_ref};
+}
+
+template <typename Key,
+          cuda::thread_scope Scope,
+          typename KeyEqual,
+          typename ProbingScheme,
+          typename StorageRef,
+          typename... Operators>
+template <typename CG>
+__device__ constexpr void
+static_multiset_ref<Key, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::initialize(
+  CG const& tile) noexcept
+{
+  this->storage_ref().initialize(tile, this->empty_key_sentinel());
+}
 
 namespace detail {
 

include/cuco/static_multiset_ref.cuh

@@ -243,6 +243,43 @@ class static_multiset_ref
   template <typename NewHash>
   [[nodiscard]] __host__ __device__ constexpr auto rebind_hash_function(NewHash const& hash) const;
 
+  /**
+   * @brief Makes a copy of the current device reference using non-owned memory
+   *
+   * This function is intended to be used to create shared memory copies of small static sets,
+   * although global memory can be used as well.
+   *
+   * @note This function synchronizes the group `tile`.
+   * @note By-default the thread scope of the copy will be the same as the scope of the parent ref.
+   *
+   * @tparam CG The type of the cooperative thread group
+   * @tparam NewScope The thread scope of the newly created device ref
+   *
+   * @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
+   *
+   * @return Copy of the current device ref
+   */
+  template <typename CG, cuda::thread_scope NewScope = thread_scope>
+  [[nodiscard]] __device__ constexpr auto make_copy(
+    CG const& tile,
+    bucket_type* const memory_to_use,
+    cuda_thread_scope<NewScope> scope = {}) const noexcept;
+
+  /**
+   * @brief Initializes the set storage using the threads in the group `tile`.
+   *
+   * @note This function synchronizes the group `tile`.
+   *
+   * @tparam CG The type of the cooperative thread group
+   *
+   * @param tile The cooperative thread group used to initialize the set
+   */
+  template <typename CG>
+  __device__ constexpr void initialize(CG const& tile) noexcept;
+
  private:
   impl_type impl_;