What it does

CLOG is the upstream name of this lossless LC framework component; LC does not provide an expansion, and no logarithmic transform is performed. It splits each chunk into a fixed 32 subchunks (not configurable — tied to the 32-lane warp used to parallelize the per-subchunk reduction). For each subchunk, it determines how many leading zero bits all values share and stores only the remaining low bits. Equivalently, it computes the minimum bit width required by the subchunk and packs every value at that width.

Unlike RREStage/RAREStage/RZEStage/RAZEStage, CLOG has no auxiliary bitmap and no per-element full/dropped decision — every element in a subchunk shares the same packed width, decided purely from that subchunk's own data. T must be unsigned (uint8/16/32/64 only — CLOG has no signed word-size variants).

HCLOGStage is the auto-selecting sibling: it additionally tries a TCMS(zigzag) reinterpretation per subchunk and picks whichever needs fewer bits, at the cost of one flag bit per subchunk.

Stage settings

clog->setChunkSize(16384); // bytes; 4096, 8192, or 16384 (default 16384)

clog->setWordSize(1); // word granularity in bytes: 1, 2, 4, or 8 (default 1)

Same chunk-size restriction and rationale as RREStage (static __shared__ budget under the 48 KB cap). word_size selects the LC CLOG_1 / CLOG_2 / CLOG_4 / CLOG_8 variant.

Alignment requirement

Requires input to be a multiple of chunk_size bytes. The pipeline pads automatically when an upstream byte-oriented stage uses a matching block size.

Graph capture

Forward (compress) is CUDA-graph capturable. The inverse (decompress) path is not — it reads the stream header (original size, per-chunk sizes) with blocking device-to-host copies before launching the decode kernel. This mirrors RREStage/RZEStage/RAREStage/RAZEStage.

Typical pipeline

auto* clog = p.addStage<CLOGStage>();
clog->setWordSize(1);
p.connect(clog, upstream);
p.finalize();

Stream layout (forward output)

Identical container to RREStage/RZEStage — CLOG's internal per-chunk header (a short original-size tag followed by the 32-entry bit-width table) lives inside each chunk's own compressed bytes (accounted for by that chunk's stored size), not in the container header:

[uint32_t: original byte count]
[uint32_t: num_chunks]
[uint32_t x num_chunks: per-chunk compressed sizes (high bit set -> chunk stored raw)]
[compressed chunk data ...]

A chunk is stored verbatim (high-bit flag) when CLOG fails to shrink it — notably including any chunk where the per-subchunk bit-width table overhead outweighs the savings, which is common for high-entropy data spanning the full width of T (see RandomBytesRoundTrip in the test suite).

Acknowledgements

The GPU kernels in CLOGStage (modules/coders/lc_common/lc_clog_components.cuh, d_CLOGencode/d_CLOGdecode<T, CLogMode::PLAIN>) are a faithful port of d_CLOG.h from the LC framework (Burtscher et al., Texas State University, BSD-3-Clause), sharing the encode/decode device code with HCLOGStage via a single template parameterized on whether the TCMS fallback is enabled.

‍Noushin Azami, Alex Fallin, Brandon Burtchell, Andrew Rodriguez, Benila Jerald, Yiqian Liu, Anju Mongandampulath Akathoott, and Martin Burtscher. LC framework for synthesizing high-speed parallel lossless and error-bounded lossy data compression and decompression algorithms for CPUs and GPUs. https://github.com/burtscher/LC-framework

See THIRD_PARTY.md for the full license text.