Add preallocated memory example (#745)

This PR adds an example demonstrating how to use cuco hash tables with
preallocated memory. It also fixes several bugs identified during the
process.

Author: PointKernel

README.md

@@ -223,6 +223,7 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
 - [Device-ref APIs for individual operations](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/device_ref_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzNWWtvG8cV_SsTFmkpa8WHUTcAJQpRbaclksqBpCQILGM93B2SA-8ru7OkGEL_vefemX1RdEylXyoDJrkzc5_nvmZ3vUIVhU6Tojd5v-vpsDcZe71IJstSLlVv0gvKUPa8XpGWeUC_hy_uE_FCvE6zba6XKyP6wYl4OXo5OsN_f_fE9c-zN7Mr8frdzY_vbq7uZu-uB3SAD_2gA5UUKhRlEqpcmJUSV5kM8OFWPPGzykka8XIwEn3acN9za_e9k3Omsk1LEcutSFIjykKBjC7EQkdKqIdAZUboRARpnEVaJoESG21WzMrRYXHEr45IOjcS-yVOZPi1aO8U0tSi09_KmGwyHG42m4FksQdpvhxGdnMx_GH2-u317dsziF4f-ymJYF6Rq99KnUPx-VbIDJIFcg55I7kRaS7kMldYMylJvsm10cnSE0W6MBuZK6YT6sLkel6ajvEqOaF_ewPMJxMY7upWzG7ve-KfV7ezW4_p_DK7-_e7n-7EL1c3N1fXd7O3t-LdDZx1_WZGrsKv78TV9a_i-9n1G08omA6s1EOWkxYQVZNZVWhteKtUR4xFasUqMhXohQ5EBSOxTNcqT6CWyFQeaws4CBkynUjH2kjDz54ox6yG98l98hedBFEZKnERlEE6LOhI4McyGwTl6nJ_RyixIxwuyiQgyjK6PLhuyixSnSWzysvCDEO1hgD-WgUmzQerQ1vUgwpKIu5nKVy6PbxLG5VL0Bj-rjO_-nF4awGYKGB28ESdWJruiQDCh2rReaZTYEDJuPOQjVswveELC8pvOVicgrla-OpBwq8KdrTr81yrhXijYnjEQF4FvxTkRxcf9qgoNBhc_Tgr2PM6CfVah6WMRJqRluzPSH8CPuHH3AwX2NIkg9nn90uQTpZRdZCAR2dFAFjPFWEIDGMgHfHiZAlSyLLWkuRj8ve9BiCkJOLAbDNVNAJ8iwSixB1lD2cACiRSJaw1p1SUlXmWIsARV9F2IGaGtlHy0YlRgCtHbrFKN2ybGE5g8k5KmRixQaLBnjDlHRUzGS1TBPsqrkUadrxkvTBjExTik9oWOC6NyGRRtOMMAiA-Q-QUwyYxlg0FRqOryWQuY_EfmYk7mKHyJDZVFoSNVE7o6xz4Xm1nwCzEyEpo7OArNitYRKxlVCqfzAr7JwhxoymtQQDwmUwgMi92CP5MZ44iCSsfogqRMxU-JfxjZQMkXZlvG5s0RuCN73XyQThQiJtK6bZB2ISblQ5WDoAFikgUdaC3T5B0naslAHkHKvyN850jmyMTMg_rxtQR3qditT-GDu9ECi2KNNDseK5zCmWJeHyeRVLGPgtB9E1qEH54NIc_KumQlh35_bNk0n3r2npmg6DiPhDvECqHINsCqjMp87CiVnXFcUxLY8WtVj8vclEGASrUoozAttruIsooRBsJekGQSWSsCEeeqH85iLee1BhtPavhhVTq-8soncvI98U61SE1H9gMUFrWfYoyhzAqu1_4q0KshtARZ5ooakHmiHOoGpNJoX9HiNVIOOJYE1vkwCMOIIRedHx3cp_s6JgsKT3BZlPgAcUqnIUPgwdxKuZRGnyyP17YH290PHg4p5TYlTtIy4Q0n4oRNYQUqUgPfSJ6USt1InZWyuHQgtGVkj_CpD2gF6JPj_q1P96D9IeTmiLT5FTkaObKlDlq10eTl-ojEdBcI4j4QueFgcKxbZWWeq0SDtCNLFrkaohT8H1cyKiwhOgM7ZYR2WqL2oE-j6EtKlkd0AYOe9QbTSaZ1Pmuq4DXhgo_eWzrJMTpqbPsOYv0mn60Au1gjLE9KxKP9ov7IH-cTsUy1yF7ct-tvJE-wOunLCRw2aiyDiaWT7igUMQ6uArDfhtbXgUJGhAeD9TQ79K85fiVXFOLgZaKE3yV9j18D3IVq8Sw63QOsjka3yxFp4IUzOajDj5N1P9RZW2lzKPLGtpMGOaPy9qzSxoR3bYJfLHUPDNDH8i8tct8W7E6qfdAZj2YAP-H3HRE9pk1UczRW3Qwx31tlUDQ3tY4cF1c0-PZukoD2UbTzGdqDtVx6vjQkPZP7Aqrs0BkkEJVliB--5ntvEl8dvtXzX6m1019TxTyXFxKmxmcR3hTqz9hB1VkWDZC6rTJPzSTTSaWFPHuoKHT93lur_WSX2B2V74Ntssdq3B2WSgETvh4XtFn5RVC3pdIH-OKRIxWP9_6aY6xE0wj-YBadf78RMZpBxEGw-mkT9isUFXSLENQBKkpBeF585TLeP3UJcO3cWa2oohSQAXzPySn8RsWtfcH9D1f4zsmHDzUiYow21j4DyjUqvRR0FBSRiEFeKEiDLI4hCy-cnw4GcJJW5GoNQVlEJQ0xtEFC31SWkJalpTrWP6AhiK6DBCKRCSk-5UIbOOpOBuzdlav_f02qdUn3G4nzXU7MwxtrqUqxqllrtpd3V43UDGp081UvBr9bTQaNcRfAyeG50o3YR9odsVuN_JGj57Yjb0xPgaDgdhpTz8yDOyEPpl0rgUuYJVLJxmz7tcJ4Lx9qOLab20dcAD2T7z2eRtunhh1z3eZsnFrti7vHcfYoWSftXvcZl4ZLo0BA7iSLrw0YDmXhbthgpG_FlEq0f5KkqvOoI1HaNW3q9StDV6dH_ReQ31qFwOlo1ojMWzT6chGM3oZUJhwgeMJ6L5XUUNU2CCysXY2ZmfTh23NGJQttFXA5CsCqwpRnQrbUjXXCbuKwxcbYXuyjpbd07h5fBYNFnR3KJqOoUP5ju0LVKCEmpTxuztWhAh8ZO5neTqHPS84hdLzUC1kGRl_JYuVX923MWnQfmz89S9l4KgkTc7STcJXnHFp-B40P9SucHpCSUk3RT2GwyXctiMNWp9Vtx2OxSeVUzqs3OfgzXWGqyO-uQbZLrXQ1Iw43BSlrZlBuZ69aAgXPj-civcfhO9Kj-_3eRE7UTBdVRZUbMXX4uWJmNLAIloGuYJyzLEAsOl6je6e6ua3aZWetMGH8wJscdmZufrjSr-9uOTy5VMIQoOXr_5x3t3Dpc4uu7TeBONp-_CZGJ8gPJsnTGhvGr64uKjpea29l5eX_cZBR8yUnb9DmfTPkdjLic8k8owhusvcIujPsKsvS6g4908GS2XQorUHZWDIiIsL5MLr7j1JdXSCJdrQJvd-9IEe2QSRhNHnA_fL8cpdbetW95lxS8cPRi0tVJrut_5_hLOKoHcQN6JdPJ1EN8pgQlpzE2N4JISq0I7VJVOoVpfWNPRfaBWIGPpDxEbdLbTMf0TRbwi06v4TEtZoVgEfQve7fBvF98lVK528wa9I-M1F1a1XIsbyk_LbrzX6rSLDx3bP5fxoNaAxg2kCYtSbkuXJ1o07uHHDFLvYWuR9pOUBJdtTMaZcy7-Z-Ee6jJHRRtKcmpeqqe_opVG57LzHGoGPny76XTWtJ7yWIdrfwW_PCd6hqsCZ9a_2XGueclWiZTeE88Xosu82OmX21sfVejWp1EGBGc4qRSVoLxnc2hucr8SdzEGlvkPO6Q6MorU9vtFlLGj2zqtrGifpyN6x9LwevWPF7Js3b457yToIxi9flWMsp5mxr5V7Z5BjGpyejr8RZzIPVtMi9r8ZibMzFCOD_-xblLNIxnN-1xzpeYtmEAQRHq7ti2E8gAuTT71Hr1pHZuisw-69xw_877_T6GKv))
 - [Custom data types, key equality operators and hash functions](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/custom_type_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJytVwtv2zYQ_is3DdvkVpbtbEMLxw7gJS1mrHCGOG1R1IVCU7RNRCI1kbKTGf7vO5KSLT_apMAcIIl55PG7u-8eXHuKKcWlUF7389rjsdftBF5CxLwgc-Z1PVrExAs8JYucmu-tFxMBL-BSZo85ny80-LQBZ-2zThN__RbA6MPwajiAy-ubv69vBrfD61FoDthD7zhlQrEYChGzHPSCwSAjFP-UkgA-sNyggbOwDb7ZMPFK2cRrnFstj7KAlDyCkBoKxVANVzDjCQP2QFmmgQugMs0STgRlsOJ6Ya8q9Vg48KlUIqea4H6CJzL8NqvvBKK30M1noXXWbbVWq1VILOxQ5vNW4jar1rvh5ZvR-E0ToW-PvRcJuhdy9k_BczR8-ggkQ2SUTBFvQlYgcyDznKFMS4N8lXPNxTwAJWd6RXJm9cRc6ZxPC73nvAon2l_fgO4jAh03GMNwPPHgj8F4OA6sno_D2z-v39_Cx8HNzWB0O3wzhusbDNboamhChd_ewmD0Cf4ajq4CYOg6vIo9ZLmxAqFy41YWOx-OGduDMZMOlsoY5TNOoaIRzOWS5QLNgozlKXeEQ5Cx1ZPwlGui7dqRcfaq1kRMxI9c0KSIGfRoQWVLmSM0SkkW0mJxcbgjJq1ZIajRSpKLukwv8kLpVsyWeEO0ZFTLPFyc2sI1ywlKW1QWwkQlqlae2K9zIhS6I33iQCLnSIXktHCrJHTWtVrwXrG8GbMZFxjme_YI-jFDgmDoC6qB4imZRrgemXVYo-uQUvrXs0gDOa9_m54bjQBRtJBKRxH-U_ojOlTjN2C9ee7eSv9DA7pA1g-bAKb422rYnJ-yYkmSgp20w0pOWDLbs0Q9bclO0XNsqe3es2ZmrVHftsYpgwVRC8DAJibND-2ywrVDsb28qO6SmSOM3_AP40kxQfTPcN9w_2EBLAveGguMLnIB9yE5B4Psm-gMcbAgkeSrEM2VbschzqmUyTMwJgsVHPGxlOULdWiBucTeBJUlqCAk0O-b3aHj7tbp6Cqs3Vz4S8njxkTYk1ahqVSgdNztKv4vBhFEkUYZ4bmCPrxu_9Jut0u6oGvGTANLNZZ9xUx2s0QZCSmwEDt4LM30o8Vf7UCA_UO71s3O5vzkSUel7dn-McXKsyWiy5wRzUzxrsqG6UhcZIU2QWu5XLH2bO-z36Ipm2Pv6IMrHd1uSu7xgqMq5Dsf7-06qm29kooXfrsRuAOmnHa7WS5pQngauRhZA2wtRhGi6B34JTg29-KiRADw-UuNVr41hTd2TN6pXR-6m29OaMbVDVK_0ah5U1asJkiXzA0CnXY7QBaASqTGbqNMTzLdZs6XTLiw1Vy9JUYII6lttyuVUzMBcNyMvZdiLcHD90KujLqVmTmSBOOGuaeRdvZCVIopYVokhjeRJIYZoWWEX7d_CrcBNVD7pf27Hrd2C-xBI6RejeKBMckQ-2K9CSrnfu1TKqlovT4m-PfpsF5anyL7c_Tsl5pnw09QP8kjZOMU3d3rBPXCil6oJdTQhkBVMXCVAOktbcyNp3Eg25HA6bHVm5k4xdWKBWiivSuYYI6HLsZ-LQmDvYx8uatANWbeMAQfF1iK3XnXzHc57VaNa7BunUjZE4n93Sn91OebKb-XkLucfurzZM4fJno9oysL9-a2nulHF6bomjGexf4pd18xbaZOge6eAQbQxnrP8xhXAVMzweL8uc8JRVJWZ0RFhVJ3NWMCUXDndN6FMMR7bJu9q4XyM_9yB7FkSvyCDeIBB_YA7rbIUWra3Ywkit2FFb9KsfJreoI9ftQYFuz8EFr6-dZ5FueteSTZsjQjPIHYzFyGaKZdIRLMWje0GBO380FYHv6_kJTqBtviWLkfnzm7EwetFAMWbWWH6WCEcuYf3VYHwERsFk6QrxxIpjUOTg3l3IA5A3_vbrPLFl5MLw29Hr6wxgWl-DL6Ad5a4hhu2TKowgla652bmcXoKpWb4QNX8Dlt3qf4Ys13r25PLCntnP1edFAsM-2e5F4Tb-zTly87r6BJcrroqzR61YZmE0cd3bTtIGZxMyHp1L7TEz6t6aSUJri4dI9qXMB-KO69TVDJsaDuydE_3uaL_fkPoxeK0w==))
 - [Key histogram](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/count_by_key_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyVWQtPI0cS_it9PkVnL34AuU1OBqMlsLlYyUEEbKIIVpN2T9tuMQ_fTA_GQfz3-6p6eh7GLHusxNr9qKqux1dfN0-dXOe5SZO8M7596piwMz7odyKZLAq50J1xRxWh7PQ7eVpkir6P3t0l4p04S1ebzCyWVnRVTxzuH347wK_3fXHx2_R8eirOLq9-vbw6vZleXgxpA2_6xSid5DoURRLqTNilFqcrqfBfOdMXv-mMrBGHw33RpQV3nXLurtM7YimbtBCx3IgktaLINcSYXMxNpIV-VHplhUmESuNVZGSitFgbu2RVpRw2R_xRCklnVmK9xI4Vvs2bK4W0len0s7R2NR6N1uv1ULLZwzRbjCK3OB_9Mj37eHH9cQDTq22fkgjuFZn-b2EyHHy2EXIFy5Scwd5IrkWaCbnINOZsSpavM2NNsuiLPJ3btcw0ywlNbjMzK2zLed5OnL-5AO6TCRx3ei2m13cd8cPp9fS6z3J-n978dPnpRvx-enV1enEz_XgtLq8QrIvzKYUK334Upxd_iJ-nF-d9oeE6qNKPq4xOAVMNuVWHzofXWrfMmKfOrHyllZkbJXwaiUX6oLMExxIrncXGJRyMDFlOZGJjpeWxF4djVaO75C75u0lUVIRaHKtCpaOctqgglquhKpYn2ytmo1mUqnv3O4DvC6XLhc1loYSgcCRtGhu1e25eJIqMk1Fr3i6zIrejUD_A0OBBK5tmw-WuJTaTSQ7nxMMXZsbSblkEjaGevxgziW2NmRTx1jJuDbIjc9YxeucS8AMXhkqLxAazTXCvN4F-lIgiOcOtmGVGz8W5juF_mGo1opBT1MpqcCcUuYGK01-nOcfZJKF5MGEhI5EipmX0InOPbETUMjuaYwk-sw4So9LE6kdLUqUzaDDbDGBQLaAvzFAPWb4UqGqbLjIZC0oegYV5jSTT19VLWJosIm8HZS2bolATM00JSKHQZJs_mkpxtAcjyU4Wf9epswu5M0cR2c1KNwz4APTR4oagp_QnVSFZHlaOJBxbFdkqBTqgKKPNUEwtLSPkQkA1cp3LPl-ma_ZRjKiy-NJKmVixBkphTZjyCq9MRosUSLGMK5NGrbC7oE7ZBTk7j8rN-T0XpA_wohOSiQAoVVBIWYNJVoUVORBLAz4bJ7YrSdH4gQrq2vylxdmn81PB9QWX_6Vbq_4jV-IGLvNJBEd6b8OfOiPZrQ0_683UIs5TVm8sBRXeXC_hPfEgo0IHFAJKI6SDNYSfcAv0jMeU1TTZEvgpMTgCy7ws7JeFIiZbcgvE57t_BnUD-MBSb03yWZRJIa78QZqHxL4UCoxalgmYowNFUSv1tgWS_TO9QABuIIU_MViWYgEfrho5jNwmSPC2lKSIA15BQmxqURsYmuH8fitlAJ86b-5NC-s2F-yx12VQjwHWWCfMGeEOg9yzGnkJ8AA0Oc_VidLn4klkrClcjW9lzBsjddCAYkGwiNKZjIJAPKQmbKFYlxKsDAS1ti_9-NQis99a6-NeefOtDY002_Jh7y55os0F4ZHzxhW3ITER6E_jcWPo2Ovt1247IbojgiBfggWEcELlpcbG8fgGjr9GVhNiUxCC3H05IjwQwp9HESiJKE1pQUZYPhEL_H9u4uEjkqFSe9TcZcLHrfNOGgDwzlX_NHyEiD3kKjpSyN9K5ZU33VQrwyZin1WhVlByXdJ0XHm9J56c30cjj-PUKbi4fKUxsStbCaXngdshC6y5zaMUvjR5kOg1SfwMfWW6DJ3AANUQUGfoEp0Yj1fSZE-k-xamfO6Lg2fmm_gxc5hXSapMY-PmJoNbGUBbUICFftXLs-_tlZKfhY6ARE2BfEyUakSbNh6U6byDE3xTmY51Yt2p_TY-MiESJVYox2NHZ-iwnFjfHlJiuanSmhx8VwcOkU-eyFuDk1wjScLnIy-WfDXXVi0DGYbdAy8hRoPLNkGagalBRSQfdehd9Vx6DMHcmzSzjeefXVbglMTPwVUb4JIXSoFkzoso2lTYIuBzhzbUWTU4t8u4UgrhmTuqpF2wUpT5Ad6QhA7AWnno2GA7DRvV1G0WUG94XcTdl9FzZ6WkaCS8mCCfq9TYGQVX3q9HgcuxqajrA_HOEcnxOJPrYJUSdcgCJXPb3YYcH4dtUY04vnDB23FF3Dh0UIxMNEmX4NjDG-JwQVzILqVlrxAXSOddoG5vr_x8Rtna6538i4JTZrRJTFzE9dWI2YTM1BJtWtki466cx8H3-zWIQiZDkE_qo3qKVbSnSus-xiu7EZTiKKuM-AddZ1BE7j5Gn7MHfAbpw6BJdAS659rkkPqgpwo58bQiCqmV5zqCkdiEtF2WetgBzKcSTYy1YlW4sGYlrQqllUSp-CQMyXS5EppMZHLuTSixtqShFOtjbDnpDg5cSNxxtyU4SlPJaO_nLV6Cj1ybIKB6ZrrZ2R201GoqjjER7_f_sb_vIByCfqoYJVdpkyo4R-Qeyiq6TA7ZoYK8EBZ8R6arCBSxjvIkMifbut3KkG-2NvS4End07bvOj3w_rRl7vz5OjPKio9OtImcOiMzYsgT2c22Ql1awAxJ3aaEzorqQCjq3Q_dqQet8DbcujBzU0uGu5qsO6D17BqSwfK0pKbmP1lB8fOnrpcy37cYtYoNfqH93gaksqe6lJdD48Vje64AbDAor8KzZ5d9-r__1ixvnGhISdHt-d3OG-W7Xz9x-BuUpfRR0OTcMcBVlCkxIXhaEQQbU2TDazobekXhu5PtpBPQjd5YYz32lYhB1N3LwyG5-PXYe1E-2mV-3WWJnZatT9GZkgEQzmZePNKihb9AkZSjmkiTuqAaaDdwsMabh-1rw1DGN2mhgEnGDnIDzPknXCaWsDB_oGawv1pqvwPRiRskJJTJSBVN2zzC8ia7ewvGYglZxRzpkpWIiXvV58-h0ES4UAW9N2XYYTM061Dn3gYY_hpVDaDORGyJq9dX8qWnSW0S96co31zpNFSw_vQTo5_9LBgPz0y6Q_ho51J05HkAAXIdsyqn_9LUmRNAjs2CVpTPU6DETORoP9VwWkQ2AGcvAv3KxaMh-rsP4b20RvyRNBsgpfpuMC8sPmNmuGzD3QRCydJ2LP7e49p_1c42H0VIJ8JYdU3P8e51RL26kQOA4Lj354VPJ27c0-MrbqiPHeiifIeDw_XdH7TWOMPE0_0wamLLX3DwQBz1kfD3Cgpo30-N67sTF5_j4uBLfb2w9OTnp-ivsTkSsL6HbCDMkItHtNWDmSoNng3e4dzYqKnpoiyKHEYidbtCg-lbxRmcCOaIMqTpTw4Tb_c-uR-3e7tiGF-DenV4V4SLqDhDA6G5Db-2Nliw_3ETapVb3REHbsOi2lY8W9NCDm5raYh6VpwPFQibbDoepRCy-0GkawMxNULt3SwpBaUH5-MavMDoz802DMhIx3W2Z2-XtajdfCA9Ar3c6ZttfGjddDLb7a4NGOkWN223ZcV13xNlfoDwtqporMf8tJ9YPjm6A2jjDGAYtqgJk6drd-v52ByGdI39BLDUTucRIp9-h66KJdFb_kaqTPCh1cPi-OMB0urLuL1idAeRP1N7ewfdiQHeJCV8gxGAAELD45d5cB5GMZ_xnrcjMGjKVUhEGH9zfoDCA5pXcd577fh6A05qHHzvPn_nf_wAAaV-z))
