diff --git a/README.md b/README.md index 2808f2838..1bb305a5c 100644 --- a/README.md +++ b/README.md @@ -267,4 +267,5 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection `cuco::experimental::roaring_bitmap` implements a Roaring bitmap following the [Roaring bitmap format specification](https://github.com/RoaringBitmap/RoaringFormatSpec). #### Examples: -- [Host-bulk APIs](https://github.com/NVIDIA/cuCollections/blob/dev/examples/roaring_bitmap/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJy9WAtPGzkQ_itzi1RtIMmGlEcbHteUlCq6HlSUe0hQrZxdJ7GyWW9tL5BD_Pcb2_uEpdDHXZDIrj3-5uFvxuPcOpJKyXgsncHFrcNCZ7DZdiISz1Iyo87ACdKQOG1H8lQE-t1bv4xhHT59HP3dOWYRPeLJSrDZXJ3TGzWA4hXcoAX9Xn-7g_922nDy53g0HsLR6dnH07Ph-fj0BF7A8Ph4_GE8PH_3qQvDKAKzUoKgkoorGnZLVR9YQGNJO-OQxopNGRUDGCYkmNNOv9vTct5lfBmvsTiI0pDCfpAG3BOcCBbP_AlTS5J0g3R--EAmVSxiauUpQZiS3XmSHN5HCoknVegF-C-k08NHJ1msmifVKqG-VVATUHORSuWF9Ard869ooLjozptEIj5jAYmaJ9OYXVEhSVSFqMpNcaPkSiq6rC2fSiUoqY8x3jCIQxjG2pDVZPR465YTb7QamCOAP0mjhU9vyDKJKIbdTk8Eo1MY0SWyDYOhqIRUIsuAT0HNKdR3Cy4djXLpQMT5Ik0MMAw_jqWhhYEcx7iQScg0wTVFYRKCuubwst9BILBgEkgcAo8p7GxVhsFNuFBkgkunXCyJasFU8KW2xuBfnFmT3hrpYyPyKaHBZ3euVCIHnjdjap5OugFfejXZ_K1c00JaJ1wyjNpKW2MUIH-DBTDrv3Y3dxb9RHmViliauYALgRHXqZFGSFM4IUsardraZQykMkJTHkX8GrWC2fCBUdGBC-vsNZrKUyXSWHYnLDaTro1S63sc8iYRn3jbm7u7JHztaSNCoojXqKz10JT_z44HRuSbnjFtZ6u0w9LjJ9qxs-U1qmsVJH4Tc0XhvMpjQb-kDLfabv2SLDBHEoVVGjqjoz-OTv3R6V8nH06HI__sdHg2PnnvYwk9Hw3PhwdYVxWHCQVJVZEoWBsx95MIcw6LBhajGNkDv9HVOT5jDk84jywXXbR6MLD5jqTDRH2R5YqvOeUnRM3R9FtEBZKiphmNqc5lf0FXEg7g4rPbgs4h2NI0GNRq236uEgwAaOYbJfQmEXhiYL3UytECJn2JNvpX-ZI2VGZTrLQv-746bOVAAJ4HR1i40MMvKcUUM_ZgVtvUKshwi6y8yxlR5kkO8dPod_ZuOPr9XXcZrumhjh7L1RgXsoA0OWVM38ultQtugxgsMNq9Pfzah82e_ujnjQPzUokLGLhuksq5PyG4w4tWgX1XU4LABrRE28cT3D5vbCy-hvkSefwk7m4F99XzcJswbVl8hGBN8XQN6ITOWOy22lYFjUO3lYPfAY0k_REy7mx9Axkb64Gh4n_BRCxAJRcfqLaF7ylaGveepqUWs7S8MdubRpHdbXx_nb9_x44_pWtY17XZ-wFd90R6N_0MrZRtFNmuiTxqMbNJy-4ZynTi9jXEV8zt3bzKFsAGsB_MDLOlz82M25IjRLHAJxIbdOXq_lmrwfNEt-X-lKBwUePb2MFlz7BEg_ShVBQvjE9uxaVTupJ5cHuX69df-kV_Y1KcoiqTKfrd5mbWxJpR9_5R1c6EuBwMkOpErKwuTHX3Fy3VxdzmCIou516aFQEVAvb30YVjgmKhPlS1HA7o4ft69JhZh5GIMtszX0xQzNBd4cZ7qoy9INk_tDhKDZoeQY4YsLJtt89m8oGPrTI8eIfiga46pldeYqeNtad-BhbRy1lSXh5M556wOKZhA2kmK4Xn9iSdTqlwC2sqys8odt7GK9xYnmnXcybOuEehKyhOUZFgaPyASLUfzIlYP3RzWwS59hNuZMy8a9V1dR3DHUKK1hRn0EHEJXUrlmSlN7tPFI5nnX0RAUtfrPZUsCXyl0RoQe0OUvYr9XHX7vBzzd6rbLztlqrHgkswoSV2y1GoUwSPIEUY7gGGsbJvrYImWaNVa7wavCdg91AzV-I3tXcse3-oEqBeJnQreFgaYSuE4Z2OvuZyRdVbvLF0shtLxSMyw7XZncRar-XvX8atCvloDWqXVuTTVS_zm5OOnlZkNYtHA6jXmT4XF_ilzEERBz3Op-4DpVU7MsOwPlS6A_OThFrd3llOZnlfU4MTJvt1K7QkGtX00GtsGuKV2LT0eSevu3h_ND4rKly1GQdNKT9kAu1uXLVXuCnTIKBSQv1zYNv85j6pAN_AMtd8kcuLdY7-4rsAvxXNHlZ1tOZLnJNTxFZxtNxW8UzDIKvftqJhmEiUzIkZyUy4V8grG5pL_Ao9GMAmTq7pw7FUVhwZzXe5xv3KLmvIAsNGrBwMy4COUf3ed-mUbUH2uXSecRls3V_44KDKvLP-xCGbGqo6bQc7zgQrpSh_GXTiqyDY7G-nmzhtDcNJp4OAB8HGxuYudIgI5gdy6e_2oNPB0qrwn9LtQdiJyHJifkuM2KSCGQRBhIP6DEI8HNBcWzh37Xweq3RtHuuVc_fZ_P0LZpjzEQ==)) \ No newline at end of file +- [Host-bulk APIs using unordered indices](https://github.com/NVIDIA/cuCollections/blob/dev/examples/roaring_bitmap/host_bulk_from_indices_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyNVA1rGkEQ_SvDFRqTnnoa2sAmppWmASEkIQ2lUMuxtzfq0r29636IqeS_d-7LnGlKq7h6b2bfzHs77jawaK3MtQ3Yt20g04CNwkBxvfR8iQELhE95EAY290aUz8OjuYYj-Hx78bV_KRV-zIsHI5crd48bx2D3CD1xCONo_C6E6y-zi9kUPt7c3d7cTe9nN9fwGqaXl7Or2fT-0-cBTJWCapMFgxbNGtPBU5UrKVBb7M9S1E4uJBoG04KLFfbHg6jMG871XL-SWiifIpwJL_KhybmRehkn0mW8GAi_On-elfKhdelQ0CK1O-8G3cp464Yprql4vEbhcjNYvZSyyq3rJnRTJMUM8qyCh0e1cx8W5BpU2xKvfsQLk2ex1CkVsjFueFYopHbr3MRIXMAFZnRAznCHFhIvFfW7BA77GqFkAq9zk6LBFBrOQeMQSYSMS907nOstYeBtyUJZuIndQ4EwgdITxsgPxjzlH49jd1r2DlCrZWzPkbOnzedtuW20iaJoRJ9xCCchjELYIVGFHHeQ8WPDz73LodExKQGA8hgZw02BRmZ09Fwxtq-4W5-xrpG9VnyCy1JyuHMDddo7PPxvVT89VSdVJGNctX6yr2f3e_Rc1d_YkzxX5-Wce-Vsr-EfWPkL674A2pHNSbPUTzk7LS1QaQlbrjb-XFxnQpvipafCYbp1xmMIC64sfdUP3bUT2Ff0JyUXznNFI9Q0UwuhEIhydMF6Qe5b2H9Ndvsmu6aa7qshFLl3cHYG88BrSarbU2SElHjjVG1eCRzM5_rg9IX9dKtIrigvheTBvcgQV4F_8dQ62u1VtJTJVbHiFdII7ZCUNAadN3oXfQ8RMBhR8HGu6XoVeVbQtWCebuFAr4UYjd_6EYXzwtVXdNCnihPx5s3oBPrciNXEZvFJBP0-2edocTQTmPYVz5Lq3lYy6XAKIRSBaypEfATQ0eofwWPYxukvtxen0Q0ev1fv35rBEmA=)) +- [Host-bulk APIs using serialized bitmap](https://github.com/NVIDIA/cuCollections/blob/dev/examples/roaring_bitmap/host_bulk_from_serialized_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJy9WAtP4zgQ_itzQUIptE1heeyWx22XLqvq9mDFcg8JVpGbuK3VNM7aDtBD_Pcb23mWsOzrrkg0scfz_GY803tHUikZj6XTv7p3WOj0t9pOROJpSqbU6TtBGhKn7UieikC_exvXMWzAxw_DvzunLKInPFkKNp2pS3qn-lC8ghu0YLu3vdvBf3ttOPtzNBwN4OT84sP5xeBydH4G6zA4PR29Hw0u337swiCKwJyUIKik4oaG3VLUexbQWNLOKKSxYhNGRR8GCQlmtLPd7Wk67zq-jtdYHERpSOEwSAPuCU4Ei6f-mKkFSbpBOjt-RJMqFjG19JQgTMnuLEmOVzmFxJMq9AL8F9LJ8ZObLFbNm2qZUN8KqBGomUil8kJ6g-b5NzRQXHRnTSQRn7KARM2bacxuqJAkqrKo0k0wUHIpFV3Ujk-kEpTU1xhvWMQldGNtyUoycrwNi4nXWgzMkIE_TqO5PxF84WMkGYnYPzT06R1ZJBHFMFjysWB0AkO6QPShcxSVkEpEHfAJqBmFevTg2tFcrx2IOJ-niREEgw8jaWBiWI5iPMgkZJLgliIxCUHdcnix3UFGYJlJIHEIPKawt1NZBjfhQpExHp1wsSCqBdoIrY3hf3VhVXpjqE8NyceEBp_cmVKJ7HvelKlZOu4GfOHVaPO38kwLYZ5wydCLS62NEYB4DubArP3a3NxYtBPpVSpiafYCLgRGQKdKGiFs4YwsaLRsa5PRkcoQTXgU8VuUCgYAfSOiA1fW2FtUladKpLHsjllsNl3rpdb3GOSNIz72drf290n4ytNKhEQRr1FY67Eq_58ej5TIg54hbW-n1MPC4yfqsbfjNYprFSB-HXNF4bKKY0E_pwxDbUO_IHPMkURh1YbO8OSPk3N_eP7X2fvzwdC_OB9cjM7e-VhSL4eDy8ER1lnFYUxBUlUkCtZKrAVJhDmHRQSLU4zogd_o8hKfMafHnEcWiy5q3e_b_EfQYaKuZ7nia0z5CVEzVP0euQJJUdKUxlTnsj-nSwlHcPXJbUHnGGyp6vdrte4wFwmGAWjkGyH0LhF4g2D91MJRAyZ9iTr6N_mRNlR2U6y8L7Z9ddzKGQF4HpxgIUMLP6cUU8zog1ltU6sAwz2i8iFHRJknOYufBr-Lt4Ph72-7i3BNL3X0Wi7GmJA5pMkoo_pBTq1NcBvIYI7e7h3g1yFs9fRHP28emZeKX8Cw6yapnPljghGetwreDzUhyNgwLbkd4o1unzc351_i-QJx_Czf_Qrfl1_Ht4mnLYtPAKzJn65hOqZTFrutthVB49Bt5cwfgEaS_ggY93a-AYyN9cBA8b9AIhagEouPRNvC9xwsjXnPw1KTWVjemfCmUWSjje-v8vfviPhzsgZ1WVu9H5C1QtK72864lbSNJLs1kic1ZjZp2YqiTCfutmbxBXV7dy-zA7AJ7Aczw4T0azPjvsQIUSzwicQ2T7m6n9Zi8D7Rbbo_IUhc1Pg2dnDZMyxQIX0pFcUL_ZNrce2UpmQW3D_k8vWXftHfmBTnKMpkin63uZk1tWbVXb2q2hkRl_0-Qp2IpZWFqe7-oqm6mNscmaLJuZXmRECFgMNDNOGUIFmoL1VNhwt6eVWOXjPn0BNRpntmi3GKWXoozHhHldEXJDbKxVVquOkVxIhhVrbx9tlsPrKxVboHZyoe6KpjeuUFdtpYe-p3YOG9HCXlMGE6-YTFMfbuj0EzXiq8t8fpZEKFW2hTEX5BsfM2VmFgeSZd7xk_Y4xCV1DcoiJB1_gBkeowmBGxcezmughy6yfc0Jh914rr6jqGEUKI1gRnrIOIS-pWNMlKbzZPFIZnnX3hAeP0laEDG5j1TzkSDIG5EjZAq4C7X6npwU9vbzI02ZRDfjhmLTDnSIS61Ewo8g8xUx_IXK3cU5n97fzvNbt6olaD8C5rDKs3oEuwdkkcDKJQVwP0jSIM4YaIqUC0VQQn6ylrPWZDoAlYuOoklfhN7ThpR6Uq1usVUXe9x6USthiaFNNA02lbEfUGh7NONpxVLCJTPJuNX1Z7Tb_6O4QVIZ8st-1Si3y7amU-JGrvaUFWsnjSgfqcaenxgF_SlOjV63ziPhJa1SNTDEthBbbm1xi1vH-wMMpAUxODG6bQ6a5vQTRXMy6ssUmI07-ZXvKhRQ8s_nB0URTz6txhMs4PmUC9G08dFGbKNAiolFD_HNmJprklLJhvYkVvnlnzeynnvv5dDL-Vm72X69ya51Unh4i9sFBze2FlEvrZVWWLN7qJRMmMmJVMhZU7qxLQnOJX6EEftnBzTVeLUlhxOzaPrY3xyuZSRIFBI1ZMhmVA-6g-4l47ZQeUfa6dr5h7W6sHH93JmXXWnjhkEwNVp-1gc53gpSDKH0Wd-CYItrZ30y3ctorhptNBhkfB5ubWPnSICGZHcuHv96DTwYqp8J_SnVDYichibH5Gjdi4wjMIgggX9XWL_HBBY23uPLTzfSy-tX2sV87DJ_P3L2TISXk=)) \ No newline at end of file diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt index 597f192de..89b997fe7 100644 --- a/benchmarks/CMakeLists.txt +++ b/benchmarks/CMakeLists.txt @@ -94,4 +94,5 @@ ConfigureBench(BLOOM_FILTER_BENCH ################################################################################################### # - roaring_bitmap benchmarks --------------------------------------------------------------------- ConfigureBench(ROARING_BITMAP_BENCH - roaring_bitmap/contains_bench.cu) \ No newline at end of file + roaring_bitmap/build_bench.cu + roaring_bitmap/contains_bench.cu) diff --git a/benchmarks/roaring_bitmap/build_bench.cu b/benchmarks/roaring_bitmap/build_bench.cu new file mode 100644 index 000000000..b0660be1e --- /dev/null +++ b/benchmarks/roaring_bitmap/build_bench.cu @@ -0,0 +1,145 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include +#include + +#include + +#include +#include +#include +#include +#include +#include + +using namespace cuco::benchmark; +using namespace cuco::utility; + +enum class build_mode { indices, sorted_indices, sorted_unique_indices }; + +template +void roaring_bitmap_build(nvbench::state& state, nvbench::type_list) +{ + using index_type = cuda::std::uint32_t; + using bitmap_type = cuco::experimental::roaring_bitmap; + + auto const num_inputs = state.get_int64("NumInputs"); + thrust::device_vector indices(num_inputs); + + [[maybe_unused]] key_generator generator{}; + if constexpr (Mode == build_mode::sorted_unique_indices) { + thrust::sequence(indices.begin(), indices.end()); + } else { + generator.generate(dist_from_state(state), indices.begin(), indices.end()); + if constexpr (Mode == build_mode::sorted_indices) { + thrust::sort(indices.begin(), indices.end()); + } + } + + state.add_element_count(num_inputs); + state.add_global_memory_reads(num_inputs, "InputSize"); + + state.exec(nvbench::exec_tag::sync | nvbench::exec_tag::timer, + [&](nvbench::launch& launch, auto& timer) { + timer.start(); + if constexpr (Mode == build_mode::indices) { + [[maybe_unused]] auto bitmap = bitmap_type::from_indices( + indices.begin(), indices.end(), {}, cuda::stream_ref{launch.get_stream()}); + timer.stop(); + } else if constexpr (Mode == build_mode::sorted_indices) { + [[maybe_unused]] auto bitmap = bitmap_type::from_sorted_indices( + indices.begin(), indices.end(), {}, cuda::stream_ref{launch.get_stream()}); + timer.stop(); + } else { + [[maybe_unused]] auto bitmap = bitmap_type::from_sorted_unique_indices( + indices.begin(), indices.end(), {}, cuda::stream_ref{launch.get_stream()}); + timer.stop(); + } + }); +} + +template +void roaring_bitmap_from_indices(nvbench::state& state, nvbench::type_list types) +{ + roaring_bitmap_build(state, types); +} + +template +void roaring_bitmap_from_sorted_indices(nvbench::state& state, nvbench::type_list types) +{ + roaring_bitmap_build(state, types); +} + +template +void roaring_bitmap_from_sorted_unique_indices(nvbench::state& state, + nvbench::type_list types) +{ + roaring_bitmap_build(state, types); +} + +void roaring_bitmap_from_indices_array_containers(nvbench::state& state) +{ + using index_type = cuda::std::uint32_t; + using bitmap_type = cuco::experimental::roaring_bitmap; + + constexpr cuda::std::int64_t num_containers = 1 << 16; + auto const cardinality = state.get_int64("ContainerCardinality"); + auto const num_inputs = num_containers * cardinality; + thrust::device_vector indices(num_inputs); + + thrust::tabulate( + indices.begin(), indices.end(), [cardinality] __device__(cuda::std::int64_t index) { + auto const container = static_cast(index / cardinality); + auto const lower = static_cast(index % cardinality); + return (container << 16) | lower; + }); + thrust::reverse(indices.begin(), indices.end()); + + state.add_element_count(num_inputs); + state.add_global_memory_reads(num_inputs, "InputSize"); + + state.exec(nvbench::exec_tag::sync | nvbench::exec_tag::timer, + [&](nvbench::launch& launch, auto& timer) { + timer.start(); + [[maybe_unused]] auto bitmap = bitmap_type::from_indices( + indices.begin(), indices.end(), {}, cuda::stream_ref{launch.get_stream()}); + timer.stop(); + }); +} + +NVBENCH_BENCH_TYPES(roaring_bitmap_from_indices, + NVBENCH_TYPE_AXES(nvbench::type_list)) + .set_name("roaring_bitmap_from_indices_unique") + .set_type_axes_names({"Distribution"}) + .add_int64_power_of_two_axis("NumInputs", {20, 24, 28}) + .add_int64_axis("Multiplicity", {1}); + +NVBENCH_BENCH(roaring_bitmap_from_indices_array_containers) + .set_name("roaring_bitmap_from_indices_array_containers") + .add_int64_axis("ContainerCardinality", {1, 8, 64, 512, 4096}); + +NVBENCH_BENCH_TYPES(roaring_bitmap_from_indices, + NVBENCH_TYPE_AXES(nvbench::type_list)) + .set_name("roaring_bitmap_from_indices_uniform") + .set_type_axes_names({"Distribution"}) + .add_int64_power_of_two_axis("NumInputs", {20, 24, 28}) + .add_int64_axis("Multiplicity", {2, 8, 32}); + +NVBENCH_BENCH_TYPES(roaring_bitmap_from_sorted_indices, + NVBENCH_TYPE_AXES(nvbench::type_list)) + .set_name("roaring_bitmap_from_sorted_indices") + .set_type_axes_names({"Distribution"}) + .add_int64_power_of_two_axis("NumInputs", {20, 24, 28}) + .add_int64_axis("Multiplicity", {2, 8, 32}); + +NVBENCH_BENCH_TYPES(roaring_bitmap_from_sorted_unique_indices, + NVBENCH_TYPE_AXES(nvbench::type_list)) + .set_name("roaring_bitmap_from_sorted_unique_indices") + .set_type_axes_names({"Distribution"}) + .add_int64_power_of_two_axis("NumInputs", {20, 24, 28}) + .add_int64_axis("Multiplicity", {1}); diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index a709c1685..8a1457e65 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -40,4 +40,7 @@ ConfigureExample(HYPERLOGLOG_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hype 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") ConfigureExample(BLOOM_FILTER_PERSISTING_L2_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/bloom_filter/persisting_l2_example.cu") -ConfigureExample(ROARING_BITMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/roaring_bitmap/host_bulk_example.cu") +ConfigureExample(ROARING_BITMAP_HOST_BULK_FROM_INDICES_EXAMPLE + "${CMAKE_CURRENT_SOURCE_DIR}/roaring_bitmap/host_bulk_from_indices_example.cu") +ConfigureExample(ROARING_BITMAP_HOST_BULK_FROM_SERIALIZED_EXAMPLE + "${CMAKE_CURRENT_SOURCE_DIR}/roaring_bitmap/host_bulk_from_serialized_example.cu") diff --git a/examples/roaring_bitmap/host_bulk_from_indices_example.cu b/examples/roaring_bitmap/host_bulk_from_indices_example.cu new file mode 100644 index 000000000..39ea136e5 --- /dev/null +++ b/examples/roaring_bitmap/host_bulk_from_indices_example.cu @@ -0,0 +1,40 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include +#include +#include + +#include + +/** + * @file host_bulk_from_indices_example.cu + * @brief Demonstrates building a roaring_bitmap from unordered indices. + */ +int main() +{ + using index_type = cuda::std::uint32_t; + + thrust::device_vector indices{0x00010002, 7, 1, 0x00010000, 7, 3, 0x00010002}; + + auto bitmap = + cuco::experimental::roaring_bitmap::from_indices(indices.begin(), indices.end()); + + thrust::device_vector queries{1, 2, 3, 7, 0x00010000, 0x00010001, 0x00010002}; + thrust::device_vector results(queries.size()); + bitmap.contains(queries.begin(), queries.end(), results.begin()); + + thrust::host_vector expected{true, false, true, true, true, false, true}; + thrust::host_vector actual = results; + bool const success = actual == expected; + + std::cout << "unique indices: " << bitmap.size() << '\n'; + std::cout << "serialized bytes: " << bitmap.size_bytes() << '\n'; + std::cout << "success: " << std::boolalpha << success << '\n'; + + return success ? 0 : 1; +} diff --git a/examples/roaring_bitmap/host_bulk_example.cu b/examples/roaring_bitmap/host_bulk_from_serialized_example.cu similarity index 92% rename from examples/roaring_bitmap/host_bulk_example.cu rename to examples/roaring_bitmap/host_bulk_from_serialized_example.cu index d64986c61..46b7bdce3 100644 --- a/examples/roaring_bitmap/host_bulk_example.cu +++ b/examples/roaring_bitmap/host_bulk_from_serialized_example.cu @@ -20,7 +20,7 @@ #include /** - * @file host_bulk_example.cu + * @file host_bulk_from_serialized_example.cu * @brief Demonstrates usage of the roaring_bitmap "bulk" lookup host APIs. * * In this example we load two 32-bit bitmaps and one 64-bit bitmap (portable format) from the @@ -94,8 +94,14 @@ bool check(std::string const& bitmap_file_path) file.close(); // Create roaring bitmap from the file - cuco::experimental::roaring_bitmap roaring_bitmap( - thrust::raw_pointer_cast(buffer.data())); + auto roaring_bitmap = [&] { + auto const* data = thrust::raw_pointer_cast(buffer.data()); + if constexpr (cuda::std::is_same_v) { + return cuco::experimental::roaring_bitmap::from_serialized(data); + } else { + return cuco::experimental::roaring_bitmap{data}; + } + }(); // Generate query keys (all should be contained in the bitmap) auto keys = generate_keys(); diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl index 61326fdd3..ee41f099c 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl @@ -5,7 +5,12 @@ #pragma once +#include + #include +#include +#include +#include #include namespace cuco::experimental { @@ -18,6 +23,62 @@ roaring_bitmap::roaring_bitmap(cuda::std::byte const* bitmap, { } +template +roaring_bitmap roaring_bitmap::from_serialized( + cuda::std::byte const* bitmap, Allocator const& alloc, cuda::stream_ref stream) +{ + static_assert(cuda::std::is_same_v, + "roaring_bitmap::from_serialized currently supports only uint32_t"); + return roaring_bitmap{bitmap, alloc, stream}; +} + +template +roaring_bitmap::roaring_bitmap(storage_type&& storage) + : storage_{cuda::std::move(storage)} +{ +} + +template +template +roaring_bitmap roaring_bitmap::from_indices(InputIt first, + InputIt last, + Allocator const& alloc, + cuda::stream_ref stream) +{ + static_assert(cuda::std::is_same_v, + "Building a roaring_bitmap from indices currently supports only uint32_t"); + detail::roaring_bitmap_builder builder{ + first, last, detail::roaring_bitmap_builder_input_order::unsorted, alloc, stream}; + auto storage = cuda::std::move(builder).build(); + return roaring_bitmap{cuda::std::move(storage)}; +} + +template +template +roaring_bitmap roaring_bitmap::from_sorted_indices( + InputIt first, InputIt last, Allocator const& alloc, cuda::stream_ref stream) +{ + static_assert(cuda::std::is_same_v, + "Building a roaring_bitmap from indices currently supports only uint32_t"); + detail::roaring_bitmap_builder builder{ + first, last, detail::roaring_bitmap_builder_input_order::sorted, alloc, stream}; + auto storage = cuda::std::move(builder).build(); + return roaring_bitmap{cuda::std::move(storage)}; +} + +template +template +roaring_bitmap roaring_bitmap::from_sorted_unique_indices( + InputIt first, InputIt last, Allocator const& alloc, cuda::stream_ref stream) +{ + static_assert(cuda::std::is_same_v, + "Building a roaring_bitmap from indices currently supports only uint32_t"); + detail::roaring_bitmap_builder builder{ + first, last, detail::roaring_bitmap_builder_input_order::sorted_unique, alloc, stream}; + auto storage = cuda::std::move(builder).build(); + return roaring_bitmap{cuda::std::move(storage)}; +} + template template void roaring_bitmap::contains(InputIt first, @@ -74,4 +135,4 @@ typename roaring_bitmap::ref_type roaring_bitmap::re { return ref_type{storage_.ref()}; } -} // namespace cuco::experimental \ No newline at end of file +} // namespace cuco::experimental diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh new file mode 100644 index 000000000..8b489bcc0 --- /dev/null +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh @@ -0,0 +1,450 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace cuco::experimental::detail { + +enum class roaring_bitmap_builder_input_order { unsorted, sorted, sorted_unique }; + +/** + * @brief One-shot builder for a 32-bit Roaring bitmap. + * + * Construction computes the temporary workspace requirement. Calling `build` allocates temporary + * storage, executes the build, and returns the serialized bitmap storage. + * + * @tparam InputIt Random access iterator with `cuda::std::uint32_t` value type + * @tparam Allocator Allocator type used for temporary and final storage + */ +template +class roaring_bitmap_builder { + using input_type = typename cuda::std::iterator_traits::value_type; + static_assert(cuda::std::is_same_v, + "roaring_bitmap factories require an input iterator with uint32_t value_type"); + + public: + using storage_type = + roaring_bitmap_storage; ///< Generated bitmap storage type + + /** + * @brief Prepares a build for the specified input range and ordering. + * + * @param first Beginning of the input range + * @param last End of the input range + * @param input_order Ordering and uniqueness guarantees for the input range + * @param alloc Allocator used for temporary and final storage + * @param stream Stream used for all allocation and build work + */ + explicit roaring_bitmap_builder(InputIt first, + InputIt last, + roaring_bitmap_builder_input_order input_order, + Allocator const& alloc, + cuda::stream_ref stream) + : first_{first}, + num_indices_{std::max(0, cuco::detail::distance(first, last))}, + num_container_slots_{static_cast( + std::min(num_indices_, + static_cast( + roaring_bitmap_metadata::max_num_containers)))}, + input_order_{input_order}, + alloc_{alloc}, + stream_{stream}, + workspace_bytes_{compute_workspace_bytes()} + { + } + + /** + * @brief Executes the prepared build. + * + * @return Storage containing the serialized bitmap + */ + [[nodiscard]] storage_type build() && + { + if (num_indices_ == 0) { + return write_serialized_bitmap(first_, empty_build_state(), nullptr, nullptr, nullptr); + } + + switch (input_order_) { + case roaring_bitmap_builder_input_order::unsorted: return build_unsorted(); + case roaring_bitmap_builder_input_order::sorted: return build_sorted(); + case roaring_bitmap_builder_input_order::sorted_unique: return build_sorted_unique(); + } + CUCO_FAIL("Invalid roaring_bitmap input order"); + } + + private: + struct sorted_indices_result { + cuda::std::uint32_t* indices; + cuda::std::uint32_t* available_buffer; + }; + + template + [[nodiscard]] auto allocate_temporary_buffer(cuda::std::size_t size) const + { + using allocator_type = typename std::allocator_traits::template rebind_alloc; + using deleter_type = cuco::detail::custom_deleter; + + allocator_type allocator{alloc_}; + auto* const data = size == 0 ? nullptr : allocator.allocate(size, stream_); + // Destruction enqueues the deallocation on the build stream, after all previously submitted + // work that uses this buffer. + return std::unique_ptr{data, deleter_type{size, allocator, stream_}}; + } + + [[nodiscard]] cuda::std::size_t compute_workspace_bytes() const + { + if (num_indices_ == 0) { return 0; } + + // CUB only uses these pointers to instantiate the requested algorithm while temporary storage + // is null; its size-query path returns before launching work or dereferencing them. + auto* const index_buffer_a = static_cast(nullptr); + auto* const index_buffer_b = static_cast(nullptr); + auto* const container_starts = static_cast(nullptr); + auto* const payload_offsets = static_cast(nullptr); + auto* const num_selected = static_cast(nullptr); + auto* const state = static_cast(nullptr); + auto const counting_begin = cuda::counting_iterator{0}; + + // After the last CUB operation, the same allocation becomes the array/bitset container queue. + cuda::std::size_t result = num_container_slots_ * sizeof(cuda::std::uint32_t); + cuda::std::size_t required_bytes = 0; + + CUCO_CUDA_TRY(cub::DeviceScan::ExclusiveSum(nullptr, + required_bytes, + payload_offsets, + payload_offsets, + num_container_slots_, + stream_.get())); + result = std::max(result, required_bytes); + + required_bytes = 0; + // Container discovery consumes the caller's iterator only when no normalization is required. + // Otherwise it consumes one of the materialized uint32_t buffers allocated by build(). + if (input_order_ == roaring_bitmap_builder_input_order::sorted_unique) { + CUCO_CUDA_TRY(cub::DeviceSelect::If(nullptr, + required_bytes, + counting_begin, + container_starts, + num_selected, + num_indices_, + is_container_start{first_, state}, + stream_.get())); + } else { + CUCO_CUDA_TRY(cub::DeviceSelect::If(nullptr, + required_bytes, + counting_begin, + container_starts, + num_selected, + num_indices_, + is_container_start{index_buffer_a, state}, + stream_.get())); + } + result = std::max(result, required_bytes); + + if (input_order_ != roaring_bitmap_builder_input_order::sorted_unique) { + required_bytes = 0; + if (input_order_ == roaring_bitmap_builder_input_order::sorted) { + CUCO_CUDA_TRY(cub::DeviceSelect::Unique(nullptr, + required_bytes, + first_, + index_buffer_a, + num_selected, + num_indices_, + stream_.get())); + } else { + CUCO_CUDA_TRY(cub::DeviceSelect::Unique(nullptr, + required_bytes, + index_buffer_a, + index_buffer_b, + num_selected, + num_indices_, + stream_.get())); + } + result = std::max(result, required_bytes); + } + + if (input_order_ == roaring_bitmap_builder_input_order::unsorted) { + required_bytes = 0; + cub::DoubleBuffer indices{index_buffer_a, index_buffer_b}; + CUCO_CUDA_TRY(cub::DeviceRadixSort::SortKeys(nullptr, + required_bytes, + indices, + num_indices_, + 0, + sizeof(cuda::std::uint32_t) * 8, + stream_.get())); + result = std::max(result, required_bytes); + } + + return result; + } + + [[nodiscard]] storage_type build_unsorted() + { + auto indices_a = + allocate_temporary_buffer(static_cast(num_indices_)); + auto indices_b = + allocate_temporary_buffer(static_cast(num_indices_)); + auto container_starts = allocate_temporary_buffer(num_container_slots_); + auto state = allocate_temporary_buffer(1); + auto workspace = allocate_temporary_buffer(workspace_bytes_); + + auto const sorted = sort_indices(first_, indices_a.get(), indices_b.get(), workspace.get()); + // Deduplication writes into the inactive radix-sort buffer. Once it completes, the sorted input + // buffer is no longer needed and is reused for payload offsets. + deduplicate_indices(sorted.indices, sorted.available_buffer, state.get(), workspace.get()); + return serialize_sorted_unique_indices(sorted.available_buffer, + container_starts.get(), + sorted.indices, + state.get(), + workspace.get()); + } + + [[nodiscard]] storage_type build_sorted() + { + auto unique_indices = + allocate_temporary_buffer(static_cast(num_indices_)); + auto container_starts = allocate_temporary_buffer(num_container_slots_); + auto payload_offsets = allocate_temporary_buffer(num_container_slots_); + auto state = allocate_temporary_buffer(1); + auto workspace = allocate_temporary_buffer(workspace_bytes_); + + deduplicate_indices(first_, unique_indices.get(), state.get(), workspace.get()); + return serialize_sorted_unique_indices(unique_indices.get(), + container_starts.get(), + payload_offsets.get(), + state.get(), + workspace.get()); + } + + [[nodiscard]] storage_type build_sorted_unique() + { + auto container_starts = allocate_temporary_buffer(num_container_slots_); + auto payload_offsets = allocate_temporary_buffer(num_container_slots_); + auto state = allocate_temporary_buffer(1); + auto workspace = allocate_temporary_buffer(workspace_bytes_); + + CUCO_CUDA_TRY(cudaMemcpyAsync(&state->num_keys, + &num_indices_, + sizeof(num_indices_), + cudaMemcpyHostToDevice, + stream_.get())); + return serialize_sorted_unique_indices( + first_, container_starts.get(), payload_offsets.get(), state.get(), workspace.get()); + } + + template + [[nodiscard]] sorted_indices_result sort_indices(SourceIt first, + cuda::std::uint32_t* indices_a, + cuda::std::uint32_t* indices_b, + cuda::std::byte* workspace) const + { + CUCO_CUDA_TRY(cub::DeviceTransform::Transform( + first, indices_a, num_indices_, cuda::std::identity{}, stream_.get())); + + cub::DoubleBuffer indices{indices_a, indices_b}; + auto workspace_bytes = workspace_bytes_; + CUCO_CUDA_TRY(cub::DeviceRadixSort::SortKeys(workspace, + workspace_bytes, + indices, + num_indices_, + 0, + sizeof(cuda::std::uint32_t) * 8, + stream_.get())); + return {indices.Current(), indices.Alternate()}; + } + + template + void deduplicate_indices(SourceIt first, + cuda::std::uint32_t* unique_indices, + roaring_bitmap_build_state* state, + cuda::std::byte* workspace) const + { + auto workspace_bytes = workspace_bytes_; + CUCO_CUDA_TRY(cub::DeviceSelect::Unique(workspace, + workspace_bytes, + first, + unique_indices, + &state->num_keys, + num_indices_, + stream_.get())); + } + + template + [[nodiscard]] storage_type serialize_sorted_unique_indices(SourceIt first, + cuda::std::int64_t* container_starts, + cuda::std::uint32_t* payload_offsets, + roaring_bitmap_build_state* state, + cuda::std::byte* workspace) const + { + analyze_containers(first, container_starts, payload_offsets, state, workspace); + auto const host_state = read_build_state(state); + auto* const container_indexes = reinterpret_cast(workspace); + return write_serialized_bitmap( + first, host_state, container_starts, payload_offsets, container_indexes); + } + + template + void analyze_containers(SourceIt first, + cuda::std::int64_t* container_starts, + cuda::std::uint32_t* payload_offsets, + roaring_bitmap_build_state* state, + cuda::std::byte* workspace) const + { + auto const counting_begin = cuda::counting_iterator{0}; + auto workspace_bytes = workspace_bytes_; + CUCO_CUDA_TRY(cub::DeviceSelect::If(workspace, + workspace_bytes, + counting_begin, + container_starts, + &state->num_containers, + num_indices_, + is_container_start{first, state}, + stream_.get())); + + compute_container_payload_sizes<<>>( + payload_offsets, num_container_slots_, container_starts, state); + workspace_bytes = workspace_bytes_; + CUCO_CUDA_TRY(cub::DeviceScan::ExclusiveSum(workspace, + workspace_bytes, + payload_offsets, + payload_offsets, + num_container_slots_, + stream_.get())); + + // No later CUB operation needs temporary storage, so reuse it as the container work queue. + auto* const container_indexes = reinterpret_cast(workspace); + CUCO_CUDA_TRY( + cudaMemsetAsync(&state->num_array_containers, + 0, + sizeof(state->num_array_containers) + sizeof(state->num_bitset_containers), + stream_.get())); + collect_container_indexes<<>>( + container_indexes, num_container_slots_, container_starts, state); + + compute_roaring_bitmap_build_size<<<1, 1, 0, stream_.get()>>>( + state, container_starts, payload_offsets); + CUCO_CUDA_TRY(cudaPeekAtLastError()); + } + + [[nodiscard]] roaring_bitmap_build_state read_build_state( + roaring_bitmap_build_state const* state) const + { + using metadata_type = roaring_bitmap_metadata; + + roaring_bitmap_build_state host_state{}; + CUCO_CUDA_TRY(cuco::detail::memcpy_async( + &host_state, state, sizeof(host_state), cudaMemcpyDeviceToHost, stream_)); + // The serialized allocation size depends on device-computed container cardinalities. This is + // the only synchronization required before final storage can be allocated. +#if CCCL_MAJOR_VERSION > 3 || (CCCL_MAJOR_VERSION == 3 && CCCL_MINOR_VERSION >= 1) + stream_.sync(); +#else + stream_.wait(); +#endif + + CUCO_EXPECTS(host_state.num_containers >= 0 && + host_state.num_containers <= metadata_type::max_num_containers, + "Invalid generated container count"); + CUCO_EXPECTS(host_state.num_keys >= 0, "Invalid generated index count"); + CUCO_EXPECTS(host_state.num_array_containers + host_state.num_bitset_containers == + static_cast(host_state.num_containers), + "Invalid generated container indexes"); + return host_state; + } + + template + [[nodiscard]] storage_type write_serialized_bitmap(SourceIt first, + roaring_bitmap_build_state const& host_state, + cuda::std::int64_t* container_starts, + cuda::std::uint32_t* payload_offsets, + cuda::std::uint32_t* container_indexes) const + { + using metadata_type = typename storage_type::metadata_type; + + auto storage = storage_type{ + metadata_type::from_no_run_build(host_state.size_bytes, + static_cast(host_state.num_keys), + static_cast(host_state.num_containers)), + alloc_, + stream_}; + + auto const header_items = std::max(1, host_state.num_containers); + write_roaring_bitmap_header<<>>( + storage.data(), first, host_state, container_starts, payload_offsets); + + if (host_state.num_containers > 0) { + constexpr cuda::std::uint32_t block_size = 256; + constexpr cuda::std::uint32_t warps_per_block = block_size / 32; + constexpr cuda::std::uint32_t bitset_blocks_per_container = + metadata_type::bitset_container_bytes / sizeof(cuda::std::uint64_t) / block_size; + auto const array_blocks = + (host_state.num_array_containers + warps_per_block - 1) / warps_per_block; + auto const bitset_blocks = host_state.num_bitset_containers * bitset_blocks_per_container; + // Array indexes grow from the front of the queue and bitset indexes from the back. Their + // internal order is irrelevant because each entry names its destination container. + auto* const bitset_containers = + container_indexes + num_container_slots_ - host_state.num_bitset_containers; + + write_roaring_containers + <<>>(storage.data(), + first, + host_state, + container_starts, + payload_offsets, + container_indexes, + bitset_containers); + } + CUCO_CUDA_TRY(cudaPeekAtLastError()); + return storage; + } + + [[nodiscard]] static constexpr roaring_bitmap_build_state empty_build_state() noexcept + { + return {0, 0, 2 * sizeof(cuda::std::uint32_t), 0, 0}; + } + + InputIt first_; + cuda::std::int64_t num_indices_; + cuda::std::size_t num_container_slots_; + roaring_bitmap_builder_input_order input_order_; + Allocator alloc_; + cuda::stream_ref stream_; + cuda::std::size_t workspace_bytes_; +}; + +} // namespace cuco::experimental::detail diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh new file mode 100644 index 000000000..a044f3c51 --- /dev/null +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh @@ -0,0 +1,282 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include + +#include +#include +#include + +namespace cuco::experimental::detail { + +/** + * @brief Device-computed scalar state shared by the Roaring construction kernels. + */ +struct roaring_bitmap_build_state { + cuda::std::int64_t num_keys; ///< Number of sorted unique indices + cuda::std::int64_t num_containers; ///< Number of high-16-bit containers + cuda::std::uint32_t size_bytes; ///< Size of the serialized bitmap + cuda::std::uint32_t num_array_containers; ///< Number of array containers + cuda::std::uint32_t num_bitset_containers; ///< Number of bitset containers +}; + +/** + * @brief Predicate selecting the first index in each high-16-bit container. + * + * @tparam KeyIt Random access iterator over sorted unique indices + */ +template +struct is_container_start { + KeyIt keys; ///< Sorted unique indices + roaring_bitmap_build_state const* state; ///< Device build state + + /** + * @brief Tests whether an index begins a container. + * + * @param index Index in the normalized input range + * @return `true` if `index` begins a container + */ + __device__ bool operator()(cuda::std::int64_t index) const noexcept + { + auto const num_keys = state->num_keys; + if (index >= num_keys) { return false; } + if (index == 0) { return true; } + return (keys[index] >> 16) != (keys[index - 1] >> 16); + } +}; + +template +is_container_start(KeyIt, roaring_bitmap_build_state const*) -> is_container_start; + +/** + * @brief Computes the encoded payload size of a container. + */ +struct container_payload_size { + cuda::std::int64_t const* container_starts; ///< Starting input index of each container + roaring_bitmap_build_state const* state; ///< Device build state + + /** + * @brief Returns the encoded payload size for one container slot. + * + * @param index Container slot index + * @return Payload size in bytes, or zero for an unused slot + */ + __device__ cuda::std::uint32_t operator()(cuda::std::int64_t index) const noexcept + { + using metadata_type = roaring_bitmap_metadata; + + auto const num_containers = state->num_containers; + if (index >= num_containers) { return 0; } + + auto const begin = container_starts[index]; + auto const end = index + 1 < num_containers ? container_starts[index + 1] : state->num_keys; + auto const cardinality = static_cast(end - begin); + return cardinality <= metadata_type::max_array_container_card + ? cardinality * sizeof(cuda::std::uint16_t) + : metadata_type::bitset_container_bytes; + } +}; + +template +CUCO_KERNEL void compute_container_payload_sizes(cuda::std::uint32_t* payload_sizes, + cuda::std::int64_t num_container_slots, + ContainerStartIt container_starts, + roaring_bitmap_build_state const* state) +{ + auto const index = cuco::detail::global_thread_id(); + if (index >= num_container_slots) { return; } + // Slots beyond the selected container count contribute zero to the fixed-size scan. + payload_sizes[index] = container_payload_size{container_starts, state}(index); +} + +template +CUCO_KERNEL void compute_roaring_bitmap_build_size(roaring_bitmap_build_state* state, + ContainerStartIt container_starts, + PayloadOffsetIt payload_offsets) +{ + using metadata_type = roaring_bitmap_metadata; + + if (blockIdx.x != 0 || threadIdx.x != 0) { return; } + + auto const num_containers = state->num_containers; + if (num_containers == 0) { + state->size_bytes = 2 * sizeof(cuda::std::uint32_t); + return; + } + + auto const last = num_containers - 1; + auto const begin = container_starts[last]; + auto const cardinality = static_cast(state->num_keys - begin); + auto const payload_size = cardinality <= metadata_type::max_array_container_card + ? cardinality * sizeof(cuda::std::uint16_t) + : metadata_type::bitset_container_bytes; + auto const payload_begin = 2 * sizeof(cuda::std::uint32_t) + + static_cast(num_containers) * + (2 * sizeof(cuda::std::uint16_t) + sizeof(cuda::std::uint32_t)); + + state->size_bytes = payload_begin + payload_offsets[last] + payload_size; +} + +template +CUCO_KERNEL void collect_container_indexes(cuda::std::uint32_t* container_indexes, + cuda::std::int64_t num_container_slots, + ContainerStartIt container_starts, + roaring_bitmap_build_state* state) +{ + using metadata_type = roaring_bitmap_metadata; + + auto const index = cuco::detail::global_thread_id(); + if (index >= state->num_containers || index >= num_container_slots) { return; } + + auto const begin = container_starts[index]; + auto const end = + index + 1 < state->num_containers ? container_starts[index + 1] : state->num_keys; + auto const cardinality = static_cast(end - begin); + if (cardinality <= metadata_type::max_array_container_card) { + // The queue does not need to preserve container order: writers use the stored container index + // to find the final payload offset. + auto const output_index = atomicAdd(&state->num_array_containers, 1); + container_indexes[output_index] = static_cast(index); + } else { + // Bitset indexes grow backward so both work lists share one num_container_slots-sized buffer. + auto const output_index = atomicAdd(&state->num_bitset_containers, 1); + container_indexes[num_container_slots - output_index - 1] = + static_cast(index); + } +} + +template +__device__ cuda::std::int64_t lower_bound_low_bits(KeyIt keys, + cuda::std::int64_t first, + cuda::std::int64_t last, + cuda::std::uint32_t value) +{ + while (first < last) { + auto const middle = first + (last - first) / 2; + auto const lower = static_cast(keys[middle]); + if (lower < value) { + first = middle + 1; + } else { + last = middle; + } + } + return first; +} + +template +CUCO_KERNEL void write_roaring_bitmap_header(cuda::std::byte* bitmap, + KeyIt keys, + roaring_bitmap_build_state state, + ContainerStartIt container_starts, + PayloadOffsetIt payload_offsets) +{ + using metadata_type = roaring_bitmap_metadata; + + auto const index = cuco::detail::global_thread_id(); + + if (index == 0) { + misaligned_store(bitmap, metadata_type::serial_cookie_no_runcontainer); + misaligned_store(bitmap + sizeof(cuda::std::uint32_t), + static_cast(state.num_containers)); + } + + if (index >= state.num_containers) { return; } + + auto const begin = container_starts[index]; + auto const end = index + 1 < state.num_containers ? container_starts[index + 1] : state.num_keys; + auto const cardinality = static_cast(end - begin); + auto const key = static_cast(keys[begin] >> 16); + auto const card_minus_one = static_cast(cardinality - 1); + + auto* const key_cards = bitmap + 2 * sizeof(cuda::std::uint32_t); + misaligned_store(key_cards + index * 2 * sizeof(cuda::std::uint16_t), key); + misaligned_store(key_cards + (index * 2 + 1) * sizeof(cuda::std::uint16_t), card_minus_one); + + auto* const offsets = key_cards + state.num_containers * 2 * sizeof(cuda::std::uint16_t); + auto const payload_begin = 2 * sizeof(cuda::std::uint32_t) + + static_cast(state.num_containers) * + (2 * sizeof(cuda::std::uint16_t) + sizeof(cuda::std::uint32_t)); + reinterpret_cast(offsets)[index] = payload_begin + payload_offsets[index]; +} + +template +CUCO_KERNEL void write_roaring_containers(cuda::std::byte* bitmap, + KeyIt keys, + roaring_bitmap_build_state state, + ContainerStartIt container_starts, + PayloadOffsetIt payload_offsets, + ContainerIndexIt array_containers, + ContainerIndexIt bitset_containers) +{ + using metadata_type = roaring_bitmap_metadata; + + constexpr cuda::std::uint32_t warps_per_block = BlockSize / cuco::detail::warp_size(); + constexpr cuda::std::uint32_t bitset_words = + metadata_type::bitset_container_bytes / sizeof(unsigned long long); + constexpr cuda::std::uint32_t bitset_blocks_per_container = bitset_words / BlockSize; + static_assert(BlockSize % cuco::detail::warp_size() == 0); + static_assert(bitset_words % BlockSize == 0); + + auto const payload_begin = 2 * sizeof(cuda::std::uint32_t) + + static_cast(state.num_containers) * + (2 * sizeof(cuda::std::uint16_t) + sizeof(cuda::std::uint32_t)); + auto const array_blocks = (state.num_array_containers + warps_per_block - 1) / warps_per_block; + auto const block = static_cast(blockIdx.x); + + // Array containers use one warp each. Remaining blocks are divided into four 256-word pieces of + // a bitset container. + if (block < array_blocks) { + auto const warp_index = block * warps_per_block + threadIdx.x / cuco::detail::warp_size(); + if (warp_index >= state.num_array_containers) { return; } + + auto const lane = static_cast(threadIdx.x) % cuco::detail::warp_size(); + auto const container_index = static_cast(array_containers[warp_index]); + auto const begin = container_starts[container_index]; + auto const end = container_index + 1 < state.num_containers + ? container_starts[container_index + 1] + : state.num_keys; + auto const cardinality = static_cast(end - begin); + auto* const container = bitmap + payload_begin + payload_offsets[container_index]; + + for (auto index = lane; index < cardinality; index += cuco::detail::warp_size()) { + auto const value = static_cast(keys[begin + index]); + misaligned_store(container + index * sizeof(cuda::std::uint16_t), value); + } + } else { + auto const bitset_block = block - array_blocks; + auto const bitset_index = bitset_block / bitset_blocks_per_container; + if (bitset_index >= state.num_bitset_containers) { return; } + + auto const quadrant = bitset_block % bitset_blocks_per_container; + auto const word = quadrant * BlockSize + threadIdx.x; + auto const container_index = static_cast(bitset_containers[bitset_index]); + auto const begin = container_starts[container_index]; + auto const end = container_index + 1 < state.num_containers + ? container_starts[container_index + 1] + : state.num_keys; + auto* const container = bitmap + payload_begin + payload_offsets[container_index]; + auto const word_begin = word * 64; + auto const word_end = word_begin + 64; + auto const first = lower_bound_low_bits(keys, begin, end, word_begin); + auto const last = lower_bound_low_bits(keys, first, end, word_end); + + // One thread constructs one 64-bit word entirely in registers before issuing one final store. + unsigned long long mask = 0; + for (auto index = first; index < last; ++index) { + auto const value = static_cast(keys[index]); + mask |= 1ULL << (value - word_begin); + } + misaligned_store(container + word * sizeof(unsigned long long), mask); + } +} + +} // namespace cuco::experimental::detail diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh index 0d8e465a3..5c2af5024 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh @@ -65,7 +65,7 @@ class roaring_bitmap_storage_ref { * @param bitmap Pointer to the serialized bitmap in a device-accessible memory location */ __device__ roaring_bitmap_storage_ref(cuda::std::byte const* bitmap) - : roaring_bitmap_storage_ref{bitmap, metadata_type{bitmap}} + : roaring_bitmap_storage_ref{bitmap, metadata_type::from_serialized(bitmap)} { } @@ -212,6 +212,8 @@ class roaring_bitmap_storage { typename std::allocator_traits::template rebind_alloc; /// Reference type for this storage using ref_type = roaring_bitmap_storage_ref; + /// Metadata type for this storage + using metadata_type = typename ref_type::metadata_type; /** * @brief Copy constructor @@ -220,13 +222,6 @@ class roaring_bitmap_storage { */ roaring_bitmap_storage(roaring_bitmap_storage const& other) = default; - /** - * @brief Move constructor - * - * @param other The roaring_bitmap_storage to move from - */ - roaring_bitmap_storage(roaring_bitmap_storage&& other) = default; - /** * @brief Copy assignment operator * @@ -235,14 +230,6 @@ class roaring_bitmap_storage { */ roaring_bitmap_storage& operator=(roaring_bitmap_storage const& other) = default; - /** - * @brief Move assignment operator - * - * @param other The roaring_bitmap_storage to move from - * @return Reference to this roaring_bitmap_storage - */ - roaring_bitmap_storage& operator=(roaring_bitmap_storage&& other) = default; - ~roaring_bitmap_storage() = default; /** @@ -256,7 +243,7 @@ class roaring_bitmap_storage { Allocator const& alloc, cuda::stream_ref stream) : allocator_{alloc}, - metadata_{bitmap}, + metadata_{metadata_type::from_serialized(bitmap)}, data_{allocator_.allocate(metadata_.size_bytes, stream), cuco::detail::custom_deleter{ metadata_.size_bytes, allocator_, stream}}, @@ -266,6 +253,67 @@ class roaring_bitmap_storage { data_.get(), bitmap, metadata_.size_bytes, cudaMemcpyHostToDevice, stream.get())); } + /** + * @brief Constructs storage for generated bitmap data + * + * @param metadata Metadata describing the generated bitmap + * @param alloc Allocator for device memory allocation + * @param stream CUDA stream for memory operations + */ + roaring_bitmap_storage(metadata_type const& metadata, + Allocator const& alloc, + cuda::stream_ref stream) + : allocator_{alloc}, + metadata_{metadata}, + data_{allocator_.allocate(metadata_.size_bytes, stream), + cuco::detail::custom_deleter{ + metadata_.size_bytes, allocator_, stream}}, + ref_{data_.get(), metadata_} + { + assert(metadata_.valid); + } + + /** + * @brief Move constructor. + * + * Rebuilds the cached reference because small run-container bitmaps store computed offsets + * directly in the metadata object. + * + * @param other Storage to move from + */ + roaring_bitmap_storage(roaring_bitmap_storage&& other) noexcept + : allocator_{std::move(other.allocator_)}, + metadata_{std::move(other.metadata_)}, + data_{std::move(other.data_)}, + ref_{data_.get(), metadata_} + { + } + + /** + * @brief Move assignment operator. + * + * Rebuilds the cached reference because small run-container bitmaps store computed offsets + * directly in the metadata object. + * + * @param other Storage to move from + * @return Reference to this storage + */ + roaring_bitmap_storage& operator=(roaring_bitmap_storage&& other) noexcept + { + allocator_ = std::move(other.allocator_); + metadata_ = std::move(other.metadata_); + data_ = std::move(other.data_); + ref_ = ref_type{data_.get(), metadata_}; + return *this; + } + + /** + * @brief Returns a mutable pointer to serialized storage + * + * @return Pointer to serialized storage + */ + cuda::std::byte* data() noexcept { return data_.get(); } + /** * @brief Returns a reference to the stored bitmap * @@ -273,9 +321,16 @@ class roaring_bitmap_storage { */ ref_type ref() const noexcept { return ref_; } + /** + * @brief Returns the allocator used to manage storage + * + * @return Allocator instance + */ + allocator_type allocator() const noexcept { return allocator_; } + private: allocator_type allocator_; - typename ref_type::metadata_type metadata_; + metadata_type metadata_; std::unique_ptr> data_; ref_type ref_; @@ -351,7 +406,7 @@ class roaring_bitmap_storage { buckets_h_{}, metadata_{ [bitmap](std::vector& bucket_metadata) { - return typename ref_type::metadata_type{bitmap, bucket_metadata}; + return ref_type::metadata_type::from_serialized(bitmap, bucket_metadata); }(bucket_metadata_)}, data_{allocator_.allocate(metadata_.size_bytes, stream), cuco::detail::custom_deleter{ @@ -384,6 +439,13 @@ class roaring_bitmap_storage { */ ref_type ref() const noexcept { return ref_; } + /** + * @brief Returns the allocator used to manage storage + * + * @return Allocator instance + */ + allocator_type allocator() const noexcept { return allocator_; } + private: allocator_type allocator_; bucket_allocator_type bucket_allocator_; diff --git a/include/cuco/detail/roaring_bitmap/util.cuh b/include/cuco/detail/roaring_bitmap/util.cuh index 87841f664..6f7f03ae5 100644 --- a/include/cuco/detail/roaring_bitmap/util.cuh +++ b/include/cuco/detail/roaring_bitmap/util.cuh @@ -32,6 +32,12 @@ __host__ __device__ __forceinline__ T misaligned_load(cuda::std::byte const* ptr return value; } +template +__host__ __device__ __forceinline__ void misaligned_store(cuda::std::byte* ptr, T value) +{ + cuda::std::memcpy(ptr, &value, sizeof(T)); +} + __host__ __device__ __forceinline__ bool check_bit(cuda::std::byte const* bitmap, cuda::std::uint32_t index) { @@ -53,6 +59,12 @@ struct roaring_bitmap_metadata { */ template <> struct roaring_bitmap_metadata { + /// Serialization cookie for bitmaps without run containers + static constexpr cuda::std::uint32_t serial_cookie_no_runcontainer = 12346; + /// Serialization cookie for bitmaps with run containers + static constexpr cuda::std::uint32_t serial_cookie = 12347; + /// Maximum number of containers in a 32-bit bitmap + static constexpr cuda::std::int32_t max_num_containers = 1 << 16; /// Maximum number of elements in an array container before converting to bitmap static constexpr cuda::std::uint32_t max_array_container_card = 4096; /// Threshold for omitting container offsets in serialized format @@ -81,19 +93,54 @@ struct roaring_bitmap_metadata { /// Whether container offsets are stored in the serialized data bool offsets_in_serialized_data = true; + roaring_bitmap_metadata() = default; + + /** + * @brief Creates metadata for a generated bitmap without run containers + * + * @param bitmap_size_bytes Size of the serialized bitmap in bytes + * @param bitmap_num_keys Number of unique keys in the bitmap + * @param bitmap_num_containers Number of containers in the bitmap + * @return Metadata describing the generated bitmap + */ + [[nodiscard]] static roaring_bitmap_metadata from_no_run_build( + cuda::std::size_t bitmap_size_bytes, + cuda::std::size_t bitmap_num_keys, + cuda::std::int32_t bitmap_num_containers) noexcept + { + roaring_bitmap_metadata metadata; + metadata.size_bytes = bitmap_size_bytes; + metadata.num_keys = bitmap_num_keys; + metadata.key_cards = 2 * sizeof(cuda::std::uint32_t); + metadata.container_offsets = + metadata.key_cards + bitmap_num_containers * 2 * sizeof(cuda::std::uint16_t); + metadata.num_containers = bitmap_num_containers; + metadata.has_run = false; + metadata.valid = true; + metadata.offsets_in_serialized_data = true; + return metadata; + } + /** - * @brief Constructs metadata from a serialized bitmap + * @brief Creates metadata from a serialized bitmap * * @param bitmap Pointer to the beginning of the serialized bitmap + * @return Metadata parsed from the serialized bitmap */ - __host__ __device__ roaring_bitmap_metadata(cuda::std::byte const* bitmap) + [[nodiscard]] __host__ __device__ static roaring_bitmap_metadata from_serialized( + cuda::std::byte const* bitmap) + { + roaring_bitmap_metadata metadata; + metadata.parse_serialized(bitmap); + return metadata; + } + + private: + __host__ __device__ void parse_serialized(cuda::std::byte const* bitmap) { - constexpr cuda::std::uint32_t serial_cookie_no_runcontainer = 12346; - constexpr cuda::std::uint32_t serial_cookie = 12347; // constexpr cuda::std::uint32_t frozen_cookie = 13766; // not implemented - constexpr cuda::std::int32_t max_containers = 1 << 16; - constexpr cuda::std::uint32_t cookie_mask = 0xFFFF; - constexpr cuda::std::uint32_t cookie_shift = 16; + constexpr cuda::std::uint32_t cookie_mask = 0xFFFF; + constexpr cuda::std::uint32_t cookie_shift = 16; cuda::std::byte const* buf = bitmap; @@ -118,7 +165,7 @@ struct roaring_bitmap_metadata { cuda::std::memcpy(&num_containers, buf, sizeof(cuda::std::uint32_t)); buf += sizeof(cuda::std::uint32_t); } - if (num_containers < 0 or num_containers > max_containers) { + if (num_containers < 0 or num_containers > max_num_containers) { valid = false; NV_IF_TARGET( NV_IS_HOST, @@ -183,6 +230,13 @@ struct roaring_bitmap_metadata { buf = container_ptr; } + if (num_containers == 0) { + // There is no last container from which to derive the end of the serialized stream. + size_bytes = static_cast(cuda::std::distance(bitmap, buf)); + valid = true; + return; + } + cuda::std::uint32_t card = 0; for (cuda::std::int32_t i = 0; i < num_containers; i++) { cuda::std::byte const* card_ptr = @@ -238,6 +292,8 @@ struct roaring_bitmap_metadata { /// Whether the metadata is valid bool valid = false; + roaring_bitmap_metadata() = default; + /** * @brief Metadata for individual buckets in a 64-bit roaring bitmap * @@ -267,13 +323,37 @@ struct roaring_bitmap_metadata { }; /** - * @brief Constructs metadata from a serialized 64-bit bitmap with bucket metadata + * @brief Creates metadata from a serialized 64-bit bitmap with bucket metadata * * @param bitmap Pointer to the beginning of the serialized bitmap * @param bucket_metadata Vector to store metadata for each bucket + * @return Metadata parsed from the serialized bitmap + */ + [[nodiscard]] __host__ static roaring_bitmap_metadata from_serialized( + cuda::std::byte const* bitmap, std::vector& bucket_metadata) + { + roaring_bitmap_metadata metadata; + metadata.parse_serialized(bitmap, bucket_metadata); + return metadata; + } + + /** + * @brief Creates metadata from a serialized 64-bit bitmap + * + * @param bitmap Pointer to the beginning of the serialized bitmap + * @return Metadata parsed from the serialized bitmap */ - __host__ roaring_bitmap_metadata(cuda::std::byte const* bitmap, - std::vector& bucket_metadata) + [[nodiscard]] __host__ __device__ static roaring_bitmap_metadata from_serialized( + cuda::std::byte const* bitmap) + { + roaring_bitmap_metadata metadata; + metadata.parse_serialized(bitmap); + return metadata; + } + + private: + __host__ void parse_serialized(cuda::std::byte const* bitmap, + std::vector& bucket_metadata) { cuda::std::size_t byte_offset = 0; cuda::std::byte const* bitmap_ptr = bitmap; @@ -287,7 +367,8 @@ struct roaring_bitmap_metadata { cuda::std::uint32_t bucket_key; cuda::std::memcpy(&bucket_key, bitmap_ptr + byte_offset, sizeof(cuda::std::uint32_t)); byte_offset += sizeof(cuda::std::uint32_t); // skip bucket key - roaring_bitmap_metadata bucket_meta{bitmap_ptr + byte_offset}; + auto const bucket_meta = + roaring_bitmap_metadata::from_serialized(bitmap_ptr + byte_offset); if (!bucket_meta.valid) { valid = false; return; @@ -300,12 +381,7 @@ struct roaring_bitmap_metadata { valid = true; } - /** - * @brief Constructs metadata from a serialized 64-bit bitmap - * - * @param bitmap Pointer to the beginning of the serialized bitmap - */ - __host__ __device__ roaring_bitmap_metadata(cuda::std::byte const* bitmap) + __host__ __device__ void parse_serialized(cuda::std::byte const* bitmap) { cuda::std::size_t byte_offset = 0; cuda::std::byte const* bitmap_ptr = bitmap; @@ -314,7 +390,8 @@ struct roaring_bitmap_metadata { for (cuda::std::size_t i = 0; i < num_buckets; ++i) { byte_offset += sizeof(cuda::std::uint32_t); // skip bucket key - roaring_bitmap_metadata bucket_meta{bitmap_ptr + byte_offset}; + auto const bucket_meta = + roaring_bitmap_metadata::from_serialized(bitmap_ptr + byte_offset); if (!bucket_meta.valid) { valid = false; return; diff --git a/include/cuco/roaring_bitmap.cuh b/include/cuco/roaring_bitmap.cuh index d01b990f4..b6305be6b 100644 --- a/include/cuco/roaring_bitmap.cuh +++ b/include/cuco/roaring_bitmap.cuh @@ -19,20 +19,21 @@ namespace cuco::experimental { * * The `roaring_bitmap` provides host-side bulk membership queries over a bitmap stored in the * [Roaring bitmap format specification](https://github.com/RoaringBitmap/RoaringFormatSpec). - * The serialized bytes are copied to device-accessible storage upon construction, and queries are + * It can be constructed by copying an existing serialized bitmap to device-accessible storage or + * built on the GPU from an unordered sequence of potentially duplicate indices. Queries are * executed on the GPU. * * In addition to bulk host APIs such as `contains`/`contains_async`, this container exposes a * non-owning reference object via `ref()` that can be used for device-side per-thread queries. * - * @tparam T Key type. Must be `cuda::std::uint32_t` or `cuda::std::uint64_t`. + * @tparam T Index type. Must be `cuda::std::uint32_t` or `cuda::std::uint64_t`. * @tparam Allocator Allocator type used to manage device-accessible storage for the serialized * bytes. */ template > class roaring_bitmap { public: - using value_type = T; ///< Key type + using value_type = T; ///< Index type using storage_type = detail::roaring_bitmap_storage; ///< Storage implementation using allocator_type = typename storage_type::allocator_type; ///< Allocator type using ref_type = roaring_bitmap_ref; ///< Non-owning reference type @@ -41,6 +42,13 @@ class roaring_bitmap { * @brief Constructs a `roaring_bitmap` by copying the serialized bytes to device-accessible * storage. * + * @note Construction of 32-bit bitmaps through this constructor is deprecated. Use + * `from_serialized` instead. The constructor remains available without a compiler + * deprecation attribute to avoid breaking existing code. + * @note `bitmap` must remain valid until `stream` completes the copy. The bitmap can be used + * immediately by work submitted to the same stream; use an explicit dependency before + * accessing it from another stream. + * * @param bitmap Pointer to the beginning of the serialized bitmap in host memory * @param alloc Allocator used to allocate device-accessible storage * @param stream CUDA stream used for device memory operations during construction @@ -49,6 +57,102 @@ class roaring_bitmap { Allocator const& alloc = {}, cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}); + /** + * @brief Creates a 32-bit `roaring_bitmap` by copying serialized bytes to device-accessible + * storage. + * + * @note This factory currently supports only `cuda::std::uint32_t` bitmaps. + * @note `bitmap` must remain valid until `stream` completes the copy. The bitmap can be used + * immediately by work submitted to the same stream; use an explicit dependency before + * accessing it from another stream. + * + * @param bitmap Pointer to the beginning of the serialized bitmap in host memory + * @param alloc Allocator used to allocate device-accessible storage + * @param stream CUDA stream used for device memory operations during construction + * + * @return Bitmap containing a copy of the serialized input + */ + [[nodiscard]] static roaring_bitmap from_serialized(cuda::std::byte const* bitmap, + Allocator const& alloc = {}, + cuda::stream_ref stream = cuda::stream_ref{ + cudaStream_t{nullptr}}); + + /** + * @brief Creates a `roaring_bitmap` from an unordered sequence of indices. + * + * @note The input must remain valid until the construction stream completes. + * @note This function may synchronize `stream` once to determine the exact serialized allocation + * size. Serialization remains stream-ordered and may still be in progress when the function + * returns. The object can be used immediately by work submitted to the same stream; use an + * explicit dependency before accessing it from another stream. + * + * @tparam InputIt Device-accessible random access input iterator whose value type is + * `cuda::std::uint32_t` + * + * @param first Beginning of the sequence of indices + * @param last End of the sequence of indices + * @param alloc Allocator used for permanent and temporary device storage + * @param stream CUDA stream used for device memory operations and kernel launches + * + * @return Bitmap containing the unique input indices + */ + template + [[nodiscard]] static roaring_bitmap from_indices(InputIt first, + InputIt last, + Allocator const& alloc = {}, + cuda::stream_ref stream = cuda::stream_ref{ + cudaStream_t{nullptr}}); + + /** + * @brief Creates a `roaring_bitmap` from sorted indices that may contain duplicates. + * + * @note The input range must be nondecreasing. This precondition is not checked. + * @note The input must remain valid until the construction stream completes. + * @note This function may synchronize `stream` once to determine the exact serialized allocation + * size. Final serialization remains stream ordered. + * + * @tparam InputIt Device-accessible random access input iterator whose value type is + * `cuda::std::uint32_t` + * + * @param first Beginning of the sorted sequence + * @param last End of the sorted sequence + * @param alloc Allocator used for permanent and temporary device storage + * @param stream CUDA stream used for device memory operations and kernel launches + * + * @return Bitmap containing the unique input indices + */ + template + [[nodiscard]] static roaring_bitmap from_sorted_indices( + InputIt first, + InputIt last, + Allocator const& alloc = {}, + cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}); + + /** + * @brief Creates a `roaring_bitmap` from sorted unique indices. + * + * @note The input range must be strictly increasing. This precondition is not checked. + * @note The input must remain valid until the construction stream completes. + * @note This function may synchronize `stream` once to determine the exact serialized allocation + * size. Final serialization remains stream ordered. + * + * @tparam InputIt Device-accessible random access input iterator whose value type is + * `cuda::std::uint32_t` + * + * @param first Beginning of the sorted unique sequence + * @param last End of the sorted unique sequence + * @param alloc Allocator used for permanent and temporary device storage + * @param stream CUDA stream used for device memory operations and kernel launches + * + * @return Bitmap containing the input indices + */ + template + [[nodiscard]] static roaring_bitmap from_sorted_unique_indices( + InputIt first, + InputIt last, + Allocator const& alloc = {}, + cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}); + /** * @brief Copy constructor * @@ -82,18 +186,18 @@ class roaring_bitmap { ~roaring_bitmap() = default; ///< Destructor /** - * @brief Bulk membership query for keys in `[first, last)`. + * @brief Bulk membership query for indices in `[first, last)`. * * @note This function synchronizes the given stream. For asynchronous execution use * `contains_async`. * - * @tparam InputIt Device-accessible random access input iterator of keys convertible to `T` + * @tparam InputIt Device-accessible random access input iterator of indices convertible to `T` * @tparam OutputIt Device-accessible random access output iterator whose `value_type` is * constructible from `bool` * - * @param first Beginning of the sequence of keys - * @param last End of the sequence of keys - * @param contained Output iterator where results are written; `true` iff the corresponding key + * @param first Beginning of the sequence of indices + * @param last End of the sequence of indices + * @param contained Output iterator where results are written; `true` iff the corresponding index * is present in the bitmap * @param stream CUDA stream used for device memory operations and kernel launches */ @@ -104,14 +208,14 @@ class roaring_bitmap { cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}}) const; /** - * @brief Asynchronously performs a bulk membership query for keys in `[first, last)`. + * @brief Asynchronously performs a bulk membership query for indices in `[first, last)`. * - * @tparam InputIt Device-accessible random access input iterator of keys convertible to `T` + * @tparam InputIt Device-accessible random access input iterator of indices convertible to `T` * @tparam OutputIt Device-accessible random access output iterator to `bool` * - * @param first Beginning of the sequence of keys - * @param last End of the sequence of keys - * @param contained Output iterator where results are written; `true` iff the corresponding key + * @param first Beginning of the sequence of indices + * @param last End of the sequence of indices + * @param contained Output iterator where results are written; `true` iff the corresponding index * is present in the bitmap * @param stream CUDA stream used for device memory operations and kernel launches */ @@ -123,14 +227,14 @@ class roaring_bitmap { cudaStream_t{nullptr}}) const noexcept; /** - * @brief Number of keys stored in the bitmap. + * @brief Number of indices stored in the bitmap. * - * @return Count of keys in the bitmap + * @return Count of indices in the bitmap */ [[nodiscard]] cuda::std::size_t size() const noexcept; /** - * @brief Checks whether the bitmap contains no keys. + * @brief Checks whether the bitmap contains no indices. * * @return `true` iff `size() == 0` */ @@ -169,9 +273,11 @@ class roaring_bitmap { [[nodiscard]] ref_type ref() const noexcept; private: + explicit roaring_bitmap(storage_type&& storage); + storage_type storage_; ///< Storage type }; } // namespace cuco::experimental -#include \ No newline at end of file +#include diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index ae98b1216..bb44b587d 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -155,4 +155,5 @@ ConfigureTest(BLOOM_FILTER_TEST ################################################################################################### # - roaring_bitmap --------------------------------------------------------------------------------- ConfigureTest(ROARING_BITMAP_TEST - roaring_bitmap/contains_test.cu) + roaring_bitmap/contains_test.cu + roaring_bitmap/build_test.cu) diff --git a/tests/roaring_bitmap/build_test.cu b/tests/roaring_bitmap/build_test.cu new file mode 100644 index 000000000..c846930d0 --- /dev/null +++ b/tests/roaring_bitmap/build_test.cu @@ -0,0 +1,449 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace { + +using index_type = cuda::std::uint32_t; +using bitmap_type = cuco::experimental::roaring_bitmap; + +std::vector copy_serialized(bitmap_type const& bitmap) +{ + std::vector bytes(bitmap.size_bytes()); + CUCO_CUDA_TRY(cudaMemcpy(bytes.data(), bitmap.data(), bytes.size(), cudaMemcpyDeviceToHost)); + return bytes; +} + +void require_contains(bitmap_type const& bitmap, + std::vector const& queries, + std::vector const& expected) +{ + thrust::device_vector query_indices(queries); + thrust::device_vector results(queries.size()); + bitmap.contains(query_indices.begin(), query_indices.end(), results.begin()); + thrust::host_vector host_results = results; + + REQUIRE(host_results.size() == expected.size()); + for (cuda::std::size_t i = 0; i < expected.size(); ++i) { + REQUIRE(host_results[i] == expected[i]); + } +} + +struct allocation_counts { + cuda::std::size_t allocations = 0; + cuda::std::size_t deallocations = 0; +}; + +template +class tracking_allocator { + public: + using value_type = T; + + tracking_allocator() : counts_{std::make_shared()} {} + + explicit tracking_allocator(std::shared_ptr counts) : counts_{counts} {} + + template + tracking_allocator(tracking_allocator const& other) noexcept : counts_{other.counts()} + { + } + + value_type* allocate(cuda::std::size_t size, cuda::stream_ref stream) + { + ++counts_->allocations; + value_type* data; + CUCO_CUDA_TRY(cudaMallocAsync(&data, size * sizeof(value_type), stream.get())); + return data; + } + + void deallocate(value_type* data, cuda::std::size_t, cuda::stream_ref stream) + { + ++counts_->deallocations; + CUCO_CUDA_TRY(cudaFreeAsync(data, stream.get())); + } + + [[nodiscard]] std::shared_ptr counts() const noexcept { return counts_; } + + private: + template + friend class tracking_allocator; + + std::shared_ptr counts_; +}; + +template +bool operator==(tracking_allocator const& lhs, tracking_allocator const& rhs) noexcept +{ + return lhs.counts() == rhs.counts(); +} + +template +bool operator!=(tracking_allocator const& lhs, tracking_allocator const& rhs) noexcept +{ + return not(lhs == rhs); +} + +} // namespace + +TEST_CASE("roaring_bitmap builds an empty bitmap", "[roaring_bitmap]") +{ + thrust::device_vector indices; + auto bitmap = bitmap_type::from_indices(indices.begin(), indices.end()); + auto sorted_bitmap = bitmap_type::from_sorted_indices(indices.begin(), indices.end()); + auto sorted_unique_bitmap = + bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.empty()); + REQUIRE(bitmap.size() == 0); + REQUIRE(bitmap.size_bytes() == 2 * sizeof(cuda::std::uint32_t)); + REQUIRE(copy_serialized(bitmap) == copy_serialized(sorted_bitmap)); + REQUIRE(copy_serialized(bitmap) == copy_serialized(sorted_unique_bitmap)); +} + +TEST_CASE("roaring_bitmap matches RoaringFormatSpec no-run serialization", "[roaring_bitmap]") +{ +#ifndef CUCO_ROARING_DATA_DIR + SKIP( + "CUCO_ROARING_DATA_DIR is not defined. Configure with -DCUCO_DOWNLOAD_ROARING_TESTDATA=ON to " + "run this test."); +#else + std::vector host_indices; + for (index_type index = 0; index < 100000; index += 1000) { + host_indices.push_back(index); + } + for (index_type index = 100000; index < 200000; ++index) { + host_indices.push_back(3 * index); + } + for (index_type index = 700000; index < 800000; ++index) { + host_indices.push_back(index); + } + + thrust::device_vector indices{host_indices}; + auto bitmap = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + std::string const path = std::string{CUCO_ROARING_DATA_DIR} + "/bitmapwithoutruns.bin"; + REQUIRE(std::filesystem::exists(path)); + std::ifstream file{path, std::ios::binary | std::ios::ate}; + REQUIRE(file.is_open()); + auto const size = static_cast(file.tellg()); + std::vector expected(size); + file.seekg(0); + file.read(reinterpret_cast(expected.data()), size); + + REQUIRE(bitmap.size() == host_indices.size()); + REQUIRE(copy_serialized(bitmap) == expected); +#endif +} + +TEST_CASE("roaring_bitmap treats reversed input ranges as empty", "[roaring_bitmap]") +{ + thrust::device_vector indices{1, 2, 3}; + + auto bitmap = bitmap_type::from_indices(indices.end(), indices.begin()); + auto sorted_bitmap = bitmap_type::from_sorted_indices(indices.end(), indices.begin()); + auto sorted_unique_bitmap = + bitmap_type::from_sorted_unique_indices(indices.end(), indices.begin()); + + REQUIRE(bitmap.empty()); + REQUIRE(copy_serialized(bitmap) == copy_serialized(sorted_bitmap)); + REQUIRE(copy_serialized(bitmap) == copy_serialized(sorted_unique_bitmap)); +} + +TEST_CASE("roaring_bitmap builds multiple array containers per block", "[roaring_bitmap]") +{ + std::vector const cardinalities{ + 1, 31, 32, 33, 255, 256, 257, 4094, 4095, 4096}; + std::vector host_indices; + + for (cuda::std::size_t container = 0; container < cardinalities.size(); ++container) { + for (cuda::std::uint32_t lower = 0; lower < cardinalities[container]; ++lower) { + host_indices.push_back((static_cast(container) << 16) | lower); + } + } + + thrust::device_vector indices(host_indices); + auto bitmap = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + auto const expected_size = + 2 * sizeof(cuda::std::uint32_t) + + cardinalities.size() * (2 * sizeof(cuda::std::uint16_t) + sizeof(cuda::std::uint32_t)) + + host_indices.size() * sizeof(cuda::std::uint16_t); + + REQUIRE(bitmap.size() == host_indices.size()); + REQUIRE(bitmap.size_bytes() == expected_size); + require_contains(bitmap, host_indices, std::vector(host_indices.size(), true)); + + std::vector absent; + absent.reserve(cardinalities.size()); + for (cuda::std::size_t container = 0; container < cardinalities.size(); ++container) { + absent.push_back((static_cast(container) << 16) | cardinalities[container]); + } + require_contains(bitmap, absent, std::vector(absent.size(), false)); +} + +TEST_CASE("roaring_bitmap writes array and bitset containers in one grid", "[roaring_bitmap]") +{ + constexpr cuda::std::uint32_t max_array_cardinality = 4096; + constexpr cuda::std::uint32_t bitset_bytes = 8192; + std::vector const cardinalities{1, 2, 3, 4, 5, 6, 7, 8, 9, 4097, 5000}; + std::vector host_indices; + + for (cuda::std::size_t container = 0; container < cardinalities.size(); ++container) { + for (cuda::std::uint32_t lower = 0; lower < cardinalities[container]; ++lower) { + host_indices.push_back((static_cast(container) << 16) | lower); + } + } + + thrust::device_vector indices(host_indices); + auto bitmap = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + auto expected_size = + 2 * sizeof(cuda::std::uint32_t) + + cardinalities.size() * (2 * sizeof(cuda::std::uint16_t) + sizeof(cuda::std::uint32_t)); + for (auto const cardinality : cardinalities) { + expected_size += cardinality <= max_array_cardinality + ? cardinality * sizeof(cuda::std::uint16_t) + : bitset_bytes; + } + + REQUIRE(bitmap.size() == host_indices.size()); + REQUIRE(bitmap.size_bytes() == expected_size); + require_contains(bitmap, host_indices, std::vector(host_indices.size(), true)); + + std::vector absent; + absent.reserve(cardinalities.size()); + for (cuda::std::size_t container = 0; container < cardinalities.size(); ++container) { + absent.push_back((static_cast(container) << 16) | cardinalities[container]); + } + require_contains(bitmap, absent, std::vector(absent.size(), false)); +} + +TEST_CASE("roaring_bitmap selects array and bitset containers at the format threshold", + "[roaring_bitmap]") +{ + SECTION("array container") + { + std::vector host_indices(4096); + for (index_type i = 0; i < host_indices.size(); ++i) { + host_indices[i] = 4095 - i; + } + thrust::device_vector indices(host_indices); + auto bitmap = bitmap_type::from_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.size() == 4096); + require_contains(bitmap, {0, 4095, 4096}, {true, true, false}); + } + + SECTION("bitset container") + { + std::vector host_indices(4097); + for (index_type i = 0; i < host_indices.size(); ++i) { + host_indices[i] = 4096 - i; + } + thrust::device_vector indices(host_indices); + auto bitmap = bitmap_type::from_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.size() == 4097); + require_contains(bitmap, {0, 4096, 4097}, {true, true, false}); + } +} + +TEST_CASE("roaring_bitmap writes a bitset after an odd-sized array container", "[roaring_bitmap]") +{ + std::vector host_indices; + host_indices.reserve(4098); + host_indices.push_back(1); + for (index_type i = 0; i < 4097; ++i) { + host_indices.push_back(0x00010000 + i); + } + std::reverse(host_indices.begin(), host_indices.end()); + + thrust::device_vector indices(host_indices); + auto bitmap = bitmap_type::from_indices(indices.begin(), indices.end()); + + auto const bytes = copy_serialized(bitmap); + cuda::std::uint32_t second_offset; + std::memcpy(&second_offset, bytes.data() + 20, sizeof(second_offset)); + REQUIRE(second_offset % alignof(cuda::std::uint64_t) == 2); + require_contains( + bitmap, {1, 2, 0x00010000, 0x00011000, 0x00011001}, {true, false, true, true, false}); +} + +TEST_CASE("roaring_bitmap accepts transformed input and removes duplicates", "[roaring_bitmap]") +{ + auto const first = cuda::make_transform_iterator( + cuda::counting_iterator{0}, + cuda::proclaim_return_type([] __device__(cuda::std::uint64_t index) { + return static_cast((31 - index) / 2); + })); + + auto bitmap = bitmap_type::from_indices(first, first + 32); + REQUIRE(bitmap.size() == 16); + require_contains(bitmap, {0, 15, 16}, {true, true, false}); +} + +TEST_CASE("roaring_bitmap factories produce identical serialized bytes", "[roaring_bitmap]") +{ + thrust::device_vector unordered{0x00010002, 7, 1, 0x00010000, 7, 3, 1}; + thrust::device_vector sorted{1, 1, 3, 7, 7, 0x00010000, 0x00010002}; + thrust::device_vector sorted_unique{1, 3, 7, 0x00010000, 0x00010002}; + auto const original = thrust::host_vector{unordered}; + + auto from_indices = bitmap_type::from_indices(unordered.begin(), unordered.end()); + auto from_sorted = bitmap_type::from_sorted_indices(sorted.begin(), sorted.end()); + auto from_sorted_unique = + bitmap_type::from_sorted_unique_indices(sorted_unique.begin(), sorted_unique.end()); + + REQUIRE(copy_serialized(from_indices) == copy_serialized(from_sorted)); + REQUIRE(copy_serialized(from_indices) == copy_serialized(from_sorted_unique)); + REQUIRE(thrust::host_vector{unordered} == original); +} + +TEST_CASE("roaring_bitmap accepts a transformed sorted unique range", "[roaring_bitmap]") +{ + auto const first = cuda::make_transform_iterator( + cuda::counting_iterator{0}, + cuda::proclaim_return_type( + [] __device__(cuda::std::uint64_t index) { return static_cast(2 * index); })); + + auto bitmap = bitmap_type::from_sorted_unique_indices(first, first + 16); + REQUIRE(bitmap.size() == 16); + require_contains(bitmap, {0, 2, 30, 31}, {true, true, true, false}); +} + +TEST_CASE("roaring_bitmap factories use the supplied allocator", "[roaring_bitmap]") +{ + SECTION("from_indices") + { + auto counts = std::make_shared(); + tracking_allocator allocator{counts}; + using tracked_bitmap = cuco::experimental::roaring_bitmap; + thrust::device_vector indices{9, 4, 9, 1, 7}; + + { + auto bitmap = tracked_bitmap::from_indices(indices.begin(), indices.end(), allocator); + REQUIRE(bitmap.size() == 4); + REQUIRE(bitmap.allocator() == allocator); + REQUIRE(counts->allocations == 6); + REQUIRE(counts->deallocations == 5); + } + REQUIRE(counts->deallocations == counts->allocations); + } + + SECTION("from_sorted_indices") + { + auto counts = std::make_shared(); + tracking_allocator allocator{counts}; + using tracked_bitmap = cuco::experimental::roaring_bitmap; + thrust::device_vector indices{1, 4, 7, 9, 9}; + + { + auto bitmap = tracked_bitmap::from_sorted_indices(indices.begin(), indices.end(), allocator); + REQUIRE(bitmap.size() == 4); + REQUIRE(bitmap.allocator() == allocator); + REQUIRE(counts->allocations == 6); + REQUIRE(counts->deallocations == 5); + } + REQUIRE(counts->deallocations == counts->allocations); + } + + SECTION("from_sorted_unique_indices") + { + auto counts = std::make_shared(); + tracking_allocator allocator{counts}; + using tracked_bitmap = cuco::experimental::roaring_bitmap; + thrust::device_vector indices{1, 4, 7, 9}; + + { + auto bitmap = + tracked_bitmap::from_sorted_unique_indices(indices.begin(), indices.end(), allocator); + REQUIRE(bitmap.size() == 4); + REQUIRE(bitmap.allocator() == allocator); + REQUIRE(counts->allocations == 5); + REQUIRE(counts->deallocations == 4); + } + REQUIRE(counts->deallocations == counts->allocations); + } +} + +TEST_CASE("roaring_bitmap build serialization is stream ordered", "[roaring_bitmap]") +{ + cudaStream_t stream; + CUCO_CUDA_TRY(cudaStreamCreate(&stream)); + + { + thrust::device_vector indices{5, 4, 3, 2, 1}; + auto bitmap = + bitmap_type::from_indices(indices.begin(), indices.end(), {}, cuda::stream_ref{stream}); + + thrust::device_vector queries{1, 5, 6}; + thrust::device_vector results(queries.size()); + bitmap.contains_async( + queries.begin(), queries.end(), results.begin(), cuda::stream_ref{stream}); + CUCO_CUDA_TRY(cudaStreamSynchronize(stream)); + + thrust::host_vector host_results = results; + REQUIRE(host_results[0]); + REQUIRE(host_results[1]); + REQUIRE_FALSE(host_results[2]); + } + + CUCO_CUDA_TRY(cudaStreamDestroy(stream)); +} + +TEST_CASE("roaring_bitmap build supports cross-stream event handoff", "[roaring_bitmap]") +{ + cudaStream_t build_stream; + cudaStream_t consume_stream; + cudaEvent_t ready; + CUCO_CUDA_TRY(cudaStreamCreate(&build_stream)); + CUCO_CUDA_TRY(cudaStreamCreate(&consume_stream)); + CUCO_CUDA_TRY(cudaEventCreateWithFlags(&ready, cudaEventDisableTiming)); + + { + thrust::device_vector indices{1, 3, 5, 7}; + auto bitmap = bitmap_type::from_sorted_unique_indices( + indices.begin(), indices.end(), {}, cuda::stream_ref{build_stream}); + + CUCO_CUDA_TRY(cudaEventRecord(ready, build_stream)); + CUCO_CUDA_TRY(cudaStreamWaitEvent(consume_stream, ready)); + + thrust::device_vector queries{1, 2, 7}; + thrust::device_vector results(queries.size()); + bitmap.contains_async( + queries.begin(), queries.end(), results.begin(), cuda::stream_ref{consume_stream}); + CUCO_CUDA_TRY(cudaStreamSynchronize(consume_stream)); + + thrust::host_vector host_results = results; + REQUIRE(host_results[0]); + REQUIRE_FALSE(host_results[1]); + REQUIRE(host_results[2]); + } + + CUCO_CUDA_TRY(cudaStreamSynchronize(build_stream)); + CUCO_CUDA_TRY(cudaEventDestroy(ready)); + CUCO_CUDA_TRY(cudaStreamDestroy(consume_stream)); + CUCO_CUDA_TRY(cudaStreamDestroy(build_stream)); +} diff --git a/tests/roaring_bitmap/contains_test.cu b/tests/roaring_bitmap/contains_test.cu index 642a1da2c..a9b992866 100644 --- a/tests/roaring_bitmap/contains_test.cu +++ b/tests/roaring_bitmap/contains_test.cu @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include #include #include @@ -20,9 +21,12 @@ #include #include #include +#include #include namespace { +using bitmap_type = cuco::experimental::roaring_bitmap; + template bool check(std::string const& bitmap_file_path) { @@ -124,19 +128,45 @@ TEST_CASE("roaring_bitmap run container without offsets", "[roaring_bitmap]") thrust::universal_host_pinned_vector buffer(bytes.size()); std::memcpy(thrust::raw_pointer_cast(buffer.data()), bytes.data(), bytes.size()); - cuco::experimental::roaring_bitmap roaring_bitmap( - thrust::raw_pointer_cast(buffer.data())); + bitmap_type original{thrust::raw_pointer_cast(buffer.data())}; + bitmap_type moved{std::move(original)}; thrust::device_vector keys{1, 2, 3, 4}; thrust::device_vector contained(keys.size(), false); - roaring_bitmap.contains(keys.begin(), keys.end(), contained.begin()); + moved.contains(keys.begin(), keys.end(), contained.begin()); thrust::host_vector contained_h = contained; REQUIRE(contained_h[0]); REQUIRE(contained_h[1]); REQUIRE(contained_h[2]); REQUIRE_FALSE(contained_h[3]); + + thrust::device_vector empty_indices; + auto assigned = + bitmap_type::from_sorted_unique_indices(empty_indices.begin(), empty_indices.end()); + assigned = std::move(moved); + assigned.contains(keys.begin(), keys.end(), contained.begin()); + + contained_h = contained; + REQUIRE(contained_h[0]); + REQUIRE(contained_h[1]); + REQUIRE(contained_h[2]); + REQUIRE_FALSE(contained_h[3]); +} + +TEST_CASE("roaring_bitmap parses a serialized empty bitmap", "[roaring_bitmap]") +{ + thrust::device_vector indices; + auto source = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + std::vector bytes(source.size_bytes()); + CUCO_CUDA_TRY(cudaMemcpy(bytes.data(), source.data(), bytes.size(), cudaMemcpyDeviceToHost)); + + auto bitmap = bitmap_type::from_serialized(bytes.data()); + REQUIRE(bitmap.empty()); + REQUIRE(bitmap.size() == 0); + REQUIRE(bitmap.size_bytes() == bytes.size()); } TEST_CASE("roaring_bitmap bulk contains from RoaringFormatSpec testdata", "[roaring_bitmap]")