Add pre-allocated example

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/eJydWI9v4kYW_lfecdId7PIjsO31RAA13eyqqG22WtKtquNkDfYAo9ge1zMOoVH-9_vejG0MCdvrEinA-P385ntv3vDYMtIYpVPTGv_nsaWi1njYbcUi3RRiI1vjVlhEotVtGV3kIX8fvFqm9Ire6myfq83WUjvs0Ohi9HWXbj7Nr-dX9PbDx58_fLy6nX-46bOsk_9RhTI1MqIijWROdivpKhMh3sonXfokcw6ERv0LarPAslU-W7Y6l87KXheUiD2l2lJhJMwoQ2sVS5IPocwsqZRCnWSxEmkoaafs1rkq7bhw6LfSiF5ZAXkBjQzf1k1JErYOnV9ba7PxYLDb7frChd3X-WYQe2Ez-HH-9t3N4l0Poddqv6QxkKVc_l6oHImv9iQyRBaKFeKNxY50TmKTSzyzmiPf5cqqdNMlo9d2J3Lp7ETK2FytCnsEXhUn8m8KAD6RArirBc0XyxZ9d7WYL7rOzq_z2-8__HJLv159_Hh1czt_t6APH7FZN9dz3ip8e09XN7_RD_Ob6y5JQAdX8iHLOQuEqhhWGXkMF1IehbHWPiyTyVCtVUgVg2ij72WeIi3KZJ4ozzUEGTk7sUqUFdatPUvOuRos02X6d5WGcRFJmoRFqAeGVcIgEVk_LLazU4lIQCIaiDwX-9OHWiMMaN_LYJPrIjP9Z_rQVXrWXFIaCEuROMnBK7-_3zreAR4RxzoUQD9IZKLzfSAfBLCSiM0LrnIl13SNhynsQBK5GsYGnDuk4tkKe73aIEXyHlCQt3soplumfemFzFbvDOEfs4iLQjwzemQGCAvLvKndOJtY48rhnKxKZN87CcGmlWSz6yJ2uxwKEJ52IId05bgTqWXP4l4rBLxPRYLt9668Ye-GKxvqUATRdX7HjHDBedFe6MFRKdKW6b3KdZrI1JpD1t-i6uVx7lz9MBrV0LKXrMgzzUHqNN73ae6S5Y6hUivBMVduxgEGriXYW2cevICthNPZoTtAJtJOonIm4o1GhW6TAy8HA_pOGORr95k0y7QwnNYPco-WMWV_l9XaJxEXEmsFFt-MAn7g1BfIETnHdM8CPh2ZZHZPJtYWJtmYw4ZL0T8K7uQ-MJUiOV-9ISx6J6fSzvJBfkoXD-vyVUXxE5gCtbXaFB7EJUxE47FRf8jANiyG3P0UopvS8OLinxcXF5WJWyAAiJSjhyNKLEWKgg51JCsUslyv8B4YdNAElllnSlzT43GM6EQelCKTYbdcj-RaFLENtsJsg3WRhhzfBLDMZjW6xuocBeUNlq_K7qoI76QNSpGJX8yEytlG12_MrEvD59Zyua4tTt0Wg9zyyNl4XAnV2qi6I81mLIfCZBEXALj0V18u5C9R5NY4HtstelaEPUArDHxr-HJjjiY44kQcWO125UtsvcCLLzFzunGzipw_n--qyzQoQQgCEgXK_sB2VI6K0M7RNixNOSC_jYm4k0Hz4eSlBOo998QbzmZtv1Lq-OLq1jU1e3zquIAbATVAdsfZ5BkLq1zHY1_n3nUzuL570O7MOIFKzVm7PPKFzoTCxSkchPhvuVFcEnFp5yK8o7RIVihnnFnKysR4UWpzqeOE50PfdY5OCfn3MkZHpapguZ-G4B06OB9LiBoHCM9o61wnJ6deGeNRcEdV5Q1xGbXh7rGZVy3TWHh8CY3uMRT9SFjR7jzxnInobJGnRz6rrav67-PzTvx0lrJNXef-8aXO_Dn9l0vt8bzKC4T8jLQPkN0Eze4wOdswPue6gTzz-eloaPKz0DzFpItN-cNPkWe23x2yQbCJ9Qo5B-QmDAVVwJ4DsHrzXdkyrarX9Igil7WMK21awdMdyzwbBjlXZSoMnFzb16QrhEbYZZDkez6CJ69keJDnnJwy6_GUqmrFtlv3IbmPfbNPw_Z5oIzMLYFmPUcW4oPLUBsL1Jvx-qtR5yxSrPtnWJ2DqtTGJkKEc8glCBUFHjKdm7IivFwDpHe4HpVglEZ4EpMnOdR-OJNpKT-PHvoP9NoD4764eA6h9P3H9uEQZyS6zsgrGnXOwfgekfOdw3eqfWkRXDsB9gySa078szieQ9Fp_imGLHUGwVjrO1wXsgrCI9yc489C50Rxg8MABdEqmj5_YOQaTt-6nm-KMMR9zw37kAHTUo3yxNGwJl3YrLAdtyNrapdW_9Ywi-m63enQo-8Mpe-1m364jyBl3A3c5AOD3fPjyOyxcQ75lsxNGXFLG24DEUXtYaVe3rl0jssjTMfiQUZ-A55KLvC5luBiUW8dj91WZjQc_7-NqNl1JhOeTYejf89woHtPHMm1i3yBYt7i6sKV3gDXuRuNq3I2OjmtB9Mgeu1n1KU3o7_q5s34Gd05tZrytbcGq099VZzYSjRKv9GGevxrgAzdbW_tPKDr_esrTgQ8GXnqGRTdm1HVDD1ZgP8Wt6ygsal1Pj_JJMz27zEELPbJSsftf5yKdunoC09Net2G0U7nsjLvA5OHyWXkwijHF21F3BhfXLw1HAzT7wVw4ntwxe3TGGg6PfFRs9ydyliyNJnQsrXwBTR24DhPIpc-g2WLRfwpnkax5yjJGFf2l229Fyp-UckHWk4pfAN7anVb5e09P_yQ10rvw3A4-roY4rHOrP-Vr9WDtWn4-vXwG-qJPNxOTRJ8c0G9HjK0PTcl4ZLci0Wycj_9xWrVsBmGYYzFe_9jHRZsXqR3radu9Rwd7eg5Srr19F_39z8wqt1Q))
 
 ### `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,150 @@
+/*
+ * 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 <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 Key   = int;
+using Value = uint32_t;
+
+// Sentinel values for empty slots
+Key constexpr empty_key_sentinel     = -1;
+Value constexpr empty_value_sentinel = 0xffffffff;
+
+// 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>>>();
+  cudaDeviceSynchronize();
+
+  // Step 2: Insert some key-value pairs
+  insert_kernel<<<2, 32>>>();
+  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;
+  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.