+- [Pre-allocated memory](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/preallocated_memory_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJylWAtv4zYS_itzOuDOzvoRe9vrQbGNpskuarRNiji7xeJ8EGiKtolIokpScVwj__1mSEmWnTi923OAJBbn8c3HmeFQu8AIY6TKTBD-axfIOAgHnSBh2apgKxGEAS9iFnQCowrN6Xv_bJ7BGVypfKvlam2hxdswPB9-24Gbz9Pr6SVc3d79ent3eT-9vemRrJP_WXKRGRFDkcVCg10LuMwZxz_lSgc-C01AYNg7hxYJzINybR60L5yVrSogZVvIlIXCCDQjDSxlIkA8cZFbkBlwleaJZBkXsJF27VyVdhwc-FIaUQvLUJ6hRo7flk1JYLaGTp-1tXnY7282mx5zsHtKr_qJFzb9n6dXH25mH7oIvVb7lCXILGjxeyE1Br7YAssRGWcLxJuwDSgNbKUFrllFyDdaWpmtOmDU0m6YFs5OLI3VclHYA_IqnBh_UwDpYxkSdzmD6WwewA-Xs-ms4-z8Nr3_8fbTPfx2eXd3eXM__TCD2zvcrJvrKW0VfvsIlzdf4KfpzXUHBFKHrsRTrikKhCqJVhF7DmdCHMBYKg_L5ILLpeRQZRCs1KPQGYYFudCp9LmGIGNnJ5GptMy6Zy-Cc67682ye_VVmPCliASNecNU3pMKjlOU9XqwnxxIxQ4m4z7Rm28mrS86tOVZUCiGi5UcRrbQqctN7YRuVpTqwKRWyL1jqJPtnfu-_dzmJ1LEkUZzhzkSpSJXeRuKJIY8CcXvBhZZiCde4mKEdlEQeDPGG-bgP02cy2uvWBiEWj0gTeLv7Qrunkii9gFmrjQH8RRlGBcNeGD0wg-wzSzlVu3E28RlVFcVkZSp63gnHTFsIMrssEpcBnGExwAYTR7hS3bDMkmf2qCQC3mYsxdTwrrxh74aqHtVREYtA6QfKFgfOi3a5J0dmGLbIHqVWWSoya_ZRf48dQRzGTp0BjcY1teQlL3SuCKTKkm0Ppi5Y6iYyswLzz5WicYRhHqa4t8485gXaSimcDXYOlImVk6icsWSlsHrX6T5n-334gRmM125zYeZZYSiszywpBIyhQIfvh5G9qBZ-ElvsM2MCclGqzzBGjDmBR1Ly4Yg0t1swibJoknQcN1Smfil6ENvIVIrgTHYHaNE7PpZ2lvfyY6ASCUNM8zDMilRoTBVfLCNnYBKGqcxa7QriL5hGaHMpV4VneJ45XSP_EJFtuOPUNiVCH8Pg_Pzv5-fnlYl7pAf5ky53XBYlgmXYCbiKRcVOrtUC_0YGW2-KlkmHwHIVhglCZzoqRUaDTvk8FktWJDZaM7OOlkXGCd8IOZtMataNVRqrzRssP5XdRcEfhI1KkZF_mDOpyUbH7-SkA4OX1rRY1hbHbv8x88WBszCshGptLMkDzSaWfdWSiAOAifa_fhzkr1H0WWHX2NBi3APsk5HvG19vzKUJno0siaxyu_I1tl7Ji68xc7xxkyo5fz3dcudZVJIQRcAK7An7bMeykjH2euwpFsYEyG9jyh5E1FwcvRZAvec-8QaTScs_KXV8cXXqmprsnn1BNgA1SHbn4OhFFlaxhqFvAt51E1zPLbTaEwqgUnPWLg58YcvCwsXjO-L4GwOG8wsAKm3N-ANgH1lgOeOBJq1IjReFFpU6jgY0LbjO0S4p_1Ek2G6hKlhqthzzDts7nVmIGk8XGu6WWqVHR2KJ8QDcQVV5Q1RGLXS3a8ZVyzQe7F5jo3NIRS9mlrXazzSgIjpb6OzAZ7V1VXPevWzTzydTtqnr3O9ea9tv6b9earvTKq8k5BvSHiC5iZrdYXSyYbzlusE85fPzwUTlB6VphiMybsoffvw8sf3uBI6iVaIWGHMEbvyQqIq0aySs3nxXtpRW1Wd8kCIXtYwrbVigpweSeTEpUqzSVBw4ufKQdIXQgF2CBN_zETx4JUM3AIrJKZMejbeyVmy55x6S-7dnthlvnSbKCG0B06zrkgXo4DLQwgfQndDzs2H7JFOk-2dcnaKq1MZNRBGKQQtMqDjylCltyorwcg2SPuC9qiSjNEJjmjiKofZDkYxL-Wn81HuCd54Y98Xh2UPp-X9b-0OcmOg4I2cwbJ-i8SMip8uK71Tb0iLm2hGxJ5hcUuBv8niKRaf5pxyS1AkGE6Ue8C6RVxQe8OYcv0mdE8WrHw5QKFqh6dE_xFzD6ZXr-abgHC-K7iaAMphpmcLyxKNhCaqweWHbbkeW0Cqt_qVhFkfvVrsNO98ZSt_LehzFkPHi4CYfNNg5PY5Mdo1zyLdkasqIW1i-jlgctwaVenkhUxpvnWg6YU8i9hvwXOYCnWspo3G33Dqaya3IYRD-t42o2XVGI5pNB8N_TvBA956uPl3dRlefri-j-7svLcJ17eKYYWmv8ZZDdd9ucO28D8Oquo1Kj8vDNPK-djvswPvh_-n1ffiiGCjwuiBq542cP3ZdZcxaYBv1aWCgSy8ZBHcXxaXzgD3xH99QXJhFQ5-YBkvy_bBqlT6VcHfWeEGLGlt-IrxfRMrz7UccGGbbdKGS1t-OFTtw8IUmLLVsoYt2u31RefM4xX7MGTpU5ayjLEsas46DX7NDrP1eIG10o64K4RgEjMdHPuqScEc4PrIwGsE8mPlqCx1XzhPTwocwD0jEH_lZnPiEBpHg5f91Wx-ZTF5V8kDLkYaua89BJyjfA-j968Ige-R8MPy2GOCyyq1_lxh00dqYv3s3-A66TPP12KTRd-fQ7WKEtutGKrxudxOWLtwLxkQuGjY55wk-fPSvBPGB1UX2EDx3qnVsfwfrWP_B87_dz38Ak3P-7g==))
 
 ### `static_multimap`
 

