diff --git a/README.md b/README.md index 0357c97e8..aaec1e279 100644 --- a/README.md +++ b/README.md @@ -223,19 +223,19 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::static_map` is a fixed-size hash table using open addressing with linear probing. See the Doxygen documentation in `static_map.cuh` for more detailed information. #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyNVw1v2zYQ_Ss3DcOcVv4KEBRw4mBe0mJGC6eI0xZFXSg0RcdEJFEjKaee4f--O5Ky5cRA2wKxTd7H47t3R2kTGWGMVIWJBt82kUyjQT-OMlY8VOxBRIOIVymL4sioSnP63X01K-AVXKlyreXD0kKLn8Bp77TXxj9nMUw-j6_HI7i6uf14czu6G99MOuTgnD5ILgojUqiKVGiwSwGjknH8CDsxfBaa0MBppwctMphFYW8WnZy7KGtVQc7WUCgLlREYRhpYyEyA-MFFaUEWwFVeZpIVXMCTtEuXKsRxcOBrCKLmlqE9Q48Sfy2alsDsDjr9W1pbDrrdp6enDnOwO0o_dDNvbLofxldvJ9O3bYS-c_tUZEgvaPFvJTUefL4GViIyzuaIN2NPoDSwBy1wzypC_qSllcVDDEYt7BPTwsVJpbFazit7QF6NE8_fNED6WIHEjaYwns4i-Hs0HU9jF-fL-O6fm0938GV0ezua3I3fTuHmFos1uR5TqfDXOxhNvsL78eQ6BoHUYSrxo9R0CoQqiVaReg6nQhzAWCgPy5SCy4XkUMsIHtRK6AKPBaXQufSCQ5Cpi5PJXFpm3dqLw7lU3VkxK36XBc-qVMAFr7jqGnLhSc7KDq-Wl4cWdqkrY7upWGGUZCW4VbpDRi9MsDIsO74lrdAMHbtcVQUVJalXfmJvNSsMspH_xMFgcoEaPb67i9J5fjqeM3vowo1NU7E4WJMKJSFYfrDouDYuXveV1-hfrneWaJ3Mq-wxET8YVlkgq357rqVYwLXIsT6IyQqskqGqhm7ZFwI1RxFQchQNRh_HZt_8d2hKu24ZUNio02yNkl8p3w0LrXIX0DmjOJxRRfOClJUq0o8qiVAnlUw-ovRRItqiNF2OhSy8_MEIS_gexXoPwalIFhYbXxatlZLpyazY4A4modjvxRq7fIgx7fl-9TPLKrFbpfVuF97mpV2DyZT1R9GCWkQU1vc4fdcr_D6LaFEWIkNSVhTJdIgJ7Bb_C8xSVRk6IY8iQ52ik6n4MuSxS2aJE5x2AlsIFOeVpjGBQ5A-y8pCyiyjMzr8nIpEDQuCICZIQFJDcENsCO2-O50_13N7h2rvEawDmkmVzwmE47XrTKFkUhuaXfO6GIK6GlWRDgZG_icS20hSVDlBMhj4rPdnr9fbB7_CkY3jCzjNVonszpkJwwxt_4BMMRQCo04mD1Zhzn1g2k38Lsbudc7Oj4LYRx_6TS5k1tqh6jbjnDSxkfQrTtUGErq7VGZRHQ2L67XgJdPuO_nSh3GKdtw2SKv5ddr0R6GoQ6DJhoh3HeXUCTvUcTDY1XbzssrbQyOXcHOsuNtt44A4KJB76hs_kurmcQfxUiVTP5oGg4PBeoHCuwzFdzzuCHU39s6pjt1qmHbm4gGb8SRu-ndEkdJS79D_MKnT7y6th_iLiUMjPk8dlg-Te2XOkbDnfNBEUEH_m00v7iHzm37cx49OpwMbGcvtdldfbzfcYcrZI6ryxUXR8gWvrV7cPBcNSV9uMGUQCD6lDQalVjxjMk-0sJUuErsuxYUXA6WnOsW-7y8vQyKAb6EfjxYlnHQIR5n7DkkSipK03CnlCWzqwAAeBhxHsKFM3-T3Ogl-3Z7XvtuTRveNXW6kPssCjY55aixsETKihwCPsOUM4mD3GpqCCOGmSCNdX3SpLJDfn-s7SM0ZH1NaCPwObx-P0lGKkoB-DKdODlsnHIPRkE5mjOKS0axvSum-meC-PvwCuxLjwT0FTaxK6JJDsu7xRhSm-NPi8xk-_MWH7mgAw-HRoX5fM0aBfrkXm8Frw8bZ8bldLtb-viIGqDhNcjGh5f45vL4lGrTPlcrCgEbnkCXxLsPDjnDPa61jaJ5hDLiPCrcGLhfQep4QFRyuBoWX68UFjvlpxTk-AP8G79yJ6Hwh3gzjROew9eGC3ulSwxV8a6LXEHy40vuXq6hYcd4_Pav6uK1K69-8ojZmHPLXr_tvoM00Xw5NnrzpQbuNl5vFPxZPI9J2xvK5ex3L5LwRk3Oe4eLKvzvhAt5VxWO0jet9bMCDfZR3tP3u_v8PKKbiNA==)) +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyNVw1v2zYQ_Ss3DcOcVv4K0BVw4mBe0mJGC6eI0xZFXSg0RdtEJFEjKaee4f--O5Ky5dTA2gKxTR7v3r17d5S2kRHGSFWYaPB1G8k0GvTjKGPFsmJLEQ0iXqUsiiOjKs3pd_fFrIAXcK3KjZbLlYUWP4Pz3nmvjX_-iGHyaXwzHsH17d2H27vR_fh20qED7tB7yUVhRApVkQoNdiVgVDKOH2Enhk9CExo47_SgRQazKOzNorML52WjKsjZBgploTIC3UgDC5kJEN-5KC3IArjKy0yyggt4knblQgU_Dg58CU7U3DK0Z3iixF-LpiUwu4dO_1bWloNu9-npqcMc7I7Sy27mjU33_fj6zWT6po3Q98c-FhnSC1r8U0mNic83wEpExtkc8WbsCZQGttQC96wi5E9aWlksYzBqYZ-YFs5PKo3Vcl7ZI_JqnJh_0wDpYwUSN5rCeDqL4K_RdDyNnZ_P4_u_bz_ew-fR3d1ocj9-M4XbOyzW5GZMpcJfb2E0-QLvxpObGARSh6HE91JTFghVEq0i9RxOhTiCsVAelikFlwvJoZYRLNVa6ALTglLoXHrBIcjU-clkLi2zbu2H5Fyo7qyYFb_KgmdVKuCSV1x1DR3hSc7KDq9WV88tUtaVVmhmlb5q7tiVroztpmKN_pO14GjQWZ0ywZqx7PSWwT2B4jq9azUrDJKRd36AlTN7fIQbm6ZicbQmFdZSsPxo0ZFknL_uCy-uP53oV2idzKvsMRHfGZZHIB1-e66lWMCNyJFYxGQF0muoHEHmBwZRLOQBtULeYPRhbA5de4-mtOuWARWJAss2qNW18jJeaJU7h-4wVtUZVdToJIlUUeFVScVwNc7kI2oWa6stasrFWMjC6xaMsITvUWwOEFz5ZWGxY2XRWiuZns2KLe5gEPL9TmywPYfo014cVj-xrBL7VVrvduFNXtoNmExZn4oWpG1RWN-c9F2v8fssokVZiAxJWZMn0yEmUOb-F5iVqjI8hDyKDGWEh0zFVyGOXTFLnOCYEqh9UJxXmvobpxd9lpWFlFlGOTr8nIpEnQaCICZIQFJDcNNnCO2-y87n9dzeoTqcCNYBzaTK5wTC8dp1plAyqQ0NnXldDEHtiKpIBwMj_xWJbQQpqpwgGXT8qvd7r9c7OL_GWYtzBzgNRYnszpkJUwhtf4NMMRQCo0ajE6zCmAfHtJv4XfTd67y6OAni4H3oN7mQWWuPqtv0c9bERtKvOFUbSOjuNphFtTcsrteCl0y77-RLH8Yp2nHbIK3m12nTp0Jeh0AjCRHvO8qpE_ao42Cwr-32xyrvjo1cwO2p4u52jQRxUCD31Dd-JNXN4xLxUiVTP5oGg6O5d4nCuwrFdzzuCXVX7f5Q7bvVMO3MxRKb8Sxunu-IIqWl3vH546BOv_uwHuJPBg6N-Dx0WD4O7pU5R8Ke80ETQQX9b7e9uIfMb_txHz86nQ5sZSx3u319vR1VOGWDQc4eUZP1gE_qK6YVyu1suKqoUMv97mVDzldbDNe0LrXiGZN5ooWtdJHYTSkuvRAoNNUo9j1_dRXCAHwNvXiyICHLIZxk7RskSShI0nIZyjPY1o4BPAw4jWBLkb7Kb3UQ_Lq7qM_uzhqdN3axkfYsCxQ61qmpsD3IiG5uj7DlDOJg9xKaYgjupkgjXV10oSyQ3__XdpCZMz6lsuD4Ld48HqWjFOUA_RjOnRR2TjQGvSGdzBjFJaM535TRQzPAQ538AjsS_cEDOU2sSuiCQ7Ie8DYUpvjd4kMVPrHFx8fRAIbDkwP9oWaMHP10Hzad14aN3PFhWy42_q4iBqg4TXIxoOX-4bm-IRq0z5XKwnDGwyFK4o8MvSbqyrhHqdYpNM8wBtwnhVsDlwtoPQ-ICg7XgsKL9fISR_y04hyfWn-Bty4jyi_4m6Gf6AJ23l3QO11ouIKvOvTugA9W-vBGFBVrzvvnr6o-bqvS-telqI0Rh_zly_5raDPNV0OTJ6970G7jxWbxj8VsRNrOWD5371CZnDd8cs4zXFz7Fx5cwHuqeIx2cb2PDXi0j_KOdt_c__8A99LEAQ==)) - [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/eJzNWftvG8cR_lc2LNJSFp9G3QCUKES1nZZIKgeS4iCwjPPybkkufK_c7ZFiCP3v_WZ270XRFpX-Uhkwydvdec83M3u7Tq7yXCdx3pl82HV00JmMe51QxstCLlVn0vGLQHZ6nTwpMp9-D1_cxeKFeJ2k20wvV0Z0_RPxcvRy1Md_f--Jq_ezN7NL8frd9c_vri9vZ--uBnSAD_2kfRXnKhBFHKhMmJUSl6n08eFWeuK9ykga8XIwEl3acNdxa3edkzOmsk0KEcmtiBMjilyBjM7FQodKqHtfpUboWPhJlIZaxr4SG21WzMrRYXHEb45IMjcS-yVOpPi1aO4U0lSi09_KmHQyHG42m4FksQdJthyGdnM-_Gn2-u3Vzds-RK-O_RKHMK_I1O-FzqD4fCtkCsl8OYe8odyIJBNymSmsmYQk32Ta6HjZE3myMBuZKaYT6Nxkel6YlvFKOaF_cwPMJ2MY7vJGzG7uOuKflzezmx7T-XV2--93v9yKXy-vry-vbmdvb8S7azjr6s2MXIVfP4jLq9_Ej7OrNz2hYDqwUvdpRlpAVE1mVYG14Y1SLTEWiRUrT5WvF9oXZRiJZbJWWQy1RKqySNuAg5AB0wl1pI00_OyRcsxqeBffxX_RsR8WgRLnfuEnw5yO-F4k04FfrC72dwQSO4Lhooh9oizDi4PrpkhD1Voyq6zIzTBQawjgrZVvkmywOrRF3Su_IOJemsCl28O7tFGZBI3hHzr1yh-Ht-YIE4WYHTxSJ5KmfcKH8IFatJ7pBDGgZNR6yMbNmd7whQ3K7zlZnIKZWnjqXsKvCna06_NMq4V4oyJ4xEBeBb_k5EeXH_aoyDUYXP48y9nzOg70WgeFDEWSkpbsz1B_RnzCj5kZLrClBoPZl_dLkI6XYXmQAo_OCh9hPVcUQ2AYIdKRL04WP4Esay1JPiZ_16kDhJREHphtqvJagO8BIErcEno4A1AikSpBpTlBUVpkaYIER16F24GYGdpG4KNjoxCunLn5KtmwbSI4gck7KWVsxAZAgz1BwjtKZjJcJkj2VVSJNGx5yXphxibIxWe1zXFcGpHKPG_mGQRAfgbAFMMmMZYNJUatq0llJiPxH5mKW5ih9CQ2lRaEjVRG0dc68KPazhCzECMtoLELX7FZwSJiLcNCeWRW2D9GihtNsAYBwGcygci82CL4ns4cRRJWPkQVIqcqeEz459IGAF2ZbWub1EbgjR90_FG4oBDXpdJNg7AJNyvtr1wA5igiYdgKvX2CpOtcLRGQt6DC3xjvHNkMSMg8rBsTR3ifitX-GDq8ExCa54mv2fFc5xTKEvH4Mou4iDwWguibxCD98GgOf5TSAZYd-f2zZNJ969p6ZpOg5D4Q75Aqh0K2EajOpMzDilrWFccxKYwVt1z9ssh54fuoUIsiBNtyu8soo5BtJOg5hUwsI0Vx1BPVLxfijSdVjDaeVeEFKPW8ZZjMZeh5Yp3ogJoPbEZQWtZdyjIXYVR2n_grU6wKoSPO1FnUCJkjzqFqTCa5_gMpVkXCEcfq3CIHHnEAKfSi5buTu3hHx2RB8ASbTREPKFbBLLgf3ItTMQ8T_7P98cL-eKOjwf0ZQWJbbj8pYtJ8KkbUEFKmAh66RPS8UupE7KyUw6ENRldKvhaT9oBeiC496lb--ADSH08qikyTocjRzJQpMtSuTyYr1CcioLlGEPGFznIDhSPbKi31WsWcoBuZN8hVIU7J92khw9wSojO0W4Zkqy1qB_o8Dm1RyuoCbeBij3qjySSVOtu1Feg1Q4WfPDR1EuL01Fn2jEV6TT8aiXYwx9ieJYkH-8V9kD9Op2KZ6YA9ue9W3kgf4PVLGlBw2ayyDiaWj7igUETavwyCbjO2emVI0IDwcKCG_pBkDcev5JpaDLRUDPAl7Pfw3c9UpGLDrtMZyGZofNMEnQogmM1HHXwSq_-jytqAzKPLGtpMGObrZe3ZJY2IbpsEniw1z0ToA8hbucyzFasFvQeQ9SAA_g_YdAT6zOos5uzNWzHHfW0JIGhvqzhwXVzd49m6SgPZRtPMZyoO5XHq-NCQdk_sCquzQGaQQiVKEL99ZDurgc9u_6bez_Ta0PdIoZ7LS2mRwXmENzX6E3ZQSYZlo0id1vhDM9lkYkkR71Y0tPq-nttrveTlmN2VZ5PtYscq9C9yhcQJHs5K-qy8Qsp7EvAxLklEaPWzrZdkGDvBNJT3qFVnzwcyhh1kGAyn4y7FZhlVBc0yFIogNaUkPKufchmvnjowfBulZivyMEGoYP6H5DR-w6L2_oC-Z2t8x4SDhzpWIWYbG_4DSrUSPnIaSoowoATPVYhBFoeA4ivHh8EQTtqKWK0pKX2_oDGOLljok2AJsCwJ61h-n4YiugwQikSkSPdKEdjGU9Efs3ZWr_39FtSqE263k-aqiQxDi7VUxRha5qrZ1e11AyWTCm6m4tXob6PRqCb-GnFieK50E_aBZlfsdqPe6KEnduPeGB-DwUDsdE8_cBjYCX0yaV0LnMMqF04yZt2tAOCseajk2m1sHXACdk96zfM23Xpi1D7fZsrGrdg63DuOsYuSfdbucZN5abgkQhjAlXThpRGWc5m7GyYY-VsRJhLtryS5KgStPUKrnl2lbm3w6uyg92rqU7voKx1WGolhk05LNprRC5_ShAscT0B3nZIassImkc21_pidTR-2NeOgbERbGZh8RWBVIapTYVuq-jphV3J4shG2J6ts2T3Om4dn0WBBd4ey6Rg6hHdsX0QFSqhJOH53x4oQgo_MvDRL5rDnOUMoPQ_UQhah8VYyX3nlfRuTBu2H2l__UgaOipO4n2xivuKMCsP3oNmhdoXhCSUl2eTVGA6XcNsOGLQ-K287HIvPKiM4LN3nwpvrDFdHfHMNsl1qRFM94nBTlDRmBuV69rwmnHv8cCo-fBSeKz2e1-VF7DwR_QsxT5JQ7Fx1FlR0xbfi5YmY0uAiGoa5hJLMOUeA0zUb3UFVTXDdMj1qhw_jA2xy0Zq9uuNSz7385DLmUSpCk5ev_nHW3sMlzy47eK-T8rR5uC_GJ0jT-gkT2puKz8_PK3q9xt6Li4tu7agjZsvW3yFE_XMk9rDxmUSeMUy3mdtI-jPsqksTKtLdk8FSGbRqzYEZMWTE-Tkw8ap9X1IenWCJNjTJfRh9pEcWKOIg_HICP5233N02bnefmb90_GD20kKp6f4I8LU4Kwn2DsaNaBZRJ9G1MpiU1tzMGB4NoSq0Y3XJFKrRrdWN_RMtAxFDn4jcqLqGhvmPKP41gUb9f0TCGs0q4EHobptvrfg-uXKlhRv8qoTfYJRdeyliJD8rr_l6o9soNnxs91zOD1YDGjeYJkKMelSyPNm6dgc3cJhmF1sbeZ9oeUBgeyrGhLX8m4l_oksZGW4kzatZoeo6j54aFczOfawR-HjJotv2wdNJWpvoOXsh6Z77nj59oO4wZv_V0m0UoCOQxZWohtOAJeeji66j5Sy5tz4u18-eZPFQJStmTGtszJT7IHVjb5i-EbcyA4PqjjujOzpCkeZ4SZfFoNk5K6-RnBIjewfU6XXoHTBm86x-s92J174_fvmqGGM5SY197d3pQ46pf3o6_k70ZeavpnnkfTcS_T6KpMF_9i1PP5TRnN-Fh3reoOn7foiHa_viGg8QWvHnzkOvXAditdbhtc7DR_73Xx4ZenE=)) -- [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/eJytV41vGjcU_1febtoG7XFAtqkVASSWtBpaRaaQtqpKdTE-A1bu7NvZB2GI_33P9h0cH21SaURKgp_9_Hvv_d6HN55iSnEplNf5vPF45HXavhcTMc_JnHkdj-YR8XxPyTyj5nvzxUTAC7iS6Trj84WGGq3DReui3cBfv_kw-jC8Hg7g6ub275vbwd3wZhSYA_bQO06ZUCyCXEQsA71gMEgJxT-FxIcPLDNo4CJoQc1smHiFbOLVL62WtcwhIWsQUkOuGKrhCmY8ZsAeKUs1cAFUJmnMiaAMVlwv7FWFHgsHPhVK5FQT3E_wRIrfZtWdQPQOuvkstE47zeZqtQqIhR3IbN6M3WbVfDe8ejMav2kg9N2x9yJG90LG_sl5hoZP10BSREbJFPHGZAUyAzLPGMq0NMhXGddczH1QcqZXJGNWT8SVzvg01wfOK3Gi_dUN6D4i0HGDMQzHEw_-GIyHY9_q-Ti8-_Pm_R18HNzeDkZ3wzdjuLnFYI2uhyZU-O0tDEaf4K_h6NoHhq7Dq9hjmhkrECo3bmWR8-GYsQMYM-lgqZRRPuMUShrBXC5ZJtAsSFmWcEc4BBlZPTFPuCbarp0YZ69qTsRE_MgFjfOIQZfmVDaVOULDhKQBzRf94x0Rac5yQY1WEverMr3IcqWbEVviDeGSUS2zYHFuC9csIyhtUpkLE5WwXHliv86IUOiO5IkDsZwjFeLzwp2SwFnXbMJ7xbJGxGZcYJgf2Br0OkWCYOhzqoHiKZmEuB6addig65BS-teLUAO5rH6bXhqNAGG4kEqHIf5T-CM8VlOrw2b73L2l_sc6dIBsHrc-TPG31bC9PGfFksQ5O2uHlZyxZHZgiXrakr2i59hS2X1gzcxao75tjVMGC6IWgIGNTZof22WFG4did3le3iVTR5havXYcT4oJon-Gh7r7DwtgUfA2WGB0ngl4CMglGGTfRGeIgwWJxF-FaK50O45xTqWMn4ExXij_hI-FLFuoYwvMJfYmKC1BBQGBXs_sDhx3d05HV2Ht5qK2lDyqT4Q9aRWaSgVKR52O4v9iEEHkSZgSninowevWL61Wq6ALumbMNLBEY9lXzGQ3i5WRkBwLsYPHklSvLf5yBwLsHdu1abS3l2dPOirtzvZOKVacLRBdZYxoZop3WTZMR-IizbUJWtPlirVnd5_9Fk7ZHHtHD1zp6HQS8oAXnFShmvPxwa6T2tYtqNivteq-O2DKaaeTZpLGhCehi5E1wNZiFCGK7pFf_FNz-_0CAcDnLxVa1awpvL5n8l7t5tjdfHtGM65ukfr1esWbsmQ1QbqkbhBot1o-sgBULDV2G2V6kuk2c75kwoWt4uodMQIYSW27XaGcmgmA42bsvRRrCR5-EHJl1K3MzBHHGDfMPY20sxeiUkwJ0yIxvLEkEcwILSL8uvVTsAuogdor7N_3uI1bYI8aIXUrFPeNSYbY_c3WL537tU-hpKT15pTg36fDemlzjuzP0XNYap4NP0b9JAuRjVN0d7ftVwsreqGSUEMbAlXGwFUCpLe0MTeexoFsTwKnx1ZvZuIUlSsWoIn2vmCCOR64GNcqSegfZOTLfQWqMPOWIfgox1Lszrtmvs9pt2pcg3XrTMqeSezvTumnPt9M-YOE3Of0U58nc_440asZXVp4MLd1TT_qm6JrxngW1c65-5ppM3UKdPcMMIA21geex7gKmJoJFufPQ04okrAqI0oqFLrLGROIgnun8z6AId5j2-x9JZSf-Zd7iCRT4hdsEI84sPtwv0OOUtPuZiRW7D4o-VWIVa2ixz_gR4Vh_t4PgaVfzTrP4rwzjyRblmaExxCZmcsQzbQrRIJZ64YWY-JuPgiKw_8XkkLdYFccS_fjM2d_4qiVYsDCnazS4sy6nBXcO7muioCJyCycYV8xkUzr0Oi70WZHxqnhnps0Z1A7AGEoaysw5pmGbhefWuOcUnwi_QBvLYMMyWw9VMEEzfYuzfBidBXKzRSCK_iuNg9VfLpm--e3J5aUti9-z9solql2b3OvgTf26MuX7VfQIBld9FQSvmpBo4Ezj27YvhCxqBGTZGof7DGfVnRSSmNcXLrXNS5gYxQP3tYv5VhZD-ToJ2_7xf78BwHUjUo=)) -- [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-BuUpfRR0OTdMj3ofVcwTqhXQkLysC4NEqJNitJ0UvSPx3Ej70wggSF4toZ7bS0Uk6qbkUJK9_XoIPbafbBPAbrPSzsqOp-jpyACQZjIv32pQSt-gV8pQzCVJ3FEUNBu4WSJOw_e14KkjHLXRgCaiCDnh532SrhPKXBk-0GtYX6w134Tp4YxyFEpkpApm7p5oeBNd2YXjMcWuopB0yErFRLzq8-bR6T5cKMLfmrntMJh6dqhzbgcNfwwrh9Bm4jjE1-ob-lPTpLf4etOVb651mip0fnqJ08__lwzG56ddWP01cqhJczwABLgV2ZRT_-lrTYigR2bBKktnKNVj5nM0Huq5LCIbADqWgX_sYtGQ_VyH8d_aIn5JmgyQU_xEGReW3zGzXRdhbofgZek6F39uUe4_61cbj6alEsAuO6am-vc6o5bcSIHAUV16-cOnkr5vafCVt1VHjvxQPkPA4fvvjtprHG_iaf6ZNDBlr7l5IA56yPh6hAU1L6jH9dyJi8_x8XElvt_YenJy0vU32Z3AWN9FtxFmSHyi22vAzJUG3Qb9cM9tVFT03hZFDiMQO91gQ_Xl4o0GBY5EGVI1qIYJt_ufXavavd2RDi_APT-9KsJF1B0ggNHdht7aGy1ZfriJtEut7omJtmHRbSvfLui9Bxc2tUVAKk8HioVMth0OU4lffKHTNICZe6F2z5cUgtKC8g2OH2N0ZuabBnMkfrrbMrersst7G3IDEGyXYTs9s-0wjRsvBtt9tkEnnSZuubM0jerLadl6XZuEE17APS2quizdBLa8WT9AugFcl1x_waBFeYA8Xbtb4N_uIKRz5C-MpWYimxjp9Dt0fTSRzuo_WnWSB6UODt8XB5hOV9b9RaszgPyJ2ts7-F4M6G4x4QuFGAyABha_3BvsIJLxjP_MFZlZQ6ZSKsLgg_ubFAbQxZL7znPfzwN5WvPwZ-f5M__7H7PCY_4=)) +- [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/eJytVw9vGjcU_ypvN22D9jgg09aKABJLWg2tIlNIW1WluhifD6zc2bezD5Ihvvue7Ts4CE1TaURKgp_9_Hvv_d4fbzzFlOJSKK_3eePxyOt1fS8hYlGQBfN6Hi0i4vmekkVOzff2i5mAF3Ahs4ecL5YaGrQJZ52zbgt__e7D5MP4cjyCi6vrv6-uRzfjq0lgDthD7zhlQrEIChGxHPSSwSgjFP-UEh8-sNyggbOgAw2zYeaVspnXPLdaHmQBKXkAITUUiqEariDmCQN2T1mmgQugMs0STgRlsOZ6aa8q9Vg48KlUIuea4H6CJzL8Ftd3AtE76Oaz1Drrtdvr9TogFnYg80U7cZtV-9344s1k-qaF0HfH3osE3Qs5-6fgORo-fwCSITJK5og3IWuQOZBFzlCmpUG-zrnmYuGDkrFek5xZPRFXOufzQh84r8KJ9tc3oPuIQMeNpjCezjz4YzQdT32r5-P45s-r9zfwcXR9PZrcjN9M4eoagzW5HJtQ4be3MJp8gr_Gk0sfGLoOr2L3WW6sQKjcuJVFzodTxg5gxNLBUhmjPOYUKhrBQq5YLtAsyFieckc4BBlZPQlPuSbarj0yzl7VnomZ-JELmhQRgz4tqGwrc4SGKckCWiyHxzsi0o4LQY1WkgwfybhmOdEyP5DoZV4o3Y7YCu8OV4zihmB5aksiFxjD5LRQ50QodEYaOFjtNrxXLG9FLOYC43PHHkA_ZBhZjFlBNVA8JdMQ10OzDhu0Gbmgfz0LNZDz-rf5udEIEIZLqXQY4j8l3PBYTaMJm-1z91b675vQA7K53_owx99Ww_b8lBUrkhTspB1WcsKS-MAS9W1L9oqeY0tt94E1sbVGPW2NUwZLopaAgU1Mfh7bZYUbh2J3eVHdJTPHqEazcRxPiszWP8Nd0_2HlausVBusDLrIBdwF5BwMsifRGeJgJSHJVyGaK92OY5xzKZNnYEyWyn_Ex1KWL9WxBeYSexNUlqCCgMBgYHYHjrs7p6OrsOhy0VhJHjVnwp60Ck2JAaWjXk_xfzGIIIo0zAjPFQzgdeeXTqdT0gVdM2UaWKqxXismsFiyRBkJKbCCOngszfSDxV_tQICDY7s2re72_ORJR6Xd2cFjipVnS0QXOSOamapb1RXTSrjICm2C1na5Yu3Z3We_hXO2wKJvLohIr5eSO1RflY-wUtZwHq7tobIw2Ba7Lf2ShsNGp-nXt2e5pAnhaejiY8HbAooiRNA_8on_2NThsLwf4POXGqUa1gze3LN4r3Zz7Gq-PaEZV7dI-2az5klZMZogVTLXvbudjo8MAJVIjS1CmUZiWsSCr5hwIau5eUeKACZS2xZVKqembXPcjA2TYh3Bw3dCro26tRkUkgRjhnmnkXL2QlSK6WD6GoY2kSSCmNAyuq87PwW7YBqog9L-fWPauAV2rxFSv0Zv35hkSD3cbP3KuV_7lEoqSm8ek_v7dFgvbU4R_Tl6DsvMs-EnqJ_kIbJxju7ud_16UUUv1JJpbEOgqhi4KoD0ljbmxtM4Re1J4PTYys1MnKJqxQI00d4XSzDHAxfjRi0B_YNsfLmvPjVmXjMEHxVYht1518j3-exWjWuwZj1K2BNJ_Z3p_C0vP5HsB6m4z-anP9_M9eMEr2eym4R6vYNRqm960NAUWjNzs6hxys2XTJsRUaCbY8DA2RgfeBzjKWBuxk0cFg-5oEjK6kyoKFDqrgZCIApunc7bAMZ4j22tt7UQfuZfbiGSTIlfsCnc43Ttw-0OOUpNi4tJothtUPGqFKtGTY9_wIsas_y9HwJLu4Z1nsV5Y140thzFhCcQmTnLEMy0KESC2eoGFWPibiYIysP_F5JS3WhXFCv345tkf-KofWLAwp0MK2LFA7Mu46qPHV9XR8BEZBZOsK-cQuZNaA3dOLMj49xwz02XMTQOQBjK2sqLOaah38d30bSgFN8zP8BbyyBDMlsHVTBDs71zM7AYXaVyM3ngCj6CzasS35n5_q3siRWl3bPfii6KZabdQ9pr4Y0D-vJl9xW0SE6XA5WGrzrQauGco1u2H0QsaiUkndvXdcLnNZ2U0gQXV-4pjAvYEMWdt_UrOVbUAzn6ydt-sT__Ab92bH4=)) +- [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/eJyVWYtOI0cW_ZVar6K1Bz-A7ExWBqMQmGysZCECJlEEo065umyX6Nd2V2McxL_vube6utvGDLOMxNhdVfd9zz3VPHUKXRQmTYrO-PapY8LO-KDfiWSyKOVCd8YdVYay0-8UaZkr-j56d5eId-Iszda5WSyt6KqeONw__HaAXx_64uK36fn0VJxdXv16eXV6M728GNIBPvSLUTopdCjKJNS5sEstTjOp8F-10he_6ZysEYfDfdGlDXedau2u0ztiKeu0FLFciyS1oiw0xJhCzE2khX5UOrPCJEKlcRYZmSgtVsYuWVUlh80Rf1RC0pmV2C9xIsO3eXunkLY2nX6W1mbj0Wi1Wg0lmz1M88UocpuL0S_Ts48X1x8HML0-9imJEF6R6_-WJofjs7WQGSxTcgZ7I7kSaS7kItdYsylZvsqNNcmiL4p0blcy1ywnNIXNzay0G8HzdsL_9gaETyYI3Om1mF7fdcQPp9fT6z7L-X1689Plpxvx--nV1enFzfTjtbi8QrIuzqeUKnz7UZxe_CF-nl6c94VG6KBKP2Y5eQFTDYVVhy6G11pvmDFPnVlFppWZGyV8GYlF-qDzBG6JTOexcQUHI0OWE5nYWGn52QvnWNXoLrlL_m4SFZWhFseqVOmooCMqiGU2VOXyZHvHbDSLUnXvfgeIfal0tbG9LZQjY3UubZq_XClsOMJKbNTutXmZKDJbRhvrdpmXhR2F-gEuBA9aQfhwuWuLzWVSIGzx8IUDsbRbtkJjqOcvnpnEbjwzKSpBy3jjIYe4YB2jd640v-eWUWmZ2GC2Du71OtCPEvmlMLkds9zouTjXMTIDU61GfgrKZ9UnzkNRGKg4_XVacAWYJDQPJixlJNKMIst5jcw96hT5zO1oji34zDpIjEoTqx8tSZXOoMFsPYBBjYC-MEM9ZPlSoN9tushlLKisBDYWDcZMX1cvYWmyiLwdVM9sikK3zDSVJqVCk23eNZXCtQcjyU4Wf9dp6g5VNUd72XWmWwZ8D1zS4oZAqYon9SdZHtaBJITLyjxLgRto12g9FFNL2wjTkFCNLmBAKJbpimMUI6ssvrJSJlasgF_YE6a8wyuT0SIFhizj2qTRRtpdUqccgoKDR43o4l4I0gfg0QnJRAKUKimlrMEkWWlFASzTANaWxzaTlI0fqNWuzV9anH06PxXceQj5X3pj139kJm4QMl9ECKSPNuKpc5K9ceBnvZ6iRWEyqffdKlZLRE88yKjUAaWAygjlYA0hK8ICPeMxVTUtbgj8lBi4wDIvS_tlocjJltwS-fnwz6AZDd-z1FuTfBZVUYgr70jbSZxLocCoZVWABWZTFG2U3rZAsn-mF0jADaTwJ4bRSizgw3Ujp5EHCAnelpKUccA7SIhNLXoDj2bw3x-lCmCvi_bZtLTucMkRe10GTR9gjXXCnBHOGdSe1ahLgAegyUWuKZQ-N08iY03pan2rct560iQNKBYEiyidySgIxENqwg0U61KBVYmgofelH19aZPZbe33e62i-daBVZlsx7N0lT3S4JDxy0bjiASUmApNrPG49OvZ6-03YTogIiSAoluAHIYJQR6l1cDy-QeCvUdWE2JSEoHBfjggPhPD-KAIlEaUpbcgJyydigf_PTTx8RDHUao_ap0z4uOXvpAUA71z3T8NHiNhDrWIihfytUl5H0y1tVNhE7LMq9Aparkuajuuo98STi_to5HGcJgU3l-80pnzVKKHyPHAnZIk9t0WUIpamCBK9Iomfoa8ql6ETGKAbApoMXSIa43EmTf5Eum9hyue-OHhmJoofM4d5taTaNDZubnKElQF0Awqw0e966fveXiX5WegISNQWyG6iVSM6tPagTP4OTvBN5TrWiXVe-2PsMiESFVYox2NHZ8hZLqxvD6mw3FJlTQEmrAOHyCdPFK3BSaFRJOHzkRdLsZprq5aBDMPugZcQY8Dl6yDNweGgIpKPOvSheq4ihmTuTdrVxuvPrirgJTF3sNgWuBSlUqCf8zKK1jW2CMTcoQ1NVg027iqukkJ45lyVdApWiqo-wBuS0AHYRh06nrhZhq1u6rYbqDe8LuPuy-w5X6koWgUvJqjnujR2ZsG19-tZ4HZsK-r6RLxzRHI8zuUqyFKiDnmgZGG725Dj87AtqpXHFyF4O6_IG6cOilGJJukSHHt4Qx4uiAvZpbQcFeIC6bwL1O3tVZ_PqFp7vZN_UXKqijaJicu4uTQxm5C5WmJMK1vmPJWLOPhuvwFRyGQI8kV91Cyxis2lyrqPcWbXgkocbZUT_6CLDprI3dToc_6AzyB9eGgSHYHuuTE5pDnoqUJBPK2MQhrlhY5gJA6hbJeVHg4A86lEE2OtWRWusnlFq0JpJVEq9oQhma5dQpOJTM69CRXWVjSUcn2MIyfdwYFLiXN3W4KjNLWMzfN8xEvwmdskCOiemW5PdgctjZqaY0zE-_1_7O87CIegn2pGyV3apgouEIWHspouU0B2qKAohCXfnukqAkWso_JEFmRbt1sb8s3WgR534o6pfdf5kW-uDWPvN-7EaC9ynW4VBXNAVMaWJbCfe4OilMEOSNylhXxEd6EUdGGH7n0G7fM9vHFh5KRWAXc9X09AH9kzIIXla01FyX22huLjy1gvZbFtN24Ra_xC_7sLTG1JfS_ttlErlvc64PGCtgo8Z3bVt9_rf-3Wlk9DQoFuz59trzDX7fqV28-gO1V8gi7XhenR3KNueUKnAhaSlz1hUARNQYy2C6J3JJ5bJX8aAQApohXM82ipSUQzkBxCcqRfT5_H9ZNt8tdtd9lZNe0UvVAyAKOZLKo3OGijbzAnZSjmkiTuaAhaDdwqkabh-0bw1JGNxmjAEtGDgrDzPklXCVWtDB_oHVlfrDTfgul1GtUnlMhIlczaPcnwJrqWC8djyl1NH8nJWsVEvBrztut0Fy4VYW_D2nYYTPM61AWPglY8hnVA6DDxG-Jqze38qW3SW1y9Hco39zpNNTI_vcTo5_9LBmPz0y6c_ho51G6cD4AAbkQ25dJ_-loTIuiReZDl6Qytesxcjp6Hei7LyAaAjWXgX3SxaMh-btL4b22RvyRNBqgpfnEZl5bfbua7LsE8CsHJ0lUh_tyi2382b2w8klZKALkcmIbm3-ucxnGrBAJHc-l9ID5V1H1Lg--8rT5yxIfqGQIO33842tzjOBMv88-khSl77cMDcdBDxTdPWFD7cnrcrJ24_BwfH9fi-62jJycnXX-L3QmMzT10G2GGxCW6vRbMXGlQbVAP96qNmoretUWRwwjkTreYUHOxeGM4gR9RhdTDqWXC7f5nN6Z2H3eEwwtwr55eFeEy6hwIYHS3pbeJxoYs_7iNtEut7omFbsKiO1a9t6B3PbisqS3yUUc6UCxksh1wmErc4guTpgXMPAu1e3VJKagsqN6_8YsYnZv5usUaiZvutsydqu3y0YbcAOTaVdjOyGwHTOO2i4ebc7ZFJZ0mHrmzNI2ai2k1et2YRBBewD1tqqcs3QK2otm8fHQPcFVy8wUPLdoDxOna3QD_dgchnSN_Waw0E9HEk06_Q1dHE-m8-VNWJ3lQ6uDwfXmA5TSz7u9cnQHkT9Te3sF3YkD3iglfJsRgADSw-OXevw4iGc_4j1-RmbVkKqUiPHxwf6nCA0yx5L7z3PfrQJ6NdcSz8_yZ__0Pf05rhg==)) - [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==)) -- [Heterogeneous Lookup](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/heterogeneous_lookup_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJy1WWtz2zYW_SsY7WwrxdTLjjeNLGurOO5E09Tu2k4znThDQyQkoqZIlgCtaDXub997L8CXRdn9skpm-ADu-9wDgN62lFBKxpFqjb5sW9JvjYZOK-TRMuNL0Rq1vMznLael4iz18Ln_6jZir9hZnGxSuQw0a3sddjg4PHbYxW-z97MpO7u8-vXyanozu7zo4Vya_1F6IlLCZ1nki5TpQLBpwj242BGH_SZSdIQd9gasjRNuW3bsttU5IS2bOGMrvmFRrFmmBKiRii1kKJj45olEMxkxL14loeSRJ9ha6oBMWT3kDvvdKonnmsN8DhIJPC2qMxnXhev4C7RORv3-er3ucXK7F6fLfmgmq_7H2dn5xfV5F1wvxD5FIWSWpeLPTKYQ-HzDeAKeeXwO_oZ8zeKU8WUqYEzH6Pk6lVpGS4epeKHXPBWkx5dKp3Ke6Vrycj8h_uoESB-PIHHTaza7vm2xd9Pr2bVDej7Pbj5cfrphn6dXV9OLm9n5Nbu8gmJdvJ9hqeDpJza9-J39PLt47zABqQNT4luSYhTgqsS0Ct_k8FqImhuL2LilEuHJhfRYjiC2jB9EGkFYLBHpShqsgZM-6QnlSmqu6d1OcGSqfxvdRv-QkRdmvmBjL_PivkIRz13xpOdlweTpDJ_3F1nkoVYeTnbGlPb7OktCURvSQZop3fcFgCLsp3ztpmIhUgE4cj2udC9onv4AvroPwtNx2jwliJV-doLUIuUw2vfiLEIAuPmbF-brlEcKMr-qC1RFJBhPBV_R6_4rg8wfqWMCAULxUkQizpQbxvF9lrjiG4cyC0hrgeMf56kUC_ZerKBIYFILxYJ4jZiFiqL5uipmVCnTfXdYsNGorNhdyQk32L3WIvNr-hsUOmwdSC9gPAzjtSIqqLhQM8rZvdiQCb1JEKhcm0ZZUEU1W6TxipDmxRGSgEi_VyhD8x1SEmeaBfwBgQtmFjJVGmeDh5mnAb-kPp7_AYXNuaNJWRntLALWiTZsHaf3CwoBZdB5XmgAKYXE5kOsIgJOUCJFRMAdOMHBIQWERgRiglFFKGTDassD5g-AZZoNz39mIgWf5Er06plXAByhjG2u8oolXKZjMOug7ckdqU8FsoEgkLK7NtypOHWl7zAv4FEkws4dAx7Ki1JCwQcG9HS4gdhQ9KgrQrHCQlAfKlLepM8hfwE8qwRU52WJgDJtXcqKgJfALzAAAKrGcGeLpMpKfGhAF6ZLRJgrYmoPGg3ySvkGrmJA4jyUesOIWOJUmRx7QLZxgr0ngHpJ-W0LFyBAO6gCAs5xoDDpohSPI8gGei9z0jMQ0-uY2eQox2AdQ53HOjCADmzB_HqQ6KSp8B3l9M5YNSsLhmu9JM5IOHYBOEChLYCvJT2XhGvqZAy5GMMpa4LFCWP9PvvZxmiQi0uBBW4cjVgTSnL9SRrPhVEPlhmyM1KFPxpRDIWd541R2KCu0VYVQbnZBx5mwiUlpwxakesTosc-kJwXQm5McSvTABxzqddSCdpe8JSaCpTBJoMUMD8WKvqeKANciyQ4ZYFHiovOMWvdbXT26ezSfX9-9nF6de6-m918nl2fu2eXv_w6vZq--3jeLo13cufqsEVoEnIsDAmqCJNKoai-lEqLhryDAHgPZrfVI9UITS2ADRAilmClRvBwj_ZVSFw5lA1qvgGSPL0HuaDVtmV9lbFOb7Hahb0xKo34SmBtJzjkunZddRmhQcn_Qi5sp8Vpu9NGFFDvf4dudcw9ipJqxniWk8N3kPsFg0Kb1XM02l3a26ji5KkgsPUfgK7TKiyXQo8HkzaIdwwaPwOZK0IhpR6T3jPpwCTZ6VSBzq5-GTXpH76kXwmQ9wsDw7qBVOgsjZhdcDG8cSWHk7aJ6hUbHg3ZgfHBxP54Gz02Q-0p-ZWgM80g_m-QA_1zGVHP4BJZkClZyY0DBMFuAULS_hIQQZNLYe3F4scP1w4rnq4-XD9F5jyOwyokQSBHXBgAd1-Vz2mgnodoKBb6BYyCzl2Mgmo6gz0rmZaSFhztXUijBx122oh2NNHpUM4Rk2c27ybjeaaNAft7amC438CwMHBCohUDFuhVCwVIYUGABpVRG2BPCaV84HGlRhlRtnJBFnawCpL0mtJQpo-mQ-n1hhah4ndaWfm23aHDusPH_aJE1aVoydwg2hssjL8U2BRg7hHewffE7FX_Oh78E_YgHNYyTmeEwg5OOTVpfbqJ3poXQMIQXK2_t9WIX7HDR6dWmPJnNeTBb4u7vyVBQW4r93ul9vTd9gUzIeyjeeriDgGW7DHUoGklmWyb3m4fH_Ok521RO6eNy-pOGF0QHLYtK5UfDsDsUe5obeC1w35oGDgcHjts2DBwNICBgRkwUGr2rAQPekZP5NvRv3rHC4e9HvYO4XL4pjeEy9FR74cCXgQZZGDlzsVSRhVOWPF70LlzWGwbL_NZO6fP7SAPxLQs1AJ2R3LlGhYhL8eVXWEZrlNpgsmknRf6i820SXmP3Gx37GQzYG7zoa8Vvm1jx8tOXqYKm73swhbtfZFfc1Nw-3iS63nsdGwK8XuC2bm2K4l0alk9qFJKp2zt_9Deu_l4Y89j9kwOyyb3fWm-TpS7U6BTEdLXFWZXx2Bnv7fnKPLcWQIXyTjK6RM0N-3HzWnTrKcYAF8CsypdPWxUVngq3n4IF5t6RDDuzSFRtmjFkG0t6JXDIaCMHLMbZPwQhnahGE-FTHfB_6PXx39b6O3bt9A30DODQW6JvkJBmFjFTd6R-wPChR5jscd71S7C6iHptjtmdaWvUbtTCpSXr0Tkmxf59HxS51k_6tywgHb1n3FlIWvDz7lBqp76QOtwsoEmU1mo4fzDvXvcneHnLNqQJgAV5Iuqw5VvXVUcBLlN6nF7f7JP0OQ7KLJj19UyXXslqxkKTFhGkm5tYOh7u7pHkDBlcAKXcemoTecJOzgoOae67zJnbZAsRIBUds8QOT5PK_GUE8kLD3fH4zF-7TZfIszu2fYpvMbB3V0ZOdAxgkUnw_3TVXVXePhUuOSgRmOH9fkd1p0022nn0f4bxinloG8E97bhzGd8IygXxXTI79NMODntmL3VqXXM1hQTiDvBxiTewvtWfqDBi10lBvAO3rQc-j4jQziLFH_xaEUPnjc8PM6GMBwn2vw5pNUF1afewcHwDevy1AtO1cp9M2DdLmz-dJc2X77wuyFfzelvJKGcV3R6nhfCS3vogRdwDonuW49OPg6MWhuHbm89fqV__wPolNZF)) +- [Heterogeneous Lookup](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_map/heterogeneous_lookup_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJy1WXtz2zYS_yoY3VwrxdTLjptGlnVVHN9E05zds51mOnGHhkhIRE2RLAFa0Wl8n_12F-BLopz-c0rGIol9728XC2rbUkIpGUeqNfqybUm_NRo6rZBHy4wvRWvU8jKft5yWirPUw_v-q_uIvWIXcbJJ5TLQrO112PHg-LQLf35w2NWvs_ezKbu4vvnl-mZ6N7u-6iEDMX2UnoiU8FkW-SJlOhBsmnAPvuyKw34VKVrDjnsD1kaC-5Zdu291zkjKJs7Yim9YFGuWKQFipGILGQomvnoi0UxGzItXSSh55Am2ljogVVYOmcN-s0LiueZAz4EjgbtFlZJxXZiOn0DrZNTvr9frHieze3G67IeGWPU_zi4ur24vu2B6wfYpCiG8LBV_ZjIFx-cbxhOwzONzsDfkaxanjC9TAWs6RsvXqdQyWjpMxQu95qkgOb5UOpXzTNeCl9sJ_lcJIHw8gsBNb9ns9r7F3k1vZ7cOyfk8u_tw_emOfZ7e3Eyv7maXt-z6BpJ19X6GqYK7f7Lp1W_s59nVe4cJCB2oEl-TFL0AUyWGVfgmhrdC1MxYxMYslQhPLqTHchixZfwk0gjcYolIV9IADoz0SU4oV1JzTc_2nCNV_fvoPvqbjLww8wUbe5kX9xWyeO6KJz0vCya7FD7vL7LIQ6k8nOytSS1SruN0f0Vpv6-zJBS1JR2kmdJ9XwBcwn7K124qFiIVgDDX40r3gmbyJ_DCfRIeqGomCWKlqwRVEglrqeAretx_ZSD1E0E9EOBAvBSRiDPlhnH8mCWu-MohPwLiUQDwp3kqxYK9FyuIrgaXhWJBvEawQSogY6u6KGZEKVM2Dxjp0agM9UNZzHdYdlYj82vyGwQ6bB1IL2A8DOO1ohqumFBTytmj2JAKvUkQYVwbhC8o4Jot0nhFEPHiCKtXpN8r5CF6h4TEmWYBf0LEgZqFTJVGarAw8zQAj8TH8z8g7nnRNwkrvZ1F0C6iDVvH6eOCXEAeNJ4XEoBLYUfywVcRQTErkWItwxUYwcEgBZ2IKt84owpXSIeVljvMnwBqRA33f2YiBZvkSvTqkVcAHKGMbq7yjCVcpmNQ66DuyQOJTwWWMQQQTXpow5WKU1f6DvMCHkUi7DwwaCB5Ukoo-NC6PB1uwDdkPemKUKwwEVQmioQ3yXPIXgDPKgHReVoi6HU2L2VGwEpoDLAAAKr68GCTpMpMfGhAF4ZLRBgrarEeFBbEleINTYZB9-Wh1BtGHSFOlYmxB10yTrAPCOiZJPy-hTsHoB1EQefMcaAw6KJkjyOIBlov825lIKbXMbPBUY7BOro6j3VgAB3YhPl1J9FIk-EHiumD0Wq2BHTXWgkFFqmEYxWAAeTaAhqtpPuyU5o8GUUu-nDOmmBxxli_z362PhrkYg-3wI2jEWtCSS4_SeO5MOJBM8Pmia3CH43Ih0LPy8rIbRDXqKuKoFztEw8z4ZKQcwalyPUZtcc-NDkvhNiY5FbIABxzqddSCZoLeEpFBcJgOiABzI-Fir6nlgGmRRKMssAjwUXlmE3qPrr4dHHtvr-8-Di9uXTfze4-z24v3Yvrf_0yvZm--3jZLpV3cuPqsEVoEnIsDAmqCJNKoii_FEqLhryCAHhPZkzqkWiEphbQDRAitsFKjeDhHg1E2LhyKBvUfAUkefoAckGqLcv6LmON3mK2C31jFBrxlcDcTnDJde225zJCg5L_gVjYSovTdqeNKKDa_w7N6phrZCXRjPEsbw7fQewXDBJtdsvRaH_nbaOIs11G6NZ_ALrOq7BcCj0eTNrA3jFo_AzNXBEKKfQY9J4JBwbJklMGOvvyZdQkf_gt-UoAv18oGNYVpEJnacTshovujSsxnLSNV6_Y8GTIjowNxvfn--i5GWq7za8EnSkG8X-DHMify4hqBrfIopmSllw5QBD0FiAk6d8CIkhyya2DWPz44dZhxd3Nh9tdZM7jOKxCEhhyxIUB9O6b8j4N1MsQDcVCfwOjIHMfoyCaTlAvcqYlpwVHex_SaEGHnTeiHVV0OhRzxOSFjbuJeB5po8B-dhUMDysYFgrOiLWiwAK9qqEAKWwIUKAyagPsKaAUDzxn1FpGlK1c4IUJVkGQXlMYyvAROaReb2gTKj7nlZ1v2x06rDt8PsxKrbpkLTs3sPYGC2MvOTYFmHuEd7A9MbPqf08Hf4cZhMNexmmEL_QgybkJ6-4QvTUPoAmDc7X63lY9fsWOn51aYsqPlZA7vy2u_hIHObmtXB_kOlB322-oCWGO5qmLEwJs2WPIQdNOMtk2Pd0-P-dBz8uidowal9mdMPpCcNiyrGR-OAC1J7mhtYXXDvuxYeF4eOqwYcPCyQAWBmbBQKnZshI8aBndkW0nP_ROFw57Pewdw9fxm94Qvk5Oej8W8CLIYAdW7lwsZVRsLCv-CBJx7MOZ3M2Pre0cWUjjxRkN9cXqdpA7YQggDzAZyZVrOghZOK5MhKWrTqUAJpN2nuQvNsom3D0ysd2xxGbBXOZLv1d6bRurXXbyFFU62bdN2KK-L_L3XBVcPp_lcp47HRs-fAlgptZ2JYhOLaJH1XbSKcv63zR3Nx9t7FnMnsdhy-S-L80rhXIyhVYqQnolwuzOGOzNegeOIS-dI3CDjKO8dYLkplncnDTNXooO8CV0VaWrB43K7k7JOwzfYqBH9OJcDoGySSuWbFlBnRwPAWVkmB2O8e0V6oVk7DKZyoL_J69P_zLT27dvoWagXgaDXBO9OgI3MYubvBoPO4SbPPpij_aqXbjVw4bb7pidlV4h7ZMUKC8ficg3D3LynKjzoh31vrCAcvVfMGUha8svmUGidm2gPTjZQJGpLNRw9uHeI05m-KaJhtEEoIL9ompw5TVUFQdBrpNq3F6fHWI08Q6K6Ng9tQzXQc5qhALjluGkS-sY2t6uzgcSSAZn8DUuDbXhPGNHR2XPqc5c5pwNnAULNJX980OOz_OKPyUhWeHhZDwe4ytq8xbCTM62TuExLu5PZGRAxzAWlQzXuzvqPvNwl7nsQY3Kjuv0HdadNOtp597-A9Yp5CBvBNe24My7d8MoFwU5xHc3Ek7edsxcdW4NsznFAOIU2BjEe3jeyg8z-GV3iQE8gycth97NyBDOIcVvFa3oyfOGx6fZEJbjRJsfMlpdEH3uHR0N37AuT73gXK3cNwPW7cLgp7s0ePnC74Z8NadfN0I5r8j0PC-Eh_bAAw_gDBI9tp6dfB06am0dqr31_Dv9-x-LDrkI)) ### `static_multimap` `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/eJyVVgtv2zYQ_is3DUXlVJYfWFbEcQJ4qYcZK5whTlsUdSHQFG0TkUiVpOx6hv_7jqTkyIuBdQlggffid3cfj9wHmmnNpdDB4Ms-4Gkw6EVBRsSqJCsWDAJapiSIAi1LRe26czEXcAF3stgpvlobCGkL-t1-r40_v0Yw_Th5NxnB3f3DX_cPo8fJ_TS2Ds7pPadMaJZCKVKmwKwZjApC8VNpIvjIlEUD_bgLoTWYB5VuHrSuXZSdLCEnOxDSQKkZhuEaljxjwL5TVhjgAqjMi4wTQRlsuVm7rao4Dg58roLIhSFoT9CjwNWyaQnEHKHbv7UxxaDT2W63MXGwY6lWncwb6877yd14Ohu3EfrR7YPIsLyg2LeSK0x8sQNSIDJKFog3I1uQCshKMdQZaZFvFTdcrCLQcmm2RDEXJ-XaKL4ozUnxapyYf9MAy0cEFm40g8lsHsBvo9lkFrk4nyaPf9x_eIRPo4eH0fRxMp7B_QM2a_puYluFq99hNP0Mf06m7yJgWDrcin0vlM0CoXJbVpb6Gs4YO4GxlB6WLhjlS06hphGs5IYpgWlBwVTOPeEQZOriZDznhhgne5Gc26ozF3PxMxc0K1MGQ1pS2dHWhSZ5mRmekyKm5fr21MysValNh8pSmNgqX6hStsFdkg2jRqrzJtwwRVDrw2AKSS05b6-x1Qx5d15rFBEa65TH_wbLJTaQkdyJuTDITi7CjeRpay72WATkui3gE9slZlcw5OMN8sVcP6s2JCuZV9Yqqzx6sLwwu8QuNbOpsMzR-gbaPRel4e9NveBoXBlaU4q9MpYXoE06GGj-NzrCFE0uu6-73W5l1ungoEBLVVKD_Ya6Wf5Q9rrdCG1BZxK1Pod2z9LCfbSjgUNic-g4NFXUGpOOYSqNpyG1h5KjMR4HupZoAk9Cbm3UrR0DWYZF0UwZxOj2xaA6qgJa7uKhySRJYUksHewwuOy-il26SDjM8pRxw7qwUaNyt9i3wrULsBwX0I8q52Px9y_bcDg1ctH251pwOFSF9XQaDE4IPPRBCsLVeWy3YHU6nLYa_UHSYQEJ1MS1eTuzGMYMp5wtE6xtN7YSkzM493RcOY_jVRzBfo_1xBz2vaiHnziOASX9S0sELwa3qJSHZgLHAxH6mtXynDyx5MWRGyKrb8NuK_pxY5-KJWjYqv28bMFWeMJq2ZevkCRVOZOQlHYat2CPo9uUSsB_lXbP4RWEU-hAvxUBP1zDoVHkieMdngDkoNvcHk_pWIt0sUZ2gnl2hqfoKrBMpIj_OeKs2S0lF8yx2dYdbLl9F_q_vL66ujrUzaqbqdeyzFLs6YZBv26pvXsaeM4TrE791u2WGJksOQLzaV83_Wo2hU3D56ROpC63CLpNVq4ZfcJgxPi88B60M8fe1XjB_QjUhZTZ7bNPE6O7LbxC_w98x1i1WQMvPlv4cuf6e8S7RELaCw4cmdzEBFHmiZPjpKxxO-KGL6I3N6wQ4Bxl9aZ8CWEj2g34BKGaPW4sY2QDwyE-BWYlpXiF_wQjRIhPAhQ6e6_FW5djuzx0H7AucX27eh8XFcFk12d3GTfHRR_S0r10cLq4g6IjP2XxoWNIZolbA_EgPNHPbXTwKVdH0V4uKMFHqX3l4btPPb9dA7GhtNe_LHuoloXxD9ugjfFu6Js3vbfQJoqub3SevO1Cu41XmMEfg1uxtJ2RfOFeuxlfNGJSSjMUbvzTFAXYB_EUHKJaj7PhRI9MDA5f3f8_JDbX8g==)) +- [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/eJyVVo1u2zYQfpWbhqJyKku2sayokxjw0gwzVjhDnLYo6kKgKdomIpEaSdn1DL_7jqTkyK2BdS0QQ7wffnf33R33gWZacyl0MPy8D3gWDPtRkBOxqsiKBcOAVhkJokDLSlH7nVzMBVzArSx3iq_WBkLagUFv0O_in18jmH6YvJ2M4fb-4a_7h_Hj5H4aWwNn9I5TJjTLoBIZU2DWDMYlofhTSyL4wJRFA4O4B6FVmAe1bB50rpyXnaygIDsQ0kClGbrhGpY8Z8C-UlYa4AKoLMqcE0EZbLlZu6tqPw4OfKqdyIUhqE_QosSvZVsTiDlCt__WxpTDJNlutzFxsGOpVknulXXybnJ7N53ddRH60ey9yDG9oNjfFVcY-GIHpERklCwQb062IBWQlWIoM9Ii3ypuuFhFoOXSbIlizk_GtVF8UZmT5DU4Mf62AqaPCEzceAaT2TyA38azySxyfj5OHv-4f_8IH8cPD-Pp4-RuBvcPWKzp24ktFX79DuPpJ_hzMn0bAcPU4VXsa6lsFAiV27SyzOdwxtgJjKX0sHTJKF9yCg2NYCU3TAkMC0qmCu4JhyAz5yfnBTfEuLPvgnNXJXMxFz9zQfMqY3BNKyoTbU1oWlS54QUpY1qtR9-qZSThhilipBq1JWatKm0SKith4vU5UcY2eH-6YRRtz6toLCpDhp2XGkWExowU8bewuMRSMVK4Yy4M8pCLcCN51pmLPYaLrLapemK71OxKhsy7QWaYq2fRhuQV88JGZIVHC1aUZpfaT80E0onljsA30O07Ly17r-oPjsq1olWlWBVjGQDaZMOh5v-gIUxR5bL3stfr1WpJgiMBNVVFDVYWmrL49uv3ehHqgs4lSn0M3b4lgPvRruAOiY0hcWhqrw0mHcNUGk84atuPozISn64lqsCTkFvrdWsbPs8xKZopgxjdvehUR7VDy1Jsj1ySDJbElte2_WXvRezCRWphlKfcum4SG7UyN8K6la5cgOm4gEFUGx-Tv_--DIdTJedtf64Eh0OdWE-n4fCEkNfeSUm4Oo9tBFamw2mnVR8kHSaQQENcG7dTi-GO4TyzaYK1rcZWYnAGJ5yOa-O7eBVHsN9jPjGGfT_q408cx4Ang0tLBH8M7qMWHtoBHBsi9Dmz3TkcFuSJpa4PsXpp06zXyOlR2OtEP6rqw7DkDDuNlT9bsBV2V3P2-QukaZ3KNCSVnbkd2OOANpUS8F9p3XN4AeEUEhh0IuCHKzi0EjxxnEP2I__c5bY1pWMsUsUq2TnlmRmeoqvBMpEh_mePs3allFwwx2Sbc7Cp9hUY_PLyzZs3h6ZQTSH1WlZ5hvXcMBg05bQbpoXnPLma0EfuttTIdMkRmA_7qm3XMClsKz4HdXLqYoug12bkmtEndEaMjwu3nZ03diPjGvsRqAsp89GzTRuj2wleoP8HvqOvRq2FFx8nfLlz9T3iXSIh7RoDRyY3LUFURerOcUo2uB1xw--8ty-sEeAMZc2lfAlhy9sN-AChnjtuJKNnA9fXuPBnFaW4qH-CMSLExY-HTt9LcbdyLJeH7h02KW52qLdxXhFMfnX2lrv2qBhAVrn3DE4W1yg68hMWnzOG5Ja4DRAPwhP93EUHH3Ldinax4Ak-Pe1bDl936vmFGogNpf3BZdVHsSyNf74GXfR3Q1-96r-GLlF0faOL9HUPul1cXwb_GLyKZd2cFAv3ps35ouWTUprj4cY_QPEA6yCegkPUyHE2nMiRicHhi_v_LwFkzNY=)) ### `static_multiset` @@ -271,4 +271,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::experimental::roaring_bitmap` implements a Roaring bitmap following the [Roaring bitmap format specification](https://github.com/RoaringBitmap/RoaringFormatSpec). #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/roaring_bitmap/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJy9WA1v2zYT_iv3qsAgN7aVpB_ZnI_NjdPNWF97sN0VQ1MIlETbhGVRI6k4XpD__h5JfSbK2nV75wCxRR7vnjs-dzzqzpFUSsYT6Qw-3jkscgZHXScmySojK-oMnDCLiNN1JM9EqJ-959cJPIdLnu4FW60VuGEHjg-PX8Hk1_FoPITL6eyX6Wy4GE8nfS1qxN-xkCaSRpAlERWg1hSGKQnxK5_pwq9UaBxw3D8EVwtcO_nctdM5NVr2PIMt2UPCFWSSohomYcliCvQ2pKkClkDIt2nMSBJS2DG1NqZyPQYO_JYr4YEiKE9wRYpPy7okEFVC15-1UunA83a7XZ8Y2H0uVl5shaX3bnx5NZlf9RB6uex9EmNgQdDfMybQ8WAPJEVkIQkQb0x2wAWQlaA4p7hGvhNMsWTVBcmXakcENXoiJpVgQaYawStwov91AQwfSTBwwzmM59cOvBnOx_Ou0fNhvPhp-n4BH4az2XCyGF_NYTrDzZqMxnqr8OktDCe_wc_jyagLFEOHpuhtKrQXCJXpsNLIxnBOaQPGkltYMqUhW7IQCgLBit9QkaBbkFKxZZZqCDIyemK2ZYooM_bIOWPKu06uk2csCeMsonAWZiH3BCcCNfoBU1uS9sNsffFIJlMsZmrvKUGYkv11ml481BQRT6rIC_FfRJcXT06yRLVPqn1KfWugIaDWIpPKi-gNOuLf0FBx0V-3icR8hYSI2yezhGHsJInrKupymvlyLxXdNpYvkQ-UNMcYbxnUvElWjSFrydjxnlsi_2ASbI0K_CCLNz69JUgFimG304FgdAkjusVNxGAoilsp9dbnKdXcLSSn1oLcjDnfZKlRDMNfxrIqFuPEZnZuCXaYL5xgmuw4vDjuoSKwygyRkPQUXr-sDYObcqFMniExt0R1YCn4VqMx-j_OLKQ3RvqtEZkjcz-5Os8lJvoK6Z8FfawlXkO2eKrWdDDDUy4ZRm1f0hoLRLgBZv3X7hbOop8orzKBfNdzIRcCI45jMouRpjAhWxrvu9plDKQyQksex3ynU8hs-MCY6MFH66wucjxTIktkP2CJmXRtlDpf45AXxDzwXh2dnJDoO0-DiIgiXquxzmMo_x6ORyCKTc-Z9vplhcPS4x_E8fql12quU5L4BzykKCzqPM7Pgnzrt2SDOZLq4ge90eX7y6k_mn6YvJsOR_5sOpyNJz_6i6v5YjRcDM-nE31IBFhhqSoTBWsj5n4aY85h0cBilCB74Ge6X-BvzOGA89hy0UXUg4HNdyQdJuo3ea74mlN-StQaod-hViAZWlrRhOpc9jd0L-EcPn5yO9C7AFuaBoNGbTsrTIJRAJr5xog-PrA_wHqpjSMCJn2JGP2bYkkXarMZVtoXx7666BSKADwPLrFwoYe_ZxRTzOApjpuKDHfIyvuCEVWeFCr-MfrNroaj_171t9EzPdTTY4UZ40IekDanDPTTQlq74LaIwQajfXiKX2dwdKg_-vfBuXmoxQWMun6aybUfENzhTafUfd8wgoqN0krbGfZr9vfBwebPdL5AHn9W70lN77dfprdNpy2LTxCsLZ6uURrQFUvcTteaoEnkdgrl90Bj3SJ9PRlfv_wLZGytB4aK_w8mYgGquPjItC18n6Olce_ztNRilpa3ZnuzOLa7jc_fFc9fseOfszVs2jo6_Bu2Hogc3h7n2irZVpFXDZEnETObtOwBUKYT91ir-BO4h7ff5gvgANjfzAyzpV-aGXcVR7AHD30iJRXK1f2zNoPnSUQT5S8JCpc1vosdXP4btghIH0pl8cL4FCjyS1vdg7v7wr7-0g_6G5NiiqZMpuhnm5t5E2tG3YdHVTcX4nIwQKoTsbe2MNXd_2ipPuY2R6XocuGlWRFSIeDsDF14S1DM3Ly0HA7o4Yd29JhZh5GIc-y5LyYoZui-dONHquxtVLI_aHmUGm16BDlilFVtu_1tJh_52KnCM8T-L9RVx_TKW-y0sfY0z8AyegVLqsuD6dxTliQ0aiFNsFd4bgfZckmFW6KpGZ9R7LyNV7ixPLeu50yccY8iV1CcogKvisoPiVRn4ZqI5xdugUWQnZ9yI2PmXWuur-sY7hBStGE4Vx3GXFK3hiQvvfl9onQ87-zLCFj6YrWngm2RvyRGBI07SNWvNMddu8NfCvu0tvG2W6ofCy7BhJbYLceRThE8gvTrBqQcS2r71ilpkjdajcarxXsCdg81cyV-23t4fn-oE6BZJnQreFGBsBXC8E5HX3O5ZuoN3lh6-Y2l5hFZ4dr8TmLRa_mHl3FrQj5Zg7oVimK67mVxc9LR04asZfFkAPU60-fiAr-SOS_joMf50n1ktI4jB4b1odYd6NrH1P7u3nIyz_uGGZww2a9boS3RWk0P_YwtI7wSm5a-6OR1F--PxrOywtWbcdCU8iMmEHfrqtPSTZmFoX4j0_yc2za_vU8qlR9gmWu_yBXFutD-zVcp_Kva7GHV1NZ-iXMKitgqjshtFc8tDPL6bSsahonE6ZqYkRzCg0Je29BC4ns4hAEc4eQzfThWxsojo_0u17pf-WUNWWDYiJWDYRkw7yMb975rp2oL8s-18wWXwc7DhY8Oqtw7608SsaWhqtN19PtRrJSiet_rJDdheHT8KjvCaQsMJ50eKjwPDw6OTqBHRLg-l1v_5BB6PSytCv8p3R5EvZhsA_OGOGZBTWcYhjEO3tiXujigubZx7rvFPFbpxjzWK-f-k_n7H7m6k-4=)) +- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/roaring_bitmap/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJy9WA1v2zYT_iv3qsAgN7aVpB_ZnI_NjdPNWF97sN0VQ1MIlETbhGVRI6k4XpD__h5JfSbK2nV75wCxRR7vnjs-dzzqzpFUSsYT6Qw-3jkscgZHXScmySojK-oMnDCLiNN1JM9EqJ-959cJPIdLnu4FW60VuGEHjg-PX8Hk1_FoPITL6eyX6Wy4GE8nfS1qxN-xkCaSRpAlERWg1hSGKQnxK5_pwq9UaBxw3D8EVwtcO_nctdM5NVr2PIMt2UPCFWSSohomYcliCvQ2pKkClkDIt2nMSBJS2DG1NqZyPQYO_JYr4YEiKE9wRYpPy7okEFVC15-1UunA83a7XZ8Y2H0uVl5shaX3bnx5NZlf9RB6uex9EmNgQdDfMybQ8WAPJEVkIQkQb0x2wAWQlaA4p7hGvhNMsWTVBcmXakcENXoiJpVgQaYawStwov91AQwfSTBwwzmM59cOvBnOx_Ou0fNhvPhp-n4BH4az2XCyGF_NYTrDzZqMxnqr8OktDCe_wc_jyagLFEOHpuhtKrQXCJXpsNLIxnBOaQPGkltYMqUhW7IQCgLBit9QkaBbkFKxZZZqCDIyemK2ZYooM_bIOWPKu06uk2csCeMsonAWZiH3BCcCNfoBU1uS9sNsffFIJlMsZmrvKUGYkv11ml481BQRT6rIC_FfRJcXT06yRLVPqn1KfWugIaDWIpPKi-gNOuLf0FBx0V-3icR8hYSI2yezhGHsJInrKupymvlyLxXdNpYvkQ-UNMcYbxnUvElWjSFrydjxnlsi_2ASbI0K_CCLNz69JUgFimG304FgdAkjusVNxGAoilsp9dbnKdXcLSSn1oLcjDnfZKlRDMNfxrIqFuPEZnZuCXaYL5xgmuw4vDjuoSKwygyRkPQUXr-sDYObcqFMniExt0R1YCn4VqMx-j_OLKQ3RvqtEZkjcz-5Os8lJvoK6Z8FfawlXkO2eKrWdDDDUy4ZRm1f0hoLRLgBZv3X7hbOop8orzKBfNdzIRcCI45jMouRpjAhWxrvu9plDKQyQksex3ynU8hs-MCY6MFH66wucjxTIktkP2CJmXRtlDpf45AXxDzwXh2dnJDoO0-DiIgiXquxzmMo_x6ORyCKTc-Z9vplhcPS4x_E8fql12quU5L4BzykKCzqPM7Pgnzrt2SDOZLq4ge90eX7y6k_mn6YvJsOR_5sOpyNJz_6i6v5YjRcDM-nE31IBFhhqSoTBWsj5n4aY85h0cBilCB74Ge6X-BvzOGA89hy0UXUg4HNdyQdJuo3ea74mlN-StQaod-hViAZWlrRhOpc9jd0L-EcPn5yO9C7AFuaBoNGbTsrTIJRAJr5xog-PrA_wHqpjSMCJn2JGP2bYkkXarMZVtoXx7666BSKADwPLrFwoYe_ZxRTzOApjpuKDHfIyvuCEVWeFCr-MfrNroaj_171t9EzPdTTY4UZ40IekDanDPTTQlq74LaIwQajfXiKX2dwdKg_-vfBuXmoxQWMun6aybUfENzhTafUfd8wgoqN0krbGfZr9vfBwebPdL5AHn9W70lN77dfprdNpy2LTxCsLZ6uURrQFUvcTteaoEnkdgrl90Bj3SJ9PRlfv_wLZGytB4aK_w8mYgGquPjItC18n6Olce_ztNRilpa3ZnuzOLa7jc_fFc9fseOfszVs2jo6_Bu2Hogc3h7n2irZVpFXDZEnETObtOwBUKYT91ir-BO4h7ff5gvgANjfzAyzpV-aGXcVR7AHD30iJRXK1f2zNoPnSUQT5S8JCpc1vosdXP4btghIH0pl8cL4FCjyS1vdg7v7wr7-0g_6G5NiiqZMpuhnm5t5E2tG3YdHVTcX4nIwQKoTsbe2MNXd_2ipPuY2R6XocuGlWRFSIeDsDF14S1DM3Ly0HA7o4Yd29JhZh5GIc-y5LyYoZui-dONHquxtVLI_aHmUGm16BDlilFVtu_1tJh_52KnCM8T-L9RVx_TKW-y0sfY0z8AyegVLqsuD6dxTliQ0aiFNsFd4bgfZckmFW6KpGZ9R7LyNV7ixPLeu50yccY8iV1CcogKvisoPiVRn4ZqI5xdugUWQnZ9yI2PmXWuur-sY7hBStGE4Vx3GXFK3hiQvvfl9onQ87-zLCFj6YrWngm2RvyRGBI07SNWvNMddu8NfCvu0tvG2W6ofCy7BhJbYLceRThE8gvTrBqQcS2r71ilpkjdajcarxXsCdg81cyV-23t4fn-oE6BZJnQreFGBsBXC8E5HX3O5ZuoN3lh6-Y2l5hFZ4dr8TmLRa_mHl3FrQj5Zg7oVimK67mVxc9LR04asZfFkAPU60-fiAr-SOS_joMf50n1ktI4jB4b1odYd6NrH1P7u3nIyz_uGGZww2a9boS3RWk0P_YwtI7wSm5a-6OR1F--PxrOywtWbcdCU8iMmEHfrqtPSTZmFoX4j0_yc2za_vU8qlR9gmWu_yBXFutD-zVcp_Kva7GHV1NZ-iXMKitgqjshtFc8tDPL6bSsahonE6ZqYkRzCg0Je29BC4ns4hAEc4eQzfThWxsojo_0u17pf-WUNWWDYiJWDYRkw7yMb975rp2oL8s-18wWXwc7DhY8Oqtw7608SsaWhqtN19PtRrJSiet_rJDdheHT8KjvCaQsMJ50eKjwPDw6OTqBHRLg-l1v_5BB6PSytCv8p3R5EvZhsA_OGOGZBTWcYhjEO3tiXujigubZx7rvFPFbpxjzWK-f-k_n7H7m6k-4=)) \ No newline at end of file diff --git a/benchmarks/benchmark_utils.hpp b/benchmarks/benchmark_utils.hpp index c36c7f018..a9bd690d7 100644 --- a/benchmarks/benchmark_utils.hpp +++ b/benchmarks/benchmark_utils.hpp @@ -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. @@ -21,8 +21,7 @@ #include -#include -#include +#include #include @@ -75,13 +74,13 @@ struct lazy_discard { }; /** - * @brief An output iterator similar to `thrust::discard_iterator` but prevents the write from being + * @brief An output iterator similar to `cuda::discard_iterator` but prevents the write from being * optimized out by the compiler. */ template auto make_lazy_discard_iterator(OutputIt it) { - return thrust::tabulate_output_iterator(lazy_discard{it}); + return cuda::make_tabulate_output_iterator(lazy_discard{it}); } } // namespace cuco::benchmark diff --git a/benchmarks/bloom_filter/add_bench.cu b/benchmarks/bloom_filter/add_bench.cu index 433b6acde..b07c285d4 100644 --- a/benchmarks/bloom_filter/add_bench.cu +++ b/benchmarks/bloom_filter/add_bench.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -24,8 +24,8 @@ #include +#include #include -#include #include #include @@ -67,7 +67,7 @@ void bloom_filter_add(nvbench::state& state, state.skip("num_sub_filters too large for size_type"); // skip invalid configurations } - thrust::counting_iterator keys(0); + cuda::counting_iterator keys(0); state.add_element_count(num_keys); @@ -108,7 +108,7 @@ void arrow_bloom_filter_add(nvbench::state& state, nvbench::type_list // configurations } - thrust::counting_iterator keys(0); + cuda::counting_iterator keys(0); state.add_element_count(num_keys); diff --git a/benchmarks/bloom_filter/contains_bench.cu b/benchmarks/bloom_filter/contains_bench.cu index e472cf1ff..7a4549524 100644 --- a/benchmarks/bloom_filter/contains_bench.cu +++ b/benchmarks/bloom_filter/contains_bench.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -24,9 +24,9 @@ #include +#include #include #include -#include #include #include @@ -73,7 +73,7 @@ void bloom_filter_contains( state.skip("num_sub_filters too large for size_type"); // skip invalid configurations } - thrust::counting_iterator keys(0); + cuda::counting_iterator keys(0); thrust::device_vector result(num_keys, false); state.add_element_count(num_keys); @@ -119,7 +119,7 @@ void arrow_bloom_filter_contains(nvbench::state& state, nvbench::type_list keys(0); + cuda::counting_iterator keys(0); thrust::device_vector result(num_keys, false); state.add_element_count(num_keys); diff --git a/benchmarks/hyperloglog/hyperloglog_bench.cu b/benchmarks/hyperloglog/hyperloglog_bench.cu index 1ada00396..e58b20a32 100644 --- a/benchmarks/hyperloglog/hyperloglog_bench.cu +++ b/benchmarks/hyperloglog/hyperloglog_bench.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -24,8 +24,8 @@ #include #include +#include #include -#include #include #include @@ -40,7 +40,7 @@ template // Casting is valid since the keys generated are representable in int64_t. using T = std::int64_t; - auto cast_iter = thrust::make_transform_iterator( + auto cast_iter = cuda::make_transform_iterator( first, cuda::proclaim_return_type([] __device__(auto i) { return static_cast(i); })); auto set = cuco::static_set{n, 0.8, cuco::empty_key{-1}}; diff --git a/benchmarks/static_multiset/retrieve_bench.cu b/benchmarks/static_multiset/retrieve_bench.cu index 427cc13ed..aed508e06 100644 --- a/benchmarks/static_multiset/retrieve_bench.cu +++ b/benchmarks/static_multiset/retrieve_bench.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -22,6 +22,7 @@ #include +#include #include #include @@ -54,7 +55,7 @@ void static_multiset_retrieve(nvbench::state& state, nvbench::type_list output_match(output_size); - auto output_probe_begin = thrust::discard_iterator{}; + auto output_probe_begin = cuda::discard_iterator{}; state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { set.retrieve( diff --git a/benchmarks/static_set/retrieve_bench.cu b/benchmarks/static_set/retrieve_bench.cu index be6902e6c..1baa00684 100644 --- a/benchmarks/static_set/retrieve_bench.cu +++ b/benchmarks/static_set/retrieve_bench.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -22,6 +22,7 @@ #include +#include #include #include @@ -54,7 +55,7 @@ void static_set_retrieve(nvbench::state& state, nvbench::type_list) auto const output_size = set.count(keys.begin(), keys.end()); thrust::device_vector output_match(output_size); - auto output_probe_begin = thrust::discard_iterator{}; + auto output_probe_begin = cuda::discard_iterator{}; state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { set.retrieve( diff --git a/examples/static_map/count_by_key_example.cu b/examples/static_map/count_by_key_example.cu index dcecf4e4e..a979e0bb6 100644 --- a/examples/static_map/count_by_key_example.cu +++ b/examples/static_map/count_by_key_example.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. @@ -17,6 +17,7 @@ #include #include +#include #include #include #include @@ -113,8 +114,8 @@ int main(void) thrust::device_vector insert_keys(num_keys); // Create a sequence of keys. Eeach distinct key has key_duplicates many matches. thrust::transform( - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(insert_keys.size()), + cuda::make_counting_iterator(0), + cuda::make_counting_iterator(insert_keys.size()), insert_keys.begin(), [] __device__(auto i) -> Key { return static_cast(i % (num_keys / key_duplicates)); }); diff --git a/examples/static_map/custom_type_example.cu b/examples/static_map/custom_type_example.cu index ffd8a8d06..1c2cbd7bf 100644 --- a/examples/static_map/custom_type_example.cu +++ b/examples/static_map/custom_type_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. @@ -17,9 +17,8 @@ #include #include +#include #include -#include -#include #include #include @@ -63,8 +62,8 @@ int main(void) auto const empty_value_sentinel = custom_value_type{-1}; // Create an iterator of input key/value pairs - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{custom_key_type{i}, custom_value_type{i}}; })); @@ -81,9 +80,9 @@ int main(void) // Reproduce inserted keys auto insert_keys = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - cuda::proclaim_return_type( - [] __device__(auto i) { return custom_key_type{i}; })); + cuda::make_transform_iterator(cuda::make_counting_iterator(0), + cuda::proclaim_return_type( + [] __device__(auto i) { return custom_key_type{i}; })); thrust::device_vector contained(num_pairs); diff --git a/examples/static_map/heterogeneous_lookup_example.cu b/examples/static_map/heterogeneous_lookup_example.cu index e21575131..4a36ae3be 100644 --- a/examples/static_map/heterogeneous_lookup_example.cu +++ b/examples/static_map/heterogeneous_lookup_example.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. @@ -17,12 +17,11 @@ #include #include +#include #include #include #include #include -#include -#include #include @@ -103,8 +102,8 @@ int main() }; thrust::device_vector d_values = {36.5f, 41.2f, 27.1f, 33.8f}; - auto pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type>( [keys = d_keys.begin(), values = d_values.begin()] __device__(int i) { return cuco::pair{keys[i], values[i]}; diff --git a/examples/static_map/host_bulk_example.cu b/examples/static_map/host_bulk_example.cu index 507fff320..a2ddfec89 100644 --- a/examples/static_map/host_bulk_example.cu +++ b/examples/static_map/host_bulk_example.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2025, NVIDIA CORPORATION. + * Copyright (c) 2020-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,10 +16,9 @@ #include +#include #include #include -#include -#include #include #include @@ -64,8 +63,8 @@ int main(void) thrust::device_vector insert_values(num_keys); thrust::sequence(insert_values.begin(), insert_values.end(), 0); // Combine keys and values into pairs {{0,0}, {1,1}, ... {i,i}} - auto pairs = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto pairs = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type>( [keys = insert_keys.begin(), values = insert_values.begin()] __device__(auto i) { return cuco::pair{keys[i], values[i]}; diff --git a/examples/static_multimap/host_bulk_example.cu b/examples/static_multimap/host_bulk_example.cu index 5e0e812a8..4de8694a3 100644 --- a/examples/static_multimap/host_bulk_example.cu +++ b/examples/static_multimap/host_bulk_example.cu @@ -16,9 +16,9 @@ #include +#include #include #include -#include #include #include @@ -45,8 +45,8 @@ int main(void) // Create a sequence of pairs. Eeach key has two matches. // E.g., {{0,0}, {1,1}, ... {0,25'000}, {1, 25'001}, ...} thrust::transform( - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(pairs.size()), + cuda::make_counting_iterator(0), + cuda::make_counting_iterator(pairs.size()), pairs.begin(), [] __device__(auto i) { return cuco::pair{i % (N / 2), i}; }); diff --git a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh index 87edd1b67..8e2b49453 100644 --- a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh +++ b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -27,6 +27,7 @@ #include #include #include +#include #include #include // TODO #include once available #include @@ -36,7 +37,6 @@ #include #include #include -#include #include @@ -605,7 +605,7 @@ class bloom_filter_impl { OutputIt output_begin, cuda::stream_ref stream) const noexcept { - auto const always_true = thrust::constant_iterator{true}; + auto const always_true = cuda::constant_iterator{true}; this->contains_if_async(first, last, always_true, cuda::std::identity{}, output_begin, stream); } diff --git a/include/cuco/detail/hyperloglog/hyperloglog_impl.cuh b/include/cuco/detail/hyperloglog/hyperloglog_impl.cuh index 534c7a135..1e66a4870 100644 --- a/include/cuco/detail/hyperloglog/hyperloglog_impl.cuh +++ b/include/cuco/detail/hyperloglog/hyperloglog_impl.cuh @@ -28,6 +28,7 @@ #include #include +#include #include // TODO #include once available #include #include @@ -36,8 +37,6 @@ #include #include #include -#include -#include #include #include @@ -183,7 +182,7 @@ class hyperloglog_impl { __host__ constexpr void add_async(InputIt first, InputIt last, cuda::stream_ref stream) { this->add_if_async( - first, last, thrust::constant_iterator{true}, cuda::std::identity{}, stream); + first, last, cuda::constant_iterator{true}, cuda::std::identity{}, stream); } /** @@ -244,8 +243,8 @@ class hyperloglog_impl { int const shmem_bytes = sketch_bytes(); void const* kernel = nullptr; - if constexpr (thrust::is_contiguous_iterator_v) { - auto const ptr = thrust::raw_pointer_cast(&first[0]); + if constexpr (cuda::std::contiguous_iterator) { + auto const ptr = cuda::std::to_address(first); auto constexpr max_vector_bytes = 32; auto const alignment = 1 << cuda::std::countr_zero(reinterpret_cast(ptr) | max_vector_bytes); @@ -276,11 +275,11 @@ class hyperloglog_impl { } if (kernel != nullptr and this->try_reserve_shmem(kernel, shmem_bytes)) { - if constexpr (thrust::is_contiguous_iterator_v) { + if constexpr (cuda::std::contiguous_iterator) { CUCO_CUDA_TRY( cudaOccupancyMaxPotentialBlockSize(&grid_size, &block_size, kernel, shmem_bytes)); - auto const ptr = thrust::raw_pointer_cast(&first[0]); + auto const ptr = cuda::std::to_address(first); void* kernel_args[] = {(void*)(&ptr), (void*)(&num_items), (void*)(&stencil), diff --git a/include/cuco/detail/open_addressing/open_addressing_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_impl.cuh index 7fee3ca6b..55d423e97 100644 --- a/include/cuco/detail/open_addressing/open_addressing_impl.cuh +++ b/include/cuco/detail/open_addressing/open_addressing_impl.cuh @@ -32,10 +32,8 @@ #include #include #include +#include #include -#include -#include -#include #include #include @@ -271,7 +269,7 @@ class open_addressing_impl { template size_type insert(InputIt first, InputIt last, Ref container_ref, cuda::stream_ref stream) { - auto const always_true = thrust::constant_iterator{true}; + auto const always_true = cuda::constant_iterator{true}; return this->insert_if(first, last, always_true, cuda::std::identity{}, container_ref, stream); } @@ -294,7 +292,7 @@ class open_addressing_impl { Ref container_ref, cuda::stream_ref stream) noexcept { - auto const always_true = thrust::constant_iterator{true}; + auto const always_true = cuda::constant_iterator{true}; this->insert_if_async(first, last, always_true, cuda::std::identity{}, container_ref, stream); } @@ -494,7 +492,7 @@ class open_addressing_impl { Ref container_ref, cuda::stream_ref stream) const noexcept { - auto const always_true = thrust::constant_iterator{true}; + auto const always_true = cuda::constant_iterator{true}; this->contains_if_async( first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream); } @@ -568,7 +566,7 @@ class open_addressing_impl { Ref container_ref, cuda::stream_ref stream) const noexcept { - auto const always_true = thrust::constant_iterator{true}; + auto const always_true = cuda::constant_iterator{true}; this->find_if_async( first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream); @@ -853,8 +851,8 @@ class open_addressing_impl { offset += stride) { auto const num_items = std::min(static_cast(this->capacity()) - offset, stride); - auto const begin = thrust::make_transform_iterator( - thrust::counting_iterator{static_cast(offset)}, + auto const begin = cuda::make_transform_iterator( + cuda::counting_iterator{static_cast(offset)}, detail::open_addressing_ns::get_slot(this->storage_ref())); auto const is_filled = detail::open_addressing_ns::slot_is_filled{ this->empty_key_sentinel(), this->erased_key_sentinel()}; diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh index 04e7a3db1..b6a11b345 100644 --- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh +++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh @@ -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. @@ -24,13 +24,13 @@ #include #include +#include #include #include #include #include #include #include -#include #include #include #if defined(CUCO_HAS_CUDA_BARRIER) @@ -1084,7 +1084,7 @@ class open_addressing_ref_impl { { auto constexpr is_outer = false; auto const n = cuco::detail::distance(input_probe_begin, input_probe_end); // TODO include - auto const always_true_stencil = thrust::constant_iterator(true); + auto const always_true_stencil = cuda::constant_iterator(true); auto const identity_predicate = cuda::std::identity{}; this->retrieve_impl(block, input_probe_begin, @@ -1141,7 +1141,7 @@ class open_addressing_ref_impl { { auto constexpr is_outer = true; auto const n = cuco::detail::distance(input_probe_begin, input_probe_end); // TODO include - auto const always_true_stencil = thrust::constant_iterator(true); + auto const always_true_stencil = cuda::constant_iterator(true); auto const identity_predicate = cuda::std::identity{}; this->retrieve_impl(block, input_probe_begin, diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh index 4012c118e..de94824ac 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh @@ -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. @@ -23,12 +23,12 @@ #include #include +#include #include #include #include #include #include -#include namespace cuco::experimental::detail { @@ -79,7 +79,7 @@ class roaring_bitmap_impl { { if (this->empty()) { cub::DeviceTransform::Transform( - thrust::constant_iterator(false), + cuda::constant_iterator(false), contained, cuda::std::distance(first, last), cuda::proclaim_return_type([] __device__(auto /* dummy */) { return false; }), @@ -327,7 +327,7 @@ class roaring_bitmap_impl { { if (this->empty()) { cub::DeviceTransform::Transform( - thrust::constant_iterator(false), + cuda::constant_iterator(false), contained, cuda::std::distance(first, last), cuda::proclaim_return_type([] __device__(auto /* dummy */) { return false; }), diff --git a/include/cuco/detail/static_multimap/static_multimap.inl b/include/cuco/detail/static_multimap/static_multimap.inl index f0b3d4330..ad529aed7 100644 --- a/include/cuco/detail/static_multimap/static_multimap.inl +++ b/include/cuco/detail/static_multimap/static_multimap.inl @@ -18,10 +18,10 @@ #include #include +#include #include #include #include -#include #include diff --git a/include/cuco/detail/trie/dynamic_bitset/dynamic_bitset.inl b/include/cuco/detail/trie/dynamic_bitset/dynamic_bitset.inl index cf1381441..2067a0d10 100644 --- a/include/cuco/detail/trie/dynamic_bitset/dynamic_bitset.inl +++ b/include/cuco/detail/trie/dynamic_bitset/dynamic_bitset.inl @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025 NVIDIA CORPORATION & AFFILIATES. All rights + * SPDX-FileCopyrightText: Copyright (c) 2023-2026 NVIDIA CORPORATION & AFFILIATES. All rights * reserved. SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -21,9 +21,9 @@ #include #include +#include #include #include -#include namespace cuco { namespace experimental { @@ -227,10 +227,10 @@ constexpr void dynamic_bitset::build_ranks_and_selects( CUCO_CUDA_TRY(cub::DeviceSelect::Flagged(nullptr, temp_storage_bytes, - thrust::make_counting_iterator(0UL), + cuda::make_counting_iterator(0UL), select_markers_begin, select_begin, - thrust::make_discard_iterator(), + cuda::make_discard_iterator(), num_blocks, stream.get())); @@ -238,10 +238,10 @@ constexpr void dynamic_bitset::build_ranks_and_selects( CUCO_CUDA_TRY(cub::DeviceSelect::Flagged(thrust::raw_pointer_cast(d_temp_storage), temp_storage_bytes, - thrust::make_counting_iterator(0UL), + cuda::make_counting_iterator(0UL), select_markers_begin, select_begin, - thrust::make_discard_iterator(), + cuda::make_discard_iterator(), num_blocks, stream.get())); diff --git a/include/cuco/detail/utils.hpp b/include/cuco/detail/utils.hpp index d0d777ed6..0b201539a 100644 --- a/include/cuco/detail/utils.hpp +++ b/include/cuco/detail/utils.hpp @@ -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. @@ -24,13 +24,23 @@ namespace cuco { namespace detail { +template +inline constexpr bool has_random_access_iterator_concept = false; + +template +inline constexpr bool + has_random_access_iterator_concept> = + cuda::std::is_base_of_v; + template __host__ __device__ constexpr inline index_type distance(Iterator begin, Iterator end) { using category = typename cuda::std::iterator_traits::iterator_category; - static_assert(cuda::std::is_base_of_v, + static_assert(cuda::std::is_base_of_v || + has_random_access_iterator_concept, "Input iterator should be a random access iterator."); - // `int64_t` instead of arch-dependant `long int` return static_cast(cuda::std::distance(begin, end)); } diff --git a/include/cuco/utility/key_generator.cuh b/include/cuco/utility/key_generator.cuh index e0cbfcbd4..4b11a5a3a 100644 --- a/include/cuco/utility/key_generator.cuh +++ b/include/cuco/utility/key_generator.cuh @@ -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. @@ -20,6 +20,7 @@ #include #include +#include #include #include // TODO include instead once available #include @@ -28,8 +29,6 @@ #include #include #include -#include -#include #include #include #include @@ -291,14 +290,14 @@ class key_generator { size_t seed = this->rng_(); thrust::transform(exec_policy, - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(num_keys), + cuda::make_counting_iterator(0), + cuda::make_counting_iterator(num_keys), out_begin, detail::generate_uniform_fn{num_keys, dist, seed}); } else if constexpr (std::is_same_v) { size_t num_keys = cuda::std::distance(out_begin, out_end); - thrust::counting_iterator seq(this->rng_()); + cuda::counting_iterator seq(this->rng_()); thrust::transform(exec_policy, seq, @@ -374,7 +373,7 @@ class key_generator { if (keep_prob < 1.0) { size_t const num_keys = cuda::std::distance(begin, end); - thrust::counting_iterator seeds(rng_()); + cuda::counting_iterator seeds(rng_()); thrust::transform_if(exec_policy, seeds, @@ -469,8 +468,8 @@ generate_random_byte_sequences(std::size_t n_sequences, // generate random lengths thrust::transform(exec_pol, - thrust::counting_iterator(0), - thrust::counting_iterator(lengths.size()), + cuda::counting_iterator(0), + cuda::counting_iterator(lengths.size()), lengths.begin(), cuda::proclaim_return_type( [min_sequence_length, max_sequence_length, seed] __device__(std::size_t idx) { @@ -508,8 +507,8 @@ generate_random_byte_sequences(std::size_t n_sequences, // fill the byte buffer with random data thrust::transform(exec_pol, - thrust::counting_iterator(0), - thrust::counting_iterator(bytes.size()), + cuda::counting_iterator(0), + cuda::counting_iterator(bytes.size()), bytes.begin(), cuda::proclaim_return_type([seed] __device__(std::size_t idx) { RNG rng; diff --git a/tests/bloom_filter/merge_intersect_test.cu b/tests/bloom_filter/merge_intersect_test.cu index cc5d48c74..5638d58a8 100644 --- a/tests/bloom_filter/merge_intersect_test.cu +++ b/tests/bloom_filter/merge_intersect_test.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. @@ -20,8 +20,8 @@ #include #include +#include #include -#include #include #include @@ -42,29 +42,30 @@ void test_merge_intersect(Filter& filter_a, size_type num_keys = capacity; size_type half_keys = capacity / 2; - // Set A: [0, capacity) - auto keys_a_begin = thrust::counting_iterator{static_cast(0)}; + auto to_key = cuda::proclaim_return_type([] __device__(size_type i) { return Key(i); }); + + auto keys_a_begin = cuda::make_transform_iterator(cuda::counting_iterator{0}, to_key); auto keys_a_end = keys_a_begin + num_keys; - // Set B: [capacity/2, capacity + capacity/2) (50% overlap with A) - auto keys_b_begin = thrust::counting_iterator{static_cast(half_keys)}; - auto keys_b_end = keys_b_begin + num_keys; + auto keys_b_begin = + cuda::make_transform_iterator(cuda::counting_iterator{half_keys}, to_key); + auto keys_b_end = keys_b_begin + num_keys; - // Intersection: [capacity/2, capacity) - auto keys_intersection_begin = thrust::counting_iterator{static_cast(half_keys)}; - auto keys_intersection_end = keys_intersection_begin + half_keys; + auto keys_intersection_begin = + cuda::make_transform_iterator(cuda::counting_iterator{half_keys}, to_key); + auto keys_intersection_end = keys_intersection_begin + half_keys; - // Union: [0, capacity + capacity/2) - auto keys_union_begin = thrust::counting_iterator{static_cast(0)}; - auto keys_union_end = keys_union_begin + num_keys + half_keys; + auto keys_union_begin = + cuda::make_transform_iterator(cuda::counting_iterator{0}, to_key); + auto keys_union_end = keys_union_begin + num_keys + half_keys; - // Unique A: [0, capacity/2) - auto keys_unique_a_begin = thrust::counting_iterator{static_cast(0)}; - auto keys_unique_a_end = keys_unique_a_begin + half_keys; + auto keys_unique_a_begin = + cuda::make_transform_iterator(cuda::counting_iterator{0}, to_key); + auto keys_unique_a_end = keys_unique_a_begin + half_keys; - // Unique B: [capacity, capacity + capacity/2) - auto keys_unique_b_begin = thrust::counting_iterator{static_cast(num_keys)}; - auto keys_unique_b_end = keys_unique_b_begin + half_keys; + auto keys_unique_b_begin = + cuda::make_transform_iterator(cuda::counting_iterator{num_keys}, to_key); + auto keys_unique_b_end = keys_unique_b_begin + half_keys; // Helper to fill filters auto refill_filters = [&]() { diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu index 0e5314898..4cf24563c 100644 --- a/tests/bloom_filter/unique_sequence_test.cu +++ b/tests/bloom_filter/unique_sequence_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,9 +19,9 @@ #include #include +#include #include #include -#include #include #include @@ -67,16 +67,16 @@ void test_unique_sequence(Filter& filter, size_type num_keys) SECTION("All conditionally inserted keys should be contained") { - filter.add_if(keys.begin(), keys.end(), thrust::counting_iterator(0), is_even); + filter.add_if(keys.begin(), keys.end(), cuda::counting_iterator(0), is_even); filter.contains_if(keys.begin(), keys.end(), - thrust::counting_iterator(0), + cuda::counting_iterator(0), is_even, contained.begin()); REQUIRE(cuco::test::equal( contained.begin(), contained.end(), - thrust::counting_iterator(0), + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto const& idx_contained, auto const& idx) { return ((idx % 2) == 0) == idx_contained; }))); diff --git a/tests/bloom_filter/variable_cg_test.cu b/tests/bloom_filter/variable_cg_test.cu index 6a9997692..adb6b703f 100644 --- a/tests/bloom_filter/variable_cg_test.cu +++ b/tests/bloom_filter/variable_cg_test.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. @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include @@ -50,7 +49,7 @@ void test_variable_cg_size(Filter& filter, size_type num_keys) thrust::device_vector contained(num_keys, false); - auto const always_true = thrust::constant_iterator{true}; + auto const always_true = cuda::constant_iterator{true}; SECTION("Check if fallback kernels work for varying combinations of CG sizes.") { diff --git a/tests/dynamic_map/erase_test.cu b/tests/dynamic_map/erase_test.cu index 1ffc9673c..453559578 100644 --- a/tests/dynamic_map/erase_test.cu +++ b/tests/dynamic_map/erase_test.cu @@ -19,12 +19,11 @@ #include #include +#include #include #include #include #include -#include -#include #include #include @@ -52,8 +51,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map erase tests", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end(), 1); thrust::sequence(thrust::device, d_values.begin(), d_values.end(), 1); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [keys = d_keys.begin(), values = d_values.begin()] __device__(auto i) { return cuco::pair{keys[i], values[i]}; @@ -102,8 +101,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map erase tests", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end(), 1); thrust::sequence(thrust::device, d_values.begin(), d_values.end(), 1); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [keys = d_keys.begin(), values = d_values.begin()] __device__(auto i) { return cuco::pair{keys[i], values[i]}; diff --git a/tests/dynamic_map/find_test.cu b/tests/dynamic_map/find_test.cu index ebd88a029..5e5520af1 100644 --- a/tests/dynamic_map/find_test.cu +++ b/tests/dynamic_map/find_test.cu @@ -19,13 +19,11 @@ #include #include +#include #include #include #include #include -#include -#include -#include #include #include @@ -54,8 +52,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map find tests", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end(), 1); thrust::sequence(thrust::device, d_values.begin(), d_values.end(), 1); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [keys = d_keys.begin(), values = d_values.begin()] __device__(auto i) { return cuco::pair{keys[i], values[i]}; @@ -85,7 +83,7 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map find tests", auto empty_zip = thrust::make_zip_iterator( cuda::std::tuple{d_nonexistent_values.begin(), - thrust::constant_iterator{cuco::empty_value{-1}.value}}); + cuda::constant_iterator{cuco::empty_value{-1}.value}}); REQUIRE(cuco::test::all_of(empty_zip, empty_zip + 100, zip_equal)); thrust::device_vector d_mixed_keys(200); @@ -105,7 +103,7 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map find tests", auto second_half_empty_zip = thrust::make_zip_iterator( cuda::std::tuple{d_mixed_values.begin() + 100, - thrust::constant_iterator{cuco::empty_value{-1}.value}}); + cuda::constant_iterator{cuco::empty_value{-1}.value}}); REQUIRE(cuco::test::all_of(second_half_empty_zip, second_half_empty_zip + 100, zip_equal)); } @@ -118,8 +116,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map find tests", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end(), 1); thrust::sequence(thrust::device, d_values.begin(), d_values.end(), 1); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [keys = d_keys.begin(), values = d_values.begin()] __device__(auto i) { return cuco::pair{keys[i], values[i]}; @@ -143,9 +141,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map find tests", map.find(d_keys.begin(), d_keys.end(), d_found_values.begin()); - auto first_half_empty_zip = thrust::make_zip_iterator( - cuda::std::tuple{d_found_values.begin(), - thrust::constant_iterator{cuco::empty_value{-1}.value}}); + auto first_half_empty_zip = thrust::make_zip_iterator(cuda::std::tuple{ + d_found_values.begin(), cuda::constant_iterator{cuco::empty_value{-1}.value}}); REQUIRE( cuco::test::all_of(first_half_empty_zip, first_half_empty_zip + num_keys / 2, zip_equal)); diff --git a/tests/dynamic_map/multiplicity_test.cu b/tests/dynamic_map/multiplicity_test.cu index 5fc7eddb7..33351239f 100644 --- a/tests/dynamic_map/multiplicity_test.cu +++ b/tests/dynamic_map/multiplicity_test.cu @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include @@ -47,15 +46,15 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", // Create pairs for first submap (keys 0 to num_keys-1) auto pairs_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - cuda::proclaim_return_type>( - [] __device__(auto i) { return cuco::pair(i, i); })); + cuda::make_transform_iterator(cuda::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); // Create pairs for second submap (keys num_keys to 2*num_keys-1) auto pairs_begin_2 = - thrust::make_transform_iterator(thrust::make_counting_iterator(num_keys), - cuda::proclaim_return_type>( - [] __device__(auto i) { return cuco::pair(i, i); })); + cuda::make_transform_iterator(cuda::make_counting_iterator(num_keys), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); thrust::device_vector d_keys(num_keys); thrust::device_vector d_results(num_keys); @@ -78,8 +77,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", // Try to insert duplicates with DIFFERENT values - should still not insert and preserve // originals - auto duplicate_pairs_diff_values = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto duplicate_pairs_diff_values = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i + 999); })); map.insert(duplicate_pairs_diff_values, duplicate_pairs_diff_values + num_keys); @@ -87,10 +86,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", // Verify original values are preserved (not overwritten by duplicate insert attempts) map.find(d_keys.begin(), d_keys.end(), d_results.begin()); - REQUIRE(cuco::test::equal(d_results.begin(), - d_results.end(), - thrust::counting_iterator(0), - cuda::std::equal_to{})); + REQUIRE(cuco::test::equal( + d_results.begin(), d_results.end(), cuda::counting_iterator(0), cuda::std::equal_to{})); } SECTION("contains finds keys in any submap") @@ -129,10 +126,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", // Find keys from FIRST submap - should return correct values map.find(d_keys.begin(), d_keys.end(), d_results.begin()); - REQUIRE(cuco::test::equal(d_results.begin(), - d_results.end(), - thrust::counting_iterator(0), - cuda::std::equal_to{})); + REQUIRE(cuco::test::equal( + d_results.begin(), d_results.end(), cuda::counting_iterator(0), cuda::std::equal_to{})); // Find keys from SECOND submap - should return correct values thrust::device_vector d_keys_2(num_keys); @@ -140,7 +135,7 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", map.find(d_keys_2.begin(), d_keys_2.end(), d_results.begin()); REQUIRE(cuco::test::equal(d_results.begin(), d_results.end(), - thrust::counting_iterator(num_keys), + cuda::counting_iterator(num_keys), cuda::std::equal_to{})); // Non-existent keys should return empty_value_sentinel (-1) @@ -210,8 +205,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", REQUIRE(map.size() == 2 * num_keys); // Create pairs with same keys as first submap but different values (value = key + 100) - auto updated_pairs = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto updated_pairs = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i + 100); })); @@ -223,7 +218,7 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", map.find(d_keys.begin(), d_keys.end(), d_results.begin()); REQUIRE(cuco::test::equal(d_results.begin(), d_results.end(), - thrust::counting_iterator(100), // Values should now be key + 100 + cuda::counting_iterator(100), // Values should now be key + 100 cuda::std::equal_to{})); // Verify second submap values are unchanged @@ -232,14 +227,14 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", map.find(d_keys_2.begin(), d_keys_2.end(), d_results.begin()); REQUIRE(cuco::test::equal(d_results.begin(), d_results.end(), - thrust::counting_iterator(num_keys), + cuda::counting_iterator(num_keys), cuda::std::equal_to{})); // Test INSERT behavior: insert_or_assign with completely NEW keys should increase size - auto new_pairs = thrust::make_transform_iterator( - thrust::make_counting_iterator(2 * num_keys), - cuda::proclaim_return_type>( - [] __device__(auto i) { return cuco::pair(i, i); })); + auto new_pairs = + cuda::make_transform_iterator(cuda::make_counting_iterator(2 * num_keys), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); map.insert_or_assign(new_pairs, new_pairs + num_keys); REQUIRE(map.size() == 3 * num_keys); // Size should increase by num_keys @@ -249,14 +244,14 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", map.find(d_keys_new.begin(), d_keys_new.end(), d_results.begin()); REQUIRE(cuco::test::equal(d_results.begin(), d_results.end(), - thrust::counting_iterator(2 * num_keys), + cuda::counting_iterator(2 * num_keys), cuda::std::equal_to{})); // Test MIXED behavior: some keys exist (update), some don't (insert) // Use keys from 0 to num_keys/2 (exist in first submap) and // keys from 3*num_keys to 3*num_keys + num_keys/2 (don't exist) - auto mixed_pairs = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto mixed_pairs = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>([] __device__(auto i) { Key key = (i < num_keys / 2) ? Key(i) : Key(3 * num_keys + i - num_keys / 2); return cuco::pair(key, T(i + 500)); @@ -300,7 +295,7 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: cross-submap duplicate handling", // Values should match keys (since we inserted key=value pairs) REQUIRE(cuco::test::equal(d_retrieved_values.begin(), d_retrieved_values.end(), - thrust::counting_iterator(0), + cuda::counting_iterator(0), cuda::std::equal_to{})); } } diff --git a/tests/dynamic_map/retrieve_all_test.cu b/tests/dynamic_map/retrieve_all_test.cu index 35f057e4b..3822e5011 100644 --- a/tests/dynamic_map/retrieve_all_test.cu +++ b/tests/dynamic_map/retrieve_all_test.cu @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include @@ -51,8 +50,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map retrieve_all tests", thrust::sequence(d_keys.begin(), d_keys.end(), 0); thrust::sequence(d_values.begin(), d_values.end(), 0); - auto pairs = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto pairs = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type>( [keys = d_keys.begin(), values = d_values.begin()] __device__(auto i) { return cuco::pair{keys[i], values[i]}; @@ -85,8 +84,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map retrieve_all tests", thrust::sequence(d_keys.begin(), d_keys.end(), 0); thrust::sequence(d_values.begin(), d_values.end(), 0); - auto pairs = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto pairs = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type>( [keys = d_keys.begin(), values = d_values.begin()] __device__(auto i) { return cuco::pair{keys[i], values[i]}; diff --git a/tests/dynamic_map/unique_sequence_test.cu b/tests/dynamic_map/unique_sequence_test.cu index 13cf08d0b..d6a8766f6 100644 --- a/tests/dynamic_map/unique_sequence_test.cu +++ b/tests/dynamic_map/unique_sequence_test.cu @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include @@ -47,9 +46,9 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: unique sequence", thrust::sequence(thrust::device, d_values.begin(), d_values.end()); auto pairs_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - cuda::proclaim_return_type>( - [] __device__(auto i) { return cuco::pair(i, i); })); + cuda::make_transform_iterator(cuda::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); thrust::device_vector d_results(num_keys); thrust::device_vector d_contained(num_keys); @@ -103,8 +102,8 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: unique sequence", REQUIRE(map.size() == num_keys); // Create pairs with same keys but different values (value = key + 1) - auto updated_pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto updated_pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i + 1); })); @@ -116,14 +115,14 @@ TEMPLATE_TEST_CASE_SIG("dynamic_map: unique sequence", map.find(d_keys.begin(), d_keys.end(), d_results.begin()); REQUIRE(cuco::test::equal(d_results.begin(), d_results.end(), - thrust::counting_iterator(1), // Values should now be key + 1 + cuda::counting_iterator(1), // Values should now be key + 1 cuda::std::equal_to{})); // Insert new keys with insert_or_assign (keys from num_keys to 2*num_keys) - auto new_pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(num_keys), - cuda::proclaim_return_type>( - [] __device__(auto i) { return cuco::pair(i, i); })); + auto new_pairs_begin = + cuda::make_transform_iterator(cuda::make_counting_iterator(num_keys), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); map.insert_or_assign(new_pairs_begin, new_pairs_begin + num_keys); REQUIRE(map.size() == 2 * num_keys); diff --git a/tests/hyperloglog/spark_parity_test.cu b/tests/hyperloglog/spark_parity_test.cu index 1f9a0d386..6dfb09d3c 100644 --- a/tests/hyperloglog/spark_parity_test.cu +++ b/tests/hyperloglog/spark_parity_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -20,10 +20,8 @@ #include #include +#include #include -#include -#include -#include #include #include @@ -76,10 +74,10 @@ TEST_CASE("hyperloglog: Spark parity: deterministic cardinality estimation", "") REQUIRE(estimator_type::sketch_bytes(sd) == estimator_type::sketch_bytes(sb)); auto items_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - cuda::proclaim_return_type([repeats] __device__(auto i) { - return static_cast(i / repeats); - })); + cuda::make_transform_iterator(cuda::make_counting_iterator(0), + cuda::proclaim_return_type([repeats] __device__(auto i) { + return static_cast(i / repeats); + })); estimator_type estimator{sd}; @@ -142,7 +140,7 @@ TEST_CASE("hyperloglog: Spark parity: merging HLL instances", "") auto num_items = 1000000; auto standard_deviation = cuco::standard_deviation(0.05); - auto items_begin = thrust::make_counting_iterator(0); + auto items_begin = cuda::make_counting_iterator(0); // count lower half of input estimator_type lower{standard_deviation}; @@ -155,7 +153,7 @@ TEST_CASE("hyperloglog: Spark parity: merging HLL instances", "") // merge upper into lower so lower has seen the entire input lower.merge(upper); - auto reversed_items_begin = thrust::make_transform_iterator( + auto reversed_items_begin = cuda::make_transform_iterator( items_begin, cuda::proclaim_return_type([num_items] __device__(auto i) { return static_cast(num_items - i); })); diff --git a/tests/hyperloglog/type_deduction_test.cu b/tests/hyperloglog/type_deduction_test.cu index 409ed02b5..5e86b4ae7 100644 --- a/tests/hyperloglog/type_deduction_test.cu +++ b/tests/hyperloglog/type_deduction_test.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. @@ -20,8 +20,7 @@ #include #include -#include -#include +#include #include @@ -32,8 +31,8 @@ TEST_CASE("hyperloglog: type deduction bug with hash functions returning referen auto constexpr sketch_size_kb = 1; auto constexpr num_items = 1000; - auto first = thrust::make_transform_iterator(thrust::counting_iterator(0), - cuco::xxhash_64{}); + auto first = cuda::make_transform_iterator(cuda::counting_iterator(0), + cuco::xxhash_64{}); auto last = first + num_items; cuco::hyperloglog estimator{ diff --git a/tests/static_map/contains_test.cu b/tests/static_map/contains_test.cu index abe2f5619..e53218933 100644 --- a/tests/static_map/contains_test.cu +++ b/tests/static_map/contains_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2025, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -19,12 +19,11 @@ #include #include +#include #include #include #include #include -#include -#include #include #include @@ -38,9 +37,9 @@ void test_unique_sequence(Map& map, size_type num_keys) using Key = typename Map::key_type; using Value = typename Map::mapped_type; - auto keys_begin = thrust::counting_iterator{0}; - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto keys_begin = cuda::counting_iterator{0}; + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, i}; })); thrust::device_vector d_contained(num_keys); @@ -63,7 +62,7 @@ void test_unique_sequence(Map& map, size_type num_keys) SECTION("All conditionally inserted keys should be contained") { auto const inserted = map.insert_if( - pairs_begin, pairs_begin + num_keys, thrust::counting_iterator(0), is_even); + pairs_begin, pairs_begin + num_keys, cuda::counting_iterator(0), is_even); REQUIRE(inserted == num_keys / 2); REQUIRE(map.size() == num_keys / 2); @@ -71,7 +70,7 @@ void test_unique_sequence(Map& map, size_type num_keys) REQUIRE(cuco::test::equal( d_contained.begin(), d_contained.end(), - thrust::counting_iterator(0), + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto const& idx_contained, auto const& idx) { return ((idx % 2) == 0) == idx_contained; }))); @@ -91,11 +90,11 @@ void test_unique_sequence(Map& map, size_type num_keys) { map.contains_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator(0), + cuda::counting_iterator(0), is_even, d_contained.begin()); auto gold_iter = - thrust::make_transform_iterator(thrust::counting_iterator(0), is_even); + cuda::make_transform_iterator(cuda::counting_iterator(0), is_even); auto zip = thrust::make_zip_iterator(cuda::std::tuple{d_contained.begin(), gold_iter}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } @@ -111,7 +110,7 @@ void test_unique_sequence(Map& map, size_type num_keys) thrust::sort(thrust::device, d_values.begin(), values_end); REQUIRE(cuco::test::equal(d_values.begin(), values_end, - thrust::make_counting_iterator(0), + cuda::make_counting_iterator(0), cuda::std::equal_to{})); } } diff --git a/tests/static_map/duplicate_keys_test.cu b/tests/static_map/duplicate_keys_test.cu index 34e7b618a..397af6a4e 100644 --- a/tests/static_map/duplicate_keys_test.cu +++ b/tests/static_map/duplicate_keys_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2025, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -19,11 +19,9 @@ #include #include +#include #include #include -#include -#include -#include #include #include @@ -79,8 +77,8 @@ TEMPLATE_TEST_CASE_SIG( thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); thrust::sequence(thrust::device, d_values.begin(), d_values.end()); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i / 2, i / 2); })); @@ -112,7 +110,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::sort(thrust::device, unique_keys.begin(), unique_keys.end()); REQUIRE(cuco::test::equal(unique_keys.begin(), unique_keys.end(), - thrust::make_counting_iterator(0), + cuda::make_counting_iterator(0), cuda::std::equal_to{})); } diff --git a/tests/static_map/erase_test.cu b/tests/static_map/erase_test.cu index 1299b4815..518b4fe14 100644 --- a/tests/static_map/erase_test.cu +++ b/tests/static_map/erase_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2025, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -19,9 +19,8 @@ #include #include +#include #include -#include -#include #include @@ -35,9 +34,9 @@ void test_erase(Map& map, size_type num_keys) thrust::device_vector d_keys_exist(num_keys); - auto keys_begin = thrust::counting_iterator(1); + auto keys_begin = cuda::counting_iterator(1); - auto pairs_begin = thrust::make_transform_iterator( + auto pairs_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(key_type const& x) { return cuco::pair(x, static_cast(x)); diff --git a/tests/static_map/find_test.cu b/tests/static_map/find_test.cu index 382578351..d26c658aa 100644 --- a/tests/static_map/find_test.cu +++ b/tests/static_map/find_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2025, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -19,13 +19,12 @@ #include #include +#include #include #include #include #include #include -#include -#include #include #include @@ -49,9 +48,9 @@ void test_unique_sequence(Map& map, size_type num_keys) using Key = typename Map::key_type; using Value = typename Map::mapped_type; - auto keys_begin = thrust::counting_iterator{0}; - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto keys_begin = cuda::counting_iterator{0}; + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, i}; })); @@ -65,8 +64,8 @@ void test_unique_sequence(Map& map, size_type num_keys) SECTION("Non-inserted keys have no matches") { map.find(keys_begin, keys_begin + num_keys, d_results.begin()); - auto zip = thrust::make_zip_iterator(cuda::std::tuple{ - d_results.begin(), thrust::constant_iterator{map.empty_key_sentinel()}}); + auto zip = thrust::make_zip_iterator( + cuda::std::tuple{d_results.begin(), cuda::constant_iterator{map.empty_key_sentinel()}}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } @@ -87,7 +86,7 @@ void test_unique_sequence(Map& map, size_type num_keys) keys_begin, keys_begin + num_keys, always_false{}, map.hash_function(), d_results.begin()); CUCO_CUDA_TRY(cudaDeviceSynchronize()); auto zip = thrust::make_zip_iterator(cuda::std::tuple{ - d_results.begin(), thrust::constant_iterator{map.empty_value_sentinel()}}); + d_results.begin(), cuda::constant_iterator{map.empty_value_sentinel()}}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } @@ -100,14 +99,14 @@ void test_unique_sequence(Map& map, size_type num_keys) map.find_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator{0}, + cuda::counting_iterator{0}, is_even, d_results.begin()); REQUIRE(cuco::test::equal( d_results.begin(), d_results.end(), - thrust::make_transform_iterator(thrust::counting_iterator{0}, gold_fn), + cuda::make_transform_iterator(cuda::counting_iterator{0}, gold_fn), cuda::proclaim_return_type( [] __device__(auto const& found, auto const& gold) { return found == gold; }))); } @@ -116,7 +115,7 @@ void test_unique_sequence(Map& map, size_type num_keys) { map.find_if_async(keys_begin, keys_begin + num_keys, - thrust::counting_iterator{0}, + cuda::counting_iterator{0}, is_even, always_false{}, map.hash_function(), @@ -124,7 +123,7 @@ void test_unique_sequence(Map& map, size_type num_keys) CUCO_CUDA_TRY(cudaDeviceSynchronize()); auto zip = thrust::make_zip_iterator(cuda::std::tuple{ - d_results.begin(), thrust::constant_iterator{map.empty_value_sentinel()}}); + d_results.begin(), cuda::constant_iterator{map.empty_value_sentinel()}}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } diff --git a/tests/static_map/for_each_test.cu b/tests/static_map/for_each_test.cu index 70124327a..3ce9d9f59 100644 --- a/tests/static_map/for_each_test.cu +++ b/tests/static_map/for_each_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include @@ -37,8 +36,8 @@ void test_for_each(Map& map, size_type num_keys) REQUIRE(num_keys % 2 == 0); // Insert pairs - auto pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>([] __device__(auto i) { // use payload as 1 for even keys and 2 for odd keys return cuco::pair{i, i % 2 + 1}; @@ -67,8 +66,8 @@ void test_for_each(Map& map, size_type num_keys) counter_storage.reset(stream); map.for_each( - thrust::counting_iterator(0), - thrust::counting_iterator(2 * num_keys), // test for false-positives + cuda::counting_iterator(0), + cuda::counting_iterator(2 * num_keys), // test for false-positives [counter = counter_storage.data()] __device__(auto const slot) { auto const& [key, value] = slot; if (((key % 2 == 0)) and (value == 1)) { counter->fetch_add(1, cuda::memory_order_relaxed); } diff --git a/tests/static_map/hash_test.cu b/tests/static_map/hash_test.cu index 8e427049c..8acf90b9c 100644 --- a/tests/static_map/hash_test.cu +++ b/tests/static_map/hash_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include @@ -45,9 +44,9 @@ void test_hash_function() cuco::storage<2>>{ num_keys, cuco::empty_key{-1}, cuco::empty_value{-1}}; - auto keys_begin = thrust::counting_iterator(1); + auto keys_begin = cuda::counting_iterator(1); - auto pairs_begin = thrust::make_transform_iterator( + auto pairs_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair(i, i); })); diff --git a/tests/static_map/heterogeneous_lookup_test.cu b/tests/static_map/heterogeneous_lookup_test.cu index 80f9715d6..b168d861d 100644 --- a/tests/static_map/heterogeneous_lookup_test.cu +++ b/tests/static_map/heterogeneous_lookup_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2025, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include #include @@ -112,12 +111,12 @@ TEMPLATE_TEST_CASE_SIG("static_map heterogeneous lookup tests", custom_key_equal{}, probe}; - auto insert_pairs = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto insert_pairs = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i); })); - auto probe_keys = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto probe_keys = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto i) { return ProbeKey{i}; })); SECTION("All inserted keys-value pairs should be contained") diff --git a/tests/static_map/insert_and_find_test.cu b/tests/static_map/insert_and_find_test.cu index 6b8463b9e..cc978ff5b 100644 --- a/tests/static_map/insert_and_find_test.cu +++ b/tests/static_map/insert_and_find_test.cu @@ -1,6 +1,6 @@ /* * Copyright (c) 2022, Jonas Hahnfeld, CERN. - * Copyright (c) 2022-2025, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -20,10 +20,9 @@ #include #include +#include #include #include -#include -#include #include @@ -75,8 +74,8 @@ TEMPLATE_TEST_CASE_SIG( cuco::storage<2>>{ num_keys, cuco::empty_key{-1}, cuco::empty_value{-1}}; - auto pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, 1}; })); diff --git a/tests/static_map/insert_or_apply_test.cu b/tests/static_map/insert_or_apply_test.cu index da087a100..e14200d5a 100644 --- a/tests/static_map/insert_or_apply_test.cu +++ b/tests/static_map/insert_or_apply_test.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. @@ -20,11 +20,9 @@ #include #include +#include #include #include -#include -#include -#include #include #include @@ -42,8 +40,8 @@ void test_insert_or_apply(Map& map, size_type num_keys, size_type num_unique_key using Value = typename Map::mapped_type; // Insert pairs - auto pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>([num_unique_keys] __device__(auto i) { return cuco::pair{i % num_unique_keys, 1}; })); @@ -63,7 +61,7 @@ void test_insert_or_apply(Map& map, size_type num_keys, size_type num_unique_key REQUIRE(cuco::test::equal(d_values.begin(), d_values.end(), - thrust::make_constant_iterator(num_keys / num_unique_keys), + cuda::make_constant_iterator(num_keys / num_unique_keys), cuda::std::equal_to{})); } @@ -106,8 +104,8 @@ void test_insert_or_apply_shmem(Map& map, size_type num_keys, size_type num_uniq typename shared_map_ref_type::storage_ref_type>(extent_type{}); // Insert pairs - auto pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>([num_unique_keys] __device__(auto i) { return cuco::pair{i % num_unique_keys, 1}; })); @@ -134,7 +132,7 @@ void test_insert_or_apply_shmem(Map& map, size_type num_keys, size_type num_uniq REQUIRE(cuco::test::equal(d_values.begin(), d_values.end(), - thrust::make_constant_iterator(num_keys / num_unique_keys), + cuda::make_constant_iterator(num_keys / num_unique_keys), cuda::std::equal_to{})); } diff --git a/tests/static_map/insert_or_assign_test.cu b/tests/static_map/insert_or_assign_test.cu index 3b73fa7ad..273fb6e2f 100644 --- a/tests/static_map/insert_or_assign_test.cu +++ b/tests/static_map/insert_or_assign_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, 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. @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include @@ -37,8 +36,8 @@ void test_insert_or_assign(Map& map, size_type num_keys) using Value = typename Map::mapped_type; // Insert pairs - auto pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, i}; })); @@ -46,8 +45,8 @@ void test_insert_or_assign(Map& map, size_type num_keys) REQUIRE(initial_size == num_keys); // all keys should be inserted // Query pairs have the same keys but different payloads - auto query_pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto query_pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i * 2); })); @@ -61,8 +60,8 @@ void test_insert_or_assign(Map& map, size_type num_keys) thrust::device_vector d_values(num_keys); map.retrieve_all(d_keys.begin(), d_values.begin()); - auto gold_values_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto gold_values_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto i) { return i * 2; })); thrust::sort(thrust::device, d_values.begin(), d_values.end()); diff --git a/tests/static_map/key_sentinel_test.cu b/tests/static_map/key_sentinel_test.cu index 3f502b6bd..b58145896 100644 --- a/tests/static_map/key_sentinel_test.cu +++ b/tests/static_map/key_sentinel_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2024, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -19,9 +19,8 @@ #include #include +#include #include -#include -#include #include @@ -54,8 +53,8 @@ TEMPLATE_TEST_CASE_SIG("static_map key sentinel tests", "", ((typename T), T), ( } CUCO_CUDA_TRY(cudaMemcpyToSymbol(A, h_A, SIZE * sizeof(int))); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i); })); diff --git a/tests/static_map/rehash_test.cu b/tests/static_map/rehash_test.cu index 62e3355e5..d4b7ffa17 100644 --- a/tests/static_map/rehash_test.cu +++ b/tests/static_map/rehash_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, 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. @@ -17,8 +17,7 @@ #include #include -#include -#include +#include #include @@ -35,9 +34,9 @@ TEST_CASE("static_map rehash test", "") cuco::empty_value{-1}, cuco::erased_key{-2}}; - auto keys_begin = thrust::counting_iterator(1); + auto keys_begin = cuda::counting_iterator(1); - auto pairs_begin = thrust::make_transform_iterator( + auto pairs_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(key_type const& x) { return cuco::pair(x, static_cast(x)); diff --git a/tests/static_map/retrieve_if_test.cu b/tests/static_map/retrieve_if_test.cu index f593dff44..c91827e57 100644 --- a/tests/static_map/retrieve_if_test.cu +++ b/tests/static_map/retrieve_if_test.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. @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -118,9 +119,9 @@ TEMPLATE_TEST_CASE_SIG("static_map retrieve_if", container_type container{num_keys * 2, cuco::empty_key{-1}, cuco::empty_value{-1}}; - auto keys_begin = thrust::counting_iterator(1); - auto vals_begin = thrust::counting_iterator(1); - auto pairs_begin = thrust::make_zip_iterator(thrust::make_tuple(keys_begin, vals_begin)); + auto keys_begin = cuda::counting_iterator(1); + auto vals_begin = cuda::counting_iterator(1); + auto pairs_begin = thrust::make_zip_iterator(cuda::std::tuple{keys_begin, vals_begin}); container.insert(pairs_begin, pairs_begin + num_keys); diff --git a/tests/static_map/retrieve_test.cu b/tests/static_map/retrieve_test.cu index bdabfd411..1f2a50ccb 100644 --- a/tests/static_map/retrieve_test.cu +++ b/tests/static_map/retrieve_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,14 +19,12 @@ #include #include +#include #include #include #include #include #include -#include -#include -#include #include #include @@ -41,9 +39,9 @@ void test_unique_sequence(Map& map, size_type num_keys) using Key = typename Map::key_type; using Value = typename Map::mapped_type; - auto keys_begin = thrust::counting_iterator{0}; - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto keys_begin = cuda::counting_iterator{0}; + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, i}; })); @@ -57,8 +55,8 @@ void test_unique_sequence(Map& map, size_type num_keys) REQUIRE(count == 0); - auto const [_, output_end] = map.retrieve( - keys_begin, keys_begin + num_keys, thrust::discard_iterator{}, d_results.begin()); + auto const [_, output_end] = + map.retrieve(keys_begin, keys_begin + num_keys, cuda::discard_iterator{}, d_results.begin()); auto const size = std::distance(d_results.begin(), output_end); REQUIRE(size == 0); @@ -74,7 +72,7 @@ void test_unique_sequence(Map& map, size_type num_keys) REQUIRE(count == num_keys); auto [_, output_end] = - map.retrieve(keys_begin, keys_begin + num_keys, thrust::discard_iterator{}, output_begin); + map.retrieve(keys_begin, keys_begin + num_keys, cuda::discard_iterator{}, output_begin); auto const size = cuda::std::distance(output_begin, output_end); REQUIRE(size == num_keys); diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu index 1ef6d05b9..cac1d08cd 100644 --- a/tests/static_map/shared_memory_test.cu +++ b/tests/static_map/shared_memory_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2025, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -108,8 +109,8 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests", SECTION("Keys are all found after insertion.") { - auto pairs_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto pairs_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type>( [d_keys = d_keys.data(), d_values = d_values.data()] __device__(int idx) { return cuco::pair(d_keys[idx], d_values[idx]); diff --git a/tests/static_map/stream_test.cu b/tests/static_map/stream_test.cu index fc85ca27f..2261b33b9 100644 --- a/tests/static_map/stream_test.cu +++ b/tests/static_map/stream_test.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. @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include @@ -58,8 +57,8 @@ TEMPLATE_TEST_CASE_SIG("static_map: unique sequence of keys on given stream", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); thrust::sequence(thrust::device, d_values.begin(), d_values.end()); - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i); })); diff --git a/tests/static_multimap/count_test.cu b/tests/static_multimap/count_test.cu index 3a0a2dfa7..aa389dc8d 100644 --- a/tests/static_multimap/count_test.cu +++ b/tests/static_multimap/count_test.cu @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -37,7 +38,7 @@ void test_multiplicity_count(Map& map, size_type num_keys) using Key = typename Map::key_type; using Value = typename Map::mapped_type; - auto const keys_begin = thrust::counting_iterator{0}; + auto const keys_begin = cuda::counting_iterator{0}; SECTION("Count of empty map should be zero.") { @@ -47,8 +48,8 @@ void test_multiplicity_count(Map& map, size_type num_keys) SECTION("Count of n unique keys should be n.") { - auto const pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto const pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, i}; })); map.insert(pairs_begin, pairs_begin + num_keys); @@ -59,8 +60,8 @@ void test_multiplicity_count(Map& map, size_type num_keys) SECTION("Count of n unique keys should be n x multiplicity.") { - auto const pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto const pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i / multiplicity, i}; })); map.insert(pairs_begin, pairs_begin + num_keys * multiplicity); diff --git a/tests/static_multimap/find_test.cu b/tests/static_multimap/find_test.cu index 68e477409..747bd47fe 100644 --- a/tests/static_multimap/find_test.cu +++ b/tests/static_multimap/find_test.cu @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include @@ -41,7 +40,7 @@ void test_multimap_find(Map& map, size_type num_keys) auto zip_equal = cuda::proclaim_return_type( [] __device__(auto val) { return cuda::std::get<0>(val) == cuda::std::get<1>(val); }); - auto const keys_begin = thrust::counting_iterator{0}; + auto const keys_begin = cuda::counting_iterator{0}; SECTION("Non-inserted keys have no matches") { @@ -49,13 +48,13 @@ void test_multimap_find(Map& map, size_type num_keys) map.find(keys_begin, keys_begin + num_keys, found_vals.begin()); auto zip = thrust::make_zip_iterator(cuda::std::tuple{ - found_vals.begin(), thrust::constant_iterator{map.empty_value_sentinel()}}); + found_vals.begin(), cuda::constant_iterator{map.empty_value_sentinel()}}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } - auto const pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto const pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, i * 2}; })); @@ -67,8 +66,8 @@ void test_multimap_find(Map& map, size_type num_keys) map.find(keys_begin, keys_begin + num_keys, found_vals.begin()); - auto const gold_vals_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto const gold_vals_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type([] __device__(auto i) { return Value{i * 2}; })); auto zip = thrust::make_zip_iterator(cuda::std::tuple{found_vals.begin(), gold_vals_begin}); @@ -86,14 +85,14 @@ void test_multimap_find(Map& map, size_type num_keys) map.find_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator{0}, + cuda::counting_iterator{0}, is_even, found_results.begin()); REQUIRE(cuco::test::equal( found_results.begin(), found_results.end(), - thrust::make_transform_iterator(thrust::counting_iterator{0}, gold_fn), + cuda::make_transform_iterator(cuda::counting_iterator{0}, gold_fn), cuda::proclaim_return_type( [] __device__(auto const& found, auto const& gold) { return found == gold; }))); } diff --git a/tests/static_multimap/for_each_test.cu b/tests/static_multimap/for_each_test.cu index fc511ee58..e6eddcc81 100644 --- a/tests/static_multimap/for_each_test.cu +++ b/tests/static_multimap/for_each_test.cu @@ -21,8 +21,7 @@ #include #include -#include -#include +#include #include #include @@ -128,12 +127,12 @@ TEMPLATE_TEST_CASE_SIG( {}, cuco::storage<2>{}}; - auto unique_keys_begin = thrust::counting_iterator(0); + auto unique_keys_begin = cuda::counting_iterator(0); auto gen_duplicate_keys = cuda::proclaim_return_type( [] __device__(auto const& k) { return static_cast(k % num_unique_keys); }); - auto keys_begin = thrust::make_transform_iterator(unique_keys_begin, gen_duplicate_keys); + auto keys_begin = cuda::make_transform_iterator(unique_keys_begin, gen_duplicate_keys); - auto const pairs_begin = thrust::make_transform_iterator( + auto const pairs_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i, i}; })); diff --git a/tests/static_multimap/heterogeneous_lookup_test.cu b/tests/static_multimap/heterogeneous_lookup_test.cu index 07cd65e27..e99ef6384 100644 --- a/tests/static_multimap/heterogeneous_lookup_test.cu +++ b/tests/static_multimap/heterogeneous_lookup_test.cu @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include #include @@ -105,12 +104,12 @@ TEMPLATE_TEST_CASE( cuco::empty_key{sentinel_key}, cuco::empty_value{sentinel_value}}; - auto insert_pairs = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto insert_pairs = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair(i, i); })); - auto probe_keys = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto probe_keys = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto i) { return ProbeKey(i); })); SECTION("All inserted keys-value pairs should be contained") diff --git a/tests/static_multimap/insert_contains_test.cu b/tests/static_multimap/insert_contains_test.cu index 834795a0a..a149e0f0b 100644 --- a/tests/static_multimap/insert_contains_test.cu +++ b/tests/static_multimap/insert_contains_test.cu @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include #include @@ -33,9 +32,9 @@ void test_insert(Map& map, std::size_t num_keys) using Key = typename Map::key_type; using Value = typename Map::mapped_type; - auto const keys_begin = thrust::counting_iterator{0}; - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto const keys_begin = cuda::counting_iterator{0}; + auto pairs_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type>( [] __device__(auto i) { return cuco::pair{i, i}; })); thrust::device_vector d_contained(num_keys); @@ -63,11 +62,11 @@ void test_insert(Map& map, std::size_t num_keys) map.contains_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator(0), + cuda::counting_iterator(0), is_even, d_contained.begin()); auto gold_iter = - thrust::make_transform_iterator(thrust::counting_iterator(0), is_even); + cuda::make_transform_iterator(cuda::counting_iterator(0), is_even); auto zip = thrust::make_zip_iterator(cuda::std::tuple{d_contained.begin(), gold_iter}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } diff --git a/tests/static_multimap/insert_if_test.cu b/tests/static_multimap/insert_if_test.cu index a5cc07e62..038dbee64 100644 --- a/tests/static_multimap/insert_if_test.cu +++ b/tests/static_multimap/insert_if_test.cu @@ -18,10 +18,10 @@ #include +#include #include #include #include -#include #include #include @@ -35,11 +35,11 @@ void test_insert_if(Map& map, std::size_t size) // 50% insertion auto const pred = [] __device__(Key k) { return k % 2 == 0; }; - auto const keys_begin = thrust::counting_iterator{0}; + auto const keys_begin = cuda::counting_iterator{0}; SECTION("Count of n / 2 insertions should be n / 2.") { - auto const pairs_begin = thrust::make_transform_iterator( + auto const pairs_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i, i}; })); @@ -53,7 +53,7 @@ void test_insert_if(Map& map, std::size_t size) SECTION("Inserting the same element n / 2 times should return n / 2.") { - auto const pairs_begin = thrust::constant_iterator>{{1, 1}}; + auto const pairs_begin = cuda::constant_iterator>{{1, 1}}; auto const num = map.insert_if(pairs_begin, pairs_begin + size, keys_begin, pred); REQUIRE(num * 2 == size); diff --git a/tests/static_multimap/multiplicity_test.cu b/tests/static_multimap/multiplicity_test.cu index fa63a4beb..58104781b 100644 --- a/tests/static_multimap/multiplicity_test.cu +++ b/tests/static_multimap/multiplicity_test.cu @@ -18,12 +18,11 @@ #include +#include #include #include #include #include -#include -#include #include #include @@ -36,9 +35,9 @@ void test_multiplicity_two(Map& map, std::size_t num_items) auto const num_keys = num_items / 2; - auto const keys_begin = thrust::counting_iterator{0}; + auto const keys_begin = cuda::counting_iterator{0}; // multiplicity = 2 - auto const pairs_begin = thrust::make_transform_iterator( + auto const pairs_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i / 2, i}; })); @@ -56,7 +55,7 @@ void test_multiplicity_two(Map& map, std::size_t num_items) REQUIRE(num == num_items); auto [_, output_end] = - map.retrieve(keys_begin, keys_begin + num_keys, thrust::discard_iterator{}, output_begin); + map.retrieve(keys_begin, keys_begin + num_keys, cuda::discard_iterator{}, output_begin); std::size_t const size = cuda::std::distance(output_begin, output_end); REQUIRE(size == num_items); diff --git a/tests/static_multimap/retrieve_if_test.cu b/tests/static_multimap/retrieve_if_test.cu index 44770395f..a1918fbaf 100644 --- a/tests/static_multimap/retrieve_if_test.cu +++ b/tests/static_multimap/retrieve_if_test.cu @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include #include @@ -119,12 +118,12 @@ TEMPLATE_TEST_CASE_SIG("static_multimap retrieve_if", container_type container{num_keys * 2, cuco::empty_key{-1}, cuco::empty_value{-1}}; - auto keys_begin = thrust::counting_iterator(1); - auto keys_end = keys_begin + num_keys; - auto pairs_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(1), - cuda::proclaim_return_type>( - [] __device__(Key i) { return cuco::pair{i, i}; })); + auto keys_begin = cuda::counting_iterator(1); + auto keys_end = keys_begin + num_keys; + auto pairs_begin = + cuda::make_transform_iterator(cuda::make_counting_iterator(1), + cuda::proclaim_return_type>( + [] __device__(Key i) { return cuco::pair{i, i}; })); container.insert(pairs_begin, pairs_begin + num_keys); diff --git a/tests/static_multimap/retrieve_test.cu b/tests/static_multimap/retrieve_test.cu index 23418f015..a5ff9085e 100644 --- a/tests/static_multimap/retrieve_test.cu +++ b/tests/static_multimap/retrieve_test.cu @@ -18,12 +18,11 @@ #include +#include #include #include #include #include -#include -#include #include #include @@ -44,9 +43,9 @@ void test_retrieve(Map& map, std::size_t num_items) auto const num_gold = num_items / 2; - auto const keys_begin = thrust::counting_iterator{0}; + auto const keys_begin = cuda::counting_iterator{0}; // multiplicity = 2 - auto const pairs_begin = thrust::make_transform_iterator( + auto const pairs_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i / 2, i / 2}; })); @@ -68,7 +67,7 @@ void test_retrieve(Map& map, std::size_t num_items) keys_begin + num_items, custom_key_eq{}, map.hash_function(), - thrust::discard_iterator{}, + cuda::discard_iterator{}, output_begin); std::size_t const size = cuda::std::distance(output_begin, output_end); @@ -84,7 +83,7 @@ void test_retrieve(Map& map, std::size_t num_items) return lhs.second < rhs.second; }); - auto const gold_begin = thrust::make_transform_iterator( + auto const gold_begin = cuda::make_transform_iterator( keys_begin, cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{(i / 2) * 2, (i / 2) * 2}; })); diff --git a/tests/static_multiset/contains_test.cu b/tests/static_multiset/contains_test.cu index 28bd8b843..82397fec2 100644 --- a/tests/static_multiset/contains_test.cu +++ b/tests/static_multiset/contains_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include #include @@ -67,11 +66,11 @@ void test_unique_sequence(Set& set, size_type num_keys) { set.contains_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator(0), + cuda::counting_iterator(0), is_even, d_contained.begin()); auto gold_iter = - thrust::make_transform_iterator(thrust::counting_iterator(0), is_even); + cuda::make_transform_iterator(cuda::counting_iterator(0), is_even); auto zip = thrust::make_zip_iterator(cuda::std::tuple{d_contained.begin(), gold_iter}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } diff --git a/tests/static_multiset/count_test.cu b/tests/static_multiset/count_test.cu index 72f60b8b6..797828e26 100644 --- a/tests/static_multiset/count_test.cu +++ b/tests/static_multiset/count_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,9 +19,8 @@ #include #include +#include #include -#include -#include #include @@ -32,8 +31,8 @@ void test_unique_sequence(Set& set, size_type num_keys) { using Key = typename Set::key_type; - auto keys_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto keys_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { return Key{i}; })); SECTION("Count of empty set should be zero.") @@ -51,8 +50,8 @@ void test_unique_sequence(Set& set, size_type num_keys) } auto constexpr multiplicity = 3; - auto query_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto query_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { return Key{i / multiplicity}; })); SECTION("Count of 3n unique keys should be 3n.") @@ -70,8 +69,8 @@ void test_count_each(Set& set, size_type num_keys) thrust::device_vector d_counts(num_keys); auto const counts_begin = d_counts.begin(); - auto keys_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto keys_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { return Key{i}; })); set.clear(); @@ -101,12 +100,12 @@ void test_count_each(Set& set, size_type num_keys) set.clear(); auto constexpr multiplicity = 3; - auto duplicate_keys_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto duplicate_keys_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { return Key{i / multiplicity}; })); set.insert(duplicate_keys_begin, duplicate_keys_begin + num_keys); - auto const query_begin = thrust::counting_iterator{0}; + auto const query_begin = cuda::counting_iterator{0}; auto const query_size = num_keys / multiplicity; SECTION("Count_each with duplicates should return correct counts.") { @@ -128,8 +127,8 @@ void test_count_each_outer(Set& set, size_type num_keys) thrust::device_vector d_counts(num_keys); auto const counts_begin = d_counts.begin(); - auto keys_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto keys_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { return Key{i}; })); set.clear(); @@ -159,13 +158,13 @@ void test_count_each_outer(Set& set, size_type num_keys) set.clear(); auto constexpr multiplicity = 3; - auto duplicate_keys_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, + auto duplicate_keys_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { return Key{i / multiplicity}; })); set.insert(duplicate_keys_begin, duplicate_keys_begin + num_keys); auto const query_size = num_keys / multiplicity; - auto const query_begin = thrust::counting_iterator{0}; + auto const query_begin = cuda::counting_iterator{0}; SECTION("Count_each_outer with duplicates should return correct counts.") { diff --git a/tests/static_multiset/custom_count_test.cu b/tests/static_multiset/custom_count_test.cu index f5ade5eeb..2cded9396 100644 --- a/tests/static_multiset/custom_count_test.cu +++ b/tests/static_multiset/custom_count_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -69,8 +70,8 @@ void test_custom_count(Set& set, size_type num_keys) } }(); - auto query_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), + auto query_begin = cuda::make_transform_iterator( + cuda::make_counting_iterator(0), cuda::proclaim_return_type([] __device__(auto i) { return static_cast(i * XXX); })); SECTION("Count of empty set should be zero.") @@ -85,7 +86,7 @@ void test_custom_count(Set& set, size_type num_keys) REQUIRE(count == num_keys); } - auto const iter = thrust::counting_iterator{0}; + auto const iter = cuda::counting_iterator{0}; set.insert(iter, iter + num_keys); SECTION("Count of n unique keys should be n.") @@ -101,7 +102,7 @@ void test_custom_count(Set& set, size_type num_keys) } set.clear(); // reset the set - auto const constants = thrust::constant_iterator{1}; + auto const constants = cuda::constant_iterator{1}; set.insert(constants, constants + num_keys); // inser the same value `num_keys` times SECTION("Count of a key whose multiplicity equals n should be n.") diff --git a/tests/static_multiset/find_test.cu b/tests/static_multiset/find_test.cu index e0c8e6825..ecb763017 100644 --- a/tests/static_multiset/find_test.cu +++ b/tests/static_multiset/find_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,10 +19,10 @@ #include #include +#include #include #include #include -#include #include #include @@ -37,7 +37,7 @@ void test_unique_sequence(Set& set, size_type num_keys) { using Key = typename Set::key_type; - auto keys_begin = thrust::counting_iterator{0}; + auto keys_begin = cuda::counting_iterator{0}; thrust::device_vector d_contained(num_keys); auto zip_equal = cuda::proclaim_return_type( @@ -48,8 +48,8 @@ void test_unique_sequence(Set& set, size_type num_keys) thrust::device_vector d_results(num_keys); set.find(keys_begin, keys_begin + num_keys, d_results.begin()); - auto zip = thrust::make_zip_iterator(cuda::std::tuple{ - d_results.begin(), thrust::constant_iterator{set.empty_key_sentinel()}}); + auto zip = thrust::make_zip_iterator( + cuda::std::tuple{d_results.begin(), cuda::constant_iterator{set.empty_key_sentinel()}}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } @@ -76,14 +76,14 @@ void test_unique_sequence(Set& set, size_type num_keys) set.find_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator{0}, + cuda::counting_iterator{0}, is_even, found_results.begin()); REQUIRE(cuco::test::equal( found_results.begin(), found_results.end(), - thrust::make_transform_iterator(thrust::counting_iterator{0}, gold_fn), + cuda::make_transform_iterator(cuda::counting_iterator{0}, gold_fn), cuda::proclaim_return_type( [] __device__(auto const& found, auto const& gold) { return found == gold; }))); } diff --git a/tests/static_multiset/for_each_test.cu b/tests/static_multiset/for_each_test.cu index 6d663f439..43b22eee9 100644 --- a/tests/static_multiset/for_each_test.cu +++ b/tests/static_multiset/for_each_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -21,8 +21,7 @@ #include #include -#include -#include +#include #include #include @@ -116,10 +115,10 @@ TEMPLATE_TEST_CASE_SIG( auto set = cuco::static_multiset{num_keys, cuco::empty_key{-1}, {}, probe{}, {}, cuco::storage<2>{}}; - auto unique_keys_begin = thrust::counting_iterator(0); + auto unique_keys_begin = cuda::counting_iterator(0); auto gen_duplicate_keys = cuda::proclaim_return_type( [] __device__(auto const& k) { return static_cast(k % num_unique_keys); }); - auto keys_begin = thrust::make_transform_iterator(unique_keys_begin, gen_duplicate_keys); + auto keys_begin = cuda::make_transform_iterator(unique_keys_begin, gen_duplicate_keys); set.insert(keys_begin, keys_begin + num_keys); diff --git a/tests/static_multiset/insert_test.cu b/tests/static_multiset/insert_test.cu index c0d21848f..0df7db4d2 100644 --- a/tests/static_multiset/insert_test.cu +++ b/tests/static_multiset/insert_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,9 +19,8 @@ #include #include +#include #include -#include -#include #include #include @@ -37,7 +36,7 @@ void test_insert(Set& set) SECTION("Inserting 300 unique keys should get 300 entries in the multiset") { - auto const keys = thrust::counting_iterator{0}; + auto const keys = cuda::counting_iterator{0}; set.insert(keys, keys + num); auto const num_insertions = set.size(); @@ -46,7 +45,7 @@ void test_insert(Set& set) SECTION("Inserting one key for 300 times should get 300 entries in the multiset") { - auto const keys = thrust::constant_iterator{0}; + auto const keys = cuda::constant_iterator{0}; set.insert(keys, keys + num); auto const num_insertions = set.size(); @@ -58,7 +57,7 @@ void test_insert(Set& set) SECTION("Inserting all even values between [0, 300) should get 150 entries in the multiset") { - auto const keys = thrust::counting_iterator{0}; + auto const keys = cuda::counting_iterator{0}; set.insert_if(keys, keys + num, keys, is_even); auto const num_insertions = set.size(); @@ -67,8 +66,8 @@ void test_insert(Set& set) SECTION("Conditionally inserting one key for 150 times should get 150 entries in the multiset") { - auto const keys = thrust::constant_iterator{0}; - set.insert_if(keys, keys + num, thrust::counting_iterator{0}, is_even); + auto const keys = cuda::constant_iterator{0}; + set.insert_if(keys, keys + num, cuda::counting_iterator{0}, is_even); auto const num_insertions = set.size(); REQUIRE(num_insertions == num / 2); diff --git a/tests/static_multiset/large_input_test.cu b/tests/static_multiset/large_input_test.cu index 4cc6c19b4..37e2dedcb 100644 --- a/tests/static_multiset/large_input_test.cu +++ b/tests/static_multiset/large_input_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -18,10 +18,9 @@ #include +#include #include #include -#include -#include #include #include @@ -34,7 +33,7 @@ void test_unique_sequence(Set& set, typename Set::value_type* res_begin, std::si { using Key = typename Set::key_type; - auto const keys_begin = thrust::counting_iterator(0); + auto const keys_begin = cuda::counting_iterator(0); auto const keys_end = keys_begin + num_keys; set.insert(keys_begin, keys_end); @@ -43,7 +42,7 @@ void test_unique_sequence(Set& set, typename Set::value_type* res_begin, std::si SECTION("All inserted keys can be retrieved.") { auto const [_, res_end] = - set.retrieve(keys_begin, keys_end, thrust::make_discard_iterator(), res_begin); + set.retrieve(keys_begin, keys_end, cuda::make_discard_iterator(), res_begin); REQUIRE(static_cast(std::distance(res_begin, res_end)) == num_keys); thrust::sort(thrust::device, res_begin, res_end); diff --git a/tests/static_multiset/retrieve_if_test.cu b/tests/static_multiset/retrieve_if_test.cu index 68286a52d..19defd557 100644 --- a/tests/static_multiset/retrieve_if_test.cu +++ b/tests/static_multiset/retrieve_if_test.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. @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -114,7 +115,7 @@ TEMPLATE_TEST_CASE_SIG( container_type container{num_keys * 2, cuco::empty_key{-1}}; - auto keys_begin = thrust::counting_iterator(1); + auto keys_begin = cuda::counting_iterator(1); auto keys_end = keys_begin + num_keys; container.insert(keys_begin, keys_end); diff --git a/tests/static_multiset/retrieve_test.cu b/tests/static_multiset/retrieve_test.cu index deda194c4..acac841df 100644 --- a/tests/static_multiset/retrieve_test.cu +++ b/tests/static_multiset/retrieve_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,12 +19,10 @@ #include #include +#include #include #include #include -#include -#include -#include #include #include @@ -47,8 +45,8 @@ void test_multiplicity(Container& container, std::size_t num_keys, std::size_t m thrust::device_vector probed_keys(num_actual_keys); thrust::device_vector matched_keys(num_actual_keys); - auto const keys_begin = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto const keys_begin = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type([multiplicity] __device__(auto const& i) { return static_cast(i / multiplicity); })); @@ -77,7 +75,7 @@ void test_outer(Container& container, std::size_t num_keys) container.clear(); - auto const keys_begin = thrust::counting_iterator{0}; + auto const keys_begin = cuda::counting_iterator{0}; auto const query_size = num_keys * 2ull; thrust::device_vector probed_keys(num_keys * 2ull); diff --git a/tests/static_multiset/stream_test.cu b/tests/static_multiset/stream_test.cu index fe9db3f7b..09ceb88f3 100644 --- a/tests/static_multiset/stream_test.cu +++ b/tests/static_multiset/stream_test.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. @@ -19,11 +19,11 @@ #include #include +#include #include #include #include #include -#include #include #include @@ -92,8 +92,7 @@ TEMPLATE_TEST_CASE_SIG("static_multiset: operations on different stream than con thrust::device_vector d_results(num_keys); multiset.find(d_keys.begin(), d_keys.end(), d_results.begin(), operation_stream); - auto zip = - thrust::make_zip_iterator(cuda::std::make_tuple(d_results.begin(), d_keys.begin())); + auto zip = thrust::make_zip_iterator(cuda::std::tuple{d_results.begin(), d_keys.begin()}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, cuda::proclaim_return_type([] __device__(auto const& p) { diff --git a/tests/static_set/atomic_storage_test.cu b/tests/static_set/atomic_storage_test.cu index f67f6e494..b611214bb 100644 --- a/tests/static_set/atomic_storage_test.cu +++ b/tests/static_set/atomic_storage_test.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. @@ -17,9 +17,8 @@ #include #include +#include #include -#include -#include #include @@ -66,7 +65,8 @@ TEST_CASE("atomic_storage_test", "") {}, cuco::storage<1>{}}; - auto keys_begin = thrust::make_transform_iterator(thrust::counting_iterator{0}, build_fn{}); + auto keys_begin = cuda::make_transform_iterator(cuda::counting_iterator{0}, + cuda::proclaim_return_type(build_fn{})); set.insert_async(keys_begin, keys_begin + num_keys); auto const count = set.size(); diff --git a/tests/static_set/for_each_test.cu b/tests/static_set/for_each_test.cu index 4262c445d..7db5d0d5f 100644 --- a/tests/static_set/for_each_test.cu +++ b/tests/static_set/for_each_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include @@ -38,8 +37,8 @@ void test_for_each(Set& set, size_type num_keys) cuda::stream_ref stream{cudaStream_t{nullptr}}; // Insert keys - auto keys_begin = thrust::make_transform_iterator( - thrust::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { + auto keys_begin = cuda::make_transform_iterator( + cuda::counting_iterator{0}, cuda::proclaim_return_type([] __device__(auto i) { // generates a sequence of 1, 2, 1, 2, ... return static_cast(i); })); @@ -62,8 +61,8 @@ void test_for_each(Set& set, size_type num_keys) // count the sum of all odd keys set.for_each( - thrust::counting_iterator(0), - thrust::counting_iterator(2 * num_keys), // test for false-positives + cuda::counting_iterator(0), + cuda::counting_iterator(2 * num_keys), // test for false-positives [counter = counter_storage.data()] __device__(auto const slot) { if (!(slot % 2 == 0)) { counter->fetch_add(slot, cuda::memory_order_relaxed); } }, diff --git a/tests/static_set/heterogeneous_lookup_test.cu b/tests/static_set/heterogeneous_lookup_test.cu index e691a2110..57a1b291a 100644 --- a/tests/static_set/heterogeneous_lookup_test.cu +++ b/tests/static_set/heterogeneous_lookup_test.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. @@ -19,10 +19,9 @@ #include #include +#include #include #include -#include -#include #include #include @@ -101,11 +100,11 @@ TEMPLATE_TEST_CASE_SIG("static_set heterogeneous lookup tests", auto my_set = cuco::static_set{capacity, cuco::empty_key{sentinel_key}, custom_key_equal{}, probe}; - auto insert_keys = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto insert_keys = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto i) { return InsertKey(i); })); - auto probe_keys = thrust::make_transform_iterator( - thrust::counting_iterator(0), + auto probe_keys = cuda::make_transform_iterator( + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto i) { return ProbeKey(i); })); SECTION("All inserted keys should be contained") diff --git a/tests/static_set/insert_and_find_test.cu b/tests/static_set/insert_and_find_test.cu index 914e99621..f01ae8884 100644 --- a/tests/static_set/insert_and_find_test.cu +++ b/tests/static_set/insert_and_find_test.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. @@ -18,9 +18,9 @@ #include +#include #include #include -#include #include @@ -30,8 +30,8 @@ void test_insert_and_find(Set& set, std::size_t num_keys) using Key = typename Set::key_type; static auto constexpr cg_size = Set::cg_size; - auto const keys_begin = thrust::counting_iterator(0); - auto const keys_end = thrust::counting_iterator(num_keys); + auto const keys_begin = cuda::counting_iterator(0); + auto const keys_end = cuda::counting_iterator(num_keys); thrust::device_vector iters1(num_keys); thrust::device_vector iters2(num_keys); diff --git a/tests/static_set/large_input_test.cu b/tests/static_set/large_input_test.cu index f0fafa995..67daeafc6 100644 --- a/tests/static_set/large_input_test.cu +++ b/tests/static_set/large_input_test.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. @@ -18,10 +18,10 @@ #include +#include #include #include #include -#include #include #include #include @@ -33,8 +33,8 @@ void test_unique_sequence(Set& set, bool* res_begin, std::size_t num_keys) { using Key = typename Set::key_type; - auto const keys_begin = thrust::counting_iterator(0); - auto const keys_end = thrust::counting_iterator(num_keys); + auto const keys_begin = cuda::counting_iterator(0); + auto const keys_end = cuda::counting_iterator(num_keys); SECTION("Non-inserted keys should not be contained.") { @@ -64,7 +64,7 @@ void test_unique_sequence(Set& set, bool* res_begin, std::size_t num_keys) REQUIRE(cuco::test::equal(output_keys.begin(), output_keys.end(), - thrust::counting_iterator(0), + cuda::counting_iterator(0), cuda::std::equal_to{})); } } diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu index 9deb02984..ed187032f 100644 --- a/tests/static_set/retrieve_all_test.cu +++ b/tests/static_set/retrieve_all_test.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. @@ -18,9 +18,9 @@ #include +#include #include #include -#include #include #include @@ -52,7 +52,7 @@ void test_unique_sequence(Set& set, std::size_t num_keys) auto d_res_end = set.retrieve_all(d_res.begin()); thrust::sort(d_res.begin(), d_res_end); REQUIRE(cuco::test::equal( - d_res.begin(), d_res_end, thrust::counting_iterator(0), cuda::std::equal_to{})); + d_res.begin(), d_res_end, cuda::counting_iterator(0), cuda::std::equal_to{})); } } diff --git a/tests/static_set/retrieve_if_test.cu b/tests/static_set/retrieve_if_test.cu index 2ee582545..129fe68af 100644 --- a/tests/static_set/retrieve_if_test.cu +++ b/tests/static_set/retrieve_if_test.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. @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -113,7 +114,7 @@ TEMPLATE_TEST_CASE_SIG("static_set retrieve_if", "", ((typename Key), Key), (int container_type container{num_keys * 2, cuco::empty_key{-1}}; - auto keys_begin = thrust::counting_iterator(1); + auto keys_begin = cuda::counting_iterator(1); auto keys_end = keys_begin + num_keys; container.insert(keys_begin, keys_end); diff --git a/tests/static_set/retrieve_test.cu b/tests/static_set/retrieve_test.cu index fd3f81b04..edbcb4c20 100644 --- a/tests/static_set/retrieve_test.cu +++ b/tests/static_set/retrieve_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-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. @@ -18,10 +18,9 @@ #include +#include #include #include -#include -#include #include #include #include @@ -38,7 +37,7 @@ void test_unique_sequence(Set& set, std::size_t num_keys) thrust::device_vector keys(num_keys); thrust::device_vector matched_keys(num_keys); - auto iter = thrust::counting_iterator{0}; + auto iter = cuda::counting_iterator{0}; SECTION("Non-inserted keys should not be contained.") { @@ -63,10 +62,10 @@ void test_unique_sequence(Set& set, std::size_t num_keys) thrust::sort(keys.begin(), probe_end); thrust::sort(matched_keys.begin(), matched_end); REQUIRE(cuco::test::equal( - keys.begin(), keys.end(), thrust::counting_iterator(0), cuda::std::equal_to{})); + keys.begin(), keys.end(), cuda::counting_iterator(0), cuda::std::equal_to{})); REQUIRE(cuco::test::equal(matched_keys.begin(), matched_keys.end(), - thrust::counting_iterator(0), + cuda::counting_iterator(0), cuda::std::equal_to{})); } } diff --git a/tests/static_set/stream_test.cu b/tests/static_set/stream_test.cu index 142ebb075..47e96283f 100644 --- a/tests/static_set/stream_test.cu +++ b/tests/static_set/stream_test.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. @@ -19,11 +19,11 @@ #include #include +#include #include #include #include #include -#include #include #include @@ -72,8 +72,7 @@ TEMPLATE_TEST_CASE_SIG("static_set: operations on different stream than construc thrust::device_vector d_results(num_keys); set.find(d_keys.begin(), d_keys.end(), d_results.begin(), operation_stream); - auto zip = - thrust::make_zip_iterator(cuda::std::make_tuple(d_results.begin(), d_keys.begin())); + auto zip = thrust::make_zip_iterator(cuda::std::tuple{d_results.begin(), d_keys.begin()}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, cuda::proclaim_return_type([] __device__(auto const& p) { diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu index 0c5ef433f..e43329b6b 100644 --- a/tests/static_set/unique_sequence_test.cu +++ b/tests/static_set/unique_sequence_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2025, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -19,11 +19,10 @@ #include #include +#include #include #include #include -#include -#include #include #include @@ -38,7 +37,7 @@ void test_unique_sequence(Set& set, size_type num_keys) { using Key = typename Set::key_type; - auto keys_begin = thrust::counting_iterator{0}; + auto keys_begin = cuda::counting_iterator{0}; thrust::device_vector d_contained(num_keys); auto zip_equal = cuda::proclaim_return_type( @@ -59,8 +58,8 @@ void test_unique_sequence(Set& set, size_type num_keys) thrust::device_vector d_results(num_keys); set.find(keys_begin, keys_begin + num_keys, d_results.begin()); - auto zip = thrust::make_zip_iterator(cuda::std::tuple{ - d_results.begin(), thrust::constant_iterator{set.empty_key_sentinel()}}); + auto zip = thrust::make_zip_iterator( + cuda::std::tuple{d_results.begin(), cuda::constant_iterator{set.empty_key_sentinel()}}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } @@ -68,7 +67,7 @@ void test_unique_sequence(Set& set, size_type num_keys) SECTION("All conditionally inserted keys should be contained") { auto const inserted = set.insert_if( - keys_begin, keys_begin + num_keys, thrust::counting_iterator(0), is_even); + keys_begin, keys_begin + num_keys, cuda::counting_iterator(0), is_even); REQUIRE(inserted == num_keys / 2); REQUIRE(set.size() == num_keys / 2); @@ -76,7 +75,7 @@ void test_unique_sequence(Set& set, size_type num_keys) REQUIRE(cuco::test::equal( d_contained.begin(), d_contained.end(), - thrust::counting_iterator(0), + cuda::counting_iterator(0), cuda::proclaim_return_type([] __device__(auto const& idx_contained, auto const& idx) { return ((idx % 2) == 0) == idx_contained; }))); @@ -97,11 +96,11 @@ void test_unique_sequence(Set& set, size_type num_keys) { set.contains_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator(0), + cuda::counting_iterator(0), is_even, d_contained.begin()); auto gold_iter = - thrust::make_transform_iterator(thrust::counting_iterator(0), is_even); + cuda::make_transform_iterator(cuda::counting_iterator(0), is_even); auto zip = thrust::make_zip_iterator(cuda::std::tuple{d_contained.begin(), gold_iter}); REQUIRE(cuco::test::all_of(zip, zip + num_keys, zip_equal)); } @@ -124,14 +123,14 @@ void test_unique_sequence(Set& set, size_type num_keys) set.find_if(keys_begin, keys_begin + num_keys, - thrust::counting_iterator{0}, + cuda::counting_iterator{0}, is_even, found_results.begin()); REQUIRE(cuco::test::equal( found_results.begin(), found_results.end(), - thrust::make_transform_iterator(thrust::counting_iterator{0}, gold_fn), + cuda::make_transform_iterator(cuda::counting_iterator{0}, gold_fn), cuda::proclaim_return_type( [] __device__(auto const& found, auto const& gold) { return found == gold; }))); }