Remove the legacy multimap

README.md

@@ -233,7 +233,7 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
 `cuco::static_multimap` is a fixed-size hash table that supports storing equivalent keys. It uses double hashing by default and supports switching to linear probing. See the Doxygen documentation in `static_multimap.cuh` for more detailed information.
 
 #### Examples:
-- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_multimap/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJylVgtv2zYQ_isHDUXtVJYfaFDEjQN4bYoZK5whTlsUcaHQFG0TkUmNpOx6hv_77ijJlpsM67AWiCHe-7vvjtwFVlgrtbJB_34XyCTod8MgZWqRs4UI-gHPExaEgdW54fTdPpsqOIN3OtsauVg6aPAm9Dq9bgv_vA5h_Hn0fjSEdze3f9zcDu9GN-OIDLzRR8mFsiKBXCXCgFsKGGaM408pCeGzMJQN9KIONEhhGpSyadB8671sdQ4rtgWlHeRWoBtpYS5TAeI7F5kDqYDrVZZKpriAjXRLH6r049OBr6UTPXMM9RlaZPg1r2sCc4fU6d_Suazfbm82m4j5tCNtFu20ULbtj6N31-PJdQtTP5h9UinCC0b8mUuDhc-2wDLMjLMZ5puyDWgDbGEEypymzDdGOqkWIVg9dxtmhPeTSOuMnOXuBLwqT6y_roDwMYXADScwmkwD-HU4GU1C7-fL6O63m0938GV4ezsc342uJ3Bzi80avx9Rq_DrAwzHX-H30fh9CAKhw1Die2aoCkxVEqwiKTCcCHGSxlwXadlMcDmXHCoawUKvhVFYFmTCrGRBOEwy8X5SuZKOOX_2pDgfqj1VU_WLVDzNEwGXPOe6bcmEx6s8dXLFsojny6tTNbc0uXXtRKzRVbwW3GkTkdITFemEYShtc50rgj-uTp7Xt9hPgeR6XuoMUxbBWEU_ZuQrtf5QKocElKqx1jJpTtUO60Q6E0aPYhu7bSaQcgOkhHt7FK1ZmotCWIlIeLAQq8xtY_q0ggoRqWfuAFpd76VmX6gWBwflUpFUObbDUevBuqTft_IvNIQxqpx3XnY6nVKt3cZdgJom5w5bClU_irnrdjoh6oJNNUqLGlpd6rz_sb7TPhOqoe2zKb1WOdkIxtoVTOM0dxKVkfF8qVEFHpXekNcNTXqaIihWGIc5-rjo1IalQ6InzkWqWQJzRmSgeT_vvIh8ucgprPKUVJcVsGENuSvsW-bbBQjHGfTC0vgA_u5pG_anSt7b7rkW7PclsAWZ-v0T-l4WTjImzfO5XQHJbGPcrPXHCIYAMqhoS3V7tQiuBS4yggmW1I2NxuIcrjYblcbX0SIKYbdDPLGGXTfs4k8URYAnvXMiQnEM_qMU7usFHMahUWBWna_Yo4ifDNwlsvqq0WmGP69clEIEbTQru-JsJhY4YdXZ_TeI4xLOuMFyWrhN2OF2drlR8G_Q7iS8gMYY2tBrhiD3b2FfA3nkeYcTgBz0wWk8tWct0oWUaEkV7GycZlcmK1SC-R89TurdMnomPJsJdyC4iy68vnh5cXGx_2fCVKVceevY6XguMZDnx9GmYkajrnRM8OTU5xlCp84wakwxzJrz3BjyZSlxnzJea_eYdbE3mrRZ6NLFmwoFHqBqh5fu7vDsIdZ4mZkHsPl8Lr-jaoK3phMUhbkfQlEkhg8C1fLspfXQ9c58k_0mA3SX5S4mmuAO8zcGZV2E-em6m_9zOpMYr1Is1zZq-dSQ_ICRChI9bfbe702LgRAFZq3mktF9f9icdeI9HCI9VAwljGjSH-pl3ctvD5BoYdVLhzc9PiPCmi1KIysQwAQGg2evjIcDzGVBgnQ9vjhXRoq1-I8Qh0eQKrWCrD7KwWnZyFrU1lPDI7DEKacdS0Hlq5nwu9_nVS08sEudp_hEQ-6N4RX4QS-qs35o8b3ZOCHR4MdkBnAwPBClXC50XeKU4kuanqb4WDXHB3eg1px3e-d5F8U6c8VrPGjhvTvgr15130CLGb4c2FX8pgOtFl7KDv84rFkkrZStZv6JnspZzSfnPMXDdfGexgO8odVjsA8rOVL1RI5MDvbf_P-_AYKkJA4=))
+- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_multimap/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyVVg1v2zYQ_Ss3DUXlVJY_sKyI4wTwUg8zVjhDnLYo6kKgKdomIpEqSdn1DP_3HUnJkRcD61oghu6Ox3fvHo_cB5ppzaXQweDLPuBpMOhFQUbEqiQrFgwCWqYkiAItS0Xtd-diLuAC7mSxU3y1NhDSFvS7_V4b__wawfTj5N1kBHf3D3_dP4weJ_fT2C5wi95zyoRmKZQiZQrMmsGoIBR_Kk8EH5myaKAfdyG0AfOg8s2D1rXLspMl5GQHQhooNcM0XMOSZwzYd8oKA1wAlXmRcSIogy03a7dVlcfBgc9VErkwBOMJrijwa9mMBGKO0O2_tTHFoNPZbrcxcbBjqVadzAfrzvvJ3Xg6G7cR-nHZB5EhvaDYt5IrLHyxA1IgMkoWiDcjW5AKyEox9BlpkW8VN1ysItByabZEMZcn5doovijNCXk1Tqy_GYD0EYHEjWYwmc0D-G00m8wil-fT5PGP-w-P8Gn08DCaPk7GM7h_wGZN301sq_DrdxhNP8Ofk-m7CBhSh1ux74WyVSBUbmllqedwxtgJjKX0sHTBKF9yCrWMYCU3TAksCwqmcu4FhyBTlyfjOTfEONuL4txWnbmYi5-5oFmZMhjSksqOtktokpeZ4TkpYlqub0_DzFqV2nSoLIWJrfOFK2Ub3CXZMGqkOh_CDVMEvT4NlpDUlvPxGlvNUHfnvUYRoZGnPP43WC6xgYzkzsyFQXVyEW4kT1tzsUcSUOuWwCe2S8yuYKjHG9SLuX52bUhWMu-sXdZ5XMHywuwS-6mZLYVlTtY30O65LI31PtQbjsFVoA2l2CtjdQHapIOB5n_jQphiyGX3dbfbrcI6HRwUGKlKarDfUDfLH8petxthLOhMotfX0O5ZWbgf7WTgkNgaOg5NlbXGpGOYSuNlSO2h5BiMx4GuJYbAk5Bbm3Vrx0CWISmaKYMY3b6YVEdVQqtdPDSZJCksiZWDHQaX3VexKxcFh1WeKm5YExs1mLvFvhWuXYB0XEA_qhYfyd-_bMPhNMhl259rweFQEevlNBicCHjokxSEq_PYbsH6dDhtNfqDokMCCdTCtXW7sBjGDKecpQnWthtbicUZnHs6rhaP41UcwX6PfGIN-17Uw584jgEt_UsrBG8G91E5D80Cjgci9JzV9pw8seTFkRuiqm_Dbiv68WBfihVo2KrXeduCrfCE1bYvXyFJKjqTkJR2Grdgj6PblErAf1G75_AKwil0oN-KgB-u4dAgeeJ0hycANeg2t8dTOtWiXGyQnWBeneEpugosEynif844a3ZLyQVzara8g6Xbd6H_y-urq6tD3ay6mXotyyzFnm4Y9OuW2rungee8wOrSb91uiZHJkiMwX_Z1c12tprAZ-FzUidXVFkG3qco1o0-YjBhfF96DdubYuxovuB-BupAyu31e08Tobgvv0P8D3zFXHdbAi88Wvty5_h7xLlGQ9oIDJyY3MUGUeeLsOClr3E644YvszQ0rBDhHWb0pX0LYyHYDvkCoZo8by5jZwHCIT4FZSSle4T_BCBHikwCNLt578dbl2C4P3SesKa5v1znuGlyfzT1uDok-pKV73-BMccdDR3624vPGkMzKtd7eb-3l_Zz-4Murjp29SNCCD1D7osM3nnp-pwZiQ2mvf1n20C0L4x-xQRvh3dA3b3pvoU0UXd_oPHnbhXYbryuDfwySydJ2RvKFe9lmfNHISSnN0Ljxz1A0IOfiKThEtR_nwIkfVRccvrr__wANsNIK))
 
 ### `static_multiset`
 

