-
Notifications
You must be signed in to change notification settings - Fork 118
Add GPU construction for 32-bit Roaring bitmaps from raw indices #839
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: dev
Are you sure you want to change the base?
Changes from 7 commits
246777b
fc82b7a
9934995
e0048a9
e5ef1e3
63ae4dd
9d7c930
d4b9416
2027216
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,145 @@ | ||
| /* | ||
| * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. | ||
| * SPDX-License-Identifier: Apache-2.0 | ||
| */ | ||
|
|
||
| #include <benchmark_utils.hpp> | ||
|
|
||
| #include <cuco/roaring_bitmap.cuh> | ||
| #include <cuco/utility/key_generator.cuh> | ||
|
|
||
| #include <nvbench/nvbench.cuh> | ||
|
|
||
| #include <cuda/std/cstdint> | ||
| #include <thrust/device_vector.h> | ||
| #include <thrust/reverse.h> | ||
| #include <thrust/sequence.h> | ||
| #include <thrust/sort.h> | ||
| #include <thrust/tabulate.h> | ||
|
|
||
| using namespace cuco::benchmark; | ||
| using namespace cuco::utility; | ||
|
|
||
| enum class build_mode { indices, sorted_indices, sorted_unique_indices }; | ||
|
|
||
| template <build_mode Mode, class Dist> | ||
| void roaring_bitmap_build(nvbench::state& state, nvbench::type_list<Dist>) | ||
| { | ||
| using index_type = cuda::std::uint32_t; | ||
| using bitmap_type = cuco::experimental::roaring_bitmap<index_type>; | ||
|
|
||
| auto const num_inputs = state.get_int64("NumInputs"); | ||
| thrust::device_vector<index_type> 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<Dist>(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<index_type>(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 <class Dist> | ||
| void roaring_bitmap_from_indices(nvbench::state& state, nvbench::type_list<Dist> types) | ||
| { | ||
| roaring_bitmap_build<build_mode::indices>(state, types); | ||
| } | ||
|
|
||
| template <class Dist> | ||
| void roaring_bitmap_from_sorted_indices(nvbench::state& state, nvbench::type_list<Dist> types) | ||
| { | ||
| roaring_bitmap_build<build_mode::sorted_indices>(state, types); | ||
| } | ||
|
|
||
| template <class Dist> | ||
| void roaring_bitmap_from_sorted_unique_indices(nvbench::state& state, | ||
| nvbench::type_list<Dist> types) | ||
| { | ||
| roaring_bitmap_build<build_mode::sorted_unique_indices>(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<index_type>; | ||
|
|
||
| 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<index_type> indices(num_inputs); | ||
|
|
||
| thrust::tabulate( | ||
| indices.begin(), indices.end(), [cardinality] __device__(cuda::std::int64_t index) { | ||
| auto const container = static_cast<index_type>(index / cardinality); | ||
| auto const lower = static_cast<index_type>(index % cardinality); | ||
| return (container << 16) | lower; | ||
| }); | ||
| thrust::reverse(indices.begin(), indices.end()); | ||
|
|
||
| state.add_element_count(num_inputs); | ||
| state.add_global_memory_reads<index_type>(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<distribution::unique>)) | ||
| .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<distribution::uniform>)) | ||
| .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<distribution::uniform>)) | ||
| .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<distribution::unique>)) | ||
| .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}); | ||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,40 @@ | ||
| /* | ||
| * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. | ||
| * SPDX-License-Identifier: Apache-2.0 | ||
| */ | ||
|
|
||
| #include <cuco/roaring_bitmap.cuh> | ||
|
|
||
| #include <cuda/std/cstdint> | ||
| #include <thrust/device_vector.h> | ||
| #include <thrust/host_vector.h> | ||
|
|
||
| #include <iostream> | ||
|
|
||
| /** | ||
| * @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<index_type> indices{0x00010002, 7, 1, 0x00010000, 7, 3, 0x00010002}; | ||
|
|
||
| auto bitmap = | ||
| cuco::experimental::roaring_bitmap<index_type>::from_indices(indices.begin(), indices.end()); | ||
|
|
||
| thrust::device_vector<index_type> queries{1, 2, 3, 7, 0x00010000, 0x00010001, 0x00010002}; | ||
| thrust::device_vector<bool> results(queries.size()); | ||
| bitmap.contains(queries.begin(), queries.end(), results.begin()); | ||
|
|
||
| thrust::host_vector<bool> expected{true, false, true, true, true, false, true}; | ||
| thrust::host_vector<bool> 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; | ||
| } |
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -5,7 +5,12 @@ | |
|
|
||
| #pragma once | ||
|
|
||
| #include <cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh> | ||
|
|
||
| #include <cuda/std/cstddef> | ||
| #include <cuda/std/cstdint> | ||
| #include <cuda/std/type_traits> | ||
| #include <cuda/std/utility> | ||
| #include <cuda/stream_ref> | ||
|
|
||
| namespace cuco::experimental { | ||
|
|
@@ -18,6 +23,62 @@ roaring_bitmap<T, Allocator>::roaring_bitmap(cuda::std::byte const* bitmap, | |
| { | ||
| } | ||
|
|
||
| template <class T, class Allocator> | ||
| roaring_bitmap<T, Allocator> roaring_bitmap<T, Allocator>::from_serialized( | ||
| cuda::std::byte const* bitmap, Allocator const& alloc, cuda::stream_ref stream) | ||
| { | ||
| static_assert(cuda::std::is_same_v<T, cuda::std::uint32_t>, | ||
| "roaring_bitmap::from_serialized currently supports only uint32_t"); | ||
| return roaring_bitmap{bitmap, alloc, stream}; | ||
| } | ||
|
|
||
| template <class T, class Allocator> | ||
| roaring_bitmap<T, Allocator>::roaring_bitmap(storage_type&& storage) | ||
| : storage_{cuda::std::move(storage)} | ||
| { | ||
| } | ||
|
|
||
| template <class T, class Allocator> | ||
|
Member
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. The three factory bodies are identical apart from the enumerator; a
Collaborator
Author
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. I would keep the three named factories. The input ordering guarantee is part of the API, and a public enum or tag makes that precondition easier to misuse. The small wrappers are intentional; the detail builder already centralizes the actual implementation, so I do not think another free helper buys enough to justify another layer. Also, adding another specialized enum to the public API feels unnecessary. |
||
| template <class InputIt> | ||
| roaring_bitmap<T, Allocator> roaring_bitmap<T, Allocator>::from_indices(InputIt first, | ||
| InputIt last, | ||
| Allocator const& alloc, | ||
| cuda::stream_ref stream) | ||
| { | ||
| static_assert(cuda::std::is_same_v<T, cuda::std::uint32_t>, | ||
| "Building a roaring_bitmap from indices currently supports only uint32_t"); | ||
| detail::roaring_bitmap_builder<InputIt, Allocator> 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 <class T, class Allocator> | ||
| template <class InputIt> | ||
| roaring_bitmap<T, Allocator> roaring_bitmap<T, Allocator>::from_sorted_indices( | ||
| InputIt first, InputIt last, Allocator const& alloc, cuda::stream_ref stream) | ||
| { | ||
| static_assert(cuda::std::is_same_v<T, cuda::std::uint32_t>, | ||
| "Building a roaring_bitmap from indices currently supports only uint32_t"); | ||
| detail::roaring_bitmap_builder<InputIt, Allocator> 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 <class T, class Allocator> | ||
| template <class InputIt> | ||
| roaring_bitmap<T, Allocator> roaring_bitmap<T, Allocator>::from_sorted_unique_indices( | ||
| InputIt first, InputIt last, Allocator const& alloc, cuda::stream_ref stream) | ||
| { | ||
| static_assert(cuda::std::is_same_v<T, cuda::std::uint32_t>, | ||
| "Building a roaring_bitmap from indices currently supports only uint32_t"); | ||
| detail::roaring_bitmap_builder<InputIt, Allocator> 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 <class T, class Allocator> | ||
| template <class InputIt, class OutputIt> | ||
| void roaring_bitmap<T, Allocator>::contains(InputIt first, | ||
|
|
@@ -74,4 +135,4 @@ typename roaring_bitmap<T, Allocator>::ref_type roaring_bitmap<T, Allocator>::re | |
| { | ||
| return ref_type{storage_.ref()}; | ||
| } | ||
| } // namespace cuco::experimental | ||
| } // namespace cuco::experimental | ||
Uh oh!
There was an error while loading. Please reload this page.