Roaring bitmaps discussion and scoping #725

This issue is intended to discuss the feasability and scope of a potential cuco::roaring_bitmap implementation.

Background

Roaring Bitmaps [paper] are a compressed bitmap format designed to efficiently represent sets of 32- or 64-bit unsigned integers. They aim to outperform traditional methods such as hash- or tree-based set data structures, bitset, and run-length encoding (RLE) methods like WAH, EWAH, and Concise by offering superior speed and compression, especially on unsorted or sparse data.

For simplicity reasons I will focus on the 32-bit implementation in the following paragraphs.

The concept is drivenby the following observations comparing the abovementioned set implementations:

• Bitsets are very efficient in case the set has a high density given its domain. In other words, if the set of integers covers a significant portion of the 32-bit integer domain [0, 2^32-1]. Operations such as membership tests, set unions, or set intersections can be implemented using efficient bitwise operations.
• In case the set density is low, the bitset will be mostly empty which makes it storage inefficient. Instead, sparse sets can be better represented using dense data structures such as hash sets, trees, or sorted arrays. Roaring chooses the latter for this situation. Set operations over sorted arrays require more irregular algorithms such as binary search or merge steps.

A roaring bitmap mixes both the bitset and sorted array concept to get the "best of both worlds" as follows:
For a given set of input integers, we first partition it into 65k buckets/containers based on the 16 MSB bits of the elements.
(img1_placeholder, GH has some issues with image upload right now. I will try again later)
Each container represents either a sorted array or a bitset that stores the 16 LSB bits of each assigned element.
(img2_placeholder)
Combining the container index with the 16bit elements stored in it reconstructs the original set members.
A sorted array container is a dynamically resizeable data structure such as a std::vector and can hold up-to 4096 elements (8KiB). When this limit is exceeded, the array gets converted into a bitmap (also 8KiB to cover the 16-bit domain). It becomes directly clear that in case the container holds less than 4096 elements, storing the subset as a compact array has a better storage efficiency than the bitset equivalent.
As a final optimization, roaring offers an opt-in step called runOptimize() which converts containers that contain large runs of consecutive elements into an RLE representation for even better storage efficiency and subsequently makes the entire data structure immutable.

Discussion

The concept of a roaring bitmap fits perfectly into the scope of cuco. However, I see some challenges regarding implementation details.

Dynamic allocation issue

Roaring bitmaps heavily rely on fine-grained dynamic memory allocation which is unfortunately not easily portable to the GPU.
In fact, this feature is what makes roaring so powerful by keeping the memory footprint small during construction and transformation.
If we want to keep this feature, we would need a memory pool such as RMM to back the allocations, as well as extensive synchronization mechanisms, e.g., we need to block concurrent threads accessing one container while it is being resized.
One could argue that relying on a memory pool that pre-allocates large chunks of memory does not help with storage efficiency.
We do not plan to add RMM or any other memory pool implementation to cuco in the foreseeable future.

Proposed Workaround

Instead of starting each container as a dynamically growing array, we could instead use bitsets during construction and then convert them into sorted arrays/RLE containers on-demand, similar to the aforementioned runOptimize() function, making the data structure immutable.

I have the following workflow in mind:

1. Sort the input elements, e.g., cub::DeviceRadixSort.
2. Count the number of elements assigned to each container. (can also be done through some device-level cub or thrust primitives)
3. For each non-zero count check if a bitmap for this container has already been allocated. If not allocate it. (cudaMallocAsync or similar)
4. Merge the input elements into their assigned bitset using, e.g., atomicOr. (custom kernel)

After all elements have been added into the roaring bitmap, we can then compress the data structure by calling runOptimize()or whatever we want to name it.

There is one problem with this approach. The domain of 32-bit integers can be represented using a 512MiB bitset. So even in the worst case (only one element per container), we can still fit the uncompressed data structure in-memory. If we switch to 64-bit integers, however, then this worst-case memory footprint explodes to 2 exabyte if my napkin math is correct.

Questions

• What's the target application for this data structure in RAPIDS/Spark? Are we primarily interested in compression ratio, e.g., for better network utilization or is it about matching existing CPU implementations in Spark?
• Does it have to be a roaring bitmap or can we emulate it with say a decently sized cuco::bloom_filter which will likely give you an even better compression ratio at a moderate false-positive rate?

Resources

• arxiv paper
• screeshots taken from this video