From 096346b739da3fb1d9c3b402190c0a3a7e554440 Mon Sep 17 00:00:00 2001 From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> Date: Tue, 10 Dec 2024 15:08:51 -0800 Subject: [PATCH 1/3] Minor improvements for arrow_filter_policy (#654) --- include/cuco/bloom_filter_policies.cuh | 2 +- .../bloom_filter/arrow_filter_policy.cuh | 23 ++++++++----------- 2 files changed, 10 insertions(+), 15 deletions(-) diff --git a/include/cuco/bloom_filter_policies.cuh b/include/cuco/bloom_filter_policies.cuh index a433e7cca..0d28b166a 100644 --- a/include/cuco/bloom_filter_policies.cuh +++ b/include/cuco/bloom_filter_policies.cuh @@ -33,7 +33,7 @@ namespace cuco { * By default, cuco::xxhash_64 hasher will be used. * */ -template > +template class XXHash64 = cuco::xxhash_64> using arrow_filter_policy = detail::arrow_filter_policy; /** diff --git a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh index edb88e6e0..bfe97cfaf 100644 --- a/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/arrow_filter_policy.cuh @@ -40,8 +40,7 @@ namespace cuco::detail { * void bulk_insert_and_eval_arrow_policy_bloom_filter(device_vector const& positive_keys, * device_vector const& negative_keys) * { - * using xxhash_64 = cuco::xxhash_64; - * using policy_type = cuco::arrow_filter_policy; + * using policy_type = cuco::arrow_filter_policy; * * // Warn or throw if the number of filter blocks is greater than maximum used by Arrow policy. * static_assert(NUM_FILTER_BLOCKS <= policy_type::max_filter_blocks, "NUM_FILTER_BLOCKS must be @@ -81,14 +80,13 @@ namespace cuco::detail { * @tparam Key The type of the values to generate a fingerprint for. * @tparam XXHash64 64-bit XXHash hasher implementation for fingerprint generation. */ -template +template class XXHash64> class arrow_filter_policy { public: - using hasher = XXHash64; ///< 64-bit XXHash hasher for Arrow bloom filter policy - using word_type = std::uint32_t; ///< uint32_t for Arrow bloom filter policy - using hash_argument_type = typename hasher::argument_type; ///< Hash function input type - using hash_result_type = decltype(std::declval()( - std::declval())); ///< hash function output type + using hasher = XXHash64; ///< 64-bit XXHash hasher for Arrow bloom filter policy + using word_type = std::uint32_t; ///< uint32_t for Arrow bloom filter policy + using key_type = Key; ///< Hash function input type + using hash_value_type = std::uint64_t; ///< hash function output type static constexpr uint32_t bits_set_per_block = 8; ///< hardcoded bits set per Arrow filter block static constexpr uint32_t words_per_block = 8; ///< hardcoded words per Arrow filter block @@ -135,10 +133,7 @@ class arrow_filter_policy { * * @return The hash value of the key */ - __device__ constexpr hash_result_type hash(hash_argument_type const& key) const - { - return hash_(key); - } + __device__ constexpr hash_value_type hash(key_type const& key) const { return hash_(key); } /** * @brief Determines the filter block a key is added into. @@ -155,7 +150,7 @@ class arrow_filter_policy { * @return The block index for the given key's hash value */ template - __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const + __device__ constexpr auto block_index(hash_value_type hash, Extent num_blocks) const { constexpr auto hash_bits = cuda::std::numeric_limits::digits; // TODO: assert if num_blocks > max_filter_blocks @@ -173,7 +168,7 @@ class arrow_filter_policy { * * @return The bit pattern for the word/segment in the filter block */ - __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const + __device__ constexpr word_type word_pattern(hash_value_type hash, std::uint32_t word_index) const { // SALT array to calculate bit indexes for the current word auto constexpr salt = SALT(); From a7679e5bc54ba2afedd29e90f951b1c3eaf12fa8 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Thu, 19 Dec 2024 10:02:19 -0800 Subject: [PATCH 2/3] Replace window with bucket instead (#656) This PR completely deprecates the window concept in cuco hash tables. --- README.md | 6 +- examples/static_set/device_subsets_example.cu | 14 ++-- examples/static_set/shared_memory_example.cu | 16 ++-- include/cuco/bucket_storage.cuh | 6 -- include/cuco/detail/extent/extent.inl | 30 -------- include/cuco/extent.cuh | 76 ------------------- include/cuco/static_map_ref.cuh | 1 - include/cuco/static_multimap_ref.cuh | 1 - include/cuco/static_multiset_ref.cuh | 1 - include/cuco/static_set_ref.cuh | 1 - tests/static_map/insert_or_apply_test.cu | 4 +- tests/static_map/shared_memory_test.cu | 28 +++---- tests/static_set/shared_memory_test.cu | 28 +++---- tests/utility/extent_test.cu | 6 +- tests/utility/probing_scheme_test.cu | 4 +- 15 files changed, 53 insertions(+), 169 deletions(-) diff --git a/README.md b/README.md index 7703effa9..198a8abdc 100644 --- a/README.md +++ b/README.md @@ -208,8 +208,8 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection #### Examples: - [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/eJx9VgtvGzcM_ivcDcOc9vwKVhRwkmJekmJGC6eI0xbDPDiyTvYJkaWbHnY9I_99pHTnx9ouAWyfSJEfP_LTaZc54Zw02mWDP3eZLLJBP88U08vAliIbZDwULMszZ4Ll9Nx9MdXwAq5NtbVyWXpo8TM4752ft_HjlxzGn0Y3oyFc391_uLsfPozuxh3aEDe9l1xoJwoIuhAWfClgWDGOX7Ulh0_CEho47_SgRQ7TrLZNs7OLGGVrAqzYFrTxEJzAMNLBQioB4gsXlQepgZtVpSTTXMBG-jKmquNEOPBHHcTMPUN_hjsqfFocewLze-j0V3pfDbrdzWbTYRF2x9hlVyVn130_ur4dT27bCH2_7aNWSC9Y8XeQFgufb4FViIyzOeJVbAPGAltagTZvCPnGSi_1MgdnFn7DrIhxCum8lfPgT8hrcGL9xw5IH9NI3HACo8k0g9-Gk9Ekj3E-jx5-v_v4AJ-H9_fD8cPodgJ399is8c2IWoVPb2E4_gPejcY3OQikDlOJL5WlKhCqJFpFkTicCHECY2ESLFcJLheSQzNGsDRrYTWWBZWwK5kGDkEWMY6SK-mZj2tfFRdTdad6qn-UmqtQCLjkgZuuoy185oTv8FC-OfXwpQ3OdwuxxiizteDe2A45feWyCJpTaqa-bVdmid36jtFhXwUOWee_-aXBdgi2OtkU63TRs_sizcevcW5L9J7Ng3qaiS8MGRZYUTLPrRQLuBEr5MZb5gUy5IjRelIPJGC_KQK2m6LB8MPIHYT3gK5kjcuAQ4UzorY4bmuTJnFhzSoGjJuxMdEpkFapq4Wh3hlsXt0mJZ8EPErsj_WPOBgxyyM3mtTkHuMIAqFCnE9ie4CCnZTao_Skbq2NLM6meofrmIoyvBNbuEJY_oJIAuh24XZV-S04ZXwCbgUNo9A-qYl-2zX-nma0KLVQSMGaqSBch-rGuUxP4EoTFG5C1oTCicBNLvCyzuNL5okBPFcEDisYzoMlQeJxQ99V8FAwz6iSiJNTS0gaIAjiDMucNRCwiHb_UMM4rOYUMnFBQkcQiTxBEsA2FoOBk_-ImT-Kq8NqFjdcwavez71e7xDxGs831DpwOogkEjRnrlY--v4EyjDsHKOxpx0sYM5DYLLOkhVj9zqvLr4J4hD9Khm5kKq1R9U9jnN2jI1mNXBqWJyBeAIju0owDPrYRH2su5pa3-4Dc3EEI52JqYbPSDqpHgHuJ_4Su_CGEuyaiHnts2_I7uvWPD8fIUWJIomEMul436JdL4d-Duc5dDo5yGfakDQ_GJwcKgkE7dkTE19Te-8mdCvKYC6WOPhneRKF0AX97h2RN4pTgcwplZCgGExSJnMlVRtbhYdemp_vhj2KOUGcdGaQklEvQfnvlzM3Rr1BTzyHTwpqGCsFfwK5OOAjTda6xwFEvcRj6QCzORP-p_6YrTEdJcO7gFxskzL3-Uq2xrNMCJ22kSthricW3SLkWTTi3DZVksEsWiepmsw1jMZVFjQqfpu4OCDCqlun8c9gVwvD4OlweYmH0CRwju_KH-BtBNDA7kwxRnYBzymUFT5YDSRoXMHrFd1X8E1gD7ewTK8575-_Cn00m8qnK1rWxnxX_OXL_mtoM8vLK7eave5Bu43C9vjhsRhRtBVbzeO9Tcn5UUzOucLFdbpk4QLqVD9lz3ljR_mc2HE4sue_4v-_fg9_jw==)) - [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/eJztWH9v4zYS_SpzOhzO2cg_EnTbwpsE50u2OKOFUyTZFkVdqDRF2bzIok6k7PiCfPe-ISXL3uwWLXB_XgIklmY482bmcTj0c2SVtdoUNhr__BzpNBqfxVEuimUtlioaR7JORRRH1tSV5Ofhm3lBb-jalLtKL1eOevKEzkfn5338-SKm2Q_Tm-mErm_vvr-9mzxMb2cDXuAXfaelKqxKqS5SVZFbKZqUQuJfI4npB1UxGjofjKjHCvOokc2jk3feys7UtBY7Koyj2iqY0ZYynStST1KVjnRB0qzLXItCKtpqt_KuGjseDv3UGDELJ6AvsKLEU3aoScLtofPPyrlyPBxut9uB8LAHploO86Bsh99Nr9_P7t_3AX2_7EORI71Uqf_UukLgix2JEsikWABvLrZkKhLLSkHmDCPfVtrpYhmTNZnbikp5O6m2rtKL2h0lr8WJ-A8VkD5RIHGTe5rezyP65-R-eh97Oz9OH_51--GBfpzc3U1mD9P393R7h2LNbqZcKjx9Q5PZT_TtdHYTk0Lq4Eo9lRVHAaia06rSkMN7pY5gZCbAsqWSOtOSWhrR0mxUVSAsKlW11oFwAJl6O7leayecf_cqOO9qOC_mxV91IfM6VXQha2mGlpfIxCo3kPXq6ljDraraumGqNrCSbJR0phqw0iuVrC4kuxb5p-W5WaJanxFa1FWBZIOP_UtjECkAblSyrExd2tcq1qWpyo7MaoMqKrH2qsM3gUP_8NxuQqlUpip2magngWIoBB-0FpVWGd2oNdLo4FohmZaT35C6y1djq281fE6-n9puhx4olZXZQAFlwk4r-mbr67f3T9uVliuSYNpC8T5sCOwQtnR-13mL7FqagncZCptVZu1lYHpAAWF6UOThEDasqhzcMgqAf1Q7y4YNXq2EXfn3tWU0puD1-1STT7XnocL-JCfs4xyIkCfkA1XblaoQawXmujuVxbR_MS3K2k0ZO4iC7CfJMjcLkScJbYxOSaLaZp0cljXA7AVTjCk-tuJxYxu7dDy2-r8qcVSczItnhErs06KFAP2SLuk1W95xLogTZx1vPxI14pfLhC1hRfA6HjdvGnWv5JgtsLkcj_ljmpSicpopftFoX_W8EG0zAY-VSJNFbuRj7-SksYNkf_kF8Hr3lANdwt0lZbtL_L_R68ETauuXhYchdVA6Azp9ouOfS-odrGpMTNMnPJxSQOOfTo4ssk3wDYH12OQFMkk-kcSJHzSl4HBjetPzhDll5z4i1uJVp5eHoXjBy7x4Yet_miIH729r9we54zeB_XOMiT-yD3KjRf6fR_9zHi2MyRugiKftWCkc-ROmLd7vcSwjLx40yahEgVzQ5SWN4CUU7mcs-AU2D128Yxr-AZIiGRhZdNFjVrUMCH3wW7Vjo4VrYkQXfb8u3Y5sbhw6eKXQuPkQV4ULUwh_rjb4PI_4JXDkmBQ2Iq-VHdDDCuLmiezK1HnKXd6qHCcpFtlarho_biUcN3nMYwqHPBkp64oHGYxp_B_cpVQ4wS3e4-y4qBhigkwmLQQE0T_rYpjV6wWbbM4AkBIgwmZXPDocbZXOblGvE7_gkt6O_j4ajTqL15gLMSPh2MKu0UjQQthmYoLu35B7kVImeFzYb4TOMEuTIIXt0eDtu0-C6KxfBqFUOu_tUQ0P7ZwcYuODu5ZcsO6g8-MrUpwrAcvzqLWNaoXiBgb0z0hYf9j6rHLCqM2qTz3PTIC5P98vUIsr9vDcWowbnX1Znl8X6OXlAC9YjlTyOR2moH2hnkcxncV0HtNgEJP29A4T03h8NJIFELxmn56wm1rt1rTfb4OFWoL-J7FfMVBFyp9HBymchtGB05DpCvlaiTxrB6BujvDjkHK-eN3x8QkXzRM2-kH1zj_t0CoUP9179MqhNqI5Auj6w80EgqqryWemiouLi7Pzr5HE86-vrq563UmH8asXqhQUfx9o_DnY98g9j4Y8KaET1Ln7fIm4MV6F7nVUpJYFKyUfufeJPA9Bc7cpzJa6Hopu0KR8QDPjVGgaW0WlsEz2bqj0ApFz_9zxHtg3GUW_tj34VwrpwkS_rwItFcZLkcdNyt3W8HVlbXzna62bxb8Vb6-WAnx8d7MpLk_t2MpGBe5UmV_p2km09Wt5Fdocy_IdX7_42pL5WBdqJTa6w_bAt9PUoI_yfZUvfwgMVgk4Zbg_MdjACtb0enxfykV5zJIQ_kfcOGZFq_QRLzoixKGSreCgkLh562zXVqCpJWJRCEkVYRmrHhyUUPNGEy9Et2sZxAKT9Y5ctZ6bbduq4pRDa3G7wLOD2QGn6bF9PkRDOzWo1MUFmuF9LSVupn-hbzyAFvZgDhuRP1fZVKVcXRXExwDeRHHE3w7goK667zyiYiPl2fnb-gxiU7rwhUjUh79LeXp69hX1RSVXl3adfDWifh_HgcMfh2BU2s_FeuG_Jcn14sCmlDLHy034SgMv0N2Lx-glbuUo2ZEcGy96-cX__gbvsxSQ)) -- [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/eJzNWQ1P20ga_ivv5qRTgJAA3V67gaDlgPairaAHdFerpfJO7EkyV8fOztgJWcR_v-edD8cmadVKe9IFhRB75v143uf9GPPYMtIYlWem1f_tsaWSVv-w00pFNinFRLb6rbhMRKvTMnmpY_7e273PaJfO8_lKq8m0oHa8Q0cHRy_28ev7Dl39PLwYntH59c3765uzu-H1VZc32E3vVCwzIxMqs0RqKqaSzuYixoe_06GfpWZr6Kh7QG1ecN_y9-5bO8dWyiovaSZWlOUFlUZCjDI0Vqkk-RDLeUEqozifzVMlsljSUhVTq8rLsebQr15IPioE1gvsmOPbuL6SRFGZzq9pUcz7vd5yuewKa3Y315Ne6hab3rvh-eXV7eU-TK-2fchSwEta_lEqDcdHKxJzWBaLEexNxZJyTWKiJe4VOVu-1KpQ2aRDJh8XS6GllZMoU2g1KosGeMFO-F9fAPhEBuDObml4e9-if57dDm87Vs4vw7t_XX-4o1_Obm7Oru6Gl7d0fYNgXV0MOVT49obOrn6ln4ZXFx2SgA6q5MNcsxcwVTGsMnEY3krZMGOcO7PMXMZqrGIKNKJJvpA6g1s0l3qmHOFgZGLlpGqmClHYaxvOWVW9--w--5vK4rRMJJ3EZZz3DG-JIyOLSMtxNy6np1vW5Br6_c3m7UTgdtITWotVY2cx1aUpeolcwIJoIWMI6U63LUFEy1huv2dikXU3tOY5AIDdCxlNdF7OzcYSkU5yUGA6a3oDUxM5blxTOUIuRXNhVs6kVrEV2tt1JPzRJof3x5QjQGYi-SAQS4bGrRlpJcd0IWeIQgETpaFpvmRSco7lmaRRmX4i7CWPKt-LYUAhaVamhYI08tI5tBxqMGJmxassUQuVlCIljwAHe6GEN2vfKFiPONLZ-6FZV4y7HOsL8ONPKIFtegVzoBvctCwTabqi5VRmkCNSJphNdzAW3AHKlIhCEFblsdVpDQvWWg1TYaZscgc1YK5FDE7BSJOnpV2vskWeLgCGBFr5ijUIMvhIAyAeDCb_RGbWNyzyQLhEueMKJZnmmQWW6Z1xyvMmI-dCWxCdf8-MZb-qC7xeovKgWmVct6TuWuFWCwhHIxSteKrkwtWaki216lKu1UtpK3aWZ_v50qYjI16s5nKN-I8oq9Jb7CjC5YUVJxU32JR5qec5yh6oka66NCx4GZdklRUSSWzrmbEUgvoZuGrFe1KIrKAlyi_WJLldEZRV_F8nviiZaaybKxHFk8gwH_g1oNfH-Oj1eieEpvPhAr_WOUZvOce4qr-GDlA1MaxxCly5WzCQLk4bOpYgbL50egZ0WOm4Kmcj1CdIxNq41FrCEZPmoLyTamG3kp3CDclXtH5B8sHxhmSZyhnEGtcQjNSFJcIfpdQrhsMF9ZNcRRw5CAHiQcpPcmXjGVbNdT7CZ2TQrWbSbwCwxOWx30fGSKEjv-rEI8utYvvL7UIhEkihiFMnGpdZzIQ4CQadngZj7lgbM1XnqaHztzTRIitTgfCuPqviq19Wha0zznpyPlLbOfUXKQjCF13TBVVLbtvsN67toML2es5LhM1Wg1qu-sIQIuG_hpg5JAVzzN2o8OvUyXd67JS8256-ofiENH6uDNeDQv7IBPCpG9LvhxXHYWu1JbyCrc1-uzb3W4Hmztvvu-wAL5GskesB3yzJddp-H8OVSKMiX1Pwm0VtyZNvlvEc9CoPbtEz1xFyAb111c62p615HnlUoqhe-ix2mAVAHZ5baqS5OnXCBvR42KEXHXrZoVcd-qFDh_h-iAuHuHKIS4c_PHlWXc7mxcqXL0yaMJIHPdQeV8b4b82d5L7FF5FWKabJhUhLabjv4Lb_xoW-TBNuPwblK2YBpoynTk0xFQXXeIzsaEyocTEqJ8-6mOT5c14W1nZU_KqqrX2WbGTEN4IR8HH_8Lgx4bjpZWhRhDU1cFEc82osqVpibRAjO4hV_e-Sa_dmD3GDjiEFKUg-L9EGTHmtv1_9Xtdca6nc4GfkM8fQ-5x7pOa4c-ezkeQK4mXmo_8AQePbXxRN0nwEfke0yFXQ1g4k263Eohw9MmU5yQ2OJ8Bwwsm7MXJa5IjWfYkKng-xdIK8xJ9JBIMLZau6bwmnbXsTU0HkM3eEQvepvWPPZKA5veXB0DnQxrBw_naHhz758EyVSh6gqW13X6hZ9wH42C_D5AFf9nzbtN92qBeafd1mLZahCnHaDWi3XUG7xwqcTXatQ6taPqhv7mqJnE8ih0-uTdsVOrdnx6uEa45WNr7cmwtMrpZV9qwDf-3IifbmNvImHpfaWEMKKtHiFaHB42Nvb4ceXV1pWtb1YWX4O5Y-v6mPzo-n--xpC9fP0tTLQK5Zpvv5b5yX7lz1JTpracOhSxDtOacZqJEw7iBpffRdxtJYgrnroxvSGSc95IX4ojn_o1wYI3z_Z5nwjOp_Kde_yHTG4tt4zjtqLL-RBkMdpnwM8Vwi1RhRrh6JNGPrR_1ANYoiM0UDSRAYZr22oqxtkNKuuUkD5ANywC9x2fHkRawwSbopvR3s-opEgun_5pbZTMeGvW5hLTTWcEisg9a1bNqSgB6dONcJu7OJBuZBMF1mazygrMhnKr7W7b87Tzu065R-N9hI7885_xnsKpuGzhgPJWJiYH1KB50t2cjtvUpFcrJDDHxI5hpAj9v3rdsyjqUx39Eba3FDWDicdO9hY2vHBy_UKI7UDOfSdkhBWLk-2IQujDCgMLiOmjSTxnb8rJzx0egfobVcrg9q9pndVKCOCUrUeCztAYwfr0ljhw7kzxaJbrPNLuPOPzyuVVOU_ZNvRogTtJ8-HgHE-vvFs_f3z94vG--negpjQlJJUO20umdKdbWnbceG2uKuH7_aMKiRDm5HPSOw4pj29lRFjroYfnCBGhiNBKpUmGn9MB8LUzTNcKVhJj7JyB9D5APO88VJNdW262AyjXcqIvPH5xwEAcYcfXYHxQsz6UGt-GDuHKsHxGlmnyrlM8yEMmyxU1aRFyJlTz2ZnHWmUulR8WfyyFgOHRyvUastqkFXR4p_Iw4NJOvy9gb15XDdlwdv5m-K_fpoD0zVpjo0PApv94KT1x4cQWE3KxOHKxTfCgdm6_OOxRIjSIyCrEFY37Uhr6F87p7WbX-AVRr_fBrp7jO1UmXj7dcNGifIxw39TyFthxkaJ_D60z_xCgKqB_ITdGJMHeFowZFc68FEFLa3N48Awak62yqGVoPA2nP_nLLm15dIsZFPG-fp2oXHJic6DSe49rZ3wIuKIx-fPGuCkdsTNGh6dAlZIbA-4z5uovLUocfau2bkUyNHw8G58Yi7hl_VE001b2yOw7Wjjm1W9urJyQmfNo9en56ettdiPA7diSxqR4Y36nl3CcK4FX-tKF6vZVHqzKY7XGx1WlxB0Mn1-r9arWwRx4dHL8tD3OZnyvZfXq19UGAQ7-0dvqJ9oePpwMyiVwe0v4--Uezb2pfIZD8Vs5H9P1iqRjWZcRynuLhw_7TCBUzU2afWUyfcR_Aa94F46-mj_fkvXfg-3w==)) -- [Using shared memory as storage](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_set/shared_memory_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyVWA1vG7kR_SvEFgXknCzZQYsrFFmoznZRIal9sJ0Eh_NhTe1SEuvVUkdypegM__e-GXK1KzluXQdxIn7MDN88vhnqKXHKOW1Klwx-fUp0ngxOu0khy3kl5yoZJFmVy6SbOFPZjD73392X4p04N6ut1fOFF53sSLw_ef-Xrrj6MrmYjMX59c3P1zfju8n1VY_W8vpPOlOlU7moylxZ4RdKjFcywz9xpiu-KEuBiPe9E9GhBfdJnLtPjj6wla2pxFJuRWm8qJyCGe3ETBdKqG-ZWnmhS5GZ5arQssyU2Gi_YFfRDocjfolGzNRLrJfYscKnWXulkH4XOv0svF8N-v3NZtOTHHbP2Hm_CItd_9Pk_PLq9vIYoe-2fS4LICus-r3SFgefboVcIbJMThFvITfCWCHnVmHOG4p8Y7XX5bwrnJn5jbSK7eTaeaunld8Dr44T528vAHyyBHDjWzG5vU_ET-PbyW2X7Xyd3P3z-vOd-Dq-uRlf3U0ub8X1DZJ1dTGhVOHTP8T46hfxcXJ10RUK0MGV-raydAqEqglWlQcMb5XaC2NmQlhupTI905moGSTmZq1siWOJlbJLHbiGIHO2U-il9tLz2IvDsav-fXlf_kmXWVHlSgyzKjN9R1uy1Cnfy6rF6GCFMfCEBWuVzq2pVq4XlvTfhcT8nQnjFgA4T5dqaew2Vd8kjqdgLiyZWq1m4gKTJcCVXiE8R8eJNGkioMQFWyLYaqL2CjaxFzEVEiDeKn-jZoglTeeFmcoiTcXa6Bz7sZWMpY_AShUdv12pUi5V3DIYqG9elT6lcbC6zM0mDUOU2__1Q5jBxHLlt3CwHb6wjkE2PRK7RYimBBlVcXRfPpGPfl98RZ61dV7IojAZHYyhaB8ewBhLONWEeMCpHnrRwB0jFxaAuHRTrSZRAKqZgb95ZSoXD-ho1BXGu27cvlnobMHON2x_rTIY03_AQGFk7thNmsbMpuLFOSNyLRTdr3to9tayqFTn6LcPlD92eoWZrtgoChBUqDLP55KZr2SxO42Z_hvBiJk1S562uN7SWslsELLCBa-Xnu3CiSOpVTOOqbOf2TpE1j4OBeemQHDDNWKRAGyXBNVyHu22XGO2dtt5meKueHpu_sbdRw0ElLasshbLYd8qCTErMvMo-uL8bty4YYQwZbJHuHt5DwcD0us0mEh5XaflZlJC_mSBfO4grDGrHAkIMh_9O7p1tIiNcAQkBnpnocMTO-Rut2W2sKasbYcYEfJSPgKvyoKP9V5WI7GQVFhK7RZB8qInB0OdA2iY4jUDcgOpoPIEcscCJMutmFVlRnbhwG97zU1ayLWiOKAxpQOnl0KTomBE1i6iYb-QHpGuVsZ6x4eobRInSqoiHQaFLpaEJD-AI8r6h6MmQWGE-Ib8EGBWTcGxNCTKWNcJUhHWtY55iZJXpz5MIgMcRhjKQ8yRiezwVZGJNMFHBBFQjYwABoEQke3MbZxaz7jwbyTxz3DhR9Gep47SabhO4VeJ4ocuglbWqIYWQbWpWIOKokSaEfScmogIZlfonup14yk7GaoxIg1BNfD14nSceoVlmZf7HCMCh50iBkR1j6P9b1zb16EaBjQh4LAO9QiBOBaGKUkvKik6C2puwsVhISaBxZB72J2VeZgBel_DIkGkDehbHpsN1-xIvql0oblgSw0ODwe0bDlBPwJDEOZyrmw036yL3KwLw_n4y-X4bkDUqIWGcofk7DlD19C4CEM69A2Aj1osBSy1sYeS1N4C0rXy-MoFqDe0kntew01JfAE302MmOuHmN9569QemypF4Eqh4pZ917pPLm5vrm4H4iHvw55wlI1iCxySSTjzfl88UAbaARrrsULtQV-SQ2498j7BgX5TuE9Z5tIC1yhNUEu0ozr4GVlzp9rpTdPjc5KkldjA3eztlJsCYDxZb0G2qdeAVbAZD-M_nn3j9YbeBAEchD9RMfqfDeDo-fa5Z_lXnUEzk_bC2i7yy3ERag4zNDzLMlmP1ZFU4E6e1yfPm8guuQ4JWfGd_rShhb9z9L9zLZbUUZbWcQmfqlqSuP7Xa0e3nlvCgHYr4BsFm63A5Zf1C80N9qNfQyLW0mp4Fta2IIfcBQ029wOnJycmIUCaDgFwFEXhowieOZPQ4QYWhM2DDh4eGJ-0O8uw1BzVmF3yhOOgIuHDQm2UIMBiME2mY2DdcYLe0aVwyjMh24yyuq6wKn6LGLtJajJgnowZ4SFPeREEdDHuYI40B-pW0KC8ej0Yhp9DzVmjUSe_F07Tr5OYtPTN-Woi9cQe9lanJ4Wrm8KxUoct543ZsrJxHWn5Ha5l6w5C8ce93svHmoFnxEHoa-3pj_x_PNcLcqQ1bl3A0elG7SDLxyKOrSClsFRpopYLu49HOiWt62Zj45gmBO6DDfQp9vOJSEe9jeD_RzdhYuSK9oy5y70pyWy4kieFDw4uHA_rUPTkotOtkaloNBvXssEXYc-BA-mmcItuNYNApIZlrdJdcGyyEHAoSq-nuyrKO7h2zafA0db7URnBjDH4oR6_93d4O3d4jro-Q58oG0yVgx0PcNgUYPtCrxbfO6xIaFuxuDzUw6d7UsA3RqNO6KU_PrZr5SeIEi3ZiQkMsCGbSvvoxsWviDx7C-36Gw-EphOr930aj0eFb6TsP1w-hGuXyAtUqU63ujLsqFNakm0QZts3XX0m5zrLT93-tTjFtVj58N5YcO5-fZT_8cPqjOJY2W5y5ZfrjiTg-BnD-mMPIVX5cyOWUvzAr9LRlM8uyAoPr8BUXBvCSLB-T5249D6D35nO1Tp5_4z__AbR8tYo=)) +- [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/eJzNWQ1P20ga_ivv5qRTgJAA3V67gaDlgPairaBX6K5WS-Wd2JNkDsfOztgJWcR_v-edD8cmadVKe9IFhRB75v143uf9GPPYMtIYlWem1f_tsaWSVv-w00pFNinFRLb6rbhMRKvTMnmpY_7e273LaJfO8_lKq8m0oHa8Q0cHRy_28ev7Dl39PLwYntH59Yf31x_ObofXV13eYDe9U7HMjEyozBKpqZhKOpuLGB_-Tod-lpqtoaPuAbV5wV3L37tr7RxbKau8pJlYUZYXVBoJMcrQWKWS5EMs5wWpjOJ8Nk-VyGJJS1VMrSovx5pDv3oh-agQWC-wY45v4_pKEkVlOr-mRTHv93rL5bIrrNndXE96qVtseu-G55dXN5f7ML3a9jFLAS9p-UepNBwfrUjMYVksRrA3FUvKNYmJlrhX5Gz5UqtCZZMOmXxcLIWWVk6iTKHVqCwa4AU74X99AeATGYA7u6HhzV2L_nl2M7zpWDm_DG__df3xln45-_Dh7Op2eHlD1x8QrKuLIYcK397Q2dWv9NPw6qJDEtBBlXyYa_YCpiqGVSYOwxspG2aMc2eWmctYjVVMgUY0yRdSZ3CL5lLPlCMcjEysnFTNVCEKe23DOauqd5fdZX9TWZyWiaSTuIzznuEtcWRkEWk57sbl9HTLmlxDv7_ZvJ0I3E56Qmuxauwspro0RS-RC1gQLWQMId3ptiWIaBnL7fdMLLLuhtY8BwCweyGjic7LudlYItJJDgpMZ01vYGoix41rKkfIpWguzMqZ1Cq2Qnu7joQ_2uTw_phyBMhMJB8EYsnQuDUjreSYLuQMUShgojQ0zZdMSs6xPJM0KtN7wl7yqPK9GAYUkmZlWihIIy-dQ8uhBiNmVrzKErVQSSlS8ghwsBdKeLP2jYL1iCOdvR-adcW4zbG-AD_-hBLYplcwB7rBTcsykaYrWk5lBjkiZYLZdAdjwR2gTIkoBGFVHlud1rBgrdUwFWbKJndQA-ZaxOAUjDR5Wtr1Klvk6QJgSKCVr1iDIIOPNADiwWDyT2RmfcMiD4RLlFuuUJJpnllgmd4ZpzxvMnIutAXR-ffMWParusDrJSoPqlXGdUvqrhVutYBwNELRiqdKLlytKdlSqy7lWr2UtmJnebafL206MuLFai7XiP-Isiq9xY4iXF5YcVJxg02Zl3qeo-yBGumqS8OCl3FJVlkhkcS2nhlLIaifgatWvCeFyApaovxiTZLbFUFZxf914ouSmca6uRJRPIkM84FfA3p9jI9er3dCaDofL_BrnWP0lnOMq_pr6ABVE8Map8CVuwUD6eK0oWNUxvcoLFbPgA4rHVflbIT6BIlYG5daSzhi0hyUd1It7FayU7gh-YrWL0g-ON6QLFM5g1jjGoKRurBE-KOUesVwuKDey1XEkYMQIB6k_CRXNp5h1VznI3xGBt1qJv0GAEtcHvt9ZIwUOvKrTjyy3Cq2v9wuFCKBFIo4daJxmcVMiJNg0OlpMOaWtTFTdZ4aOn9LEy2yMhUI7-qzKr76ZVXYOuOsJ-cjtZ1Tf5GCIHzRNV1QteS2zX7j2g4qbK_nvETYbDWo5aovDCES_muImUNS5MvI36jw69TJd3rslLzbnr6h-IQ0fq4M14NC_sgE8Kkb0u-HFcdha7UlvIKtzX67NvdbgebO2--77AAvkayR6wHfLMl12n4fw5VIoyJfU_CbRW3Jk2-W8Rz0Kg9u0DPXEXIBvXHVzranrXkeeVSiqF76LHaYBUAdnltqpLk6dcIG9HjYoRcdetmhVx36oUOH-H6IC4e4cohLhz88eVZdzubFypcvTJowkgc91B5XxvhvzZ3krsUXkVYppsmFSEtpuO_gtv_Ghb5ME24_BuUrZgGmjKdOTTEVBdd4jOxoTKhxMSonz7qY5PlzXhbWdlT8qqqtfZZsZMQ3ghHwcf_wuDHhuOllaFGENTVwURzzaiypWmJtECM7iFX975Jr92YPcYOOIQUpSD4v0QZMea2_X_1e11xrqdzgZ-Qzx9D7nHuk5rhz57OR5AriZeaj_wBB49tfFE3SfAR-R7TIVdDWDiTbrcSiHD0yZTnJDY4nwHDCybsxclrkiNZ9iQqeD7F0grzEn0kEgwtlq7pvCadtexNTQeQzd4RCd9_esWcy0Jze8mDoHGhjWDh_u8NDn3x4pkolD9DUtrsv1Kz7AHzsl2HygC97vm3abzvUC82-brMWy1CFOO0GtNuuoN1jBc4mu9ahVS0f1Dd3tUTOJ5HDJ9em7Qqd27PjVcI1RysbX-7NBSZXyyp71oG_duREe3MbeROPS22sIQWVaPGK0ODxsbe3Q4-urjQt6_qwMvwdS5_f1Cfnx9Nd9rSF62dp6mUg1yzT_fw3zkt3rvoSnbW04dAliPac0wzUSBh3kLQ--i5jaSzB3PXRDemMkx7yQnzRnP9RLowRvv-zTHhG9b-U619kOmPxbTznHTWWf5AGQx2mfAzxXCLVGFGuHok0Y-tH_UA1iiIzRQNJEBhmvbairG2Q0q65SQPkA3LAL3HZ8eRFrDBJuim9Hez6ikSC6f_mltlMx4a9bmEtNNZwSKyD1rVs2pKAHp041wm7s4kG5kEwXWZrPKCsyGcqvtbtvztPO7TrlH432Ejvzzn_Gewqm4bOGA8lYmJgfUoHnS3ZyO29SkVyskMMfEjmGkCP23etmzKOpTHf0RtrcUNYOJx072Bja8cHL9QojtQM59J2SEFYuT7YhC6MMKAwuI6aNJPGdvysnPHR6B-htVyuD2r2md1UoI4JStR4LO0BjB-vSWOHDuTPFolus80u484_PK5VU5T9k29GiBO0nz4eAcT6-8Wz9_fP3i8b76d6CmNCUklQ7bS6Z0p1tadtx4ba4q4fv9owqJEObkc9I7DimPb2VEWOuhh-cIEaGI0EqlSYaf0wHwtTNM1wpWEm7rHBHUPkA87zxUk11bbrYDKNdyoi88fnHAQBxhx9dgfFCzPpQa34YO4cqwfEaWafKuUzzIQybLFTVpEXImVPPZmcdaZS6VHxZ_LIWA4dHK9Rqy2qQVdHin8jDg0k6_L2BvXlcN2XB2_mb4r9-mQPTNWmOjQ8Cm_3gpPXHhxBYTcrE4crFN8KB2br847FEiNIjIKsQVjftSGvoXzuntZtf4BVGv98GunuM7VSZePt1w0aJ8jHDf1PIW2HGRon8PrTP_EKAqoH8hN0Ykwd4WjBkVzrwUQUtrc3jwDBqTrbKoZWg8Dac_-csubXl0ixkU8b5-nahccmJzoNJ7j2tnfAi4ojn548a4KR2xM0aHp0CVkhsD7jPm6i8tShx9q7ZuRTI0fDwbnxiLuGX9UTTTVvbI7DtaOObVb26snJCZ82j16fnp6212I8Dt2JLGpHhjfqeXcJwrgVf60oXq9lUerMpjtcbHVaXEHQyfX6v1qtbBHHh0cvy0Pc5mfK9l9erX1QYBDv7R2-on2h4-nAzKJXB7S_j75R7Nval8hkPxWzkf0_WKpGNZlxHKe4uHD_tMIFTNTZfeupE-4jeI37QLz19Mn-_BdaHj4p)) +- [Using shared memory as storage](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_set/shared_memory_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyVWIFuG7kR_RViiwJyTpbsoMUViixUZ7uokNQ-2E6Cw_mwpnYpifVqqSO5UnSG_71vhlztSo5b10GciBzODN_MvBnqKXHKOW1Klwx-fUp0ngxOu0khy3kl5yoZJFmVy6SbOFPZjD73392X4p04N6ut1fOFF53sSLw_ef-Xrrj6MrmYjMX59c3P1zfju8n1VY9kWf6TzlTpVC6qMldW-IUS45XM8E_c6YovypIj4n3vRHRI4D6Je_fJ0QfWsjWVWMqtKI0XlVNQo52Y6UIJ9S1TKy90KTKzXBValpkSG-0XbCrqYXfEL1GJmXoJeYkTK3yatSWF9DvX6Wfh_WrQ7282m55kt3vGzvtFEHb9T5Pzy6vby2O4vjv2uSyArLDq90pbXHy6FXIFzzI5hb-F3AhjhZxbhT1vyPON1V6X865wZuY30irWk2vnrZ5Wfg-82k_cvy0A-GQJ4Ma3YnJ7n4ifxreT2y7r-Tq5--f15zvxdXxzM766m1zeiusbBOvqYkKhwqd_iPHVL-Lj5OqiKxSggyn1bWXpFnBVE6wqDxjeKrXnxswEt9xKZXqmM1FnkJibtbIlriVWyi51yDU4mbOeQi-1l57XXlyOTfXvy_vyT7rMiipXYphVmek7OpKlTvleVi1GBxLGwBIE1iqdW1OtXC-I9N-FwPydE8YtAHCeLtXS2G2qvklcT0FdEJlarWbiApslwJVewT1H14lp0nhAgQu6RNDVeO0VdOIsfCokQLxV_kbN4EuazgszlUWairXROc7jKClLH4GVKjp-u1KlXKp4ZDBQ37wqfUrrYlpljxANSxTb__VDmEHFcuW3MLAdvtCORVY9EjsheFMiGVVxdF8-kY1-X3xFnLV1XsiiMBldjKFoXx7AGEs41QnxgFs99KKCO0YuCCBxqVKtJlIAqpmBvXllKhcv6GjVFca7bjy-WehswcY3rH-tMijTf0BBYWTu2Eyaxsim4sU9I3ItFN2ve2j21rKoVOfotw8UPzZ6hZ2u2ChyEKlQZZ7vJTNfyWJ3GzP9N5wRM2uWvG1R3tJaydkgZIUCr0XPdu7EldSqGfvU2Y9s7SJzH7uCe5MjqHANXyQA2wVBtYxHvS3T2K3Ndl6GuCuenpu_8fRRAwGFLaushTj0WyVBZkVmHkVfnN-NGzOMELZM9ghzL-twMCC-ToOKlOU6LTOTEvQnC8RzB2GNWeWIQBD5aN9R1ZEQK2EPiAz0TkOHN3bI3W7LbGFNWesOPsLlpXwEXpVFPtZnmY3EQlJjKbVbBMqLlhwUdQ6g4RSvMyA3oApqT0ju2IBkuRWzqsxILwz4ba-ppIVcK_IDHFM65PRSaGIUrMjaRFTsF9LD09XKWJQGXaLWSTlRUhfpMChUWBKU_IAcUdY_HDUBCiuUb4gPAWbVVJd5GgJlrOsEqghyrWteouXVoQ-biAC7EZby4HPMRDb4KsnENMFHOBFQjRkBDEJCxGzn3Mat9Ywb_0ZS_hlu_Gja89RROA33Kfwq0fwwRZBkjWoYEVQ7FWtQ0ZSIMwKf0xARwewK3VO9brxlJ0M3hqfBqQa-XtyOW69kWeblfo5RAoeTIjpEfY-9_W-5ts9DNQwYQpDDOvQjOOKYGKZEveikmCxouAmFw0RMBIsl97C7K-dhBuh9DYtEIm2QvuWx2XDPjsk3lS4MF6ypweHhIC1bRjCPQBGIuZwrG9U3cjE368ZwPv5yOb4bUGrUREOxQ3D2jGFqaEyEJR3mBsBHI5YCltrYQ0pqH0HSteL4SgHUB1rBPa_hpiC-gJvTYyY6ofIba736A6fKkXgS6Hiln3Xuk8ubm-ubgfiIOvhzzpQRNMFiEpNOPN-Xz-QBjiCNdNmhcaHuyCG2H7mOILBPSvcJ8zxGwJrlCSqJcRR3XwMr7nR70ykmfB7y1BInODd7O2YmwDgfLI5g2lTrkFfQGRThP59_YvnDaQMOjkIcaJj8zoTxdHz6XGf5V52DMRH3w94u8sryEGkNIjY_iDBrjt2TWeFMnNYqz5viF9yHBEl853zNKOFsPP0v1OWyWoqyWk7BM_VIUvefmu2o-nkkPBiHIr6BsFk7TE6ZvzD80BzqNThyLa2mZ0GtK2LIc8BQ0yxwenJyMiKUSSEgV4EEHhr3KUcyepygw9AdcODDQ5Mn7Qny7DUDNWYXXFDsdARcOPDNMjgYFMaNNGzsKy5wWto0igwjst24i3KVVeFT9NhFWpMR58moAR7UlDde0ATDFuYIY4B-JS3ai8ejUcgp-LzlmlMHF23GdTLzlpkZPy3E3niC3so05HA3c3hWqjDlvPE4DlbOIyy_Y7RMvWFI3nj2O9F4s9PMeHA9jXO9sf-P5RphntSGrSIcjV70LqJMPPKoFCmErUYDrlTgfTzaOXDNLBsD3zwhUAM61BNLbhS3iliP4f1ElbGxckV8R1PkXknyWC4kkeFDkxcPB-lTz-RIod0kU6fVYFDvDlsJew4ciD-NU6S7IQy6JShzjemSe4MFkdPrJni8K1nm0b1rNgOepsmXxggejJEfytFrf3e2Q9V7xP0R9FzZoLoE7HiI26YBwwZmtfjWeZ1Cg8CuemiASfe2hm2IRp1WpTw9t3rmJ4kbLNqBCQOxIJiJ--rHxG6IP3gI79sZDoenIKr3fxuNRodvpe88XD-EbpTLC3SrTLWmM56q0FiTbhJp2DZffyXlOstO3_-1OsW2Wfnw3Vhy7Hx-lv3ww-mP4ljabHHmlumPJ-L4GMD5Y3YjV_lxIZdT_sKs0NOWzizLCiyuw1dcWMBLsnxMnrv1PoDe28_VOnn-jf_8B80ktFI=)) - [Using set as mapping table to handle large keys or indeterministic sentinels](https://github.com/NVIDIA/cuCollections/blob/dev/examples/static_set/mapping_table_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJy1WQ1v20YS_SsDHe4iJZRkKck1UGyjOqdFjRb2IXZaFLYhr8iVRITkqsulFJ3g_35vZknqw3KdHnIOHIvk7OybmTcfS60buc7z2GR5Y3CzbsRRY9ALGonKpoWa6sagERaRagSN3BQ25Ovuy9uMXtKZma9sPJ05aoYt6h_13wR08ev5h_MhnV1-_Pflx-H1-eVFh2VF_pc41FmuIyqySFtyM03DuQrxp3wS0K_aMhDqd46oyQK3jfLZbaP1XrSsTEGpWlFmHBW5hpo4p0mcaNJfQj13FGcUmnSexCoLNS1jN5OtSj0Ch34vlZixU5BXWDHH1WRbkpSrofPPzLn5oNtdLpcdJbA7xk67iRfOu7-cn_1wcfVDG9DrZZ-yBJ4lq_8oYgvDxytScyAL1Rh4E7UkY0lNrcYzZxj50sYuzqYB5Wbilspq0RPFubPxuHA7zqtwwv5tAbhPZXDc8IrOr24b9K_h1flVIHp-O7_-6fLTNf02_PhxeHF9_sMVXX5EsC4-nHOocPUjDS9-p5_PLz4EpOE6bKW_zC1bAagxu1VH3odXWu_AmBgPK5_rMJ7EIVUMoqlZaJvBLJprm8aeawAZiZ4kTmOnnNx7ZJxs1b3NbrO_xVmYFJGm47AITTfnJeEo164TFrPTfYlIQSLqKmvV6vTgI1fME73zyM1skbtupBfYe7TQoTO2s686NvC0Vqnc7r70cf5e-JcitjBy5Di6I_1FwVsa6Go6fD-2sZ7QB53CVmeV0znNzJJDz0yeqXxGMIgUnEOJMZ-LOYmyipkG6RZnKqFIOVWrvWeH3PvlIp5TXoQzViOPBoONs-45jLt3AfseC-ZzYx3Ik6zoTfcdqOo0fdar3Af7mrNsEylSNgbpaaxDxdAPQLAamljSmRRkOENKgs7tYRa1r5ZqTs2z4VVLdBuwoow_cwibIhNybf2tJm79UWh4Lm8FTBrxxEzZiNODMqxc8FalBXlpQlsMqDj2zl9iyw5dG9KZeLUymndNlJ1q2TwPoEJzolqDNOf04tzkYKY6c5VBiQk_V7b6SuHZX-YtJRqwOvSTWeKDDXydqpUuTZFEkFGS-IqrmUCCK4AmldIV6RBMg3ejAplm6i1mqLlki8zFqUbhyh3T4zGqTeEdRlHMzlRJsgoOxYpDiCggWFHEuc6pCvZ4R0MsNEkSS2HGU5MUQoFUu5mJtuhRG1fWO9FrxXt4sIiRPgr8Bu5MJ1CuuFZHKIacB0VWcGvIE4MQohDy3hXLCclShK6wGtFjH1c6FioppP4pBqbtAhr8PdGeIp1BUco4AsRVjONXKgck2J526MI4XfpWEAnz4ICSdytxxoE9uVZNcAtlXc_UIkax8H4IUX7H9QacY776aqmMHARazrRHEWfzwomh5FZz0SvNS38RqkfaccXMPAZsHUcHkIA8lZsU92MwCBSMNgQQWLnUI6iPmFpbJQjGVTFlLnoKMDrpLXnBeb5CLBnETo3KRDnH0KJYCid8CSyrFrQ9qloD30_bEkjOttLjcGMZmA0xifNIKAIIh0ogSX3vVCo_CcT7du_el9DHnsrVRNdFS4VhyfV7VnbT7t3ds1CciaNrtWegkUkFF9OZ48I1CiSHmFtJwBRcafghENftHvaLMbzCC23hZ9ur2WKsct-4S2M3kfsesdU-fmVDYYRcsDYhZMfPCzs3XJG5-nXo3JW8gEanQVQBk0uzgRNTFA1RX5cbR0uMQ5CJfMyqzVQyBWI3SzeNuG56vpt9QvTaVSbUfmFVlcluaWpr_UrH_krpmi63vCEJUO7iNHYHOdGUcTNTKHTX6Li-DlC6GvlN1hCvr5rX2BEiL4UaLRrQiD-s-b8HWj-wrMdOG_RnMx1-RsgnNUpPuoipE2KLOjTwSmxrNnqulroOcNI_qfby5iazXFif6IkDMUCKnGsiVOovu5K2lPTT9ROiVqMiZvQCPtEvKgtkqtmzwS_o8p_RqJxuRt5XPNvR2JikbMLGNltNUOZ1f-QYb0DVBSC1_BrWI46HO7vegaBJQm9I-w4puao8FKFUWcT0syCptImHNJ5PosFgqt3x0WlTQnkDRHctOjl5UsCKgO_6m5896d5z6np_SV3_OXX9v6Tu9XPqXu-o40MRCbdr-stT3H94_1y6SnWaFFlYFX9fex-V2U19kubsSfjNctkX1jqZ_eXXZPMWm1UB9Ac4LGAPEveSx0TZyw86k9hiWigpvM3gp4vtV3FWEPwPcXKxNBcZnmCNYMRCV6ObYu7N6F5A3h-IxzXPE2XXFKE_j8fCYLaQrSrX_8Mvw5AuXhPTQoNx5fgYx8sbHC3x4bHZfpEXCuigVO-x1E5KHFjS_4olT4F6Xa29Obr7U1wbwd7d10PbrOp_rfrXpeDd3S1i32ASMAc4zqmKsyYHo_I7uHo92x0c6lmxPLVgcs2qM041bPAIyT1_rpPEE7zgD7yQY86a_Uj3M_TtFBnBeFyUKRRQiMNWIF0i2BaTPnccmQKTWkBvPv1yevq-xHueqinXe95WxNC-_iOAZ8WUTZEn24akxuoKNpoOj-244kabItXy2I-TpfpNQ3liJuaXPkxUmGANT90y0G8dKXy9EJbTTPKQTnyYxTR_9j-GZ07X-45Z93pFQC_UCxzp0Hkfu2Td6wTUx-9r_L7pPDwEuwQ6VPj_35r7fdY8flrzW2j8J36_w--7b4r5G2ieqCT_Gnc8vN8LrbDupIxwh6-avhD7Vz2Dwc67Hok3Rb7P-DW-NNc5qHAI3ctAlfkeIzdwuC6z5r7OTtKZKXBgF1CL8u3WcmaS6vBnVTbV5SacvNWri_K9yfCqzg89mcRhrKUZKH9YKpPeZPKmrMzoocD0eV3iqTJTDkdyYD6YOFgeijakEKqFPzj7c9LWIYl1bfbYDJJoLG41gndGtdoTavdKH27FhTWN5s7isXd3h_9rtjqokmWERJoPmlyadt9hrf0N_YV74bHwgUM7cqdrCfhL6j8E3p1jrgJvj_5OiVEReMRhfoJ_W0QU7ZUt68dWPcnh-qc6k6wrS59f4rdNoF_ZEb87QCSPe8FmIhJ6nq53rzcbPDzJ1Xz_HP22ntT34uLXjGTNyRNJUgf-dH0UEAD2kYDIP598_D7WH-ibW8o6Y41-0OT3d1s3cS5ttlrb7Ciy-I9C-_03P88jaXLcW_vZv6VthL2gh9FhXMOAtdAjnFmaWyIVxkdqsiLdCdSJ7xH8yp1f0x3SEezvXVm5N0MhTta_5fDzAm8lD-SFPnLRlTWiitegGhdIxtKmwIwB6eg9_oiC9_TqVdzaGnQ_lgb7t1mEzVnj_isHL-7nP1_7brZsuInvNiMs_ymH3iM_uTSCBvfuONF28x1SI1uEYa__tujhsZk7_wVTow0fnISvXvW-o7ay4ewkT0ffHVG7jRLi2pLUkY7aiUrH8q1TEo-3dIZhmODmwn9PhBvoOtnnxkNQPUcW7TwHZRoPd_Lvv3bzBxQ=)) ### `static_map` @@ -256,4 +256,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries. #### Examples: -- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) +- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw)) \ No newline at end of file diff --git a/examples/static_set/device_subsets_example.cu b/examples/static_set/device_subsets_example.cu index 498c41d69..054154348 100644 --- a/examples/static_set/device_subsets_example.cu +++ b/examples/static_set/device_subsets_example.cu @@ -44,7 +44,7 @@ */ auto constexpr cg_size = 8; ///< A CUDA Cooperative Group of 8 threads to handle each subset -auto constexpr window_size = 1; ///< Number of concurrent slots handled by each thread +auto constexpr bucket_size = 1; ///< Number of concurrent slots handled by each thread auto constexpr N = 10; ///< Number of elements to insert and query using key_type = int; ///< Key type @@ -54,7 +54,7 @@ using probing_scheme_type = ///< and probing scheme (linear ///< probing v.s. double hashing) /// Type of bulk allocation storage -using storage_type = cuco::aow_storage; +using storage_type = cuco::aow_storage; /// Lightweight non-owning storage ref type using storage_ref_type = typename storage_type::ref_type; using ref_type = cuco::static_set_ref(cuco::make_window_extent(subset_sizes[i]))); + static_cast(cuco::make_bucket_extent(subset_sizes[i]))); } std::vector offsets(num + 1, 0); - // prefix sum to compute offsets and total number of windows + // prefix sum to compute offsets and total number of buckets std::size_t current_sum = 0; for (std::size_t i = 0; i < valid_sizes.size(); ++i) { current_sum += valid_sizes[i]; offsets[i + 1] = current_sum; } - // total number of windows is located at the back of the offsets array - auto const total_num_windows = offsets.back(); + // total number of buckets is located at the back of the offsets array + auto const total_num_buckets = offsets.back(); // Create a single bulk storage used by all subsets - auto set_storage = storage_type{total_num_windows}; + auto set_storage = storage_type{total_num_buckets}; // Initializes the storage with the given sentinel set_storage.initialize(empty_key_sentinel); diff --git a/examples/static_set/shared_memory_example.cu b/examples/static_set/shared_memory_example.cu index 3c9ccd76e..5fc3a1076 100644 --- a/examples/static_set/shared_memory_example.cu +++ b/examples/static_set/shared_memory_example.cu @@ -24,16 +24,16 @@ */ template -__global__ void shmem_set_kernel(typename SetRef::extent_type window_extent, +__global__ void shmem_set_kernel(typename SetRef::extent_type bucket_extent, cuco::empty_key empty_key_sentinel) { // We first allocate the shared memory storage for the `set`. - // The storage is comprised of contiguous windows of slots, + // The storage is comprised of contiguous buckets of slots, // which allow for vectorized loads. - __shared__ typename SetRef::window_type windows[window_extent.value()]; + __shared__ typename SetRef::bucket_type buckets[bucket_extent.value()]; // Next, we construct the actual storage object from the raw array. - auto storage = SetRef::storage_ref_type(window_extent, windows); + auto storage = SetRef::storage_ref_type(bucket_extent, buckets); // Now we can instantiate the set from the storage. auto set = SetRef(empty_key_sentinel, {}, {}, {}, storage); @@ -75,7 +75,7 @@ int main(void) // Inserting or retrieving this value is UB. cuco::empty_key constexpr empty_key_sentinel{-1}; // Width of vectorized loads during probing. - auto constexpr window_size = 1; + auto constexpr bucket_size = 1; // Cooperative group size auto constexpr cg_size = 1; @@ -93,7 +93,7 @@ int main(void) thrust::equal_to, probing_scheme_type, cuco::cuda_allocator, - cuco::storage>; + cuco::storage>; // Next, we can derive the non-owning reference type from the set type. // This is the type we use in the kernel to wrap a raw shared memory array as a `static_set`. using set_ref_type = typename set_type::ref_type<>; @@ -101,9 +101,9 @@ int main(void) // Cuco imposes a number of non-trivial contraints on the capacity value. // This function will take the requested capacity (1000) and return the next larger // valid extent. - auto constexpr window_extent = cuco::make_window_extent(extent_type{}); + auto constexpr bucket_extent = cuco::make_bucket_extent(extent_type{}); // Launch the kernel with a single thread block. - shmem_set_kernel<<<1, 128>>>(window_extent, empty_key_sentinel); + shmem_set_kernel<<<1, 128>>>(bucket_extent, empty_key_sentinel); cudaDeviceSynchronize(); } diff --git a/include/cuco/bucket_storage.cuh b/include/cuco/bucket_storage.cuh index eb16c8313..f895d79aa 100644 --- a/include/cuco/bucket_storage.cuh +++ b/include/cuco/bucket_storage.cuh @@ -33,10 +33,6 @@ namespace cuco { template using bucket = detail::bucket; -/// Alias for bucket -template -using window = bucket; - /** * @brief Non-owning array of buckets storage reference type. * @@ -56,7 +52,6 @@ class bucket_storage_ref : public detail::bucket_storage_base using size_type = typename base_type::size_type; ///< Storage size type using value_type = typename base_type::value_type; ///< Slot type using bucket_type = typename base_type::bucket_type; ///< Slot bucket type - using window_type = bucket_type; ///< Slot bucket type using base_type::capacity; using base_type::num_buckets; diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl index 916d75e2c..bb5145c70 100644 --- a/include/cuco/detail/extent/extent.inl +++ b/include/cuco/detail/extent/extent.inl @@ -111,24 +111,12 @@ template } } -template -[[nodiscard]] auto constexpr make_window_extent(extent ext) -{ - return make_bucket_extent(ext); -} - template [[nodiscard]] auto constexpr make_bucket_extent(SizeType size) { return make_bucket_extent(extent{size}); } -template -[[nodiscard]] auto constexpr make_window_extent(SizeType size) -{ - return make_bucket_extent(extent{size}); -} - template [[nodiscard]] auto constexpr make_bucket_extent(extent ext) { @@ -167,15 +155,6 @@ template N>(ext); } -template -[[nodiscard]] auto constexpr make_window_extent(extent ext) -{ - return make_bucket_extent(ext); -} - template [[nodiscard]] auto constexpr make_bucket_extent(SizeType size) { @@ -185,15 +164,6 @@ template dynamic_extent>(extent{size}); } -template -[[nodiscard]] auto constexpr make_window_extent(SizeType size) -{ - return make_bucket_extent(extent{size}); -} - namespace detail { template diff --git a/include/cuco/extent.cuh b/include/cuco/extent.cuh index 157aff9e4..e7059ab2a 100644 --- a/include/cuco/extent.cuh +++ b/include/cuco/extent.cuh @@ -85,10 +85,6 @@ struct extent { template struct bucket_extent; -/// Alias for bucket_extent -template -using window_extent = bucket_extent; - /** * @brief Computes valid bucket extent based on given parameters. * @@ -112,25 +108,6 @@ using window_extent = bucket_extent; template [[nodiscard]] auto constexpr make_bucket_extent(extent ext); -/** - * @brief Computes valid bucket extent based on given parameters. - * - * @deprecated Use the equivalent `make_bucket_extent` instead. - * - * @tparam CGSize Number of elements handled per CG - * @tparam BucketSize Number of elements handled per Bucket - * @tparam SizeType Size type - * @tparam N Extent - * - * @param ext The input extent - * - * @throw If the input extent is invalid - * - * @return Resulting valid extent - */ -template -[[nodiscard]] auto constexpr make_window_extent(extent ext); - /** * @brief Computes valid bucket extent/capacity based on given parameters. * @@ -153,24 +130,6 @@ template template [[nodiscard]] auto constexpr make_bucket_extent(SizeType size); -/** - * @brief Computes valid bucket extent/capacity based on given parameters. - * - * @deprecated Use the equivalent `make_bucket_extent` instead. - * - * @tparam CGSize Number of elements handled per CG - * @tparam BucketSize Number of elements handled per Bucket - * @tparam SizeType Size type - * - * @param size The input size - * - * @throw If the input size is invalid - * - * @return Resulting valid extent - */ -template -[[nodiscard]] auto constexpr make_window_extent(SizeType size); - template [[nodiscard]] auto constexpr make_bucket_extent(cuco::extent ext); @@ -199,24 +158,6 @@ template template [[nodiscard]] auto constexpr make_bucket_extent(extent ext); -/** - * @brief Computes a valid bucket extent/capacity for a given container type. - * - * @deprecated Use the equivalent `make_bucket_extent` instead. - * - * @tparam Container Container type to compute the extent for - * @tparam SizeType Size type - * @tparam N Extent - * - * @param ext The input extent - * - * @throw If the input extent is invalid - * - * @return Resulting valid `bucket extent` - */ -template -[[nodiscard]] auto constexpr make_window_extent(extent ext); - /** * @brief Computes a valid capacity for a given container type. * @@ -238,23 +179,6 @@ template template [[nodiscard]] auto constexpr make_bucket_extent(SizeType size); -/** - * @brief Computes a valid capacity for a given container type. - * - * @deprecated Use the equivalent `make_bucket_extent` instead. - * - * @tparam Container Container type to compute the extent for - * @tparam SizeType Size type - * - * @param size The input size - * - * @throw If the input size is invalid - * - * @return Resulting valid extent - */ -template -[[nodiscard]] auto constexpr make_window_extent(SizeType size); - } // namespace cuco #include diff --git a/include/cuco/static_map_ref.cuh b/include/cuco/static_map_ref.cuh index 75f654a73..d9b043ff5 100644 --- a/include/cuco/static_map_ref.cuh +++ b/include/cuco/static_map_ref.cuh @@ -89,7 +89,6 @@ class static_map_ref using hasher = typename probing_scheme_type::hasher; ///< Hash function type using storage_ref_type = StorageRef; ///< Type of storage ref using bucket_type = typename storage_ref_type::bucket_type; ///< Bucket type - using window_type = bucket_type; ///< Bucket type using value_type = typename storage_ref_type::value_type; ///< Storage element type using extent_type = typename storage_ref_type::extent_type; ///< Extent type using size_type = typename storage_ref_type::size_type; ///< Probing scheme size type diff --git a/include/cuco/static_multimap_ref.cuh b/include/cuco/static_multimap_ref.cuh index 0407d2d70..cb5a5a337 100644 --- a/include/cuco/static_multimap_ref.cuh +++ b/include/cuco/static_multimap_ref.cuh @@ -86,7 +86,6 @@ class static_multimap_ref using hasher = typename probing_scheme_type::hasher; ///< Hash function type using storage_ref_type = StorageRef; ///< Type of storage ref using bucket_type = typename storage_ref_type::bucket_type; ///< Bucket type - using window_type = bucket_type; ///< Bucket type using value_type = typename storage_ref_type::value_type; ///< Storage element type using extent_type = typename storage_ref_type::extent_type; ///< Extent type using size_type = typename storage_ref_type::size_type; ///< Probing scheme size type diff --git a/include/cuco/static_multiset_ref.cuh b/include/cuco/static_multiset_ref.cuh index 203cdc8b5..569a4551c 100644 --- a/include/cuco/static_multiset_ref.cuh +++ b/include/cuco/static_multiset_ref.cuh @@ -77,7 +77,6 @@ class static_multiset_ref using hasher = typename probing_scheme_type::hasher; ///< Hash function type using storage_ref_type = StorageRef; ///< Type of storage ref using bucket_type = typename storage_ref_type::bucket_type; ///< Bucket type - using window_type = bucket_type; ///< Bucket type using value_type = typename storage_ref_type::value_type; ///< Storage element type using extent_type = typename storage_ref_type::extent_type; ///< Extent type using size_type = typename storage_ref_type::size_type; ///< Probing scheme size type diff --git a/include/cuco/static_set_ref.cuh b/include/cuco/static_set_ref.cuh index 846ad9252..0bf2ed4e3 100644 --- a/include/cuco/static_set_ref.cuh +++ b/include/cuco/static_set_ref.cuh @@ -77,7 +77,6 @@ class static_set_ref using hasher = typename probing_scheme_type::hasher; ///< Hash function type using storage_ref_type = StorageRef; ///< Type of storage ref using bucket_type = typename storage_ref_type::bucket_type; ///< Bucket type - using window_type = bucket_type; ///< Bucket type using value_type = typename storage_ref_type::value_type; ///< Storage element type using extent_type = typename storage_ref_type::extent_type; ///< Extent type using size_type = typename storage_ref_type::size_type; ///< Probing scheme size type diff --git a/tests/static_map/insert_or_apply_test.cu b/tests/static_map/insert_or_apply_test.cu index 03d8ce5c9..451a310b2 100644 --- a/tests/static_map/insert_or_apply_test.cu +++ b/tests/static_map/insert_or_apply_test.cu @@ -101,7 +101,7 @@ void test_insert_or_apply_shmem(Map& map, size_type num_keys, size_type num_uniq cuco::storage<1>>; using shared_map_ref_type = typename shared_map_type::ref_type<>; - auto constexpr window_extent = cuco::make_window_extent(extent_type{}); + auto constexpr bucket_extent = cuco::make_bucket_extent(extent_type{}); // Insert pairs auto pairs_begin = thrust::make_transform_iterator( @@ -122,7 +122,7 @@ void test_insert_or_apply_shmem(Map& map, size_type num_keys, size_type num_uniq init, cuco::reduce::plus{}, map.ref(cuco::op::insert_or_apply), - window_extent); + bucket_extent); REQUIRE(map.size() == num_unique_keys); diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu index 0059b61dd..2f782aa57 100644 --- a/tests/static_map/shared_memory_test.cu +++ b/tests/static_map/shared_memory_test.cu @@ -30,7 +30,7 @@ #include -template +template __global__ void shared_memory_test_kernel(Ref* maps, typename Ref::key_type const* const insterted_keys, typename Ref::mapped_type const* const inserted_values, @@ -42,7 +42,7 @@ __global__ void shared_memory_test_kernel(Ref* maps, const size_t map_id = blockIdx.x; const size_t offset = map_id * number_of_elements; - __shared__ typename Ref::window_type sm_buffer[NumWindows]; + __shared__ typename Ref::bucket_type sm_buffer[NumBuckets]; auto g = cuco::test::cg::this_thread_block(); auto insert_ref = maps[map_id].make_copy(g, sm_buffer, cuco::thread_scope_block); @@ -122,9 +122,9 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests", thrust::device_vector d_refs(h_refs); // maybe_unused to silence false positive "variable set but not used" warning - [[maybe_unused]] auto constexpr num_windows = cuco::make_window_extent(extent_type{}); + [[maybe_unused]] auto constexpr num_buckets = cuco::make_bucket_extent(extent_type{}); - shared_memory_test_kernel + shared_memory_test_kernel <<>>(d_refs.data().get(), d_keys.data().get(), d_values.data().get(), @@ -152,9 +152,9 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests", thrust::device_vector d_refs(h_refs); // maybe_unused to silence false positive "variable set but not used" warning - [[maybe_unused]] auto constexpr num_windows = cuco::make_window_extent(extent_type{}); + [[maybe_unused]] auto constexpr num_buckets = cuco::make_bucket_extent(extent_type{}); - shared_memory_test_kernel + shared_memory_test_kernel <<>>(d_refs.data().get(), d_keys.data().get(), d_values.data().get(), @@ -167,19 +167,19 @@ TEMPLATE_TEST_CASE_SIG("static_map shared memory tests", } auto constexpr cg_size = 1; -auto constexpr window_size = 1; +auto constexpr bucket_size = 1; -template +template __global__ void shared_memory_hash_table_kernel(bool* key_found) { using Key = int32_t; using Value = int32_t; using slot_type = cuco::pair; - __shared__ cuco::window map[NumWindows]; + __shared__ cuco::bucket map[NumBuckets]; - using extent_type = cuco::extent; - using storage_ref_type = cuco::aow_storage_ref; + using extent_type = cuco::extent; + using storage_ref_type = cuco::aow_storage_ref; auto raw_ref = cuco::static_map_ref{cuco::empty_key{-1}, @@ -213,11 +213,11 @@ TEST_CASE("static map shared memory slots.", "") { constexpr std::size_t N = 256; // maybe_unused to silence false positive "variable set but not used" warning - [[maybe_unused]] auto constexpr num_windows = - cuco::make_window_extent(cuco::extent{}); + [[maybe_unused]] auto constexpr num_buckets = + cuco::make_bucket_extent(cuco::extent{}); thrust::device_vector key_found(N, false); - shared_memory_hash_table_kernel<<<8, 32>>>(key_found.data().get()); + shared_memory_hash_table_kernel<<<8, 32>>>(key_found.data().get()); CUCO_CUDA_TRY(cudaDeviceSynchronize()); REQUIRE(cuco::test::all_of(key_found.begin(), key_found.end(), thrust::identity{})); diff --git a/tests/static_set/shared_memory_test.cu b/tests/static_set/shared_memory_test.cu index b970e735f..fd1c34118 100644 --- a/tests/static_set/shared_memory_test.cu +++ b/tests/static_set/shared_memory_test.cu @@ -30,7 +30,7 @@ #include -template +template __global__ void shared_memory_test_kernel(Ref* sets, typename Ref::key_type const* const insterted_keys, size_t number_of_elements, @@ -41,7 +41,7 @@ __global__ void shared_memory_test_kernel(Ref* sets, const size_t set_id = blockIdx.x; const size_t offset = set_id * number_of_elements; - __shared__ typename Ref::window_type sm_buffer[NumWindows]; + __shared__ typename Ref::bucket_type sm_buffer[NumBuckets]; auto g = cuco::test::cg::this_thread_block(); auto insert_ref = sets[set_id].make_copy(g, sm_buffer, cuco::thread_scope_block); @@ -108,9 +108,9 @@ TEMPLATE_TEST_CASE_SIG( } thrust::device_vector d_refs(h_refs); - auto constexpr num_windows = cuco::make_window_extent(extent_type{}); + auto constexpr num_buckets = cuco::make_bucket_extent(extent_type{}); - shared_memory_test_kernel + shared_memory_test_kernel <<>>(d_refs.data().get(), d_keys.data().get(), elements_in_set, @@ -136,9 +136,9 @@ TEMPLATE_TEST_CASE_SIG( } thrust::device_vector d_refs(h_refs); - auto constexpr num_windows = cuco::make_window_extent(extent_type{}); + auto constexpr num_buckets = cuco::make_bucket_extent(extent_type{}); - shared_memory_test_kernel + shared_memory_test_kernel <<>>(d_refs.data().get(), d_keys.data().get(), elements_in_set, @@ -150,18 +150,18 @@ TEMPLATE_TEST_CASE_SIG( } auto constexpr cg_size = 1; -auto constexpr window_size = 1; +auto constexpr bucket_size = 1; -template +template __global__ void shared_memory_hash_set_kernel(bool* key_found) { using Key = int32_t; using slot_type = Key; - __shared__ cuco::window set[NumWindows]; + __shared__ cuco::bucket set[NumBuckets]; - using extent_type = cuco::extent; - using storage_ref_type = cuco::aow_storage_ref; + using extent_type = cuco::extent; + using storage_ref_type = cuco::aow_storage_ref; auto raw_ref = cuco::static_set_ref{cuco::empty_key{-1}, @@ -191,11 +191,11 @@ __global__ void shared_memory_hash_set_kernel(bool* key_found) TEST_CASE("static_set shared memory slots test", "") { constexpr std::size_t N = 256; - [[maybe_unused]] auto constexpr num_windows = - cuco::make_window_extent(cuco::extent{}); + [[maybe_unused]] auto constexpr num_buckets = + cuco::make_bucket_extent(cuco::extent{}); thrust::device_vector key_found(N, false); - shared_memory_hash_set_kernel<<<8, 32>>>(key_found.data().get()); + shared_memory_hash_set_kernel<<<8, 32>>>(key_found.data().get()); CUCO_CUDA_TRY(cudaDeviceSynchronize()); REQUIRE(cuco::test::all_of(key_found.begin(), key_found.end(), thrust::identity{})); diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu index 96e97e4b9..c88ae08d3 100644 --- a/tests/utility/extent_test.cu +++ b/tests/utility/extent_test.cu @@ -26,7 +26,7 @@ TEMPLATE_TEST_CASE_SIG( SizeType constexpr num = 1234; SizeType constexpr gold_reference = 314; // 157 x 2 auto constexpr cg_size = 2; - auto constexpr window_size = 4; + auto constexpr bucket_size = 4; SECTION("Static extent must be evaluated at compile time.") { @@ -43,14 +43,14 @@ TEMPLATE_TEST_CASE_SIG( SECTION("Compute static valid extent at compile time.") { auto constexpr size = cuco::extent{}; - auto constexpr res = cuco::make_window_extent(size); + auto constexpr res = cuco::make_bucket_extent(size); STATIC_REQUIRE(gold_reference == res.value()); } SECTION("Compute dynamic valid extent at run time.") { auto const size = cuco::extent{num}; - auto const res = cuco::make_window_extent(size); + auto const res = cuco::make_bucket_extent(size); REQUIRE(gold_reference == res.value()); } } diff --git a/tests/utility/probing_scheme_test.cu b/tests/utility/probing_scheme_test.cu index 994d5e443..4eb01199a 100644 --- a/tests/utility/probing_scheme_test.cu +++ b/tests/utility/probing_scheme_test.cu @@ -79,7 +79,7 @@ __global__ void generate_cg_probing_sequence(Key key, TEMPLATE_TEST_CASE_SIG( "utility probing_scheme tests", "", - ((typename Key, cuco::test::probe_sequence Probe, int32_t WindowSize), Key, Probe, WindowSize), + ((typename Key, cuco::test::probe_sequence Probe, int32_t BucketSize), Key, Probe, BucketSize), (int32_t, cuco::test::probe_sequence::double_hashing, 1), (int32_t, cuco::test::probe_sequence::double_hashing, 2), (int64_t, cuco::test::probe_sequence::double_hashing, 1), @@ -89,7 +89,7 @@ TEMPLATE_TEST_CASE_SIG( (int64_t, cuco::test::probe_sequence::linear_probing, 1), (int64_t, cuco::test::probe_sequence::linear_probing, 2)) { - auto const upper_bound = cuco::make_window_extent<1, WindowSize>(cuco::extent{10}); + auto const upper_bound = cuco::make_bucket_extent<1, BucketSize>(cuco::extent{10}); constexpr size_t seq_length{8}; constexpr Key key{42}; From e79787be2cb3de1b12e90d56355612e47395cce5 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Thu, 19 Dec 2024 12:01:03 -0800 Subject: [PATCH 3/3] Upgrade rapids-cmake to 25.02 (#652) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Upgrade `rapids-cmake` to 25.02 Co-authored-by: Daniel Jünger --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 8de18db73..9ba53cbfb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -16,7 +16,7 @@ cmake_minimum_required(VERSION 3.23.1 FATAL_ERROR) if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake) - file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-24.10/RAPIDS.cmake + file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-25.02/RAPIDS.cmake ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake) endif() include(${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake)