Promote experimental::dynamic_map to cuco::dynamic_map (#794)

This PR promotes `cuco::experimental::dynamic_map` to
`cuco::dynamic_map` and removes the legacy implementation.

**Changes:**
- Added `erase`, `find`, `contains`, `retrieve_all` methods
- Added `size()`, `capacity()`, `load_factor()` getters
- Added sentinel and hash function getters (`empty_key_sentinel()`,
`erased_key_sentinel()`, `key_eq()`, `hash_function()`)
- Updated all tests and benchmarks to use the new API
- Added `insert_or_assign`

Closes #732

Author: PointKernel

README.md

@@ -15,6 +15,8 @@ Similar to how [Thrust](https://github.com/thrust/thrust) and [CUB](https://gith
 
 ### Major Updates
 
+__02/03/2026__ Modernized `dynamic_map`: promoted `cuco::experimental::dynamic_map` to `cuco::dynamic_map` and removed the legacy implementation
+
 __01/30/2026__ Removed legacy `static_multimap` implementation and promoted `cuco::experimental::static_multimap` to `cuco::static_multimap`
 
 __10/08/2025__ Changed `cuda_allocator` to stream-ordered, requiring `cuda::stream_ref` parameter in `allocate`/`deallocate`.
@@ -244,7 +246,7 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
 
 ### `dynamic_map`
 
-`cuco::dynamic_map` links together multiple `cuco::static_map`s to provide a hash table that can grow as key-value pairs are inserted. It currently only provides host-bulk APIs. See the Doxygen documentation in `dynamic_map.cuh` for more detailed information.
+`cuco::dynamic_map` links together multiple `cuco::static_map`s to provide a hash table that can grow as key-value pairs are inserted. It supports `insert`, `insert_or_assign`, `erase`, `find`, `contains`, and `retrieve_all` operations via host-bulk APIs with kernel-based implementations for optimal performance. See the Doxygen documentation in `dynamic_map.cuh` for more detailed information.
 
 #### Examples:
 - [Host-bulk APIs (TODO)]()

benchmarks/dynamic_map/contains_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2026, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -62,7 +62,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_contains(
   state.add_element_count(num_keys);
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
-    map.contains(keys.begin(), keys.end(), result.begin(), {}, {}, launch.get_stream());
+    map.contains(keys.begin(), keys.end(), result.begin(), {launch.get_stream()});
   });
 }
 

benchmarks/dynamic_map/erase_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2026, 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,15 +56,14 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_erase(
 
   state.exec(nvbench::exec_tag::sync | nvbench::exec_tag::timer,
              [&](nvbench::launch& launch, auto& timer) {
-               // dynamic map with erase support
                cuco::dynamic_map<Key, Value> map{static_cast<size_t>(initial_size),
                                                  cuco::empty_key<Key>{-1},
                                                  cuco::empty_value<Value>{-1},
                                                  cuco::erased_key<Key>{-2}};
-               map.insert(pairs.begin(), pairs.end(), {}, {}, launch.get_stream());
+               map.insert(pairs.begin(), pairs.end(), {launch.get_stream()});
 
                timer.start();
-               map.erase(keys.begin(), keys.end(), {}, {}, launch.get_stream());
+               map.erase(keys.begin(), keys.end(), {launch.get_stream()});
                timer.stop();
              });
 }

benchmarks/dynamic_map/find_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2026, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -62,7 +62,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_find(
   state.add_element_count(num_keys);
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
-    map.find(keys.begin(), keys.end(), result.begin(), {}, {}, launch.get_stream());
+    map.find(keys.begin(), keys.end(), result.begin(), {launch.get_stream()});
   });
 }
 

benchmarks/dynamic_map/insert_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2026, 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,15 +57,12 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_insert(
 
   state.exec(
     nvbench::exec_tag::sync | nvbench::exec_tag::timer, [&](nvbench::launch& launch, auto& timer) {
-      cuco::dynamic_map<Key, Value> map{static_cast<size_t>(initial_size),
-                                        cuco::empty_key<Key>{-1},
-                                        cuco::empty_value<Value>{-1},
-                                        {},
-                                        launch.get_stream()};
+      cuco::dynamic_map<Key, Value> map{
+        static_cast<size_t>(initial_size), cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
 
       timer.start();
       for (int64_t i = 0; i < num_keys; i += batch_size) {
-        map.insert(pairs.begin() + i, pairs.begin() + i + batch_size, {}, {}, launch.get_stream());
+        map.insert(pairs.begin() + i, pairs.begin() + i + batch_size, {launch.get_stream()});
       }
       timer.stop();
     });

benchmarks/dynamic_map/retrieve_all_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2025, NVIDIA CORPORATION.
+ * Copyright (c) 2025-2026, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -53,14 +53,14 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_retrieve_all(
   cuco::dynamic_map<Key, Value> map{
     static_cast<size_t>(initial_size), cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
   map.insert(pairs.begin(), pairs.end());
-  // Prepare output buffers
-  thrust::device_vector<Key> retrieved_keys(map.get_size());
-  thrust::device_vector<Value> retrieved_values(map.get_size());
 
-  state.add_element_count(map.get_size());
+  thrust::device_vector<Key> retrieved_keys(map.size());
+  thrust::device_vector<Value> retrieved_values(map.size());
+
+  state.add_element_count(map.size());
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
-    map.retrieve_all(retrieved_keys.begin(), retrieved_values.begin(), launch.get_stream());
+    map.retrieve_all(retrieved_keys.begin(), retrieved_values.begin(), {launch.get_stream()});
   });
 }