benchmarks/static_multimap/count_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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,7 +57,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_count(
 
   state.add_element_count(num_keys);
 
-  cuco::experimental::static_multimap<Key, Value> map{
+  cuco::static_multimap<Key, Value> map{
     size, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
   map.insert(pairs.begin(), pairs.end());
 

benchmarks/static_multimap/insert_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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,15 @@ 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::experimental::static_multimap<Key, Value> map{size,
-                                                                   cuco::empty_key<Key>{-1},
-                                                                   cuco::empty_value<Value>{-1},
-                                                                   {},
-                                                                   {},
-                                                                   {},
-                                                                   {},
-                                                                   {},
-                                                                   {launch.get_stream()}};
+               cuco::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()});

benchmarks/static_multimap/query_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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,11 +62,13 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_query(
   map.insert(pairs.begin(), pairs.end());
 
   auto const output_size = map.count(keys.begin(), keys.end());
-  thrust::device_vector<pair_type> output(output_size);
+  thrust::device_vector<Key> output_probe(output_size);
+  thrust::device_vector<pair_type> output_match(output_size);
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
-    auto const count = map.count(keys.begin(), keys.end(), launch.get_stream());
-    map.retrieve(keys.begin(), keys.end(), output.begin(), launch.get_stream());
+    auto const count = map.count(keys.begin(), keys.end(), {launch.get_stream()});
+    map.retrieve(
+      keys.begin(), keys.end(), output_probe.begin(), output_match.begin(), {launch.get_stream()});
   });
 }
 