examples/CMakeLists.txt

@@ -47,6 +47,7 @@ ConfigureExample(STATIC_MAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/stati
 ConfigureExample(STATIC_MAP_DEVICE_SIDE_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/device_ref_example.cu")
 ConfigureExample(STATIC_MAP_CUSTOM_TYPE_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/custom_type_example.cu")
 ConfigureExample(STATIC_MAP_COUNT_BY_KEY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/count_by_key_example.cu")
+ConfigureExample(STATIC_MAP_PREALLOCATED_MEMORY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/preallocated_memory_example.cu")
 ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_multimap/host_bulk_example.cu")
 ConfigureExample(HYPERLOGLOG_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/host_bulk_example.cu")
 ConfigureExample(HYPERLOGLOG_DEVICE_REF_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/device_ref_example.cu")

examples/static_map/preallocated_memory_example.cu

@@ -0,0 +1,151 @@
+/*
+ * Copyright (c) 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.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cuco/static_map.cuh>
+
+#include <cuda/std/array>
+#include <cuda/std/limits>
+
+#include <cooperative_groups.h>
+
+#include <cstdio>
+#include <iostream>
+
+/**
+ * @file preallocated_memory_example.cu
+ * @brief Demonstrates usage of static_map with pre-allocated device memory.
+ *
+ * This example shows how to use a static_map with device memory that is allocated
+ * at compile time. This can be useful for cases where you want to avoid dynamic memory
+ * allocation or when working with memory-constrained environments.
+ *
+ * @note This example is for demonstration purposes only. It is not intended to show the most
+ * performant way to do the example algorithm.
+ */
+
+// Basic types
+using Value = uint32_t;
+using Key   = int;
+
+// Sentinel values for empty slots
+Key constexpr empty_key_sentinel     = -1;
+Value constexpr empty_value_sentinel = cuda::std::numeric_limits<Value>::min();
+
+// Map configuration
+std::size_t constexpr capacity = 100'000;
+
+// Type aliases for cleaner code
+using probing_scheme_type = cuco::linear_probing<1, cuco::default_hash_function<Key>>;
+using storage_type        = cuco::bucket_storage<cuco::pair<Key, Value>, 1>;
+using storage_ref_type    = typename storage_type::ref_type;
+using map_ref_type        = cuco::static_map_ref<Key,
+                                                 Value,
+                                                 cuda::thread_scope_device,
+                                                 cuda::std::equal_to<Key>,
+                                                 probing_scheme_type,
+                                                 storage_ref_type>;
+
+// Pre-allocated device memory
+__device__ auto constexpr valid_extent =
+  cuco::make_valid_extent<probing_scheme_type, cuco::storage<1>>(cuco::extent<size_t, capacity>{});
+
+__device__ cuda::std::array<typename storage_ref_type::value_type, valid_extent.value()>
+  storage_array;
+__device__ int found_count = 0;  // Track number of items found (for verification)
+
+// Helper function to create map reference from pre-allocated storage
+__device__ map_ref_type create_map()
+{
+  storage_ref_type storage_ref{valid_extent.value(), storage_array.data()};
+  return map_ref_type(cuco::empty_key{empty_key_sentinel},
+                      cuco::empty_value{empty_value_sentinel},
+                      cuda::std::equal_to<Key>{},
+                      probing_scheme_type{},
+                      cuco::cuda_thread_scope<cuda::thread_scope_device>{},
+                      storage_ref);
+}
+
+/**
+ * @brief Initialize the pre-allocated storage
+ */
+__global__ void init_kernel()
+{
+  auto map         = create_map();
+  auto const block = cooperative_groups::this_thread_block();
+
+  // Initialize storage using all threads in the block
+  map.initialize(block);
+  block.sync();
+}
+
+/**
+ * @brief Insert key-value pairs (key -> key*2)
+ */
+__global__ void insert_kernel()
+{
+  auto map        = create_map();
+  auto insert_ref = map.rebind_operators(cuco::insert);
+
+  // Each thread inserts one key-value pair
+  auto key = threadIdx.x + blockIdx.x;
+  insert_ref.insert(cuco::pair(key, key * 2));
+}
+
+/**
+ * @brief Find and verify inserted key-value pairs
+ */
+__global__ void find_kernel()
+{
+  auto map      = create_map();
+  auto find_ref = map.rebind_operators(cuco::find);
+
+  // Each thread looks up one key
+  auto key    = threadIdx.x + blockIdx.x;
+  auto result = find_ref.find(key);
+
+  // Count successful finds (no printf output)
+  if (result != find_ref.end()) {
+    auto ref = cuda::atomic_ref<int, cuda::thread_scope_device>{found_count};
+    ref.fetch_add(1, cuda::memory_order_relaxed);
+  }
+}
+
+int main()
+{
+  // Step 1: Initialize the pre-allocated storage
+  init_kernel<<<1, 128>>>();
+  CUCO_CUDA_TRY(cudaDeviceSynchronize());
+
+  // Step 2: Insert some key-value pairs
+  insert_kernel<<<2, 32>>>();
+  CUCO_CUDA_TRY(cudaDeviceSynchronize());
+
+  // Step 3: Find and verify the inserted pairs
+  find_kernel<<<2, 32>>>();
+
+  // Check results - expect to find all 64 keys (2 blocks * 32 threads)
+  int host_found_count;
+  CUCO_CUDA_TRY(cudaMemcpyFromSymbol(&host_found_count, found_count, sizeof(int)));
+  int expected_count = 2 * 32;  // Total number of keys inserted and queried
+
+  if (host_found_count == expected_count) {
+    std::cout << "Success: all keys are found" << std::endl;
+  } else {
+    std::cout << "Fail" << std::endl;
+  }
+
+  return 0;
+}

include/cuco/bucket_storage.cuh

@@ -69,7 +69,7 @@ class bucket_storage_ref {
    *
    * @return An iterator to one past the last slot
    */
-  [[nodiscard]] __device__ constexpr iterator end() noexcept;
+  [[nodiscard]] __host__ __device__ constexpr iterator end() noexcept;
 
   /**
    * @brief Returns a const_iterator to one past the last slot.
@@ -79,21 +79,21 @@ class bucket_storage_ref {
    *
    * @return A const_iterator to one past the last slot
    */
-  [[nodiscard]] __device__ constexpr iterator end() const noexcept;
+  [[nodiscard]] __host__ __device__ constexpr iterator end() const noexcept;
 
   /**
    * @brief Gets slots array.
    *
    * @return Pointer to the first slot
    */
-  [[nodiscard]] __device__ constexpr value_type* data() noexcept;
+  [[nodiscard]] __host__ __device__ constexpr value_type* data() noexcept;
 
   /**
    * @brief Gets slots array.
    *
    * @return Pointer to the first slot
    */
-  [[nodiscard]] __device__ constexpr value_type* data() const noexcept;
+  [[nodiscard]] __host__ __device__ constexpr value_type* data() const noexcept;
 
   /**
    * @brief Returns an array of slots (or a bucket) for a given index.