diff --git a/README.md b/README.md index fa5885281..2808f2838 100644 --- a/README.md +++ b/README.md @@ -260,7 +260,7 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection #### 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/eJydVgtr40YQ_iuDCq2dWH6E3h3IccC9aw7DcckloRTqYtarkbVE3lX34YsJ-e-dXUm27DhHqQxGmvd88-1Iz5FBY4SSJkr-eo5EGiWjXlQwuXJshVEScZeyqBcZ5TT3z4OzuYQzuL_99Gd8LQr8qMqtFqvcPuCTTWD3CB3ehYvhxa8x_b3vwdc_Zp9mU_h4c3d7czd9mN18hZ9hen09-zKbPvx-34dpUUDwNKDRoN5g2t-n-iI4SoPxLEVpRSZQJzAtGc8xvugPvd1gLufyJyF54VKES-64GiwLpdaLTBQWdZ-7_OrQxubaGTvgyknb98pXqhQ3lHixQW6VPm2CT8idJQQXpSoE3562MviPQ8mxf1yDUMZqZOsgFtLCmgnZ2SiRdufymRoDL-Q0IItPpQbp1otH3BqgawKj4S_D4XAMu2swGFzCZ5SomcVaDd7-dCRbVm6TfdwzGPbfjetIM8JcW7A5Qia0sZCzIgvxfDBVKwK8bySQrxLEdebxaw_jlvW0DHlchM5CHXTrlXGjLFFDGG6d3YPnTX1pYLclVaWOLQCcEXLlTRbBZOKzj3eu9yX7LoEdeMF3YXNIMWOusFANGJhMjwsK3XvKJUmbc5dNrqs63nOrw5c6tbFpkhAHLVxewjz6zRWPVFjAPWD8dj0ZtYO61B7EVTVz4mFdZkKx5hJaF4UPuVCmRZ27YmeSHPC8VbUfWKeZXHfcdmkY3fGq_hJXRNtuL3j0KYO_H3XrNMx5rpSLYEW4t13GbQNyDGTZ2Z63yVJZyV2YyuFA1QTYl7Ef8GkyqyxIgoM3rOnM0rTTVNGrM3V_iNpSqeLKW9L6oul0qsJ7kLHC4CF2Jx1l2_HgpLRi1K18c6i3rdMX6O4fS02hlTPFtuYQpvvWyPEhFwZMrlyRQpWO7Ij0VjuMS2WEFRuEsDwImIfbu8mojQodVUv7ybyCprdvvBnsfyk3tGlzZoFpBKmsb8DQivdVHS0Xf0RpbeLBgTDBsdRqyQphrOCQMsuI59px6yhWrxWmjoJPuVgKes-wCtejvq9v766GdMZKasvvC1VVQqNYUuWESvAVTW-trvxiODTNmcnR0OKhMfidRf2ehlPu4JQ1nPINOLNCEVwebV_spDriQdhp6BVeaJ1Dsp0YUFtUn1hPg24XBnXAin4Vd6u82f_Ie9xJW_SDvLJ7ekF6Us4jf9-AEMR-brU828uPF55GYoWEIT2-0GcNV-uSPmL0_usnkhvORxfv3IjUqrTVp1EUU6AJPz8ffYCYaZ5PzHrxYQhxTO8tS3-WcmAaF2y9DN9LhVi2YnLOCxJuKBHFIwH1Kx-jl16jpzfHgZ6wi17-Dr9_AWLvS3Y=)) -- [Persisting L2 access](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/persisting_l2_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzVWI1v2zYW_1cePKyzO1v-iGOnbmPAS9uDgXTJ0my34TLINEVHRChSJ1KJvbb_-x5JSZZjJ7i7DjhMLWKLfPy97w_6U0MzrbmSujH516cGjxqTfrshiLzNyS1rTBo0j0ij3dAqz6h97768kfASPl6-_bXzngt2ptJNxm9jc83WZgLVKzRpCwa9wagNP_4yfzufwdnF1eXF1ex6fvEjvIDZ-_fz8_ns-t3HAGZCgDukIWOaZfcsCrZczjllUrPOPGLS8BVn2QRmKaEx6wyCnqXr3sgb-Q2XVOQRgzc0p6q7FEol4YoLw7KA5vF0jyJihnDRZVmmsiBO0-kuiomzXBukukf-4T2jxpJND5Bo9u-cScrc7g4bbaKIraaP17g0jhI_ISFcNu8Vj1o38hMqA92X1sKA2l_HXANbkyQVDHSsHjTgHzAKaMaIYUDgB6sleC3BxMTiZXcaOJrSr3YeVBYBoRTdzDQQXYCnLNNcGy5vgUs4HwCRUQlfOKFrPw3CMqAqw5dUycgeQGrc6BDNUaMHbmJHc_bz21kBnuXoqYTB7HIe-KVKJQaLy4r1-WDmBDs1Wc4WkCrB6QZSkpGEWY0ScocyL6y7JpO6RxfAEm6sEh2vGsq0rxi1QQIxmhmtoTK4FWpJRCdhico2pdWsfXTgJCv4RwqZSmVKO3icAp4SDD1uNrDcWCszsXoNJE3xIDE2jayjchFVZ0lNoM6O4ZYMZWLWdAW0c2anECuNiS6Mu7A5-NbF4Udmzjlq3rRL7lvdmGdWzo_8D9aGIAhaC-dT58R92wgubTiskNcDQQvUWJ2Z9ZU9tAfdbC123fmztvJjlCopNvAQM-kioVBhZaOQ-6VK7S7aoLRhG9Ap1ak7tgFtMLITiKswZVLltzEIRYmwVncxXqjBopJToQxaM2O5xuoB81VdEpQwySnCkOzWJ4oN-XYZrWgmzKIqRyjuUoURLHMiUC-bUriATi346Joh24iWEe1ToGbfjN1iOJQsrCO0wOzKZcYEsaLfsUwyocGminh82kcu9_XQ7Bq9KKdX17_C5dUFjHq9HoYOoXcPDOvoB7Lu_AT_VFgFjItIeBdx-9mG_uAEPvAfrOZYmXsf0DdpbjCKWaInOywKs2kMJfgMJIragHWM5MKUrzVZ0aGfnb2wkuk64XZth7rg0dk-SNt5-vXpxYOE_imYjIZOY_d8hmFwctKDRBevR8Gwd1y-ljSDYDgc7iy6tcEY1wpQxKiBngSD8agOerTlsWU87o32QYdHW9BXozpovx_0eqUYTtLxaB_0ZLwv6fH4eAtaerzYHASj3qACfRX0Xo0fgx4Hx8PjR6DHwdG4X4G6ZguQa-tNzNnQbFIGpxhM5rXfsm0NfY8Juk4zkHkSIplGkn7vOwzWgqrbhbODbayowq4e7YZZUeqZJEthk9w1SPzvCgy2DKy9BGUhuiyqyKMMR4_axn5KWWqK4iCJONSOFrUWhPAL152CrdYeyyu-85zCfqsKPfWb0lS-JPzvj-ewXmONjMPRsMKdfi0wDiaTSY7OOxqE5mvBTv7-AP3_O8BXq7AiQn91wP0lIDaFpq9tRnbfwOVzWW0ZkdwOuVUJWW4M0yHWghATi95hmtnOpFbNWiJOJnaQc19bdtJ5vFM7XytTfrpHQFuVAEfYs4vQzrHh9dVvbsb6BzN-8mq-8LStVu14QtbhtkSFYhC6xh26xvk0qEdE6JkxGV_mhjW9eV88B9iG4qw9hW3-icnvkZj1Mlu_JyhK85TjNHPS-xaUn5aQO0_yBHSepiqzQ8oTs6sdelgJX066xai2W7Lt2O34FbyLoVeQXNLY0uwMvcVMdCgG0BTIPnTsw1VGqJtt0MbByetdcjB2yjNl-fW-8PZ1IxFFm2rzxpU7uxmaafMw-stnHdx61rv_8aS-L23lu5pOOl8WJLaRHtCw-zhLtv53bbKwsL8eHroI-YuiUATvAnZcJUYlnD6adjtVvjLvpf1-t210xSZbG7y479jbh_JkgqMzI1GoqUpZ6AO3XU_dqXebB_5Us8GXQjt_A59Mdm7p25Zo5xPdLCcQ77DySHlrb9qtYInjumy22u5EwGRkv_drSTSXGOUG8ELgSGz6q9oVwxmj-IoD8pOgrWcFXyolpjahcGKpxG77ClyT5aecZZvt1OPQa_zLufsZzTyLcq8G_Q7FMfv3RK1qNxf7t8pVvC27uzKXMcvwSo4pJoobs-WC-STpJjicKc9cNGsiXTHBiLtosqfq0RP4_2Um9iquGTN5Jl39_nIjG-0GVUnKBR6rfiVryHtK-4PjvI_bKjX-J7RGB8P8lH7_fX8MHZLR-FQn4biH1xQsYqbjUiFiUUeQZOl-VxN8WcOklApcvLfyKYkL2DblXeNLu9zHhNrZx_BpfPnd_fsTdEMSAg==)) +- [Persisting L2 access](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/persisting_l2_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJzVWAtv27YW_isHHtbZnS0_6tip2xjw0nYwkC5Zmr1wM8g0RUdEKFITqcRe2_--Q1KS5dgJdm8HXEwtYos8_M77QX9saKY1V1I3Jv_52OBRY9JvNwSRNzm5YY1Jg-YRabQbWuUZte_d59cSnsOHize_dt5xwU5Vusn4TWyu2NpMoHqFJm3BoDcYteGHn-dv5jM4Pb-8OL-cXc3Pf4BnMHv3bn42n129_RDATAhwhzRkTLPsjkXBlssZp0xq1plHTBq-4iybwCwlNGadQdCzdN1reS2_4pKKPGLwmuZUdZdCqSRccWFYFtA8nu5RRMwQLrosy1QWxGk63UUxcZZrg1R3yD-8Y9RYsukBEs3-yJmkzO3usNEmithq-nCNS-Mo8RMSwmXzTvGodS0_ojLQfW4tDKj9Vcw1sDVJUsFAx-peA_4Bo4BmjBgGBL6zWoLXEkxMLF52q4GjKf1q515lERBK0c1MA9EFeMoyzbXh8ga4hLMBEBmV8IUTuvbTICwDqjJ8SZWM7AGkxo0O0Rw1uucmdjSnP72ZFeBZjp5KGMwu5oFfqlRisLioWJ8NZk6wE5PlbAGpEpxuICUZSZjVKCG3KPPCumsyqXt0ASzhxirR8aqhTPuKURskEKOZ0RoqgxuhlkR0EpaobFNazdpHB06ygn-kkKlUprSDxyngKcHQ42YDy421MhOrV0DSFA8SY9PIOioXUXWW1ATq7BhuyVAmZk1XQDtndgqx0pjowrgLm4NvXBx-YOaMo-ZNu-S-1Y15auX8wP9kbQiCoLVwPnVO3LeN4NKGwwp53RO0QI3VqVlf2kN70M3WYtedP2krP0apkmID9zGTLhIKFVY2CrlfqtTuog1KG7YBnVKdumUb0AYjO4G4ClMmVX4Tg1CUCGt1F-OFGiwqORXKoDUzlmusHjBf1SVBCZOcIgzJbnyi2JBvl9GKZsIsqnKE4i5VGMEyJwL1simFC-jUgo-uGbKNaBnRPgVq9s3YDYZDycI6QgvMrlxmTBAr-i3LJBMabKqIh6d95HJfD82u0Ytyenn1K1xcnsOo1-th6BB6e8-wjr4n686P8IvCKmBcRMLbiNvPNvQHx_Cef2c1x8rce4--SXODUcwSPdlhUZhNYyjBJyBR1AasYyQXpnytyYoO_eTshZVM1wm3azvUBY_O9kHazuOvjy8eJPRPwWQ0dBq75xMMg-PjHiS6eH0RDHtH5WtJMwiGw-HOolsbjHGtAEWMGuhxMBiP6qAvtjy2jMe90T7o8MUW9OWoDtrvB71eKYaTdDzaBz0e70t6ND7agpYeLzYHwag3qEBfBr2X44egR8HR8OgB6FHwYtyvQF2zBci19SbmbGg2KYMTDCbzym_Ztoa-xwRdpxnIPAmRTCNJv_cNBmtB1e3C6cE2VlRhV492w6wo9UySpbBJ7hok_ncFBlsG1l6CshBdFlXkUYajR21jP6UsNUVxkEQcakeLWgtC-IXrTsFWa4_lFd95TmC_VYWe-nVpKl8S_vfHc1ivsUbG4WhY4U6_FHj4pQDH_36A_v8d4ItVWBGhvzjI_hEQmzbTVzYLu6_h4qlMtoxIbgfbqmwsN4bpEPM_xGSit5hathupVbOWfJOJHd7c15adbh7u1M7XSpOf6BHQViLAsfX0PLSza3h1-Zubq75nxk9bzWeettWqHU_IOtyWpVAMQtesQ9csHwf1iAg9Mybjy9ywpjfvs6cA21CctaewtT8y7T0Qs15a63cDRWmecpxgjntfg_ITEnLnSZ6AztNUZXYweWRetYMOK-HL6bYYz3bLtB21Hb-CdzHoCpJLGluanUG3mIMOxQCaAtmHjn24ygh18wzaODh-tUsOxk52piy53hfevm4MomhTbV7j3WsysZuhmTYPoz9_0sGtJ737t6fzfWkr39V00vmyILHN84CG3YdZsvW_a42Fhf2V8NDlx18OhSI4_9sRlRiVcPpgwu1U-cq8l_Z73La5FZtsbfCyvmNvH8qTCY7LjEShpiploQ_cdj11p95tHvhjzQafC-38rXsy2bmZb9ugnUl0s5w6vMPKI-VNvWm3giWO6LLZarsTAZOR_d6vJdFcYpQbwEuAI7Hpr2rXCmeM4isOxY-Ctp4UfKmUmNqEwimlErvtK3BNlh9zlm22k45Dr_EvZ-0nNPMsyr0a9FsUx-zfDbWq3Vbs3ypX8Ybs7sdcxizDazimmChuyZYL5pOkm-BwpjxxuayJdMkEI-5yyR6rR4_g_5eZ2Ku4ZszkmXT1-_O1bLQbVCUpF3is-mWsIe8o7Q-O8j5uq9T4n80aHQzzE_rtt_0xdEhG4xOdhOMeXk2wiJmOS4WIRR1BkqX7LU3wZQ2TUipw8c7KpyQuYNuUt43P7XIfE2pnH8On8fl39-8vIiINfw==)) ### roaring_bitmap diff --git a/benchmarks/bloom_filter/add_bench.cu b/benchmarks/bloom_filter/add_bench.cu index a89de9806..60eddf9f2 100644 --- a/benchmarks/bloom_filter/add_bench.cu +++ b/benchmarks/bloom_filter/add_bench.cu @@ -18,8 +18,6 @@ #include #include -#include -#include using namespace cuco::benchmark; // defaults, dist_from_state, rebind_hasher_t using namespace cuco::utility; // key_generator, distribution @@ -28,20 +26,21 @@ using namespace cuco::utility; // key_generator, distribution * @brief A benchmark evaluating `cuco::bloom_filter::add_async` performance */ template void bloom_filter_add(nvbench::state& state, nvbench::type_list, nvbench::enum_type, nvbench::enum_type, nvbench::enum_type, nvbench::enum_type>) { - auto constexpr words_per_block = BlockBits / cuda::std::numeric_limits::digits; + auto constexpr word_bits = WordBytes * cuda::std::numeric_limits::digits; + auto constexpr words_per_block = BlockBits / word_bits; auto constexpr pattern_bits_per_word = PatternBits / words_per_block; // Check for a valid configuration @@ -50,8 +49,7 @@ void bloom_filter_add(nvbench::state& state, state.skip("Invalid filter block size"); } else if constexpr (HorizontalLayout * VerticalLayout != words_per_block) { state.skip("Invalid vectorization layout"); - } else if constexpr ((pattern_bits_per_word <= 0) or - (pattern_bits_per_word > cuda::std::numeric_limits::digits) or + } else if constexpr ((pattern_bits_per_word <= 0) or (pattern_bits_per_word > word_bits) or (pattern_bits_per_word * words_per_block > 64)) { state.skip("Invalid pattern bits per word"); } else { @@ -60,7 +58,7 @@ void bloom_filter_add(nvbench::state& state, auto constexpr contains_horizontal_layout = 1; using policy_type = cuco::bloom_filter_policy, - Word, + WordBytes, words_per_block, PatternBits, HorizontalLayout, @@ -104,15 +102,15 @@ void bloom_filter_add(nvbench::state& state, // Default benchmark: single layout matching default `cuco::bloom_filter_policy`. NVBENCH_BENCH_TYPES(bloom_filter_add, NVBENCH_TYPE_AXES(nvbench::type_list, - nvbench::type_list, ///< Word - nvbench::enum_type_list<256>, ///< BlockBits - nvbench::enum_type_list<8>, ///< PatternBits - nvbench::enum_type_list<8>, ///< HorizontalLayout - nvbench::enum_type_list<1> ///< VerticalLayout + nvbench::enum_type_list<4>, ///< WordBytes + nvbench::enum_type_list<256>, ///< BlockBits + nvbench::enum_type_list<8>, ///< PatternBits + nvbench::enum_type_list<8>, ///< HorizontalLayout + nvbench::enum_type_list<1> ///< VerticalLayout )) .set_name("bloom_filter_add_unique_size") .set_type_axes_names( - {"Key", "Word", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) + {"Key", "WordBytes", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) .add_int64_axis("NumInputs", {defaults::BF_N}) .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); @@ -121,14 +119,14 @@ NVBENCH_BENCH_TYPES(bloom_filter_add, // NVBENCH_BENCH_TYPES( // bloom_filter_add, // NVBENCH_TYPE_AXES(nvbench::type_list, -// nvbench::type_list, ///< Word -// nvbench::enum_type_list<64, 128, 256, 512, 1024>, ///< BlockBits -// nvbench::enum_type_list<8, 16>, ///< PatternBits -// nvbench::enum_type_list<1, 2, 4, 8, 16>, ///< -// HorizontalLayout nvbench::enum_type_list<1, 2, 4, 8, 16> ///< VerticalLayout +// nvbench::enum_type_list<8, 4>, ///< WordBytes +// nvbench::enum_type_list<64, 128, 256, 512, 1024>, ///< BlockBits +// nvbench::enum_type_list<8, 16>, ///< PatternBits +// nvbench::enum_type_list<1, 2, 4, 8, 16>, ///< HorizontalLayout +// nvbench::enum_type_list<1, 2, 4, 8, 16> ///< VerticalLayout // )) // .set_name("bloom_filter_add_full_sweep_u64") // .set_type_axes_names( -// {"Key", "Word", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) +// {"Key", "WordBytes", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) // .add_int64_axis("NumInputs", {defaults::BF_N}) // .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); diff --git a/benchmarks/bloom_filter/contains_bench.cu b/benchmarks/bloom_filter/contains_bench.cu index 59d1a85b4..2bb7138dd 100644 --- a/benchmarks/bloom_filter/contains_bench.cu +++ b/benchmarks/bloom_filter/contains_bench.cu @@ -18,9 +18,6 @@ #include #include -#include -#include - using namespace cuco::benchmark; // defaults, dist_from_state, rebind_hasher_t using namespace cuco::utility; // key_generator, distribution @@ -28,20 +25,21 @@ using namespace cuco::utility; // key_generator, distribution * @brief A benchmark evaluating `cuco::bloom_filter::contains_async` performance */ template void bloom_filter_contains(nvbench::state& state, nvbench::type_list, nvbench::enum_type, nvbench::enum_type, nvbench::enum_type, nvbench::enum_type>) { - auto constexpr words_per_block = BlockBits / cuda::std::numeric_limits::digits; + auto constexpr word_bits = WordBytes * cuda::std::numeric_limits::digits; + auto constexpr words_per_block = BlockBits / word_bits; auto constexpr pattern_bits_per_word = PatternBits / words_per_block; // Check for a valid configuration @@ -50,8 +48,7 @@ void bloom_filter_contains(nvbench::state& state, state.skip("Invalid filter block size"); } else if constexpr (HorizontalLayout * VerticalLayout > words_per_block) { state.skip("Invalid vectorization layout"); // TODO check if this is correct - } else if constexpr ((pattern_bits_per_word <= 0) or - (pattern_bits_per_word > cuda::std::numeric_limits::digits) or + } else if constexpr ((pattern_bits_per_word <= 0) or (pattern_bits_per_word > word_bits) or (pattern_bits_per_word * words_per_block > 64)) { state.skip("Invalid pattern bits per word"); } else { @@ -60,7 +57,7 @@ void bloom_filter_contains(nvbench::state& state, auto constexpr add_horizontal_layout = words_per_block; using policy_type = cuco::bloom_filter_policy, - Word, + WordBytes, words_per_block, PatternBits, add_horizontal_layout, @@ -113,8 +110,7 @@ void bloom_filter_contains(nvbench::state& state, thrust::device_vector keys(num_keys); thrust::sequence(thrust::device, keys.begin(), keys.end(), 0); - state.add_global_memory_reads(num_keys * - ((words_per_block * sizeof(Word)) + sizeof(Key))); + state.add_global_memory_reads(num_keys * ((words_per_block * WordBytes) + sizeof(Key))); state.add_global_memory_writes(num_keys * sizeof(bool)); state.exec([&](nvbench::launch& launch) { @@ -126,15 +122,15 @@ void bloom_filter_contains(nvbench::state& state, // Default benchmark: single layout matching default `cuco::bloom_filter_policy`. NVBENCH_BENCH_TYPES(bloom_filter_contains, NVBENCH_TYPE_AXES(nvbench::type_list, - nvbench::type_list, ///< Word - nvbench::enum_type_list<256>, ///< BlockBits - nvbench::enum_type_list<8>, ///< PatternBits - nvbench::enum_type_list<1>, ///< HorizontalLayout - nvbench::enum_type_list<8> ///< VerticalLayout + nvbench::enum_type_list<4>, ///< WordBytes + nvbench::enum_type_list<256>, ///< BlockBits + nvbench::enum_type_list<8>, ///< PatternBits + nvbench::enum_type_list<1>, ///< HorizontalLayout + nvbench::enum_type_list<8> ///< VerticalLayout )) .set_name("bloom_filter_contains_unique_size") .set_type_axes_names( - {"Key", "Word", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) + {"Key", "WordBytes", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) .add_int64_axis("NumInputs", {defaults::BF_N}) .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); @@ -143,14 +139,14 @@ NVBENCH_BENCH_TYPES(bloom_filter_contains, // NVBENCH_BENCH_TYPES( // bloom_filter_contains, // NVBENCH_TYPE_AXES(nvbench::type_list, -// nvbench::type_list, ///< Word -// nvbench::enum_type_list<64, 128, 256, 512, 1024>, ///< BlockBits -// nvbench::enum_type_list<8, 16>, ///< PatternBits -// nvbench::enum_type_list<1, 2, 4, 8, 16>, ///< -// HorizontalLayout nvbench::enum_type_list<1, 2, 4, 8, 16> ///< VerticalLayout +// nvbench::enum_type_list<8, 4>, ///< WordBytes +// nvbench::enum_type_list<64, 128, 256, 512, 1024>, ///< BlockBits +// nvbench::enum_type_list<8, 16>, ///< PatternBits +// nvbench::enum_type_list<1, 2, 4, 8, 16>, ///< HorizontalLayout +// nvbench::enum_type_list<1, 2, 4, 8, 16> ///< VerticalLayout // )) // .set_name("bloom_filter_contains_full_sweep_u64") // .set_type_axes_names( -// {"Key", "Word", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) +// {"Key", "WordBytes", "BlockBits", "PatternBits", "HorizontalLayout", "VerticalLayout"}) // .add_int64_axis("NumInputs", {defaults::BF_N}) // .add_int64_axis("FilterSizeMB", defaults::BF_SIZE_MB_RANGE_CACHE); diff --git a/examples/bloom_filter/persisting_l2_example.cu b/examples/bloom_filter/persisting_l2_example.cu index b3ea7a4a6..e407c8422 100644 --- a/examples/bloom_filter/persisting_l2_example.cu +++ b/examples/bloom_filter/persisting_l2_example.cu @@ -48,7 +48,7 @@ int main(void) // default policy, except the final `PersistingL2Access` parameter is `true`. using policy_type = cuco::bloom_filter_policy, - std::uint32_t, + 4, 8, 8, 8, diff --git a/include/cuco/bloom_filter_policy.cuh b/include/cuco/bloom_filter_policy.cuh index 895a31faa..3803be1df 100644 --- a/include/cuco/bloom_filter_policy.cuh +++ b/include/cuco/bloom_filter_policy.cuh @@ -23,8 +23,8 @@ namespace cuco { * * @tparam Key Key type to hash. * @tparam Hash 64-bit hash functor type. Defaults to `cuco::xxhash_64`. - * @tparam Word Underlying word type of a filter block. Defaults to `std::uint32_t`. - * @tparam WordsPerBlock Words per filter block. Defaults to the number of `Word`s that fit in one + * @tparam WordBytes Size in bytes of the underlying word type. Must be `4` or `8`. Defaults to `4`. + * @tparam WordsPerBlock Words per filter block. Defaults to the number of words that fit in one * 32-byte sector. * @tparam PatternBits Fingerprint bits per key (paper's k). Defaults to `WordsPerBlock`. * @tparam AddHorizontalLayout CG size for add (paper's Theta). Defaults to `WordsPerBlock` for @@ -48,8 +48,8 @@ namespace cuco { */ template , - class Word = std::uint32_t, - std::uint32_t WordsPerBlock = 32 / sizeof(Word), + std::uint32_t WordBytes = 4, + std::uint32_t WordsPerBlock = 32 / WordBytes, std::uint32_t PatternBits = WordsPerBlock, std::uint32_t AddHorizontalLayout = WordsPerBlock, std::uint32_t AddVerticalLayout = 1, @@ -59,7 +59,7 @@ template using bloom_filter_policy = detail::bloom_filter_policy -using parametric_filter_policy = detail::bloom_filter_policy; +using parametric_filter_policy = + detail::bloom_filter_policy(sizeof(Word)), + WordsPerBlock, + PatternBits, + AddHorizontalLayout, + AddVerticalLayout, + ContainsHorizontalLayout, + ContainsVerticalLayout, + ConditionalAdd, + EarlyExitContains, + PersistingL2Access>; } // namespace cuco diff --git a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh index b2efb3c9e..65ca75f72 100644 --- a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh +++ b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh @@ -45,13 +45,10 @@ class bloom_filter_impl { using size_type = typename extent_type::value_type; using policy_type = Policy; using word_type = typename policy_type::word_type; - static_assert(sizeof(word_type) == 4 || sizeof(word_type) == 8, - "word_type must be 4 or 8 bytes wide for atomicOr"); - // atomicOr overloads resolve on canonical 32- and 64-bit unsigned integer types. - // Normalize by size so any policy-provided word_type (uint32_t, uint64_t, unsigned long, ...) - // resolves to a matching overload via the reinterpret_cast in atomic_or(). - using atomic_word_type = - cuda::std::conditional_t; + static_assert( + cuda::std::is_same_v || + cuda::std::is_same_v, + "Policy::word_type must be unsigned int or unsigned long long int for native atomicOr"); static constexpr auto thread_scope = Scope; static constexpr auto words_per_block = policy_type::words_per_block; @@ -70,14 +67,6 @@ class bloom_filter_impl { static_assert(cuda::std::has_single_bit(words_per_block) and words_per_block <= 32, "Number of words per block must be a power-of-two and less than or equal to 32"); - static_assert( - cuda::std::is_constructible_v, word_type&> && - cuda::std::is_invocable_r_v::fetch_or), - cuda::atomic_ref*, - word_type, - cuda::std::memory_order>, - "Invalid word type"); __host__ __device__ static constexpr size_t alignment() noexcept { @@ -558,18 +547,16 @@ class bloom_filter_impl { template __device__ constexpr void atomic_or(word_type* word_ptr, word_type pattern) const { - // Native atomicOr: cuda::atomic_ref::fetch_or produces consistently slower codegen here. auto const do_or = [&]() { - auto* const p = filter_access(reinterpret_cast(word_ptr)); - auto const v = static_cast(pattern); + auto* const p = filter_access(word_ptr); if constexpr (thread_scope == cuda::thread_scope_thread) { - *p |= v; + *p |= pattern; } else if constexpr (thread_scope == cuda::thread_scope_block) { - atomicOr_block(p, v); + atomicOr_block(p, pattern); } else if constexpr (thread_scope == cuda::thread_scope_device) { - atomicOr(p, v); + atomicOr(p, pattern); } else if constexpr (thread_scope == cuda::thread_scope_system) { - atomicOr_system(p, v); + atomicOr_system(p, pattern); } else { static_assert(cuco::dependent_false, "unsupported cuda::thread_scope for native atomic_or"); diff --git a/include/cuco/detail/bloom_filter/bloom_filter_policy.cuh b/include/cuco/detail/bloom_filter/bloom_filter_policy.cuh index 64288d121..0fb18a95e 100644 --- a/include/cuco/detail/bloom_filter/bloom_filter_policy.cuh +++ b/include/cuco/detail/bloom_filter/bloom_filter_policy.cuh @@ -39,7 +39,7 @@ namespace cuco::detail { * space. * * @tparam Hash 64-bit hash functor whose call operator returns `uint64_t`. - * @tparam Word Underlying word type of a filter block. Must be an atomically updatable integral. + * @tparam WordBytes Size in bytes of the underlying word type. Must be `4` or `8`. * @tparam WordsPerBlock Words per filter block. Must be a power of two and <= 32. * @tparam PatternBits Number of fingerprint bits (k in the paper). * @tparam AddHorizontalLayout CG size used for `add` (paper's Theta). Must be a power of two and @@ -66,7 +66,7 @@ namespace cuco::detail { * persisting, thrashing the cache and affecting unrelated kernels. */ template class bloom_filter_policy { + static_assert(WordBytes == 4 || WordBytes == 8, "WordBytes must be 4 or 8 for native atomicOr"); + public: - using hasher = Hash; ///< 64-bit hash functor type - using word_type = Word; ///< Underlying filter-block word type + using hasher = Hash; ///< 64-bit hash functor type + using word_type = cuda::std::conditional_t; ///< Filter-block word type using hash_result_type = uint64_t; ///< Hash function output type private: @@ -98,6 +102,7 @@ class bloom_filter_policy { static constexpr uint32_t word_bits = cuda::std::numeric_limits::digits; public: + static constexpr uint32_t word_bytes = WordBytes; ///< Size in bytes of each filter word static constexpr uint32_t words_per_block = WordsPerBlock; ///< Number of words per filter block static constexpr uint32_t pattern_bits = PatternBits; ///< Fingerprint bits per key diff --git a/tests/bloom_filter/arrow_compat_test.cu b/tests/bloom_filter/arrow_compat_test.cu index e898c7b44..abf36ea26 100644 --- a/tests/bloom_filter/arrow_compat_test.cu +++ b/tests/bloom_filter/arrow_compat_test.cu @@ -144,8 +144,7 @@ TEMPLATE_TEST_CASE_SIG("bloom_filter arrow-compatible policy bitset validation", // Apache Arrow Block-Split Bloom Filter parameters: 256-bit blocks (8 x uint32_t), 8 fingerprint // bits per key, fully horizontal add (Theta=8) and fully vertical contains (Phi=8). - using policy_type = - cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 1, 8>; + using policy_type = cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 1, 8>; cuco::bloom_filter, cuda::thread_scope_device, policy_type> filter{ sub_filters}; diff --git a/tests/bloom_filter/bulk_ref_equivalence_test.cu b/tests/bloom_filter/bulk_ref_equivalence_test.cu index 7329bc0b9..0748ad563 100644 --- a/tests/bloom_filter/bulk_ref_equivalence_test.cu +++ b/tests/bloom_filter/bulk_ref_equivalence_test.cu @@ -90,13 +90,10 @@ TEMPLATE_TEST_CASE_SIG( "", ((class Key, class Policy), Key, Policy), (int32_t, cuco::bloom_filter_policy), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 1, 1, 1, 1, 1, 1>), - (uint64_t, - cuco::bloom_filter_policy, uint32_t, 8, 12, 8, 1, 4, 2>), - (float, cuco::bloom_filter_policy, uint64_t, 4, 4, 2, 2, 1, 2>), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 8, 8, 2, 2, 1, 8>)) + (int32_t, cuco::bloom_filter_policy, 4, 1, 1, 1, 1, 1, 1>), + (uint64_t, cuco::bloom_filter_policy, 4, 8, 12, 8, 1, 4, 2>), + (float, cuco::bloom_filter_policy, 8, 4, 4, 2, 2, 1, 2>), + (int32_t, cuco::bloom_filter_policy, 4, 8, 8, 2, 2, 1, 8>)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>; @@ -172,13 +169,10 @@ TEMPLATE_TEST_CASE_SIG( "", ((class Key, class Policy), Key, Policy), (int32_t, cuco::bloom_filter_policy), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 1, 1, 1, 1, 1, 1>), - (uint64_t, - cuco::bloom_filter_policy, uint32_t, 8, 12, 8, 1, 4, 2>), - (float, cuco::bloom_filter_policy, uint64_t, 4, 4, 2, 2, 1, 2>), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 8, 8, 2, 2, 1, 8>)) + (int32_t, cuco::bloom_filter_policy, 4, 1, 1, 1, 1, 1, 1>), + (uint64_t, cuco::bloom_filter_policy, 4, 8, 12, 8, 1, 4, 2>), + (float, cuco::bloom_filter_policy, 8, 4, 4, 2, 2, 1, 2>), + (int32_t, cuco::bloom_filter_policy, 4, 8, 8, 2, 2, 1, 8>)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>; diff --git a/tests/bloom_filter/layout_equivalence_test.cu b/tests/bloom_filter/layout_equivalence_test.cu index ce1e91b53..c76b981bd 100644 --- a/tests/bloom_filter/layout_equivalence_test.cu +++ b/tests/bloom_filter/layout_equivalence_test.cu @@ -7,7 +7,7 @@ // across (ContainsH, ContainsV) layout permutations, equivalence between dynamic vs static // `cuco::extent`, and invariance under the ConditionalAdd / EarlyExitContains policy knobs // (both are optimizations that must not change results) -- all for fixed -// (Hash, Word, WordsPerBlock, PatternBits, keys). +// (Hash, WordBytes, WordsPerBlock, PatternBits, keys). #include @@ -24,16 +24,37 @@ #include #include +#include + +TEST_CASE("bloom_filter: word byte selection", "") +{ + using hash_type = cuco::xxhash_64; + using default_policy = cuco::bloom_filter_policy; + using wide_policy = cuco::bloom_filter_policy; + using legacy_default_policy = + cuco::parametric_filter_policy; + using legacy_wide_policy = + cuco::parametric_filter_policy; + + STATIC_REQUIRE((std::is_same_v)); + STATIC_REQUIRE((std::is_same_v)); + STATIC_REQUIRE((std::is_same_v)); + STATIC_REQUIRE((std::is_same_v)); + STATIC_REQUIRE(default_policy::word_bytes == 4); + STATIC_REQUIRE(wide_policy::word_bytes == 8); + STATIC_REQUIRE(default_policy::words_per_block == 8); + STATIC_REQUIRE(wide_policy::words_per_block == 4); +} TEMPLATE_TEST_CASE_SIG( "bloom_filter: bitset is invariant under (AddH, AddV) layout permutations", "", ((class AltPolicy), AltPolicy), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 1, 8, 1, 8>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 2, 4, 1, 8>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 4, 2, 1, 8>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 2, 2, 1, 8>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 4, 1, 1, 8>)) + (cuco::bloom_filter_policy, 4, 8, 8, 1, 8, 1, 8>), + (cuco::bloom_filter_policy, 4, 8, 8, 2, 4, 1, 8>), + (cuco::bloom_filter_policy, 4, 8, 8, 4, 2, 1, 8>), + (cuco::bloom_filter_policy, 4, 8, 8, 2, 2, 1, 8>), + (cuco::bloom_filter_policy, 4, 8, 8, 4, 1, 1, 8>)) { using Key = int32_t; using default_policy = cuco::bloom_filter_policy; @@ -64,11 +85,11 @@ TEMPLATE_TEST_CASE_SIG( "bloom_filter: contains results are invariant under (ContainsH, ContainsV) permutations", "", ((class AltPolicy), AltPolicy), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 8, 1>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 2, 4>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 4, 2>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 2, 2>), - (cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 1, 4>)) + (cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 8, 1>), + (cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 2, 4>), + (cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 4, 2>), + (cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 2, 2>), + (cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 1, 4>)) { using Key = int32_t; using default_policy = cuco::bloom_filter_policy; @@ -143,7 +164,7 @@ TEST_CASE("bloom_filter: bitset is invariant under ConditionalAdd", "") Key, cuco::extent, cuda::thread_scope_device, - cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 1, 8, true>>; + cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 1, 8, true>>; constexpr int32_t num_blocks = 1'000; constexpr int32_t num_keys = 400; @@ -170,16 +191,16 @@ TEST_CASE("bloom_filter: contains results are invariant under EarlyExitContains" { using Key = int32_t; // ContainsHorizontalLayout > 1 so the compare_patterns early-exit branch is actually used. - using filter_off_t = cuco::bloom_filter< - Key, - cuco::extent, - cuda::thread_scope_device, - cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 8, 1>>; + using filter_off_t = + cuco::bloom_filter, + cuda::thread_scope_device, + cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 8, 1>>; using filter_on_t = cuco::bloom_filter< Key, cuco::extent, cuda::thread_scope_device, - cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 8, 1, false, true>>; + cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 8, 1, false, true>>; constexpr int32_t num_blocks = 1'000; constexpr int32_t num_keys = 400; diff --git a/tests/bloom_filter/merge_intersect_test.cu b/tests/bloom_filter/merge_intersect_test.cu index 371241abb..f4fb03c34 100644 --- a/tests/bloom_filter/merge_intersect_test.cu +++ b/tests/bloom_filter/merge_intersect_test.cu @@ -152,11 +152,9 @@ TEMPLATE_TEST_CASE_SIG( "", ((class Key, class Policy), Key, Policy), (int32_t, cuco::bloom_filter_policy), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 1, 1, 1, 1, 1, 1>), - (int64_t, - cuco::bloom_filter_policy, uint64_t, 1, 1, 1, 1, 1, 1>), - (int64_t, cuco::bloom_filter_policy, uint64_t, 8, 8>)) + (int32_t, cuco::bloom_filter_policy, 4, 1, 1, 1, 1, 1, 1>), + (int64_t, cuco::bloom_filter_policy, 8, 1, 1, 1, 1, 1, 1>), + (int64_t, cuco::bloom_filter_policy, 8, 8, 8>)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>; diff --git a/tests/bloom_filter/persisting_l2_access_test.cu b/tests/bloom_filter/persisting_l2_access_test.cu index 326621afc..cdcb8d020 100644 --- a/tests/bloom_filter/persisting_l2_access_test.cu +++ b/tests/bloom_filter/persisting_l2_access_test.cu @@ -35,7 +35,7 @@ using shared_ref_type = cuco::bloom_filter_ref, - uint32_t, + 4, 8, 8, 8, @@ -82,7 +82,7 @@ TEST_CASE("bloom_filter: PersistingL2Access preserves global-memory results", "" cuda::thread_scope_device, cuco::bloom_filter_policy, - uint32_t, + 4, 8, 8, 8, @@ -97,7 +97,7 @@ TEST_CASE("bloom_filter: PersistingL2Access preserves global-memory results", "" STATIC_REQUIRE_FALSE((cuco::bloom_filter_policy::persisting_l2_access)); STATIC_REQUIRE((cuco::bloom_filter_policy, - uint32_t, + 4, 8, 8, 8, diff --git a/tests/bloom_filter/unique_sequence_test.cu b/tests/bloom_filter/unique_sequence_test.cu index 31c1cb051..92b546f2f 100644 --- a/tests/bloom_filter/unique_sequence_test.cu +++ b/tests/bloom_filter/unique_sequence_test.cu @@ -79,13 +79,10 @@ TEMPLATE_TEST_CASE_SIG( "", ((class Key, class Policy), Key, Policy), (int32_t, cuco::bloom_filter_policy), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 1, 1, 1, 1, 1, 1>), - (uint64_t, - cuco::bloom_filter_policy, uint32_t, 8, 12, 8, 1, 4, 2>), - (float, cuco::bloom_filter_policy, uint64_t, 4, 4, 2, 2, 1, 2>), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 8, 8, 2, 2, 1, 8>)) + (int32_t, cuco::bloom_filter_policy, 4, 1, 1, 1, 1, 1, 1>), + (uint64_t, cuco::bloom_filter_policy, 4, 8, 12, 8, 1, 4, 2>), + (float, cuco::bloom_filter_policy, 8, 4, 4, 2, 2, 1, 2>), + (int32_t, cuco::bloom_filter_policy, 4, 8, 8, 2, 2, 1, 8>)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>; diff --git a/tests/bloom_filter/variable_cg_test.cu b/tests/bloom_filter/variable_cg_test.cu index 12a5dc27e..c413f3e8d 100644 --- a/tests/bloom_filter/variable_cg_test.cu +++ b/tests/bloom_filter/variable_cg_test.cu @@ -71,10 +71,8 @@ TEMPLATE_TEST_CASE_SIG( "", ((class Key, class Policy), Key, Policy), (int32_t, cuco::bloom_filter_policy), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 1, 1, 1, 1, 1, 1>), - (int32_t, - cuco::bloom_filter_policy, uint32_t, 8, 8, 4, 2, 4, 2>)) + (int32_t, cuco::bloom_filter_policy, 4, 1, 1, 1, 1, 1, 1>), + (int32_t, cuco::bloom_filter_policy, 4, 8, 8, 4, 2, 4, 2>)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>; @@ -108,12 +106,8 @@ TEMPLATE_TEST_CASE_SIG( "bloom_filter device ref CG contains is reduced across the group", "", ((int32_t CGSize, class Key, class Policy), CGSize, Key, Policy), - (4, - int32_t, - cuco::bloom_filter_policy, uint32_t, 8, 8, 4, 2, 4, 2>), - (8, - int32_t, - cuco::bloom_filter_policy, uint32_t, 8, 8, 8, 1, 8, 1>)) + (4, int32_t, cuco::bloom_filter_policy, 4, 8, 8, 4, 2, 4, 2>), + (8, int32_t, cuco::bloom_filter_policy, 4, 8, 8, 8, 1, 8, 1>)) { using filter_type = cuco::bloom_filter, cuda::thread_scope_device, Policy>;