benchmarks/static_multimap/retrieve_bench.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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,10 +62,12 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_retrieve(
   map.insert(pairs.begin(), pairs.end());
 
   auto const output_size = map.count(keys.begin(), keys.end());
-  thrust::device_vector<pair_type> output(output_size);
+  thrust::device_vector<Key> output_probe(output_size);
+  thrust::device_vector<pair_type> output_match(output_size);
 
   state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
-    map.retrieve(keys.begin(), keys.end(), output.begin(), launch.get_stream());
+    map.retrieve(
+      keys.begin(), keys.end(), output_probe.begin(), output_match.begin(), {launch.get_stream()});
   });
 }
 

examples/static_multimap/host_bulk_example.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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.
@@ -16,12 +16,13 @@
 
 #include <cuco/static_multimap.cuh>
 
+#include <thrust/count.h>
 #include <thrust/device_vector.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/sequence.h>
 #include <thrust/transform.h>
 
-#include <limits>
+#include <iostream>
 
 int main(void)
 {
@@ -52,23 +53,22 @@ int main(void)
   // Inserts all pairs into the map
   map.insert(pairs.begin(), pairs.end());
 
-  // Sequence of probe keys {0, 1, 2, ... 49'999}
-  thrust::device_vector<key_type> keys_to_find(N);
+  // Sequence of probe keys {0, 1, 2, ... 24'999}
+  // Each key should have 2 matches in the map
+  thrust::device_vector<key_type> keys_to_find(N / 2);
   thrust::sequence(keys_to_find.begin(), keys_to_find.end(), 0);
 
-  // Counts the occurrences of keys in [0, 50'000) contained in the multimap.
-  // The `_outer` suffix indicates that the occurrence of a non-match is 1.
-  auto const output_size = map.count_outer(keys_to_find.begin(), keys_to_find.end());
+  // Check that keys are contained in the map
+  thrust::device_vector<bool> contained(N / 2);
+  map.contains(keys_to_find.begin(), keys_to_find.end(), contained.begin());
 
-  thrust::device_vector<cuco::pair<key_type, value_type>> d_results(output_size);
+  // Verify all keys are found
+  auto const num_found = thrust::count(contained.begin(), contained.end(), true);
 
-  // Finds all keys {0, 1, 2, ...} and stores associated key/value pairs into `d_results`
-  // If a key `keys_to_find[i]` doesn't exist, `d_results[i].second == empty_value_sentinel`
-  auto output_end = map.retrieve_outer(keys_to_find.begin(), keys_to_find.end(), d_results.begin());
-  auto retrieve_size = output_end - d_results.begin();
-
-  // The total number of outer matches should be `N + N / 2`
-  assert(not(output_size == retrieve_size == N + N / 2));
+  if (num_found == N / 2) {
+    std::cout << "Success! All " << N / 2 << " unique keys found in the multimap.\n";
+    std::cout << "Each key has 2 duplicate values, for a total of " << N << " pairs.\n";
+  }
 
   return 0;
 }