From 246777ba514f5a177720515d8c30edc3c8a5590d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 28 Aug 2026 16:55:49 -0700 Subject: [PATCH 1/8] Prepare Roaring bitmap storage for construction --- .../detail/roaring_bitmap/roaring_bitmap.inl | 11 ++ .../roaring_bitmap/roaring_bitmap_storage.cuh | 85 ++++++++++--- include/cuco/detail/roaring_bitmap/util.cuh | 115 +++++++++++++++--- include/cuco/roaring_bitmap.cuh | 27 ++++ 4 files changed, 200 insertions(+), 38 deletions(-) diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl index 61326fdd3..cdc5aef1f 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl @@ -6,6 +6,8 @@ #pragma once #include +#include +#include #include namespace cuco::experimental { @@ -18,6 +20,15 @@ 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 template void roaring_bitmap::contains(InputIt first, diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh index 0d8e465a3..7b839ae01 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,52 @@ 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); + } + + // For small run-container bitmaps, ref_ points into metadata_.computed_offsets. Rebuilding the + // reference after a move keeps that pointer attached to this object's metadata. + 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_} + { + } + + 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 +306,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 +391,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 +424,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..136e172ff 100644 --- a/include/cuco/roaring_bitmap.cuh +++ b/include/cuco/roaring_bitmap.cuh @@ -41,6 +41,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 +56,26 @@ 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 Copy constructor * From fc82b7a1f68ff72b98364530abac0e38574a16da Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 28 Aug 2026 16:56:20 -0700 Subject: [PATCH 2/8] Add GPU construction for 32-bit Roaring bitmaps --- .../detail/roaring_bitmap/roaring_bitmap.inl | 52 +- .../roaring_bitmap/roaring_bitmap_builder.cuh | 449 ++++++++++++++++++ .../roaring_bitmap/roaring_bitmap_kernels.cuh | 260 ++++++++++ include/cuco/roaring_bitmap.cuh | 113 ++++- tests/CMakeLists.txt | 3 +- tests/roaring_bitmap/build_test.cu | 449 ++++++++++++++++++ tests/roaring_bitmap/contains_test.cu | 36 +- 7 files changed, 1340 insertions(+), 22 deletions(-) create mode 100644 include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh create mode 100644 include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh create mode 100644 tests/roaring_bitmap/build_test.cu diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl index cdc5aef1f..ee41f099c 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl @@ -5,9 +5,12 @@ #pragma once +#include + #include #include #include +#include #include namespace cuco::experimental { @@ -29,6 +32,53 @@ roaring_bitmap roaring_bitmap::from_serialized( 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, @@ -85,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..9d3166821 --- /dev/null +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh @@ -0,0 +1,449 @@ +/* + * 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; + + /** + * @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..a67b208f5 --- /dev/null +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh @@ -0,0 +1,260 @@ +/* + * 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 { + +struct roaring_bitmap_build_state { + cuda::std::int64_t num_keys; + cuda::std::int64_t num_containers; + cuda::std::uint32_t size_bytes; + cuda::std::uint32_t num_array_containers; + cuda::std::uint32_t num_bitset_containers; +}; + +template +struct is_container_start { + KeyIt keys; + roaring_bitmap_build_state const* state; + + __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; + +struct container_payload_size { + cuda::std::int64_t const* container_starts; + roaring_bitmap_build_state const* state; + + __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 warp_size = 32; + constexpr cuda::std::uint32_t warps_per_block = BlockSize / 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 % 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 / warp_size; + if (warp_index >= state.num_array_containers) { return; } + + auto const lane = static_cast(threadIdx.x) % 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 += 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/roaring_bitmap.cuh b/include/cuco/roaring_bitmap.cuh index 136e172ff..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 @@ -76,6 +77,82 @@ class roaring_bitmap { 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 * @@ -109,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 */ @@ -131,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 */ @@ -150,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` */ @@ -196,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]") From 99349953a8422a2c2c39b803e9312b249445474f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 28 Aug 2026 16:56:46 -0700 Subject: [PATCH 3/8] Add Roaring bitmap build examples and benchmarks --- README.md | 3 +- benchmarks/CMakeLists.txt | 3 +- benchmarks/roaring_bitmap/build_bench.cu | 145 +++++++++++++++++++++++ examples/CMakeLists.txt | 1 + examples/roaring_bitmap/build_example.cu | 36 ++++++ 5 files changed, 186 insertions(+), 2 deletions(-) create mode 100644 benchmarks/roaring_bitmap/build_bench.cu create mode 100644 examples/roaring_bitmap/build_example.cu diff --git a/README.md b/README.md index fa5885281..610b8896b 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 +- [Build from unordered indices](https://github.com/NVIDIA/cuCollections/blob/dev/examples/roaring_bitmap/build_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyNVA1r2zAQ_SuHB2vS2o7jsgXUJCO0KwRKW9oyBvMIiqzGYrbs6SOkC_nvO38lTtexJUSO753evXc-a-torrXIpXbIt60jYocMXSelcmXpijvEYTamjuvo3CpW3g9OIwmn8Hh_9dW7Fim_zIsXJVaJeeIbQ2B_Cz3WhzAIP7pw-2V-NZ_B5d3D_d3D7Gl-dwvvYXZ9Pb-Zz54-P_owS1OoNmlQXHO15rF_qHIjGJeae_OYSyOeBVcEZgVlCfdCPyjzBpGM5DshWWpjDmNmWT5QOVVCrhZLYTJa-Mwm09dZMR1oEw8YLkKaaRc0ibLaDGK-xuKLNWcmV37yVkqSa9NN6KYIxBSnWRXGEpBRIXv9SG5RNViN-kDImG8W5qXgMIFSEyGohxCL-efhwlyUewHqaoQcKRofNk9LIkT0NtgEQTDEX-jCyIWhC_tIUEXOO5Fw1_BTa3KoewWTMgBQtpEQvim4Ehm2nqaEHHe1W5-QZ5Vni0ZFr7n6S74qLbutPJ_LuNfv_7ernxaroyu0EVbSR8d-9v-Hr139jX2Z5-m0nDObGt1r-H0tfvFaF0A7Mjl6FvKQs_fSBiovbsvV4q_NdSakKV72lBkeb42y3IVnmmq81DfdtQMcO_qTkjJjaYoj1IipjSAEaEMb0JZh9zUcfyb7fZO9qEZ9NYQstwbGY4gcKwW6bp8iwUgZbzpVN68MnESRPLl4Yz--1YKmmBfD8sW8ybCogH_x1D7a7RVa2qRpkdAq0hjtkJQ0ihur5B79BAEQGCK4iyQebyzPCjzM1OEUdOSasWH4wQ4RzgtTH5GOhxUn7OxsOAKPKpZMdLYYBeB52D6Di8GZ4LGX0mxZnZupWHY4GWMpBtdYCPkwgI9W_nB2bovjK3eE4-g6u-_V9zckk-PF)) +- [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==)) 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..268b64a2c 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -41,3 +41,4 @@ ConfigureExample(HYPERLOGLOG_DEVICE_REF_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyp 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_BUILD_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/roaring_bitmap/build_example.cu") diff --git a/examples/roaring_bitmap/build_example.cu b/examples/roaring_bitmap/build_example.cu new file mode 100644 index 000000000..74874dff5 --- /dev/null +++ b/examples/roaring_bitmap/build_example.cu @@ -0,0 +1,36 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include +#include +#include + +#include + +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; +} From e5ef1e3e6cdf8c678bec6acf307de0e86aec7952 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 28 Aug 2026 17:11:50 -0700 Subject: [PATCH 4/8] Rename Roaring bitmap host-bulk examples --- README.md | 4 ++-- examples/CMakeLists.txt | 6 ++++-- .../{build_example.cu => host_bulk_from_indices_example.cu} | 4 ++++ ...bulk_example.cu => host_bulk_from_serialized_example.cu} | 2 +- 4 files changed, 11 insertions(+), 5 deletions(-) rename examples/roaring_bitmap/{build_example.cu => host_bulk_from_indices_example.cu} (90%) rename examples/roaring_bitmap/{host_bulk_example.cu => host_bulk_from_serialized_example.cu} (99%) diff --git a/README.md b/README.md index 58b6736b6..29b1c2ff8 100644 --- a/README.md +++ b/README.md @@ -267,5 +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: -- [Build from unordered indices](https://github.com/NVIDIA/cuCollections/blob/dev/examples/roaring_bitmap/build_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJyNVA1r2zAQ_SuHB2vS2o7jsgXUJCO0KwRKW9oyBvMIiqzGYrbs6SOkC_nvO38lTtexJUSO753evXc-a-torrXIpXbIt60jYocMXSelcmXpijvEYTamjuvo3CpW3g9OIwmn8Hh_9dW7Fim_zIsXJVaJeeIbQ2B_Cz3WhzAIP7pw-2V-NZ_B5d3D_d3D7Gl-dwvvYXZ9Pb-Zz54-P_owS1OoNmlQXHO15rF_qHIjGJeae_OYSyOeBVcEZgVlCfdCPyjzBpGM5DshWWpjDmNmWT5QOVVCrhZLYTJa-Mwm09dZMR1oEw8YLkKaaRc0ibLaDGK-xuKLNWcmV37yVkqSa9NN6KYIxBSnWRXGEpBRIXv9SG5RNViN-kDImG8W5qXgMIFSEyGohxCL-efhwlyUewHqaoQcKRofNk9LIkT0NtgEQTDEX-jCyIWhC_tIUEXOO5Fw1_BTa3KoewWTMgBQtpEQvim4Ehm2nqaEHHe1W5-QZ5Vni0ZFr7n6S74qLbutPJ_LuNfv_7ernxaroyu0EVbSR8d-9v-Hr139jX2Z5-m0nDObGt1r-H0tfvFaF0A7Mjl6FvKQs_fSBiovbsvV4q_NdSakKV72lBkeb42y3IVnmmq81DfdtQMcO_qTkjJjaYoj1IipjSAEaEMb0JZh9zUcfyb7fZO9qEZ9NYQstwbGY4gcKwW6bp8iwUgZbzpVN68MnESRPLl4Yz--1YKmmBfD8sW8ybCogH_x1D7a7RVa2qRpkdAq0hjtkJQ0ihur5B79BAEQGCK4iyQebyzPCjzM1OEUdOSasWH4wQ4RzgtTH5GOhxUn7OxsOAKPKpZMdLYYBeB52D6Di8GZ4LGX0mxZnZupWHY4GWMpBtdYCPkwgI9W_nB2bovjK3eE4-g6u-_V9zckk-PF)) -- [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==)) +- [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/eJy9WAtP20gQ_itzRqocSOKQ8mjD45qSUkXXg4pyDwkqa2NvklUcr7u7BnKI_36zu36CKfRxFyRi787O85vZmdw6kkrJeCydwcWtw0JnsNl2IhLPUjKjzsAJ0pA4bUfyVAT63Vu_jGEdPn0c_d05ZhE94slKsNlcndMbNYDiFdygBf1ef7uD_3bacPLneDQewtHp2cfTs-H5-PQEXsDw-Hj8YTw8f_epC8MoAnNSgqCSiisadktRH1hAY0k745DGik0ZFQMYJiSY006_29N03mV8Ga-xOIjSkMJ-kAbcE5wIFs_8CVNLknSDdH74gCZVLGJq5SlBmJLdeZIc3ucUEk-q0AvwX0inh49uslg1b6pVQn0roEag5iKVygvpFZrnX9FAcdGdN5FEfMYCEjVvpjG7okKSqMqiSjfFQMmVVHRZOz6VSlBSX2O8YRGX0I21JSvJyPHWLSbeaDEwRwb-JI0W_lTwpY-RZCRi_9DQpzdkmUQUw2DJJ4LRKYzoEtGHzlFUQioRdcCnoOYU6tGDS0dzvXQg4nyRJkYQDD-OpYGJYTmO8SCTkEmCa4rEJAR1zeFlv4OMwDKTQOIQeExhZ6uyDG7ChSITPDrlYklUC7QRWhvD_-LMqvTWUB8bkk8JDT67c6USOfC8GVPzdNIN-NKr0eZv5ZkWwjzhkqEXV1obIwDxHCyAWfu1ubmxaCfSq1TE0uwFXAiMgE6VNELYwglZ0mjV1iajI5UhmvIo4tcoFQwABkZEBy6ssdeoKk-VSGPZnbDYbLrWS63vMcibRHzibW_u7pLwtaeVCIkiXqOw1kNV_j89HiiRBz1D2s5WqYeFx0_UY2fLaxTXKkD8JuaKwnkVx4J-SRmG2oZ-SRaYI4nCqg2d0dEfR6f-6PSvkw-nw5F_djo8G5-897Gkno-G58MDrLOKw4SCpKpIFKyVWAuSCHMOiwgWpxjRA7_R1Tk-Y05POI8sFl3UejCw-Y-gw0R9keWKrzHlJ0TNUfVb5AokRUkzGlOdy_6CriQcwMVntwWdQ7ClajCo1br9XCQYBqCRb4TQm0TgDYL1UwtHDZj0JeroX-VH2lDZTbHyvuz76rCVMwLwPDjCQoYWfkkpppjRB7PaplYBhltE5V2OiDJPchY_DX5n74aj3991l-GaXurotVyMMSFzSJNRRvW9nFqb4DaQwQK93dvDr33Y7OmPft44MC8Vv4Bh101SOfcnBCO8aBW872pCkLFhWnLbxxvdPm9sLL7G8yXi-Em-uxW-r57Ht4mnLYuPAKzJn65hOqEzFrutthVB49Bt5czvgEaS_ggYd7a-AYyN9cBA8b9AIhagEosPRNvC9xQsjXlPw1KTWVjemPCmUWSjje-v8_fviPhTsoZ1WZu9H5B1j6R308-4lbSNJNs1kkc1ZjZp2T1FmU7cvmbxFXV7N6-yA7AB7Aczw4T0uZlxW2KEKBb4RGKbp1zdT2sxeJ_oNt2fEiQuanwbO7jsGZaokL6UiuKF_sm1uHRKUzILbu9y-fpLv-hvTIpTFGUyRb_b3MyaWrPq3r-q2hkRl4MBQp2IlZWFqe7-oqm6mNscmaLJuZXmRECFgP19NOGYIFmoL1VNhwt6-b4cvWbOoSeiTPfMFuMUs3RXmPGeKqMvSGyUi6vUcNMriBHDrGzj7bPZfGBjq3QPzlQ80FXH9MpL7LSx9tTvwMJ7OUrKYcJ08gmLY-zdH4JmslJ4b0_S6ZQKt9CmIvyMYudtrMLA8ky63jN-xhiFrqC4RUWCrvEDItV-MCdi_dDNdRHk2k-4oTH7rhXX1XUMI4QQrQnOWAcRl9StaJKV3myeKAzPOvvCAxa-WO1xZFkifkmEGtRmkLJfqa-7NsLPVXuvEnjbLVWvBZdgQkvslqNQpwheQYowjAG6sRK3VgGTrNGqNV4N1hOwMdTIlfhN7Yxl54cqAOplQreCh6UStkIY3GnvayxXRL3FiaWTTSwVi8gMz2YzidVe098fzq0I-WgNapda5NtVK_PJSXtPC7KSxaMO1OdMn4sH_JLmoPCDXudT94HQqh6ZYlgfKt2B-YlCrW7vLCazvK-JwQ2T_boVWhLN1fTQa2wa4khsWvq8k9ddvD8anxUVrtqMg4aUHzKBejee2ivMlGkQUCmh_jmwbX5zn1Qw38Ay1zzI5cU65_7iuxh-Kzd7WdW5NQ9xTg4RW8VRc1vFMwmDrH7bioZuIlEyJ2YlU-FeIa8ENKf4FXowgE3cXNOXYymsuDKaZ7nGeGXDGqLAoBErB8MyoH1Un_sunbItyD6XzjOGwdb9gw8uqsw6a08csqmBqtN2sONMsFKK8pdCJ74Kgs3-drqJ21Yx3HQ6yPAg2NjY3IUOEcH8QC793R50OlhaFf5Tuj0IOxFZTsxvixGbVHgGQRDhor6DkB8uaKwtnLt2vo9VuraP9cq5-2z-_gWO0fmv)) \ No newline at end of file diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 268b64a2c..8a1457e65 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -40,5 +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_BUILD_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/roaring_bitmap/build_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/build_example.cu b/examples/roaring_bitmap/host_bulk_from_indices_example.cu similarity index 90% rename from examples/roaring_bitmap/build_example.cu rename to examples/roaring_bitmap/host_bulk_from_indices_example.cu index 74874dff5..39ea136e5 100644 --- a/examples/roaring_bitmap/build_example.cu +++ b/examples/roaring_bitmap/host_bulk_from_indices_example.cu @@ -11,6 +11,10 @@ #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; diff --git a/examples/roaring_bitmap/host_bulk_example.cu b/examples/roaring_bitmap/host_bulk_from_serialized_example.cu similarity index 99% rename from examples/roaring_bitmap/host_bulk_example.cu rename to examples/roaring_bitmap/host_bulk_from_serialized_example.cu index d64986c61..e5273f2bb 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 From 63ae4ddef53cc75ed056c6ff8e8692a909d96e97 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 28 Aug 2026 17:13:02 -0700 Subject: [PATCH 5/8] Document Roaring bitmap build internals --- .../roaring_bitmap/roaring_bitmap_builder.cuh | 3 +- .../roaring_bitmap/roaring_bitmap_kernels.cuh | 56 +++++++++++++------ .../roaring_bitmap/roaring_bitmap_storage.cuh | 19 ++++++- 3 files changed, 58 insertions(+), 20 deletions(-) diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh index 9d3166821..8b489bcc0 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh @@ -48,7 +48,8 @@ class roaring_bitmap_builder { "roaring_bitmap factories require an input iterator with uint32_t value_type"); public: - using storage_type = roaring_bitmap_storage; + using storage_type = + roaring_bitmap_storage; ///< Generated bitmap storage type /** * @brief Prepares a build for the specified input range and ordering. diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh index a67b208f5..a044f3c51 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh @@ -14,19 +14,33 @@ 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; - cuda::std::int64_t num_containers; - cuda::std::uint32_t size_bytes; - cuda::std::uint32_t num_array_containers; - cuda::std::uint32_t num_bitset_containers; + 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; - roaring_bitmap_build_state const* state; - + 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; @@ -39,10 +53,19 @@ struct is_container_start { 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; - roaring_bitmap_build_state const* state; - + 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; @@ -196,12 +219,11 @@ CUCO_KERNEL void write_roaring_containers(cuda::std::byte* bitmap, { using metadata_type = roaring_bitmap_metadata; - constexpr cuda::std::uint32_t warp_size = 32; - constexpr cuda::std::uint32_t warps_per_block = BlockSize / warp_size; + 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 % warp_size == 0); + static_assert(BlockSize % cuco::detail::warp_size() == 0); static_assert(bitset_words % BlockSize == 0); auto const payload_begin = 2 * sizeof(cuda::std::uint32_t) + @@ -213,10 +235,10 @@ CUCO_KERNEL void write_roaring_containers(cuda::std::byte* bitmap, // 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 / warp_size; + 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) % warp_size; + 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 @@ -225,7 +247,7 @@ CUCO_KERNEL void write_roaring_containers(cuda::std::byte* bitmap, auto const cardinality = static_cast(end - begin); auto* const container = bitmap + payload_begin + payload_offsets[container_index]; - for (auto index = lane; index < cardinality; index += warp_size) { + 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); } diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh index 7b839ae01..5c2af5024 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh @@ -273,8 +273,14 @@ class roaring_bitmap_storage { assert(metadata_.valid); } - // For small run-container bitmaps, ref_ points into metadata_.computed_offsets. Rebuilding the - // reference after a move keeps that pointer attached to this object's metadata. + /** + * @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_)}, @@ -283,6 +289,15 @@ class roaring_bitmap_storage { { } + /** + * @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_); From 9d7c9307395c3b8795d93ad65d0751c98471dde6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 28 Aug 2026 17:29:29 -0700 Subject: [PATCH 6/8] Use from_serialized in Roaring example --- README.md | 2 +- .../host_bulk_from_serialized_example.cu | 10 ++++++++-- 2 files changed, 9 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 29b1c2ff8..1bb305a5c 100644 --- a/README.md +++ b/README.md @@ -268,4 +268,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection #### Examples: - [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/eJy9WAtP20gQ_itzRqocSOKQ8mjD45qSUkXXg4pyDwkqa2NvklUcr7u7BnKI_36zu36CKfRxFyRi787O85vZmdw6kkrJeCydwcWtw0JnsNl2IhLPUjKjzsAJ0pA4bUfyVAT63Vu_jGEdPn0c_d05ZhE94slKsNlcndMbNYDiFdygBf1ef7uD_3bacPLneDQewtHp2cfTs-H5-PQEXsDw-Hj8YTw8f_epC8MoAnNSgqCSiisadktRH1hAY0k745DGik0ZFQMYJiSY006_29N03mV8Ga-xOIjSkMJ-kAbcE5wIFs_8CVNLknSDdH74gCZVLGJq5SlBmJLdeZIc3ucUEk-q0AvwX0inh49uslg1b6pVQn0roEag5iKVygvpFZrnX9FAcdGdN5FEfMYCEjVvpjG7okKSqMqiSjfFQMmVVHRZOz6VSlBSX2O8YRGX0I21JSvJyPHWLSbeaDEwRwb-JI0W_lTwpY-RZCRi_9DQpzdkmUQUw2DJJ4LRKYzoEtGHzlFUQioRdcCnoOYU6tGDS0dzvXQg4nyRJkYQDD-OpYGJYTmO8SCTkEmCa4rEJAR1zeFlv4OMwDKTQOIQeExhZ6uyDG7ChSITPDrlYklUC7QRWhvD_-LMqvTWUB8bkk8JDT67c6USOfC8GVPzdNIN-NKr0eZv5ZkWwjzhkqEXV1obIwDxHCyAWfu1ubmxaCfSq1TE0uwFXAiMgE6VNELYwglZ0mjV1iajI5UhmvIo4tcoFQwABkZEBy6ssdeoKk-VSGPZnbDYbLrWS63vMcibRHzibW_u7pLwtaeVCIkiXqOw1kNV_j89HiiRBz1D2s5WqYeFx0_UY2fLaxTXKkD8JuaKwnkVx4J-SRmG2oZ-SRaYI4nCqg2d0dEfR6f-6PSvkw-nw5F_djo8G5-897Gkno-G58MDrLOKw4SCpKpIFKyVWAuSCHMOiwgWpxjRA7_R1Tk-Y05POI8sFl3UejCw-Y-gw0R9keWKrzHlJ0TNUfVb5AokRUkzGlOdy_6CriQcwMVntwWdQ7ClajCo1br9XCQYBqCRb4TQm0TgDYL1UwtHDZj0JeroX-VH2lDZTbHyvuz76rCVMwLwPDjCQoYWfkkpppjRB7PaplYBhltE5V2OiDJPchY_DX5n74aj3991l-GaXurotVyMMSFzSJNRRvW9nFqb4DaQwQK93dvDr33Y7OmPft44MC8Vv4Bh101SOfcnBCO8aBW872pCkLFhWnLbxxvdPm9sLL7G8yXi-Em-uxW-r57Ht4mnLYuPAKzJn65hOqEzFrutthVB49Bt5czvgEaS_ggYd7a-AYyN9cBA8b9AIhagEosPRNvC9xQsjXlPw1KTWVjemPCmUWSjje-v8_fviPhTsoZ1WZu9H5B1j6R308-4lbSNJNs1kkc1ZjZp2T1FmU7cvmbxFXV7N6-yA7AB7Aczw4T0uZlxW2KEKBb4RGKbp1zdT2sxeJ_oNt2fEiQuanwbO7jsGZaokL6UiuKF_sm1uHRKUzILbu9y-fpLv-hvTIpTFGUyRb_b3MyaWrPq3r-q2hkRl4MBQp2IlZWFqe7-oqm6mNscmaLJuZXmRECFgP19NOGYIFmoL1VNhwt6-b4cvWbOoSeiTPfMFuMUs3RXmPGeKqMvSGyUi6vUcNMriBHDrGzj7bPZfGBjq3QPzlQ80FXH9MpL7LSx9tTvwMJ7OUrKYcJ08gmLY-zdH4JmslJ4b0_S6ZQKt9CmIvyMYudtrMLA8ky63jN-xhiFrqC4RUWCrvEDItV-MCdi_dDNdRHk2k-4oTH7rhXX1XUMI4QQrQnOWAcRl9StaJKV3myeKAzPOvvCAxa-WO1xZFkifkmEGtRmkLJfqa-7NsLPVXuvEnjbLVWvBZdgQkvslqNQpwheQYowjAG6sRK3VgGTrNGqNV4N1hOwMdTIlfhN7Yxl54cqAOplQreCh6UStkIY3GnvayxXRL3FiaWTTSwVi8gMz2YzidVe098fzq0I-WgNapda5NtVK_PJSXtPC7KSxaMO1OdMn4sH_JLmoPCDXudT94HQqh6ZYlgfKt2B-YlCrW7vLCazvK-JwQ2T_boVWhLN1fTQa2wa4khsWvq8k9ddvD8anxUVrtqMg4aUHzKBejee2ivMlGkQUCmh_jmwbX5zn1Qw38Ay1zzI5cU65_7iuxh-Kzd7WdW5NQ9xTg4RW8VRc1vFMwmDrH7bioZuIlEyJ2YlU-FeIa8ENKf4FXowgE3cXNOXYymsuDKaZ7nGeGXDGqLAoBErB8MyoH1Un_sunbItyD6XzjOGwdb9gw8uqsw6a08csqmBqtN2sONMsFKK8pdCJ74Kgs3-drqJ21Yx3HQ6yPAg2NjY3IUOEcH8QC793R50OlhaFf5Tuj0IOxFZTsxvixGbVHgGQRDhor6DkB8uaKwtnLt2vo9VuraP9cq5-2z-_gWO0fmv)) \ No newline at end of file +- [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/examples/roaring_bitmap/host_bulk_from_serialized_example.cu b/examples/roaring_bitmap/host_bulk_from_serialized_example.cu index e5273f2bb..46b7bdce3 100644 --- a/examples/roaring_bitmap/host_bulk_from_serialized_example.cu +++ b/examples/roaring_bitmap/host_bulk_from_serialized_example.cu @@ -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(); From d4b941682d600193d76dd957d7b5c009b88c4145 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 4 Sep 2026 18:06:49 -0700 Subject: [PATCH 7/8] Address Roaring bitmap review feedback --- benchmarks/roaring_bitmap/build_bench.cu | 76 +++-- .../detail/roaring_bitmap/roaring_bitmap.inl | 4 +- .../roaring_bitmap/roaring_bitmap_builder.cuh | 123 +++---- .../roaring_bitmap/roaring_bitmap_impl.cuh | 2 +- .../roaring_bitmap/roaring_bitmap_kernels.cuh | 311 ++++++++++-------- .../roaring_bitmap/roaring_bitmap_storage.cuh | 54 +-- include/cuco/detail/roaring_bitmap/util.cuh | 61 +++- include/cuco/detail/utility/cuda.cuh | 2 +- include/cuco/roaring_bitmap.cuh | 19 +- include/cuco/roaring_bitmap_ref.cuh | 37 ++- tests/roaring_bitmap/build_test.cu | 97 +++++- tests/roaring_bitmap/contains_test.cu | 98 ++++-- tests/roaring_bitmap/test_utils.cuh | 29 ++ 13 files changed, 563 insertions(+), 350 deletions(-) create mode 100644 tests/roaring_bitmap/test_utils.cuh diff --git a/benchmarks/roaring_bitmap/build_bench.cu b/benchmarks/roaring_bitmap/build_bench.cu index b0660be1e..54847b373 100644 --- a/benchmarks/roaring_bitmap/build_bench.cu +++ b/benchmarks/roaring_bitmap/build_bench.cu @@ -22,11 +22,24 @@ using namespace cuco::utility; enum class build_mode { indices, sorted_indices, sorted_unique_indices }; +template +auto build_roaring_bitmap(InputIt first, InputIt last, cuda::stream_ref stream) +{ + using bitmap_type = cuco::experimental::roaring_bitmap; + + if constexpr (Mode == build_mode::indices) { + return bitmap_type::from_indices(first, last, {}, stream); + } else if constexpr (Mode == build_mode::sorted_indices) { + return bitmap_type::from_sorted_indices(first, last, {}, stream); + } else { + return bitmap_type::from_sorted_unique_indices(first, last, {}, stream); + } +} + 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; + using index_type = cuda::std::uint32_t; auto const num_inputs = state.get_int64("NumInputs"); thrust::device_vector indices(num_inputs); @@ -47,19 +60,9 @@ void roaring_bitmap_build(nvbench::state& state, nvbench::type_list) 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(); - } + [[maybe_unused]] auto bitmap = build_roaring_bitmap( + indices.begin(), indices.end(), cuda::stream_ref{launch.get_stream()}); + timer.stop(); }); } @@ -82,12 +85,12 @@ void roaring_bitmap_from_sorted_unique_indices(nvbench::state& state, roaring_bitmap_build(state, types); } -void roaring_bitmap_from_indices_array_containers(nvbench::state& state) +template +void roaring_bitmap_build_container_cardinality(nvbench::state& state) { - using index_type = cuda::std::uint32_t; - using bitmap_type = cuco::experimental::roaring_bitmap; + using index_type = cuda::std::uint32_t; - constexpr cuda::std::int64_t num_containers = 1 << 16; + constexpr cuda::std::int64_t num_containers = 4096; auto const cardinality = state.get_int64("ContainerCardinality"); auto const num_inputs = num_containers * cardinality; thrust::device_vector indices(num_inputs); @@ -98,7 +101,7 @@ void roaring_bitmap_from_indices_array_containers(nvbench::state& state) auto const lower = static_cast(index % cardinality); return (container << 16) | lower; }); - thrust::reverse(indices.begin(), indices.end()); + if constexpr (Mode == build_mode::indices) { thrust::reverse(indices.begin(), indices.end()); } state.add_element_count(num_inputs); state.add_global_memory_reads(num_inputs, "InputSize"); @@ -106,12 +109,27 @@ void roaring_bitmap_from_indices_array_containers(nvbench::state& state) 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()}); + [[maybe_unused]] auto bitmap = build_roaring_bitmap( + indices.begin(), indices.end(), cuda::stream_ref{launch.get_stream()}); timer.stop(); }); } +void roaring_bitmap_from_indices_container_cardinality(nvbench::state& state) +{ + roaring_bitmap_build_container_cardinality(state); +} + +void roaring_bitmap_from_sorted_indices_container_cardinality(nvbench::state& state) +{ + roaring_bitmap_build_container_cardinality(state); +} + +void roaring_bitmap_from_sorted_unique_indices_container_cardinality(nvbench::state& state) +{ + roaring_bitmap_build_container_cardinality(state); +} + NVBENCH_BENCH_TYPES(roaring_bitmap_from_indices, NVBENCH_TYPE_AXES(nvbench::type_list)) .set_name("roaring_bitmap_from_indices_unique") @@ -119,9 +137,17 @@ NVBENCH_BENCH_TYPES(roaring_bitmap_from_indices, .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(roaring_bitmap_from_indices_container_cardinality) + .set_name("roaring_bitmap_from_indices_container_cardinality") + .add_int64_axis("ContainerCardinality", {1, 8, 64, 512, 4096, 8192, 32768}); + +NVBENCH_BENCH(roaring_bitmap_from_sorted_indices_container_cardinality) + .set_name("roaring_bitmap_from_sorted_indices_container_cardinality") + .add_int64_axis("ContainerCardinality", {1, 8, 64, 512, 4096, 8192, 32768}); + +NVBENCH_BENCH(roaring_bitmap_from_sorted_unique_indices_container_cardinality) + .set_name("roaring_bitmap_from_sorted_unique_indices_container_cardinality") + .add_int64_axis("ContainerCardinality", {1, 8, 64, 512, 4096, 8192, 32768}); NVBENCH_BENCH_TYPES(roaring_bitmap_from_indices, NVBENCH_TYPE_AXES(nvbench::type_list)) diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl index ee41f099c..6ea4b4c04 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap.inl @@ -27,9 +27,7 @@ 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}; + return roaring_bitmap{storage_type{bitmap, alloc, stream}}; } template diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh index 8b489bcc0..64706c439 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_builder.cuh @@ -10,6 +10,7 @@ #include #include #include +#include #include #include @@ -30,6 +31,9 @@ namespace cuco::experimental::detail { +/** + * @brief Ordering and uniqueness guarantees for Roaring bitmap builder input. + */ enum class roaring_bitmap_builder_input_order { unsorted, sorted, sorted_unique }; /** @@ -65,17 +69,15 @@ class roaring_bitmap_builder { 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()} + : first_{first}, input_order_{input_order}, alloc_{alloc}, stream_{stream} { + using metadata_type = roaring_bitmap_metadata; + + num_indices_ = cuco::detail::distance(first, last); + CUCO_EXPECTS(num_indices_ >= 0, "Invalid input range"); + num_container_slots_ = std::min( + num_indices_, static_cast(metadata_type::max_num_containers)); + workspace_bytes_ = compute_workspace_bytes(); } /** @@ -124,19 +126,22 @@ class roaring_bitmap_builder { // 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 container_starts = static_cast(nullptr); auto* const payload_offsets = static_cast(nullptr); - auto* const num_selected = static_cast(nullptr); + auto* const num_selected = static_cast(nullptr); auto* const state = static_cast(nullptr); - auto const counting_begin = cuda::counting_iterator{0}; + auto const counting_begin = cuda::counting_iterator{0}; + auto const payload_sizes = cuda::make_transform_iterator( + counting_begin, container_payload_size{container_starts, state}); // 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 result = + static_cast(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_sizes, payload_offsets, num_container_slots_, stream_.get())); @@ -210,9 +215,10 @@ class roaring_bitmap_builder { 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 container_starts = allocate_temporary_buffer( + static_cast(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 @@ -229,10 +235,12 @@ class roaring_bitmap_builder { { 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_); + auto container_starts = allocate_temporary_buffer( + static_cast(num_container_slots_)); + auto payload_offsets = allocate_temporary_buffer( + static_cast(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(), @@ -244,16 +252,15 @@ class roaring_bitmap_builder { [[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())); + auto container_starts = allocate_temporary_buffer( + static_cast(num_container_slots_)); + auto payload_offsets = allocate_temporary_buffer( + static_cast(num_container_slots_)); + auto state = allocate_temporary_buffer(1); + auto workspace = allocate_temporary_buffer(workspace_bytes_); + + CUCO_CUDA_TRY(cuco::detail::memcpy_async( + &state->num_indices, &num_indices_, sizeof(num_indices_), cudaMemcpyHostToDevice, stream_)); return serialize_sorted_unique_indices( first_, container_starts.get(), payload_offsets.get(), state.get(), workspace.get()); } @@ -290,17 +297,18 @@ class roaring_bitmap_builder { workspace_bytes, first, unique_indices, - &state->num_keys, + &state->num_indices, 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 + [[nodiscard]] storage_type serialize_sorted_unique_indices( + SourceIt first, + cuco::detail::index_type* 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); @@ -311,12 +319,12 @@ class roaring_bitmap_builder { template void analyze_containers(SourceIt first, - cuda::std::int64_t* container_starts, + cuco::detail::index_type* 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 const counting_begin = cuda::counting_iterator{0}; auto workspace_bytes = workspace_bytes_; CUCO_CUDA_TRY(cub::DeviceSelect::If(workspace, workspace_bytes, @@ -327,15 +335,12 @@ class roaring_bitmap_builder { is_container_start{first, state}, stream_.get())); - compute_container_payload_sizes<<>>( - payload_offsets, num_container_slots_, container_starts, state); + auto const payload_sizes = cuda::make_transform_iterator( + counting_begin, container_payload_size{container_starts, state}); workspace_bytes = workspace_bytes_; CUCO_CUDA_TRY(cub::DeviceScan::ExclusiveSum(workspace, workspace_bytes, - payload_offsets, + payload_sizes, payload_offsets, num_container_slots_, stream_.get())); @@ -351,10 +356,7 @@ class roaring_bitmap_builder { cuco::detail::default_block_size(), 0, stream_.get()>>>( - container_indexes, num_container_slots_, container_starts, state); - - compute_roaring_bitmap_build_size<<<1, 1, 0, stream_.get()>>>( - state, container_starts, payload_offsets); + container_indexes, num_container_slots_, container_starts, payload_offsets, state); CUCO_CUDA_TRY(cudaPeekAtLastError()); } @@ -377,7 +379,7 @@ class roaring_bitmap_builder { 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_indices >= 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"); @@ -387,7 +389,7 @@ class roaring_bitmap_builder { template [[nodiscard]] storage_type write_serialized_bitmap(SourceIt first, roaring_bitmap_build_state const& host_state, - cuda::std::int64_t* container_starts, + cuco::detail::index_type* container_starts, cuda::std::uint32_t* payload_offsets, cuda::std::uint32_t* container_indexes) const { @@ -395,7 +397,7 @@ class roaring_bitmap_builder { 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_indices), static_cast(host_state.num_containers)), alloc_, stream_}; @@ -409,11 +411,11 @@ class roaring_bitmap_builder { 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 warps_per_block = block_size / cuco::detail::warp_size(); constexpr cuda::std::uint32_t bitset_blocks_per_container = - metadata_type::bitset_container_bytes / sizeof(cuda::std::uint64_t) / block_size; + metadata_type::bitset_container_words / block_size; auto const array_blocks = - (host_state.num_array_containers + warps_per_block - 1) / warps_per_block; + cuco::detail::int_div_ceil(host_state.num_array_containers, 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. @@ -435,16 +437,17 @@ class roaring_bitmap_builder { [[nodiscard]] static constexpr roaring_bitmap_build_state empty_build_state() noexcept { - return {0, 0, 2 * sizeof(cuda::std::uint32_t), 0, 0}; + using metadata_type = roaring_bitmap_metadata; + return {0, 0, metadata_type::no_run_header_bytes(0), 0, 0}; } InputIt first_; - cuda::std::int64_t num_indices_; - cuda::std::size_t num_container_slots_; + cuco::detail::index_type num_indices_{}; + cuco::detail::index_type num_container_slots_{}; roaring_bitmap_builder_input_order input_order_; Allocator alloc_; cuda::stream_ref stream_; - cuda::std::size_t workspace_bytes_; + cuda::std::size_t workspace_bytes_{}; }; } // namespace cuco::experimental::detail diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh index 7128f5aac..6466fa35f 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh @@ -98,7 +98,7 @@ class roaring_bitmap_impl { template __device__ bool dispatch_contains(cuda::std::uint32_t value) const { - cuda::std::uint16_t const upper = value >> 16; + cuda::std::uint16_t const upper = storage_ref_type::metadata_type::container_key(value); cuda::std::uint16_t const lower = value & 0xFFFF; cuda::std::uint16_t key; diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh index a044f3c51..a98866d20 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh @@ -8,17 +8,20 @@ #include #include +#include #include #include -#include +#include namespace cuco::experimental::detail { +CUCO_SUPPRESS_KERNEL_WARNINGS + /** * @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 + cuco::detail::index_type num_indices; ///< 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 @@ -28,11 +31,11 @@ struct roaring_bitmap_build_state { /** * @brief Predicate selecting the first index in each high-16-bit container. * - * @tparam KeyIt Random access iterator over sorted unique indices + * @tparam IndexIt Random access iterator over sorted unique indices */ -template +template struct is_container_start { - KeyIt keys; ///< Sorted unique indices + IndexIt indices; ///< Sorted unique indices roaring_bitmap_build_state const* state; ///< Device build state /** @@ -41,24 +44,52 @@ struct is_container_start { * @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 + __device__ bool operator()(cuco::detail::index_type index) const noexcept { - auto const num_keys = state->num_keys; - if (index >= num_keys) { return false; } + auto const num_indices = state->num_indices; + if (index >= num_indices) { return false; } if (index == 0) { return true; } - return (keys[index] >> 16) != (keys[index - 1] >> 16); + return roaring_bitmap_metadata::container_key(indices[index]) != + roaring_bitmap_metadata::container_key(indices[index - 1]); } }; -template -is_container_start(KeyIt, roaring_bitmap_build_state const*) -> is_container_start; +template +is_container_start(IndexIt, roaring_bitmap_build_state const*) -> is_container_start; + +/** + * @brief Half-open range of normalized indices belonging to one container. + */ +struct roaring_bitmap_container_bounds { + cuco::detail::index_type begin; ///< First index in the container + cuco::detail::index_type end; ///< One-past-the-end index in the container +}; + +/** + * @brief Returns the normalized input range belonging to a container. + * + * @param container Container index + * @param container_starts Starting normalized input index of each container + * @param state Device build state + * @return Half-open normalized input range + */ +[[nodiscard]] __device__ inline roaring_bitmap_container_bounds container_bounds( + cuco::detail::index_type container, + cuco::detail::index_type const* container_starts, + roaring_bitmap_build_state const* state) noexcept +{ + auto const begin = container_starts[container]; + auto const end = + container + 1 < state->num_containers ? container_starts[container + 1] : state->num_indices; + return {begin, end}; +} /** * @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 + cuco::detail::index_type const* container_starts; ///< Starting index of each container + roaring_bitmap_build_state const* state; ///< Device build state /** * @brief Returns the encoded payload size for one container slot. @@ -66,77 +97,43 @@ struct container_payload_size { * @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 + __device__ cuda::std::uint32_t operator()(cuco::detail::index_type 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; + auto const bounds = container_bounds(index, container_starts, state); + auto const cardinality = static_cast(bounds.end - bounds.begin); + return metadata_type::container_payload_bytes(cardinality); } }; -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) +/** + * @brief Partitions container indexes into array and bitset work queues. + * + * The final container also computes the exact serialized bitmap size. + * + * @param container_indexes Shared array/bitset work queue + * @param num_container_slots Number of allocated container slots + * @param container_starts Starting normalized input index of each container + * @param payload_offsets Exclusive payload offsets for every container + * @param state Device build state + */ +inline CUCO_KERNEL void collect_container_indexes(cuda::std::uint32_t* container_indexes, + cuco::detail::index_type num_container_slots, + cuco::detail::index_type const* container_starts, + cuda::std::uint32_t const* payload_offsets, + 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); + auto const bounds = container_bounds(index, container_starts, state); + auto const cardinality = static_cast(bounds.end - bounds.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. @@ -148,32 +145,56 @@ CUCO_KERNEL void collect_container_indexes(cuda::std::uint32_t* container_indexe container_indexes[num_container_slots - output_index - 1] = static_cast(index); } + + if (index == state->num_containers - 1) { + state->size_bytes = + metadata_type::no_run_header_bytes(static_cast(state->num_containers)) + + payload_offsets[index] + metadata_type::container_payload_bytes(cardinality); + } } -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) +/** + * @brief Finds the first normalized index whose low 16 bits are not less than `value`. + * + * @tparam IndexIt Random access iterator over sorted unique indices + * + * @param indices Sorted unique indices + * @param first Beginning of the search range + * @param last End of the search range + * @param value Low-16-bit value to locate + * @return Position of the first matching or greater value + */ +template +__device__ cuco::detail::index_type lower_bound_low_bits(IndexIt indices, + cuco::detail::index_type first, + cuco::detail::index_type 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; + auto const begin = indices + first; + auto const found = + cuda::std::lower_bound(begin, indices + last, value, [] __device__(auto index, auto lower) { + return static_cast(index) < lower; + }); + return first + cuda::std::distance(begin, found); } -template +/** + * @brief Writes the no-run Roaring header, descriptors, and container offsets. + * + * @tparam IndexIt Random access iterator over sorted unique indices + * + * @param bitmap Output serialized bitmap + * @param indices Sorted unique indices + * @param state Completed build state + * @param container_starts Starting normalized input index of each container + * @param payload_offsets Exclusive payload offsets for every container + */ +template CUCO_KERNEL void write_roaring_bitmap_header(cuda::std::byte* bitmap, - KeyIt keys, + IndexIt indices, roaring_bitmap_build_state state, - ContainerStartIt container_starts, - PayloadOffsetIt payload_offsets) + cuco::detail::index_type const* container_starts, + cuda::std::uint32_t const* payload_offsets) { using metadata_type = roaring_bitmap_metadata; @@ -187,48 +208,58 @@ CUCO_KERNEL void write_roaring_bitmap_header(cuda::std::byte* bitmap, 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 bounds = container_bounds(index, container_starts, &state); + auto const cardinality = static_cast(bounds.end - bounds.begin); + auto const container_key = metadata_type::container_key(indices[bounds.begin]); 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); + auto* const key_cards = bitmap + metadata_type::no_run_key_cards_offset; + misaligned_store(key_cards + index * 2 * sizeof(cuda::std::uint16_t), container_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]; + auto* const offsets = bitmap + metadata_type::no_run_container_offsets_offset( + static_cast(state.num_containers)); + auto const offset = static_cast( + metadata_type::no_run_header_bytes(static_cast(state.num_containers)) + + payload_offsets[index]); + misaligned_store(offsets + index * sizeof(cuda::std::uint32_t), offset); } -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) +/** + * @brief Writes array and bitset container payloads. + * + * @tparam BlockSize Number of threads per block + * @tparam IndexIt Random access iterator over sorted unique indices + * + * @param bitmap Output serialized bitmap + * @param indices Sorted unique indices + * @param state Completed build state + * @param container_starts Starting normalized input index of each container + * @param payload_offsets Exclusive payload offsets for every container + * @param array_containers Array-container work queue + * @param bitset_containers Bitset-container work queue + */ +template +CUCO_KERNEL __launch_bounds__(BlockSize) void write_roaring_containers( + cuda::std::byte* bitmap, + IndexIt indices, + roaring_bitmap_build_state state, + cuco::detail::index_type const* container_starts, + cuda::std::uint32_t const* payload_offsets, + cuda::std::uint32_t const* array_containers, + cuda::std::uint32_t const* bitset_containers) { - using metadata_type = roaring_bitmap_metadata; + using metadata_type = roaring_bitmap_metadata; + using bitset_word_type = typename metadata_type::bitset_word_type; - 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 warps_per_block = BlockSize / cuco::detail::warp_size(); + constexpr cuda::std::uint32_t bitset_words = metadata_type::bitset_container_words; 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 payload_begin = + metadata_type::no_run_header_bytes(static_cast(state.num_containers)); auto const array_blocks = (state.num_array_containers + warps_per_block - 1) / warps_per_block; auto const block = static_cast(blockIdx.x); @@ -239,16 +270,14 @@ CUCO_KERNEL void write_roaring_containers(cuda::std::byte* bitmap, 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]; + auto const container_index = + static_cast(array_containers[warp_index]); + auto const bounds = container_bounds(container_index, container_starts, &state); + auto const cardinality = static_cast(bounds.end - bounds.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]); + auto const value = static_cast(indices[bounds.begin + index]); misaligned_store(container + index * sizeof(cuda::std::uint16_t), value); } } else { @@ -256,26 +285,24 @@ CUCO_KERNEL void write_roaring_containers(cuda::std::byte* bitmap, 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); + 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 bounds = container_bounds(container_index, container_starts, &state); + 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(indices, bounds.begin, bounds.end, word_begin); + auto const last = lower_bound_low_bits(indices, first, bounds.end, word_end); // One thread constructs one 64-bit word entirely in registers before issuing one final store. - unsigned long long mask = 0; + bitset_word_type mask = 0; for (auto index = first; index < last; ++index) { - auto const value = static_cast(keys[index]); - mask |= 1ULL << (value - word_begin); + auto const value = static_cast(indices[index]); + mask |= bitset_word_type{1} << (value - word_begin); } - misaligned_store(container + word * sizeof(unsigned long long), mask); + misaligned_store(container + word * sizeof(bitset_word_type), mask); } } diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh index 5c2af5024..82c08c421 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_storage.cuh @@ -49,10 +49,7 @@ class roaring_bitmap_storage_ref { : metadata_{metadata}, data_{bitmap}, run_container_bitmap_{bitmap + metadata_.run_container_bitmap}, - key_cards_{bitmap + metadata_.key_cards}, - container_offsets_{metadata_.offsets_in_serialized_data - ? (bitmap + metadata_.container_offsets) - : reinterpret_cast(metadata_.computed_offsets)} + key_cards_{bitmap + metadata_.key_cards} { assert(metadata.valid); } @@ -114,7 +111,9 @@ class roaring_bitmap_storage_ref { */ __host__ __device__ cuda::std::byte const* container_offsets() const noexcept { - return container_offsets_; + return metadata_.offsets_in_serialized_data + ? data_ + metadata_.container_offsets + : reinterpret_cast(metadata_.computed_offsets); } private: @@ -122,7 +121,6 @@ class roaring_bitmap_storage_ref { cuda::std::byte const* data_; cuda::std::byte const* run_container_bitmap_; cuda::std::byte const* key_cards_; - cuda::std::byte const* container_offsets_; }; /** @@ -246,8 +244,7 @@ class roaring_bitmap_storage { metadata_{metadata_type::from_serialized(bitmap)}, data_{allocator_.allocate(metadata_.size_bytes, stream), cuco::detail::custom_deleter{ - metadata_.size_bytes, allocator_, stream}}, - ref_{data_.get(), metadata_} + metadata_.size_bytes, allocator_, stream}} { CUCO_CUDA_TRY(cudaMemcpyAsync( data_.get(), bitmap, metadata_.size_bytes, cudaMemcpyHostToDevice, stream.get())); @@ -267,73 +264,52 @@ class roaring_bitmap_storage { metadata_{metadata}, data_{allocator_.allocate(metadata_.size_bytes, stream), cuco::detail::custom_deleter{ - metadata_.size_bytes, allocator_, stream}}, - ref_{data_.get(), metadata_} + metadata_.size_bytes, allocator_, stream}} { assert(metadata_.valid); } /** - * @brief Move constructor. - * - * Rebuilds the cached reference because small run-container bitmaps store computed offsets - * directly in the metadata object. + * @brief Move constructor * * @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_} - { - } + roaring_bitmap_storage(roaring_bitmap_storage&& other) noexcept = default; /** - * @brief Move assignment operator. - * - * Rebuilds the cached reference because small run-container bitmaps store computed offsets - * directly in the metadata object. + * @brief Move assignment operator * * @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; - } + roaring_bitmap_storage& operator=(roaring_bitmap_storage&& other) noexcept = default; /** * @brief Returns a mutable pointer to serialized storage * * @return Pointer to serialized storage */ - cuda::std::byte* data() noexcept { return data_.get(); } + [[nodiscard]] cuda::std::byte* data() noexcept { return data_.get(); } /** * @brief Returns a reference to the stored bitmap * * @return Reference to the bitmap storage */ - ref_type ref() const noexcept { return ref_; } + [[nodiscard]] ref_type ref() const noexcept { return ref_type{data_.get(), metadata_}; } /** * @brief Returns the allocator used to manage storage * * @return Allocator instance */ - allocator_type allocator() const noexcept { return allocator_; } + [[nodiscard]] allocator_type allocator() const noexcept { return allocator_; } private: allocator_type allocator_; metadata_type metadata_; std::unique_ptr> data_; - ref_type ref_; }; /** @@ -437,14 +413,14 @@ class roaring_bitmap_storage { * * @return Reference to the bitmap storage */ - ref_type ref() const noexcept { return ref_; } + [[nodiscard]] 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_; } + [[nodiscard]] allocator_type allocator() const noexcept { return allocator_; } private: allocator_type allocator_; diff --git a/include/cuco/detail/roaring_bitmap/util.cuh b/include/cuco/detail/roaring_bitmap/util.cuh index 6f7f03ae5..728a23df3 100644 --- a/include/cuco/detail/roaring_bitmap/util.cuh +++ b/include/cuco/detail/roaring_bitmap/util.cuh @@ -59,6 +59,8 @@ struct roaring_bitmap_metadata { */ template <> struct roaring_bitmap_metadata { + using bitset_word_type = cuda::std::uint64_t; ///< Word type used by bitset containers + /// 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 @@ -71,6 +73,61 @@ struct roaring_bitmap_metadata { static constexpr cuda::std::int32_t no_offset_threshold = 4; /// Fixed size of a bitset container in bytes static constexpr cuda::std::uint32_t bitset_container_bytes = 8192; + /// Number of words in a bitset container + static constexpr cuda::std::uint32_t bitset_container_words = + bitset_container_bytes / sizeof(bitset_word_type); + /// Byte offset of key/cardinality descriptors in a no-run bitmap + static constexpr cuda::std::uint32_t no_run_key_cards_offset = 2 * sizeof(cuda::std::uint32_t); + + /** + * @brief Returns the byte offset of container offsets in a no-run bitmap. + * + * @param bitmap_num_containers Number of containers + * @return Byte offset of the container offset table + */ + [[nodiscard]] __host__ __device__ static constexpr cuda::std::uint32_t + no_run_container_offsets_offset(cuda::std::uint32_t bitmap_num_containers) noexcept + { + return no_run_key_cards_offset + bitmap_num_containers * 2 * sizeof(cuda::std::uint16_t); + } + + /** + * @brief Returns the byte size of a no-run bitmap header. + * + * @param bitmap_num_containers Number of containers + * @return Header size in bytes + */ + [[nodiscard]] __host__ __device__ static constexpr cuda::std::uint32_t no_run_header_bytes( + cuda::std::uint32_t bitmap_num_containers) noexcept + { + return no_run_container_offsets_offset(bitmap_num_containers) + + bitmap_num_containers * sizeof(cuda::std::uint32_t); + } + + /** + * @brief Returns the encoded payload size for a container cardinality. + * + * @param cardinality Number of indices in the container + * @return Payload size in bytes + */ + [[nodiscard]] __host__ __device__ static constexpr cuda::std::uint32_t container_payload_bytes( + cuda::std::uint32_t cardinality) noexcept + { + return cardinality <= max_array_container_card ? cardinality * sizeof(cuda::std::uint16_t) + : bitset_container_bytes; + } + + /** + * @brief Returns the high-16-bit container key for an index. + * + * @param index Bitmap index + * @return Container key + */ + [[nodiscard]] __host__ __device__ static constexpr cuda::std::uint16_t container_key( + cuda::std::uint32_t index) noexcept + { + return static_cast(index >> 16); + } /// Total size of the bitmap in bytes cuda::std::size_t size_bytes = 0; @@ -111,9 +168,9 @@ struct roaring_bitmap_metadata { 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.key_cards = no_run_key_cards_offset; metadata.container_offsets = - metadata.key_cards + bitmap_num_containers * 2 * sizeof(cuda::std::uint16_t); + no_run_container_offsets_offset(static_cast(bitmap_num_containers)); metadata.num_containers = bitmap_num_containers; metadata.has_run = false; metadata.valid = true; diff --git a/include/cuco/detail/utility/cuda.cuh b/include/cuco/detail/utility/cuda.cuh index 75cdea974..db416d45c 100644 --- a/include/cuco/detail/utility/cuda.cuh +++ b/include/cuco/detail/utility/cuda.cuh @@ -32,7 +32,7 @@ using index_type = cuda::std::int64_t; ///< CUDA thread index type /// Default block size /// CUDA warp size -[[nodiscard]] __device__ constexpr cuda::std::int32_t warp_size() noexcept { return 32; } +[[nodiscard]] __host__ __device__ constexpr cuda::std::int32_t warp_size() noexcept { return 32; } /** * @brief Returns the global thread index in a 1D scalar grid diff --git a/include/cuco/roaring_bitmap.cuh b/include/cuco/roaring_bitmap.cuh index b6305be6b..2828de064 100644 --- a/include/cuco/roaring_bitmap.cuh +++ b/include/cuco/roaring_bitmap.cuh @@ -42,9 +42,8 @@ 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 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. @@ -52,16 +51,16 @@ class roaring_bitmap { * @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 + * + * @throw cuco::logic_error If the serialized bitmap header is invalid or unsupported */ roaring_bitmap(cuda::std::byte const* 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. + * @brief Creates a `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. @@ -71,6 +70,8 @@ class roaring_bitmap { * @param stream CUDA stream used for device memory operations during construction * * @return Bitmap containing a copy of the serialized input + * + * @throw cuco::logic_error If the serialized bitmap header is invalid or unsupported */ [[nodiscard]] static roaring_bitmap from_serialized(cuda::std::byte const* bitmap, Allocator const& alloc = {}, @@ -95,6 +96,8 @@ class roaring_bitmap { * @param stream CUDA stream used for device memory operations and kernel launches * * @return Bitmap containing the unique input indices + * + * @throw cuco::logic_error If `[first, last)` is not a valid range */ template [[nodiscard]] static roaring_bitmap from_indices(InputIt first, @@ -120,6 +123,8 @@ class roaring_bitmap { * @param stream CUDA stream used for device memory operations and kernel launches * * @return Bitmap containing the unique input indices + * + * @throw cuco::logic_error If `[first, last)` is not a valid range */ template [[nodiscard]] static roaring_bitmap from_sorted_indices( @@ -145,6 +150,8 @@ class roaring_bitmap { * @param stream CUDA stream used for device memory operations and kernel launches * * @return Bitmap containing the input indices + * + * @throw cuco::logic_error If `[first, last)` is not a valid range */ template [[nodiscard]] static roaring_bitmap from_sorted_unique_indices( diff --git a/include/cuco/roaring_bitmap_ref.cuh b/include/cuco/roaring_bitmap_ref.cuh index baa756ef4..cfb183ce3 100644 --- a/include/cuco/roaring_bitmap_ref.cuh +++ b/include/cuco/roaring_bitmap_ref.cuh @@ -25,14 +25,15 @@ namespace cuco::experimental { * "Standard 32-bit Roaring Bitmap" format; for 64-bit bitmaps, the "portable" format is * supported. * - * @tparam T Key type stored in the bitmap. Must be `cuda::std::uint32_t` or `cuda::std::uint64_t`. + * @tparam T Index type stored in the bitmap. Must be `cuda::std::uint32_t` or + * `cuda::std::uint64_t`. */ template class roaring_bitmap_ref { using impl_type = detail::roaring_bitmap_impl; public: - using value_type = T; ///< Key type stored in the bitmap + using value_type = T; ///< Index type stored in the bitmap using storage_ref_type = typename impl_type::storage_ref_type; ///< Implementation storage ref /** @@ -55,17 +56,17 @@ class roaring_bitmap_ref { __device__ roaring_bitmap_ref(cuda::std::byte const* bitmap); /** - * @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 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 */ @@ -76,14 +77,14 @@ class roaring_bitmap_ref { 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 */ @@ -95,23 +96,23 @@ class roaring_bitmap_ref { cudaStream_t{nullptr}}) const noexcept; /** - * @brief Device-side membership query for a single key. + * @brief Device-side membership query for a single index. * - * @param value Key to test for membership + * @param value Index to test for membership * * @return `true` iff `value` is contained in the bitmap */ __device__ bool contains(T value) const; /** - * @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]] __host__ __device__ 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` */ @@ -137,4 +138,4 @@ class roaring_bitmap_ref { } // namespace cuco::experimental -#include \ No newline at end of file +#include diff --git a/tests/roaring_bitmap/build_test.cu b/tests/roaring_bitmap/build_test.cu index c846930d0..bffa3d7a4 100644 --- a/tests/roaring_bitmap/build_test.cu +++ b/tests/roaring_bitmap/build_test.cu @@ -3,6 +3,8 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "test_utils.cuh" + #include #include @@ -10,9 +12,12 @@ #include #include #include +#include #include #include #include +#include +#include #include @@ -129,16 +134,7 @@ TEST_CASE("roaring_bitmap matches RoaringFormatSpec no-run serialization", "[roa "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); - } + auto const host_indices = cuco::test::make_roaring_bitmap_without_runs_indices(); thrust::device_vector indices{host_indices}; auto bitmap = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); @@ -157,18 +153,68 @@ TEST_CASE("roaring_bitmap matches RoaringFormatSpec no-run serialization", "[roa #endif } -TEST_CASE("roaring_bitmap treats reversed input ranges as empty", "[roaring_bitmap]") +TEST_CASE("roaring_bitmap rejects reversed input ranges", "[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_THROWS_AS(bitmap_type::from_indices(indices.end(), indices.begin()), cuco::logic_error); + REQUIRE_THROWS_AS(bitmap_type::from_sorted_indices(indices.end(), indices.begin()), + cuco::logic_error); + REQUIRE_THROWS_AS(bitmap_type::from_sorted_unique_indices(indices.end(), indices.begin()), + cuco::logic_error); +} - 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 format boundary indices", "[roaring_bitmap]") +{ + auto constexpr max_index = cuda::std::numeric_limits::max(); + thrust::device_vector indices{0, max_index}; + + auto bitmap = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.size() == 2); + require_contains(bitmap, {0, 1, max_index - 1, max_index}, {true, false, false, true}); +} + +TEST_CASE("roaring_bitmap removes duplicates across a container boundary", "[roaring_bitmap]") +{ + thrust::device_vector indices{ + 0x0000FFFE, 0x0000FFFF, 0x0000FFFF, 0x00010000, 0x00010000, 0x00010001}; + + auto bitmap = bitmap_type::from_sorted_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.size() == 4); + require_contains(bitmap, + {0x0000FFFD, 0x0000FFFE, 0x0000FFFF, 0x00010000, 0x00010001, 0x00010002}, + {false, true, true, true, true, false}); +} + +TEST_CASE("roaring_bitmap builds a full container", "[roaring_bitmap]") +{ + constexpr cuda::std::uint32_t num_indices = 1 << 16; + thrust::device_vector indices(num_indices); + thrust::sequence(indices.begin(), indices.end(), index_type{0x12340000}); + + auto bitmap = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.size() == num_indices); + require_contains( + bitmap, {0x1233FFFF, 0x12340000, 0x1234FFFF, 0x12350000}, {false, true, true, false}); +} + +TEST_CASE("roaring_bitmap builds the maximum number of containers", "[roaring_bitmap]") +{ + constexpr cuda::std::uint32_t num_containers = 1 << 16; + thrust::device_vector indices(num_containers); + thrust::tabulate(indices.begin(), indices.end(), [] __device__(index_type container) { + return (container << 16) | index_type{7}; + }); + + auto bitmap = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.size() == num_containers); + require_contains(bitmap, + {7, 8, 0x7FFF0007, 0x80000007, 0xFFFF0007, 0xFFFF0008}, + {true, false, true, true, true, false}); } TEST_CASE("roaring_bitmap builds multiple array containers per block", "[roaring_bitmap]") @@ -268,6 +314,21 @@ TEST_CASE("roaring_bitmap selects array and bitset containers at the format thre REQUIRE(bitmap.size() == 4097); require_contains(bitmap, {0, 4096, 4097}, {true, true, false}); } + + SECTION("sorted bitset container with duplicates") + { + std::vector host_indices; + host_indices.reserve(4099); + for (index_type i = 0; i < 4097; ++i) { + host_indices.push_back(i); + if (i == 2048 || i == 4096) { host_indices.push_back(i); } + } + thrust::device_vector indices(host_indices); + auto bitmap = bitmap_type::from_sorted_indices(indices.begin(), indices.end()); + + REQUIRE(bitmap.size() == 4097); + require_contains(bitmap, {0, 2048, 4096, 4097}, {true, true, true, false}); + } } TEST_CASE("roaring_bitmap writes a bitset after an odd-sized array container", "[roaring_bitmap]") diff --git a/tests/roaring_bitmap/contains_test.cu b/tests/roaring_bitmap/contains_test.cu index a9b992866..b3e00da5f 100644 --- a/tests/roaring_bitmap/contains_test.cu +++ b/tests/roaring_bitmap/contains_test.cu @@ -3,6 +3,8 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "test_utils.cuh" + #include #include #include @@ -25,38 +27,30 @@ #include namespace { +using index_type = cuda::std::uint32_t; using bitmap_type = cuco::experimental::roaring_bitmap; template bool check(std::string const& bitmap_file_path) { - auto generate_keys = []() -> thrust::device_vector { + auto generate_indices = []() -> thrust::device_vector { if constexpr (cuda::std::is_same_v) { - std::vector keys; - for (cuda::std::uint32_t k = 0; k < 100000; k += 1000) { - keys.push_back(k); - } - for (int k = 100000; k < 200000; ++k) { - keys.push_back(3 * k); - } - for (int k = 700000; k < 800000; ++k) { - keys.push_back(k); - } - return thrust::device_vector(keys.begin(), keys.end()); + auto const indices = cuco::test::make_roaring_bitmap_without_runs_indices(); + return thrust::device_vector(indices.begin(), indices.end()); } else if constexpr (cuda::std::is_same_v) { - std::vector keys; - for (cuda::std::uint64_t k = 0x00000ull; k < 0x09000ull; ++k) { - keys.push_back(k); + std::vector indices; + for (cuda::std::uint64_t index = 0x00000ull; index < 0x09000ull; ++index) { + indices.push_back(index); } - for (cuda::std::uint64_t k = 0x0A000ull; k < 0x10000ull; ++k) { - keys.push_back(k); + for (cuda::std::uint64_t index = 0x0A000ull; index < 0x10000ull; ++index) { + indices.push_back(index); } - keys.push_back(0x20000ull); - keys.push_back(0x20005ull); + indices.push_back(0x20000ull); + indices.push_back(0x20005ull); for (cuda::std::uint64_t i = 0; i < 0x10000ull; i += 2ull) { - keys.push_back(0x80000ull + i); + indices.push_back(0x80000ull + i); } - return thrust::device_vector(keys.begin(), keys.end()); + return thrust::device_vector(indices.begin(), indices.end()); } else { static_assert(cuco::dependent_false, "KeyType must be uint32_t or uint64_t"); return {}; @@ -73,13 +67,13 @@ bool check(std::string const& bitmap_file_path) file.read(reinterpret_cast(thrust::raw_pointer_cast(buffer.data())), file_size); file.close(); - cuco::experimental::roaring_bitmap roaring_bitmap( + auto roaring_bitmap = cuco::experimental::roaring_bitmap::from_serialized( thrust::raw_pointer_cast(buffer.data())); - auto keys = generate_keys(); - thrust::device_vector contained(keys.size(), false); + auto indices = generate_indices(); + thrust::device_vector contained(indices.size(), false); - roaring_bitmap.contains(keys.begin(), keys.end(), contained.begin()); + roaring_bitmap.contains(indices.begin(), indices.end(), contained.begin()); bool const all_contained = thrust::all_of(contained.begin(), contained.end(), ::cuda::std::identity{}); @@ -128,7 +122,7 @@ 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()); - bitmap_type original{thrust::raw_pointer_cast(buffer.data())}; + auto original = bitmap_type::from_serialized(thrust::raw_pointer_cast(buffer.data())); bitmap_type moved{std::move(original)}; thrust::device_vector keys{1, 2, 3, 4}; @@ -155,18 +149,52 @@ TEST_CASE("roaring_bitmap run container without offsets", "[roaring_bitmap]") REQUIRE_FALSE(contained_h[3]); } -TEST_CASE("roaring_bitmap parses a serialized empty bitmap", "[roaring_bitmap]") +TEST_CASE("roaring_bitmap_storage_ref keeps computed offsets when copied", "[roaring_bitmap]") +{ + auto const bytes = make_run_container_no_offsets_bitmap(); + using storage_ref = cuco::experimental::detail::roaring_bitmap_storage_ref; + using metadata_type = storage_ref::metadata_type; + + storage_ref original{bytes.data(), metadata_type::from_serialized(bytes.data())}; + storage_ref copy = original; + + REQUIRE(copy.container_offsets() == + reinterpret_cast(copy.metadata().computed_offsets)); +} + +TEST_CASE("roaring_bitmap parses serialized round trips", "[roaring_bitmap]") { - thrust::device_vector indices; - auto source = bitmap_type::from_sorted_unique_indices(indices.begin(), indices.end()); + SECTION("empty 32-bit 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)); + 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()); + auto bitmap = bitmap_type::from_serialized(bytes.data()); + REQUIRE(bitmap.empty()); + REQUIRE(bitmap.size() == 0); + REQUIRE(bitmap.size_bytes() == bytes.size()); + } + + SECTION("nonempty 32-bit bitmap") + { + thrust::device_vector indices{0, 7, 0x00010000, 0xFFFFFFFF}; + 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.size() == indices.size()); + + thrust::device_vector queries{0, 1, 7, 0x00010000, 0xFFFFFFFF}; + thrust::device_vector contained(queries.size()); + bitmap.contains(queries.begin(), queries.end(), contained.begin()); + thrust::host_vector result = contained; + REQUIRE(result == (thrust::host_vector{true, false, true, true, true})); + } } TEST_CASE("roaring_bitmap bulk contains from RoaringFormatSpec testdata", "[roaring_bitmap]") diff --git a/tests/roaring_bitmap/test_utils.cuh b/tests/roaring_bitmap/test_utils.cuh new file mode 100644 index 000000000..d0dd97b08 --- /dev/null +++ b/tests/roaring_bitmap/test_utils.cuh @@ -0,0 +1,29 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include + +#include + +namespace cuco::test { + +inline std::vector make_roaring_bitmap_without_runs_indices() +{ + std::vector indices; + for (cuda::std::uint32_t index = 0; index < 100000; index += 1000) { + indices.push_back(index); + } + for (cuda::std::uint32_t index = 100000; index < 200000; ++index) { + indices.push_back(3 * index); + } + for (cuda::std::uint32_t index = 700000; index < 800000; ++index) { + indices.push_back(index); + } + return indices; +} + +} // namespace cuco::test From 2027216da1ba687374222f54e0e196dd1ccaedc7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Fri, 4 Sep 2026 18:44:19 -0700 Subject: [PATCH 8/8] Fix Roaring kernel linkage for Clang builds --- include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh index a98866d20..2a4f3600c 100644 --- a/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh +++ b/include/cuco/detail/roaring_bitmap/roaring_bitmap_kernels.cuh @@ -121,7 +121,7 @@ struct container_payload_size { * @param payload_offsets Exclusive payload offsets for every container * @param state Device build state */ -inline CUCO_KERNEL void collect_container_indexes(cuda::std::uint32_t* container_indexes, +static CUCO_KERNEL void collect_container_indexes(cuda::std::uint32_t* container_indexes, cuco::detail::index_type num_container_slots, cuco::detail::index_type const* container_starts, cuda::std::uint32_t const* payload_offsets,