diff --git a/README.md b/README.md index 444113097..633757272 100644 --- a/README.md +++ b/README.md @@ -208,6 +208,7 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::static_set` is a fixed-size container that stores unique elements in no particular order. See the Doxygen documentation in `static_set.cuh` for more detailed information. #### Examples: +- [Heterogeneous Lookup](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_set/heterogeneous_lookup_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVw1v4zYS_StzOvRgJ7KdZLto4XygvuwWZ3SRFEm2RbEOFFoaW8RKpJai7PgM__ebISXZTtbda23AFsnhzJs3M-RoHZRYllKrMhh-WgcyCYanYZAJNa_EHINhEFeJCMKg1JWJeTw4mig4gmtdrIycpxY6cRfOTs7ehnDz2_jdeATXt3e_3t6NHsa3N32WdfIfZIyqxAQqlaABmyKMChHTX70Swm9oGAic9U-gwwKToF6bBN1zp2WlK8jFCpS2UJVIamQJM5kh4HOMhQWpINZ5kUmhYoSltKkzVetxcOCPWomeWkHygnYUNJrtSoKwLXT-pNYWw8FguVz2hYPd12Y-yLxwOfgwvn5_c_--R9DbbR9VRsyCwS-VNOT4dAWiIGSxmBLeTCxBGxBzg7RmNSNfGmmlmodQ6pldCoNOTyJLa-S0snvkNTjJ_10Bok8oIm50D-P7SQD_Ht2P70On5_fxw39uPz7A76O7u9HNw_j9PdzeUbBu3o05VDT6GUY3f8Av45t3ISBRR6bwuTDsBUGVTCsmnsN7xD0YM-1hlQXGciZjaDII5nqBRpFbUKDJpc81Apk4PZnMpRXWzb1yzpkaTNRE_VOqOKsShIu4ivWg5C1xVKLtx1V69VIiESSRDJgVNY8WEpdXuwI2NVVpBwkuyEy0wNhq03-pRWrajiJ304MjH9KfXKqlaNHoOSrUVRllWn-uigifBfGDhKdNgJ-mRuIM3mFO3lkjLJaQ6iUHm3M3FWUK5AIIogO8GrAuOepc1FRgUokMEmGFT2bvE3zGVbm1Q9WA8MClUKPgrOCIJK1prquiMoWmbKUkyVZ9GFsW40qSyiJx79KwdAjJeE4EOPUUNtKVC2VhSVVDMol2Eo0xkc0JqU3zbbxawjwFH0s0vQRnUpERwg5UE-QWqeJipUQHu9StU6SBMFexhXwVeck1zbejDkd4OKQQ808b4yNHUxeGEPHDmn82sN7wVo8HtoiuU4w_lyBnrWVgGDJBRZlFFqei9OVEnkpKfpVQspRcpTRR69qPjzBGrPzKgP-iqM6wiLwkj7iUYKp1BpooFZR0nW6HqH9zFlnI0jKEZmDSsuv3sB7nO9A5Yiuj3InXu9rR4Jz9RPsfQ7eRD0rYcBD-TwhfZ5MRRTT-BipRURCNF4VLz_wnGj-eewRE_QB-lqa0EDPlzHiGam5TT7g7M0TuGeXFTm23X8r_YqcL_7hs1Ncz3YaJmchK3DXzkKJqEwrpkCaTwoiYqtXLcEl0WAt5IwntyTn9XcC-xXM4Ppbdxr89TJ_k4w4eGr3Cwhs2DaQmYqbCJigAh3LXMUdim_Nv1Y87N2aVil1RE_1u4vV50WawtJy3CT5vTwxLHIkcHuC22cJVaVcFvipA1k7HclOBfvg3SnAnF6smpb5SCA7ogTzbzTEn93j-57VBGw7Vw9cwHKgE1vIKULvfsc-51ObhjZALbE7qvWh5kXoHp9DJIxzDTu69SB2W_KupQ7iou5Gqs9Ay6U6Uw-urY5sgTbT5BhA11BCWqaSioamFMNJ3Ka5U-bamC93M3fVM_cWP1MzQbcaaq5K9_IX0sUcttNoTFzapisrCpffMX77D4d7te_F1167Wk4BbJur_Qth7nApFX_9M54oxqz2Rza55Fzlgggmiw7LPds0NnRPVS2aUT0c3EcJTHfInXnN0ACpdzVN_kS3q5mWZar6BndOGuiCsjTCJ3_eYOGoudU490jV1aaQq13xezahrknT7eIZZ0Y9eWCvXCNVMjxxMz3eN57w28NQ7ffIBXYhMJtRb0F2mMONhRQkpuSUmbdRJZBkoZMAijrm9e3I11Tt9fGJdWxvbWwPzwq4iYidq1V5C7_Q10awpKlylerJ5otPtz9F2fDH6WqbG5xK4m-OQN-3c2k_gMzUk9sLngzuvr9YugEdwtgk9m1NN_L49-Y46J5HQAcx51JzZhz619saV9WunNuG3dDRdyLpx9NtbvNmM9AsTFUZPKZAXp-H2OL1at49btZvNoQSte5W6JaFmETPMyYHyRSj8jsjtuDxQeW2sr9YnIRCosxDehPD9xoWKO2xJnbixnR1l_SnSMdLphrsW-tRC0u28mxBfKjSrg5YP13y2oprGP63uPT-dnaiucDfYr_BGlF_36Eo8BIgbo6vOVlm3paDeW_rFrft-6BwPtwaa9ca8azu2cP9FFzLdkK1422s4KmJO64sLlqHfSTAhHcHuHVZfDXzfbOi9nPsdeikx29f3QC3i-PTsbXVKy7qw_t0-6JH6y_j4-PQH6AkTp5dlHv1wAr0elbftuYqjV4BeJvKpe-HP5HRHZxzHGU0u_Cs6TVB_oD4Hm7BZpxTfWydqg82j-_4PrwCRKw==)) - [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_set/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJx9VgtvIjcQ_ivTraqSu-UVKTqJPFSapCq6EzmF3J1OpSLGa1grxqZ-wFHEf--MvcvjHk0kYD3jmW9mvs_rbeaEc9Jol_X-2mayyHrdPFNMzwObi6yX8VCwLM-cCZbTc_vVWMMruDXLjZXz0kODn8F55_y8iR8XOQw_Du4Gfbh9eHz_8Nh_GjwMW7QhbnonudBOFBB0ISz4UkB_yTh-VZYcPgpLaOC81YEGOYyzyjbOzi5jlI0JsGAb0MZDcALDSAczqQSIL1wsPUgN3CyWSjLNBaylL2OqKk6EA5-rIGbqGfoz3LHEp9mxJzC_h05_pffLXru9Xq9bLMJuGTtvq-Ts2u8Gt_fD0X0Toe-3fdAK2wtW_BOkxcKnG2BLRMbZFPEqtgZjgc2tQJs3hHxtpZd6noMzM79mVsQ4hXTeymnwJ82rcWL9xw7YPqaxcf0RDEbjDH7vjwajPMb5NHj68-HDE3zqPz72h0-D-xE8POKwhncDGhU-_QH94Wd4Oxje5SCwdZhKfFlaqgKhSmqrKFIPR0KcwJiZBMstBZczyaGmEczNSliNZcFS2IVMhEOQRYyj5EJ65uPaN8XFVO2xHuufpeYqFAKueOCm7WgLnzjhWzyUN197FAw9ivYsaE6Rmbo5tvvSBufbhVhhlslKcG9sq_yeizJznJb6vtHhXAWSrPV1fmlwHIItTjbFOl30bL9K_Pgt8rZE78k0qJeJ-MKwwwIrSuaplWIGd2KBvfGWeYEdctTRiqmHJuC8KQKOm6JB__3AHYT3hK5kjcuApEKOqA3SbWUSE2fWLGLAuBkHE50CaZWmWhiancHhVWNS8kXAs8T5WP-MxIhZnrnRpCb3HCkIhApxvojNAQpOUmqP0pO6sTKyOBvrLa5jKsrwVmzgGmH5S2oSQLsN94ul34BTxifgVhAZhfZJTfTbrvD3OKNFqYXCFqyYCsK1qG7kZXoCV5qgcBN2TSicOG5ygZdVHl8yTx3Ac0UgWcFwHiwJEo8b-l4GDwXzjCqJODmNhKQBgiBOsMxJDQGLaHYPNQzDYkohUy9I6AgiNU-QBHCMRa_n5L9i4o_i6rCYxA3XcNH5tdPpHCLe4vmGWgdOB5HEBk2Zq5SPvr-AMgwnx4jWtIMFzHkITNZJsmLsTuvi8rsgDtGvk5ELqRp7VO3jOGfH2IirgdPAIgfiCYzdVYJh0Oc66nM11TT6ZheYixSM7UydqvsZm06qR4B7xl_hFG4owbaOmFc--4Fsvx3NbneEFCWKTSSUScf7EW07OXRzOM-h1cpB7mhD0nyvd3JoJBC0Z9-Y-Jrae9ehG1EGUzFH4p_lSRRCF_S7c9S8QWQFdk6phATFYJIymSup2jgqPPQSf34Y9ijmCHHSmUFKRr0E5X9cztQYdYOeeA6fFFR3rBT8BeTsgI80WekeCYh6icfSAWZ9JvxP_TFbbTpKhncBOdskZe7zlWyFZ5kQOm0jV8JcMRbdIuRJNCJv6yrJYGaNk1R15goGvTOIX0h0WRBb_Ga7q-FgyY3T4GewrVRh8Gi4usITaBQ4xxflT_BHzF5jbo0xRnYJuxTKCh-sBlIzruDdii4r-BqwhytYplecd88vQhfNZunT_SxrYr5r_vp19w00meXltVtM3nSg2URVe_zwWIkomootpvHSpuT0KCbnXOHiKt2wcAFFql-yXV7bUTsndmRGtvs7_v8HAzB-mg==)) - [Device-ref APIs for individual operations](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_set/device_ref_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJztWG1v4zYS_itTFcU5G_klQfdaeJPgfMkWNVokRZJtUdSFSlO0xYssqiJlxw3y3-8ZUrLsTbZogft4CZBYmuEzz7xwOPRTZJW12hQ2Gv_yFOk0Gp_EUS6KZS2WKhpHsk5FFEfW1JXk5-GbWUFv6NKU20ovM0c9eUSno9PTPv68jen6x-nVdEKXN7c_3NxO7qc31wNe4Bd9r6UqrEqpLlJVkcsUTUoh8a-RxPSjqpgNnQ5G1GOFWdTIZtHRO4-yNTWtxJYK46i2CjDa0kLnitSjVKUjXZA0qzLXopCKNtpl3lSD4-nQzw2ImTsBfYEVJZ4W-5ok3I46_2TOlePhcLPZDISnPTDVcpgHZTv8fnr5_vrufR_Ud8s-FDnCS5X6vdYVHJ9vSZRgJsUcfHOxIVORWFYKMmeY-abSThfLmKxZuI2olMdJtXWVntfuIHgtT_i_r4DwiQKBm9zR9G4W0b8nd9O72OP8NL3_9ubDPf00ub2dXN9P39_RzS2SdX015VTh6RuaXP9M302vr2JSCB1MqceyYi9AVXNYVRpieKfUAY2FCbRsqaReaEltGdHSrFVVwC0qVbXSoeBAMvU4uV5pJ5x_98I5b2o4K2bF57qQeZ0qOpO1NEPLS2RilRvIOrv4WCMV0EiHi7qQjCzyi325y6raumGq1rCSrJV0phpkr6nkZols5a8LLfKqUGSDF_aNgacguFbJsjJ1aV-qgF6qFgew2iCLSqy86vBNqKF_-dpuqFZqoSo2mahHgWQoOB-05pVWC7pSK4TRwbRCMC0HvynqLl4NVt9q2Jz8MLXdDt1TKiuzhgLShJ1W9M3G529nnzaZlhlJVNpc8T5sCtjBben8rvOIbFqagncZEruozMrLUOmBBYTpXpKHQ2BYVTmYZRYg_6C2loENXmXCZv59bZmNKXj9LtTkQ-3rUGF_khP2YQZGiBPigaxtS1WIlULlulu1iGn3YlqUtZsydxQCop8ky9zMRZ4ktDY6JYlsm1Wyn9ZAsxegmFN8iOJ5Yxu7dDy2-g-VOCqOZsUTXCW2adFCwH5J5_SyWt5xLIgDZx1vPxI1_JfLhJGwIlgdj5s3jbpXclwtwFyOx_wxTUpROc174KzRvuh5IdpmgjpWIk3muZEPvaOjBgfB_ueX4OvNUw52CXeXlHGX-H-lV4NH5NYvCw9D6qh0ADp9pMOfc-rtrWogpukjHo4psPFPRweIjIl6g2M9hjxDJMkHkjjwgyYV7G5Mb3q-YI7ZuPeItXjV8fm-K17wPCueGf1vl8je-5va_cXa8ZvA_r2KiT_CR3GjRf6_jv7ndTQ3Jm-Iwp-2Y6Uw5E-YNnl_VmML8uJBE4xKFIgFnZ_TCFZC4n7Bgl-BuW_iHZfhXyhSBAMjiy56XFVtBYQ--J3aMmjhGh_RRd-vSrclmxuHDl4pNG4-xFXhwhTCn6s1Ps8ifgkeOSaFtchrZQd0n0HcPJHNTJ2n3OWtynFSYpGtZdbYcZlw3OQxjykc8mSkrCseZDCm8X_ULqXCCW7xnmdXi4opJohk0lKAE_2TzofrejVnyOYMQFGCRNjsikeHg63S4Rb1KvELzunt6B-j0ahDvMRciBkJxxZ2jUaA5sI2ExN0v0DsRUoLwePAbiN0wCxNghTYo8Hbd6-S6NDPg1Aqnfd2rIb7OEf73PjgriUnrDvo_PiKEOdKAHkWtdjIVkhuqID-CQnrD1sfVQ4YtVH1oeeZCTR35_sZcnHBFp5axLjR2aXl6WWCnp_3-KLKEUo-p8MUtEvU0yimk5hOYxoMYtK-vMPENB4fjFyBBK_ZhSfspla7hfb7bTBXS5T_UexXDFSR8ufRXginYXTgMCx0hXhlIl-0A1A3R_hxSDmfvO74eMVE84SNvpe909cNWoXkpzuLXjnkRjRHAF1-uJpAUHU5-cRUcXZ2dnL6NYJ4-vXFxUWvO-kwfvVCloLinxONP0X7DrHn0ZAnJXSCOnefThE3xovQvQ6S1FZBpuQD9z6R58Fp7jaF2VDXQ9ENmpAP6No4FZrGRlEpLBd7N1R6gci5f255D-yajKLf2h78G4VwYWLfZYGWCuOlyOMm5G5j-LqyMr7ztehm_h_F26stAT6-u9kUl6d2bGVQgTvVwq907STa2rW8Cm2OZfmWr198bVl4X-cqE2vdcbvn22lq0Ef5vsqXPzgGVAJPGe5PTDZUBWt6Pb4v5aI8rJLg_ke1cVgVrdJHddEVQhwy2Qr2Eombt15s2ww0uYQvCi6pIixj1b2DEmoeNPFCdLu2glhgFr0DU63lZtvyDY37EdojDjp0F7d9et4NDjhKD8H5BA291CBNZ2fohHe1lLiWfkbfeOst58EMGJE_VBmqUq6uCuIzAG-iOOKvBnBKV90XHlGxlvLk9G19ArEpXfg2JOrD3rk8Pj75ivqiktm5XSVfjajfx1ng8MfBE5X2c7Ga-69Icj3fw5RS5ni5Dt9n4AVae_EQPcetHPk6kGPXRc-_-t__AoyhE5s=)) - [One single storage for multiple sets](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_set/device_subsets_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzNWYtOG0kW_ZU7XmllwNiYbDYZA9awQLLWRJANZEajIeopd5ftWtrdnqpuGw_i3_fcerTbj0TJalZaI2PcXXUf5577qOapYaQxKs9Mo_frU0MljV631UhFNi7FWDZ6jbhMRKPVMHmpY_7e2b_PaJ8u8tlSq_GkoGa8R8dHxy8O8etli65_GlwOzuni5sP7mw_nd4Ob6zZvsJveqVhmRiZUZonUVEwknc9EjA9_p0U_Sc3W0HH7iJq84L7h79039k6slGVe0lQsKcsLKo2EGGVopFJJ8jGWs4JURnE-naVKZLGkhSomVpWXY82hX7yQfFgIrBfYMcO3UX0liaIynV-Topj1Op3FYtEW1ux2rsed1C02nXeDi6vr26tDmF5t-5ilgJe0_L1UGo4PlyRmsCwWQ9ibigXlmsRYS9wrcrZ8oVWhsnGLTD4qFkJLKydRptBqWBZr4AU74X99AeATGYA7v6XB7X2D_nF-O7htWTk_D-7-efPxjn4-__Dh_PpucHVLNx8QrOvLAYcK397Q-fUv9OPg-rJFEtBBlXycafYCpiqGVSYOw1sp18wY5c4sM5OxGqmYAo1onM-lzuAWzaSeKkc4GJlYOamaqkIU9tqWc1ZV5z67z_6isjgtE0mncRnnHcNb4sjIItJy1I7LSX9zVSKwKukIrcWyv_PWqMxiVizStfvFRJem6CRyDiOiuYyLXLcnu5YgqGUsd98zscjaW1blOTCA6XMZjXVezszWEpGOc7BgMl23GfYmcrR2TeWIuhTrC7NyKrWKrdDOvuPhDzY_vD-mHAI1E8lHgXBKQOfWDLWSI7qUUwSigInS0CRfMC85zfJM0rBMHwh7yQAQDizuxTCgkDQt00JBGnnpHF2ONkgxteJVlqi5SkqRkkeA4z1Xwpt1aBSsRyjp_P3ArIrGXY71BSjyB5TANr2EOdANelqiiTRd0mIiM8gRKXPMZjxIC_oAZUpEIQir8tjqtIYFa62GiTATNrmFMjDTAoSIYaTJ09KuV9k8T-cAQwKtfMkaBBl8pAEQDwbzfywz6xsWeSBcrtxxkZLM9MwCywzPOOt5k5EzoS2Izr8NY9mv6gKvlyg-KFgZly6p21a41QLC0RB1K54oOXflpmRLrbqUy_VC2qKd5dlhvrAZyYgXy5lcIf4DKqv0FjuKcIVhxUnFDTZlVupZjsoHaqTLNg0KXsZVWWWFRB7bkmYshaB-Cq5a8Z4UIitogQqMNUluVwRlFf9XuS9KZhrr5mJE8TgyzAd-ndHrE3x0Op1TQt_5eIlfqxyjt5xjXNhfQweomhjWOAGu3DAYSBenLR3DMn5AbbF6zqhb6bgup0OUKEjE2rjUWsIRk-agvJNqYbeSncItyde0ekHy0cmWZJnKKcQa1xOM1IUlwu-l1EuGwwX1QS4jjhyEAPEg5Ue5tPEMq2Y6H-IzMmhYU-k3AFjiKtrrIWOk0JFfdeqR5W6x--V2oRAJpFDEqROFInoaDOr3gzF3rI2ZqvPU0MVbGmuRlalAeJefVfHVL6vC1hlnPTkfqemc-pMUBOHztmmDqiV3bvYb1_ZQYTsd5yXCZqtBLVd9YQiR8F9DzBySgWbuXgVhq86__onT8253Bof6EzJ5Ux-uB538kQlAVLel1wsrTsLWakt4BXPXu-7K3G_Fmjtwr-cSBNREvkauDfyXktAcez1MWSKNinxFxG-WtiNbvlnGJu5VNtyic66C5GJ662qebVI7sz3ywERRvQCunLbTTY03130n7Iyeui160SIM5a9a9H2LuvjexYUurnRxqfv9syfW1XRWLH0Rw8gJI3niQwVyxYz_1txP7ht8EcmVYqyci7SUhrsPbvtvXO7LNOEmZFDEYhZgynji1BQTUXClx-yO9oRKF6N-8tCLkZ4_Z2VhbUfdr2rbymfJRkZ8IxgBHw-7J2tzjpthBhZFWFMDFyUyr4aTqjHWxjGy41jVBa-4gm93EjfuGFKQgvzzEm3AlNf62_Vvdc21xsptfko-eQy9z7lTao479z8bSa4jXmY-_DcQNL4JRtE4zYfgd0TzXAVtzUCy_UositITU5bz3OCcAgzHnL9bg6dFjmjVnajgKRFLx0hN_JlEMLhQtrb7xtBv2puYDSKfvEOUu4fmnj2cgeb0lsdD50ATI8PF2z0e_eTjhiqVPEJT0-6-VNP2I_CxXwbJI74c-OZpv-1RJ7T8us1aLEIh4rQ7o_1mBe0BK3A22bUOrWr5WX1zW0vkfBI5fHJtmq7WuT17XiVcc7Sy8eUOXWB-tayyhx74awdPNDm3kTfx0NTEGlJQiUavCG0eHwcHe_Tk6sq6ZW0fVoa_Zenzq_rk_Hi-z553cP08Tb0M5Jplup8CR3npDlhforOWNhy6BNE2Oc1ADYVxJ0rro280lsYSzF2d4ZDOOPIhL8QXzfkf5cII4fs_y4QNqv-pXP8i0xmLb-M576ix_IM0GO0w62OU5xKpRohy9WxkPbZ-4A9UoygyEzSQBIFh1msrytoGKc2am3SGfEAO-CUuO569iCXmSTerN4NdX5FIMP1f3DLX03HNXrewFhprOCTWQWtbNu1IQI9OnOuE3dlGA1MhmC6zFR5QVuRTFd_o5l-dpy3ad0q_O9tK7885_xnsKpsGzhgPJWJiYH1KR60d2cjtvUpFcrJDDHxIZhpAj5r3jdsyjqUx39Eba_GasHBEad_DxsaeD16oURypKU6nzZCCsHJ1vAldGGFAYXAdNVlPGtvxs3LKB6S_h9ZytTqu2Yd3E4E6JihRo5G0xzB-ziaNHTqQPzskus02u4w7BfG4Vk1R9k--GSFO0N5_OgaI9feLjfffNt4v197P9RTGhKSSoNppdU-W6mr7TceG2uK2H7-aMGgtHdyOekZgxQkdHKiKHHUx_PgCNTAaClQpP8fHwhTr6sOw60rDVDzIyAmRjzjUF6e7BuTtibdZB5opvleRnD8-5zzIMWJmsKsobJhXj2qFCTPpSD0ihlP73CmfYl6UYYudwIq8ECmjEIjGE22l0OPlz-yRsew6OlnhWVtUA7WOIf9GhNYwrss7OKsvh-O-cHgjf1Xs1Sd7mqo21YHhIXmXD5zU9lgJarsZmjiMoShXGDCLNzsZy4sgL3KSzsLqtiVCDd8L9yRv98Ot0vjH1ygCPn8rRWZ1gLXMXh0tnza0P4dUHmRopkDqD_8sLGyvntaP0Z0xiYTjBsdwpQVTUtje3D4WBJfqLKuYWQ0HK7_9E8yaV1-iw1aObR2zaxee1tnQWnOC63FzD4yo2PHp2fMlGLmetCE3g6Ynl6QVAqtz79M2Ks8teqq9a0Y-r-Wme3jNj3hqD79r-FV90lQzyPaIXDv-2AZmr56envIJ9Ph1v99vrsR4HNpjWdSOEW_UZscJwrg9f60oXq9lUerMJjpcbLQaXDnQ3fXqX16NbB7H3eOXZRe3-Wmz_X9Y4xAUOIsPDrqv6FDoeHJmptGrIzo8RC8pDm1BTGRymIrp0P6TLFXDmsw4jlNcnLv_aOECpuzsofHcCvcRvLX7QLzx_Mn-_Af5iUl4)) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 12b508404..9df526f95 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -33,17 +33,23 @@ endfunction(ConfigureExample) ### Example sources ############################################################################### ################################################################################################### -ConfigureExample(STATIC_SET_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/host_bulk_example.cu") ConfigureExample(STATIC_SET_DEVICE_REF_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/device_ref_example.cu") ConfigureExample(STATIC_SET_DEVICE_SUBSETS_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/device_subsets_example.cu") -ConfigureExample(STATIC_SET_SHARED_MEMORY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/shared_memory_example.cu") +ConfigureExample(STATIC_SET_HETEROGENEOUS_LOOKUP_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/heterogeneous_lookup_example.cu") +ConfigureExample(STATIC_SET_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/host_bulk_example.cu") ConfigureExample(STATIC_SET_MAPPING_TABLE_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/mapping_table_example.cu") -ConfigureExample(STATIC_MULTISET_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_multiset/host_bulk_example.cu") +ConfigureExample(STATIC_SET_SHARED_MEMORY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_set/shared_memory_example.cu") + ConfigureExample(STATIC_MAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/host_bulk_example.cu") ConfigureExample(STATIC_MAP_DEVICE_SIDE_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/device_ref_example.cu") ConfigureExample(STATIC_MAP_CUSTOM_TYPE_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/custom_type_example.cu") ConfigureExample(STATIC_MAP_COUNT_BY_KEY_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_map/count_by_key_example.cu") + +ConfigureExample(STATIC_MULTISET_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_multiset/host_bulk_example.cu") + ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/static_multimap/host_bulk_example.cu") + ConfigureExample(HYPERLOGLOG_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/host_bulk_example.cu") ConfigureExample(HYPERLOGLOG_DEVICE_REF_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/device_ref_example.cu") + ConfigureExample(BLOOM_FILTER_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/host_bulk_example.cu") diff --git a/examples/static_set/heterogeneous_lookup_example.cu b/examples/static_set/heterogeneous_lookup_example.cu new file mode 100644 index 000000000..ec0fd3405 --- /dev/null +++ b/examples/static_set/heterogeneous_lookup_example.cu @@ -0,0 +1,125 @@ +/* + * Copyright (c) 2025, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include +#include + +#include + +/** + * @file heterogeneous_lookup_example.cu + * + * @brief Demonstrates how to use hash set as a lookup table of the original data with string keys + * + * @note This example is for demonstration purposes only. It is not intended to show the most + * performant way to do the example algorithm. + */ + +/** + * @brief User-defined key equal to compare two keys + */ +struct my_equal { + my_equal(cuda::std::string_view* data) : _data{data} {} + /** + * @brief Checks if two keys are identical based on their indices in the + * original data array + */ + __device__ constexpr bool operator()(int32_t lhs, int32_t rhs) const + { + return this->operator()(_data[lhs], rhs); + } + + __device__ constexpr bool operator()(cuda::std::string_view lhs_str, int32_t rhs) const + { + auto rhs_str = _data[rhs]; + + // First check if lengths are the same + if (lhs_str.size() != rhs_str.size()) return false; + + // Then compare each character + for (size_t i = 0; i < lhs_str.size(); ++i) { + if (lhs_str[i] != rhs_str[i]) return false; + } + + return true; + } + cuda::std::string_view* _data; +}; + +/** + * @brief User-defined hash function to hash the original data based on its index + * + * @tparam T Original key type + */ +struct my_hasher { + my_hasher(cuda::std::string_view* data) : _data{data} {} + __device__ uint32_t operator()(int32_t index) const + { + auto str = _data[index]; + return this->operator()(str); + } + + __device__ uint32_t operator()(cuda::std::string_view str) const + { + uint32_t hash = 0; + // Naive string hash function + hash = str[0] + str.size(); + + return hash; + } + cuda::std::string_view* _data; +}; + +int main(void) +{ + // The original key type is a string, which is variable length and larger than 8 bytes + using Key = std::string; + + auto input = + thrust::device_vector{"apple", "apple", "banana", "cherry", "apple"}; + + auto const size = input.size(); + + // The actual key type is an index type, `int32_t` is large enough to cover the whole input range + // and 4-byte atomic CAS is more efficient than the 8-byte one. + using ActualKey = int32_t; + // `-1` is a valid sentinel value since one will never access `data[-1]` + ActualKey constexpr empty_key_sentinel = -1; + + auto const data_ptr = input.data().get(); + auto set = cuco::static_set{cuco::extent{size * 2}, // about 50% load factor + cuco::empty_key{empty_key_sentinel}, + my_equal{data_ptr}, + cuco::linear_probing<1, my_hasher>{my_hasher{data_ptr}}}; + + // The actual keys are indices of elements + auto const actual_keys = thrust::device_vector{0, 1, 2, 3, 4}; + set.insert(actual_keys.begin(), actual_keys.end()); + + auto query = thrust::device_vector{"lychee", "cherry", "apple"}; + auto const query_size = query.size(); + + auto contained = thrust::device_vector(query_size); + set.contains(query.begin(), query.end(), contained.begin()); + + for (auto const& it : contained) { + std::cout << it << "\n"; + } + + return 0; +} \ No newline at end of file