Support for 32- and 64-bit cuco::experimental::roaring_bitmap lookups (#741)

Closes #725

Author: sleeepyjack

CMakeLists.txt

@@ -88,6 +88,12 @@ target_include_directories(cuco INTERFACE
 target_link_libraries(cuco INTERFACE CCCL::CCCL CUDA::toolkit)
 target_compile_features(cuco INTERFACE cxx_std_17 cuda_std_17)
 
+###################################################################################################
+# - Optionally download RoaringFormatSpec test data -----------------------------------------------
+
+option(CUCO_DOWNLOAD_ROARING_TESTDATA "Download RoaringFormatSpec test data" ON)
+include(${CMAKE_CURRENT_LIST_DIR}/cmake/roaring_testdata.cmake)
+
 ###################################################################################################
 # - optionally build tests ------------------------------------------------------------------------
 

README.md

@@ -259,4 +259,11 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
 `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries.
 
 #### Examples:
-- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw))
+- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw))
+
+### roaring_bitmap
+
+`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/eJy9WA1v2zYT_iv3qsAgN7aVpB_ZnI_NjdPNWF97sN0VQ1MIlETbhGVRI6k4XpD__h5JfSbK2nV75wCxRR7vnjs-dzzqzpFUSsYT6Qw-3jkscgZHXScmySojK-oMnDCLiNN1JM9EqJ-959cJPIdLnu4FW60VuGEHjg-PX8Hk1_FoPITL6eyX6Wy4GE8nfS1qxN-xkCaSRpAlERWg1hSGKQnxK5_pwq9UaBxw3D8EVwtcO_nctdM5NVr2PIMt2UPCFWSSohomYcliCvQ2pKkClkDIt2nMSBJS2DG1NqZyPQYO_JYr4YEiKE9wRYpPy7okEFVC15-1UunA83a7XZ8Y2H0uVl5shaX3bnx5NZlf9RB6uex9EmNgQdDfMybQ8WAPJEVkIQkQb0x2wAWQlaA4p7hGvhNMsWTVBcmXakcENXoiJpVgQaYawStwov91AQwfSTBwwzmM59cOvBnOx_Ou0fNhvPhp-n4BH4az2XCyGF_NYTrDzZqMxnqr8OktDCe_wc_jyagLFEOHpuhtKrQXCJXpsNLIxnBOaQPGkltYMqUhW7IQCgLBit9QkaBbkFKxZZZqCDIyemK2ZYooM_bIOWPKu06uk2csCeMsonAWZiH3BCcCNfoBU1uS9sNsffFIJlMsZmrvKUGYkv11ml481BQRT6rIC_FfRJcXT06yRLVPqn1KfWugIaDWIpPKi-gNOuLf0FBx0V-3icR8hYSI2yezhGHsJInrKupymvlyLxXdNpYvkQ-UNMcYbxnUvElWjSFrydjxnlsi_2ASbI0K_CCLNz69JUgFimG304FgdAkjusVNxGAoilsp9dbnKdXcLSSn1oLcjDnfZKlRDMNfxrIqFuPEZnZuCXaYL5xgmuw4vDjuoSKwygyRkPQUXr-sDYObcqFMniExt0R1YCn4VqMx-j_OLKQ3RvqtEZkjcz-5Os8lJvoK6Z8FfawlXkO2eKrWdDDDUy4ZRm1f0hoLRLgBZv3X7hbOop8orzKBfNdzIRcCI45jMouRpjAhWxrvu9plDKQyQksex3ynU8hs-MCY6MFH66wucjxTIktkP2CJmXRtlDpf45AXxDzwXh2dnJDoO0-DiIgiXquxzmMo_x6ORyCKTc-Z9vplhcPS4x_E8fql12quU5L4BzykKCzqPM7Pgnzrt2SDOZLq4ge90eX7y6k_mn6YvJsOR_5sOpyNJz_6i6v5YjRcDM-nE31IBFhhqSoTxdRGzP40xqzDsoHlKEH-wM90v8DfmMUB57Flo4u4BwOb8Ug7TNVv8mzxNav8lKg1gr9DvUAytLWiCdXZ7G_oXsI5fPzkdqB3AbY4DQaN6nZWmASjADT3jRF9gGCHgBVTG0cETPoSMfo3xZIu1GYzrLUvjn110SkUAXgeXGLpQg9_zygmmcFTHDgVHe6Ql_cFJ6pMKVT8YwScXQ1H_73qb6NneqinxwozxoU8IG1OGeinhbR2wW0Rgw1G-_AUv87g6FB_9O-Dc_NQiwsYdf00k2s_ILjDm06p-75hBBUbpZW2M-zY7O-Dg82f6XyBTP6s3pOa3m-_TG-bTlsYnyBYWzxdozSgK5a4na41QZPI7RTK74HGukn6ejK-fvkXyNhaEQwV_x9MxBJUcfGRaVv6PkdL497naanFLC1vzfZmcWx3G5-_K56_Ysc_Z2vYtHV0-DdsPRA5vD3OtVWyrSKvGiJPImY2adkDoEwn7rFW8SdwD2-_zRfAAbC_mRlmS780M-4qjmAXHvpESiqUqztobQbPk4gmyl8SFC5rfBd7uPw3bBGQPpbK4oXxKVDk17a6B3f3hX39pR_0NybFFE2ZTNHPNjfzNtaMug-Pqm4uxOVggFQnYm9tYaq7_9FSfcxtjkrR5cJLsyKkQsDZGbrwlqCYuXtpORzQww_t6DGzDiMR59hzX0xQzNB96caPVNn7qGR_0PIoNdr0CHLEKKsad_vbTD7ysVOFZ4gdYKirjumWt9hrY-1pnoFl9AqWVNcH07unLElo1EKaYK_w3A6y5ZIKt0RTMz6j2Hsbr3BjeW5dz5k44x5FrqA4RQVeFpUfEqnOwjURzy_cAosgOz_lRsbMu9ZcX9cx3CGkaMNwrjqMuaRuDUleevMbRel43tuXEbD0xWpPBdsif0mMCBq3kKpfaY67doe_FPZpbeNtt1Q_FlyCCS2xX44jnSJ4BOkXDkg5ltT2rVPSJG-0Go1Xi_cE7B5q5kr8tjfx_AZRJ0CzTOhW8KICYSuE4Z2OvuZyzdQbvLP08jtLzSOywrX5rcSi1_IPr-PWhHyyBnUrFMV03cvi7qSjpw1Zy-LJAOp1ps_FBX4lc17GQY_zpfvIaB1HDgzrQ6070LWPqf3dveVknvcNMzhhsl-3QluitZoe-hlbRngpNk190cvrPt4fjWdlhas346Ap5UdMIO7WVaelmzILQ_1Opvk5t21-e59UKj_AMtd-lSuKdaH9m69S-Fe12cOqqa39GucUFLFVHJHbKp5bGOT121Y0DBOJ0zUxIzmEB4W8tqGFxPdwCAM4wsln-nCsjJVHRvttrnW_8usassCwESsHwzJg3kg2bn7XTtUW5J9r5wuug52HCx8dVLl31p8kYkukqtN19PtRrJOiet_rJDdheHT8KjvCaQsLJ50eqjsPDw6OTqBHRLg-l1v_5BB6PSysCv8p3RxEvZhsA_OGOGZBTWcYhjEO3tiXujigmbZx7rvFPNboxjxWK-f-

benchmarks/CMakeLists.txt

@@ -28,7 +28,7 @@ function(ConfigureBench BENCH_NAME)
     target_include_directories(${BENCH_NAME} PRIVATE
                                              "${CMAKE_CURRENT_SOURCE_DIR}")
     target_compile_options(${BENCH_NAME} PRIVATE --compiler-options=-Wall --compiler-options=-Wextra
-      --compiler-options=-Werror -Wno-deprecated-gpu-targets --expt-extended-lambda)
+      --compiler-options=-Werror -Wno-deprecated-gpu-targets --expt-extended-lambda -lineinfo)
     # Add GCC-specific warning suppression only for GCC
     if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
       target_compile_options(${BENCH_NAME} PRIVATE -Xcompiler -Wno-subobject-linkage)
@@ -104,3 +104,8 @@ ConfigureBench(HYPERLOGLOG_BENCH
 ConfigureBench(BLOOM_FILTER_BENCH
   bloom_filter/add_bench.cu
   bloom_filter/contains_bench.cu)
+
+###################################################################################################
+# - roaring_bitmap benchmarks ---------------------------------------------------------------------
+ConfigureBench(ROARING_BITMAP_BENCH
+  roaring_bitmap/contains_bench.cu)

benchmarks/roaring_bitmap/contains_bench.cu

@@ -0,0 +1,96 @@
+/*
+ * 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 <filesystem>
+#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
+  auto const file_size = std::filesystem::file_size(bitmap_file);
+
+  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::experimental::roaring_bitmap<T> roaring_bitmap(thrust::raw_pointer_cast(buffer.data()));
+
+  thrust::device_vector<T> items(num_items);
+
+  key_generator gen{};
+  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);

cmake/roaring_testdata.cmake

@@ -0,0 +1,41 @@
+# =============================================================================
+# 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")
+
+message(STATUS "Roaring Bitmap test data downloaded to: ${CUCO_ROARING_DATA_DIR}")
+
+# Define macro only when data is available
+add_compile_definitions(CUCO_ROARING_DATA_DIR="${CUCO_ROARING_DATA_DIR}")

examples/CMakeLists.txt

@@ -52,3 +52,4 @@ ConfigureExample(STATIC_MULTIMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/
 ConfigureExample(HYPERLOGLOG_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/hyperloglog/host_bulk_example.cu")
 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(ROARING_BITMAP_HOST_BULK_EXAMPLE "${CMAKE_CURRENT_SOURCE_DIR}/roaring_bitmap/host_bulk_example.cu")