@@ -301,44 +301,6 @@ class open_addressing_ref_impl {
301301 */
302302 [[nodiscard]] __host__ __device__ constexpr iterator end () noexcept { return storage_ref_.end (); }
303303
304- /* *
305- * @brief Makes a copy of the current device reference using non-owned memory.
306- *
307- * This function is intended to be used to create shared memory copies of small static data
308- * structures, although global memory can be used as well.
309- *
310- * @tparam CG The type of the cooperative thread group
311- *
312- * @param g The cooperative thread group used to copy the data structure
313- * @param memory_to_use Array large enough to support `capacity` elements. Object does not take
314- * the ownership of the memory
315- */
316- template <typename CG>
317- __device__ void make_copy (CG const & g, bucket_type* const memory_to_use) const noexcept
318- {
319- auto const num_buckets = static_cast <size_type>(this ->extent ());
320- #if defined(CUCO_HAS_CUDA_BARRIER)
321- #pragma nv_diagnostic push
322- // Disables `barrier` initialization warning.
323- #pragma nv_diag_suppress static_var_with_dynamic_init
324- __shared__ cuda::barrier<cuda::thread_scope::thread_scope_block> barrier;
325- #pragma nv_diagnostic pop
326- if (g.thread_rank () == 0 ) { init (&barrier, g.size ()); }
327- g.sync ();
328-
329- cuda::memcpy_async (
330- g, memory_to_use, this ->storage_ref ().data (), sizeof (bucket_type) * num_buckets, barrier);
331-
332- barrier.arrive_and_wait ();
333- #else
334- bucket_type const * const buckets_ptr = this ->storage_ref ().data ();
335- for (size_type i = g.thread_rank (); i < num_buckets; i += g.size ()) {
336- memory_to_use[i] = buckets_ptr[i];
337- }
338- g.sync ();
339- #endif
340- }
341-
342304 /* *
343305 * @brief Makes a copy of the current device reference using non-owned memory.
344306 *
@@ -354,7 +316,7 @@ class open_addressing_ref_impl {
354316 template <typename CG>
355317 __device__ void make_copy (CG const & g, value_type* const memory_to_use) const noexcept
356318 {
357- auto const num_slots = static_cast <size_type>( this ->extent () );
319+ auto const num_slots = this ->capacity ( );
358320#if defined(CUCO_HAS_CUDA_BARRIER)
359321#pragma nv_diagnostic push
360322// Disables `barrier` initialization warning.
0 commit comments