Fix run-container bitmaps without serialized offsets and add a minimal repro test

Author: sleeepyjack

include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh

@@ -174,12 +174,16 @@ class roaring_bitmap_impl<cuda::std::uint32_t> {
   __device__ bool contains_container(cuda::std::uint16_t lower, cuda::std::uint32_t index) const
   {
     cuda::std::uint32_t offset;
-    cuda::std::byte const* offset_ptr =
-      storage_ref_.container_offsets() + index * sizeof(cuda::std::uint32_t);
-    if (offsets_aligned_) {
-      offset = aligned_load<cuda::std::uint32_t>(offset_ptr);
+    if (storage_ref_.metadata().offsets_in_serialized_data) {
+      cuda::std::byte const* offset_ptr =
+        storage_ref_.container_offsets() + index * sizeof(cuda::std::uint32_t);
+      if (offsets_aligned_) {
+        offset = aligned_load<cuda::std::uint32_t>(offset_ptr);
+      } else {
+        offset = misaligned_load<cuda::std::uint32_t>(offset_ptr);
+      }
     } else {
-      offset = misaligned_load<cuda::std::uint32_t>(offset_ptr);
+      offset = storage_ref_.metadata().computed_offsets[index];
     }
     cuda::std::byte const* container = storage_ref_.data() + offset;
     if (storage_ref_.metadata().has_run and check_bit(storage_ref_.run_container_bitmap(), index)) {
@@ -381,4 +385,4 @@ class roaring_bitmap_impl<cuda::std::uint64_t> {
   storage_ref_type storage_ref_;
 };
 
-}  // namespace cuco::experimental::detail
+}  // namespace cuco::experimental::detail

tests/roaring_bitmap/contains_test.cu

@@ -21,11 +21,13 @@
 #include <cuda/std/cstdint>
 #include <cuda/std/type_traits>
 #include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
 #include <thrust/logical.h>
 #include <thrust/universal_vector.h>
 
 #include <catch2/catch_test_macros.hpp>
 
+#include <cstring>
 #include <filesystem>
 #include <fstream>
 #include <string>
@@ -90,8 +92,61 @@ bool check(std::string const& bitmap_file_path)
     thrust::all_of(contained.begin(), contained.end(), ::cuda::std::identity{});
   return all_contained;
 }
+
+std::vector<cuda::std::byte> make_run_container_no_offsets_bitmap()
+{
+  std::vector<cuda::std::byte> bytes;
+  auto append_u16 = [&bytes](cuda::std::uint16_t value) {
+    bytes.push_back(static_cast<cuda::std::byte>(value & 0xFF));
+    bytes.push_back(static_cast<cuda::std::byte>((value >> 8) & 0xFF));
+  };
+  auto append_u32 = [&bytes](cuda::std::uint32_t value) {
+    bytes.push_back(static_cast<cuda::std::byte>(value & 0xFF));
+    bytes.push_back(static_cast<cuda::std::byte>((value >> 8) & 0xFF));
+    bytes.push_back(static_cast<cuda::std::byte>((value >> 16) & 0xFF));
+    bytes.push_back(static_cast<cuda::std::byte>((value >> 24) & 0xFF));
+  };
+
+  constexpr cuda::std::uint32_t serial_cookie = 12347;
+  constexpr cuda::std::uint16_t key           = 0;
+  constexpr cuda::std::uint16_t card_minus_1  = 2;  // card = 3
+  constexpr cuda::std::uint16_t num_runs      = 1;
+  constexpr cuda::std::uint16_t run_start     = 1;
+  constexpr cuda::std::uint16_t run_length    = 2;  // 3 values total
+
+  append_u32(serial_cookie);                            // 1 container encoded in cookie
+  bytes.push_back(static_cast<cuda::std::byte>(0x01));  // run container bitmap (1 byte)
+  append_u16(key);
+  append_u16(card_minus_1);
+  append_u16(num_runs);
+  append_u16(run_start);
+  append_u16(run_length);
+
+  return bytes;
+}
 }  // namespace
 
+TEST_CASE("roaring_bitmap run container without offsets", "[roaring_bitmap]")
+{
+  auto const bytes = make_run_container_no_offsets_bitmap();
+  thrust::universal_host_pinned_vector<cuda::std::byte> buffer(bytes.size());
+  std::memcpy(thrust::raw_pointer_cast(buffer.data()), bytes.data(), bytes.size());
+
+  cuco::experimental::roaring_bitmap<cuda::std::uint32_t> roaring_bitmap(
+    thrust::raw_pointer_cast(buffer.data()));
+
+  thrust::device_vector<cuda::std::uint32_t> keys{1, 2, 3, 4};
+  thrust::device_vector<bool> contained(keys.size(), false);
+
+  roaring_bitmap.contains(keys.begin(), keys.end(), contained.begin());
+
+  thrust::host_vector<bool> contained_h = contained;
+  REQUIRE(contained_h[0]);
+  REQUIRE(contained_h[1]);
+  REQUIRE(contained_h[2]);
+  REQUIRE_FALSE(contained_h[3]);
+}
+
 TEST_CASE("roaring_bitmap bulk contains from RoaringFormatSpec testdata", "[roaring_bitmap]")
 {
 #ifndef CUCO_ROARING_DATA_DIR
@@ -131,4 +186,4 @@ TEST_CASE("roaring_bitmap bulk contains from RoaringFormatSpec testdata", "[roar
     REQUIRE(check<cuda::std::uint64_t>(path));
   }
 #endif
-}
+}