Support for 32- and 64-bit cuco::experimental::roaring_bitmap lookups #741
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change | ||||
|---|---|---|---|---|---|---|
| @@ -0,0 +1,97 @@ | ||||||
| /* | ||||||
| * Copyright (c) 2025, NVIDIA CORPORATION. | ||||||
| * | ||||||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||||||
| * you may not use this file except in compliance with the License. | ||||||
| * You may obtain a copy of the License at | ||||||
| * | ||||||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||||||
| * | ||||||
| * Unless required by applicable law or agreed to in writing, software | ||||||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||||||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||||||
| * See the License for the specific language governing permissions and | ||||||
| * limitations under the License. | ||||||
| */ | ||||||
|
|
||||||
| #include <benchmark_defaults.hpp> | ||||||
| #include <benchmark_utils.hpp> | ||||||
|
|
||||||
| #include <cuco/roaring_bitmap.cuh> | ||||||
| #include <cuco/utility/key_generator.cuh> | ||||||
|
|
||||||
| #include <nvbench/nvbench.cuh> | ||||||
|
|
||||||
| #include <cuda/std/cstddef> | ||||||
| #include <cuda/std/cstdint> | ||||||
| #include <thrust/device_vector.h> | ||||||
| #include <thrust/universal_vector.h> | ||||||
|
|
||||||
| #include <fstream> | ||||||
| #include <string> | ||||||
|
|
||||||
| using namespace cuco::benchmark; // defaults | ||||||
| using namespace cuco::utility; // key_generator, distribution | ||||||
|
|
||||||
| template <typename T> | ||||||
| void roaring_bitmap_contains(nvbench::state& state, nvbench::type_list<T>) | ||||||
| { | ||||||
| auto const num_items = state.get_int64("NumInputs"); | ||||||
| auto const bitmap_file = state.get_string_or_default("BitmapFile", {}); | ||||||
|
|
||||||
| std::ifstream file(bitmap_file, std::ios::binary); | ||||||
| if (!file.is_open()) { state.skip("Bitmap file not found"); } | ||||||
|
|
||||||
| // Get file size | ||||||
| file.seekg(0, std::ios::end); | ||||||
| std::streamsize file_size = file.tellg(); | ||||||
| file.seekg(0, std::ios::beg); | ||||||
sleeepyjack marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||||||
|
|
||||||
| thrust::universal_host_pinned_vector<cuda::std::byte> buffer(file_size); | ||||||
|
|
||||||
| file.read(reinterpret_cast<char*>(thrust::raw_pointer_cast(buffer.data())), file_size); | ||||||
| file.close(); | ||||||
|
|
||||||
| cuco::roaring_bitmap<T> roaring_bitmap(thrust::raw_pointer_cast(buffer.data())); | ||||||
|
|
||||||
| thrust::device_vector<T> items(num_items); | ||||||
|
|
||||||
| key_generator gen{}; | ||||||
|
Contributor
There was a problem hiding this comment. Just caffeinating the engine ;)
Suggested change
Collaborator
Author
There was a problem hiding this comment. Ultimately, I would like to use the same default seed for all cuco benchmarks and add an option to adjust it via nvbench. We currently use |
||||||
| gen.generate(distribution::unique{}, items.begin(), items.end()); | ||||||
|
|
||||||
| thrust::device_vector<bool> contained(items.size(), false); | ||||||
|
|
||||||
| state.add_element_count(items.size()); | ||||||
| state.add_global_memory_reads<T>(items.size(), "InputSize"); | ||||||
|
|
||||||
| auto& summ = state.add_summary("BitmapSizeMB"); | ||||||
| summ.set_string("hint", "BitmapSize"); | ||||||
| summ.set_string("short_name", "BitmapSizeMB"); | ||||||
| summ.set_string("description", "Bitmap size in MB"); | ||||||
| summ.set_float64("value", static_cast<double>(file_size) / (1024 * 1024)); | ||||||
|
|
||||||
| state.exec([&](nvbench::launch& launch) { | ||||||
| roaring_bitmap.contains_async( | ||||||
| items.begin(), items.end(), contained.begin(), {launch.get_stream()}); | ||||||
| }); | ||||||
| } | ||||||
|
|
||||||
| NVBENCH_BENCH_TYPES(roaring_bitmap_contains, | ||||||
| NVBENCH_TYPE_AXES(nvbench::type_list<nvbench::uint32_t>)) | ||||||
| .set_name("roaring_bitmap_contains") | ||||||
| .add_int64_power_of_two_axis("NumInputs", {32}) | ||||||
| // Default benchmark is only available if the Roaring bitmap testdata has been downloaded | ||||||
| #ifdef CUCO_ROARING_DATA_DIR | ||||||
| .add_string_axis("BitmapFile", {std::string(CUCO_ROARING_DATA_DIR) + "/bitmapwithruns.bin"}) | ||||||
| #endif | ||||||
| .set_max_noise(cuco::benchmark::defaults::MAX_NOISE); | ||||||
|
|
||||||
| NVBENCH_BENCH_TYPES(roaring_bitmap_contains, | ||||||
| NVBENCH_TYPE_AXES(nvbench::type_list<nvbench::uint64_t>)) | ||||||
| .set_name("roaring_bitmap_contains") | ||||||
| .add_int64_power_of_two_axis("NumInputs", {31}) | ||||||
| // Default benchmark is only available if the Roaring bitmap testdata has been downloaded | ||||||
| #ifdef CUCO_ROARING_DATA_DIR | ||||||
| .add_string_axis("BitmapFile", {std::string(CUCO_ROARING_DATA_DIR) + "/portable_bitmap64.bin"}) | ||||||
| #endif | ||||||
| .set_max_noise(cuco::benchmark::defaults::MAX_NOISE); | ||||||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -51,6 +51,8 @@ HOST_COMPILER=${CXX:-g++} # $CXX if set, otherwise `g++` | |
| CUDA_ARCHS=native # detect system's GPU architectures | ||
| CXX_STANDARD=17 | ||
|
|
||
| EXTRA_CMAKE_OPTIONS=() | ||
|
|
||
| function usage { | ||
| echo "cuCollections build script" | ||
| echo "Usage: $0 [OPTIONS]" | ||
|
|
@@ -62,9 +64,9 @@ function usage { | |
| echo " --prefix: Build directory prefix (Defaults to <repo_root>/build)" | ||
| echo " -i/--infix: Build directory infix (Defaults to local)" | ||
| echo " -d/--debug: Debug build" | ||
| echo " -p/--parallel: Build parallelism (Defaults to $PARALLEL_LEVEL if set, otherwise the system's number of CPU cores)" | ||
| echo " --cuda: CUDA compiler (Defaults to $CUDACXX if set, otherwise nvcc)" | ||
| echo " --cxx: Host compiler (Defaults to $CXX if set, otherwise g++)" | ||
| echo " --arch: Target CUDA arches, e.g. \"60-real;70;80-virtual\" (Defaults to the system's native GPU archs)" | ||
| echo " --std: CUDA/C++ standard (Defaults to 17)" | ||
| echo " -v/-verbose/--verbose: Enable shell echo for debugging" | ||
|
|
@@ -103,6 +105,9 @@ function usage { | |
| echo " Enables verbose mode for detailed output and builds with C++17 standard." | ||
| echo " Build files will be written to <repo_root>/build/local and symlinked to <repo_root>/build/latest." | ||
| echo | ||
| echo "Pass-through:" | ||
| echo " -- [CMake args...] Anything after -- is forwarded to CMake" | ||
| echo | ||
| exit 1 | ||
| } | ||
|
|
||
|
|
@@ -126,6 +131,7 @@ while [ "${#args[@]}" -ne 0 ]; do | |
| --arch) CUDA_ARCHS="${args[1]}"; args=("${args[@]:2}");; | ||
| --std) CXX_STANDARD="${args[1]}"; args=("${args[@]:2}");; | ||
| -v | -verbose | --verbose) VERBOSE=1; args=("${args[@]:1}");; | ||
| --) EXTRA_CMAKE_OPTIONS+=("${args[@]:1}"); break;; | ||
|
||
| -h | -help | --help) usage ;; | ||
| *) echo "Unrecognized option: ${args[0]}"; usage ;; | ||
| esac | ||
|
|
@@ -200,8 +206,14 @@ echo "-- BUILD_TESTS: ${BUILD_TESTS}" | |
| echo "-- BUILD_EXAMPLES: ${BUILD_EXAMPLES}" | ||
| echo "-- BUILD_BENCHMARKS: ${BUILD_BENCHMARKS}" | ||
|
|
||
| if [ ${#EXTRA_CMAKE_OPTIONS[@]} -gt 0 ]; then | ||
| echo "-- EXTRA_CMAKE_OPTIONS: ${EXTRA_CMAKE_OPTIONS[*]}" | ||
| else | ||
| echo "-- EXTRA_CMAKE_OPTIONS: (none)" | ||
| fi | ||
|
|
||
| # configure | ||
| cmake -S .. -B $BUILD_DIR $CMAKE_OPTIONS "${EXTRA_CMAKE_OPTIONS[@]}" | ||
| echo "========================================" | ||
|
|
||
| if command -v sccache >/dev/null; then | ||
|
|
||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,39 @@ | ||
| # ============================================================================= | ||
| # Copyright (c) 2025, NVIDIA CORPORATION. | ||
| # | ||
| # Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except | ||
| # in compliance with the License. You may obtain a copy of the License at | ||
| # | ||
| # http://www.apache.org/licenses/LICENSE-2.0 | ||
| # | ||
| # Unless required by applicable law or agreed to in writing, software distributed under the License | ||
| # is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express | ||
| # or implied. See the License for the specific language governing permissions and limitations under | ||
| # the License. | ||
| # ============================================================================= | ||
|
|
||
| # Only act if enabled | ||
| if(NOT CUCO_DOWNLOAD_ROARING_TESTDATA) | ||
| return() | ||
| endif() | ||
|
|
||
| set(CUCO_ROARING_DATA_DIR "${CMAKE_BINARY_DIR}/data/roaring_bitmap") | ||
|
|
||
| file(MAKE_DIRECTORY "${CUCO_ROARING_DATA_DIR}") | ||
|
|
||
| set(ROARING_FORMATSPEC_BASE "https://raw.githubusercontent.com/RoaringBitmap/RoaringFormatSpec/5177ad9") | ||
|
|
||
| rapids_cmake_download_with_retry("${ROARING_FORMATSPEC_BASE}/testdata/bitmapwithoutruns.bin" | ||
| "${CUCO_ROARING_DATA_DIR}/bitmapwithoutruns.bin" | ||
| "d719ae2e0150a362ef7cf51c361527585891f01460b1a92bcfb6a7257282a442") | ||
|
|
||
| rapids_cmake_download_with_retry("${ROARING_FORMATSPEC_BASE}/testdata/bitmapwithruns.bin" | ||
| "${CUCO_ROARING_DATA_DIR}/bitmapwithruns.bin" | ||
| "1f1909bfdd354fa2f0694fe88b8076833ca5383ad9fc3f68f2709c84a2ab70e3") | ||
|
|
||
| rapids_cmake_download_with_retry("${ROARING_FORMATSPEC_BASE}/testdata64/portable_bitmap64.bin" | ||
| "${CUCO_ROARING_DATA_DIR}/portable_bitmap64.bin" | ||
| "b5a553a759167f5f9ccb3fa21552d943b4c73235635b753376f4faf62067d178") | ||
|
|
||
| # Define macro only when data is available | ||
| add_compile_definitions(CUCO_ROARING_DATA_DIR="${CUCO_ROARING_DATA_DIR}") |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.