|
FZGPUModules 2.0
GPU-accelerated modular compression pipelines
|
Header: modules/coders/raze/raze_stage.h
Class: fz::RAZEStage — no template parameters
Category: Coder (lossless)
Common instantiation:
Zero-Adaptive Reduction Encoding — the RAZE lossless component of the LC framework. RAZE is the auto-k generalization of RZEStage: rather than RZE's binary "word is exactly zero, or it's kept in full" test, it histograms how many top bits of each word (on its own — not compared against a predecessor) are zero across the whole chunk, picks one global cut keep (0 <= keep < word_size*8) that maximizes total bit savings, then:
bits - keep bits are all zero store only their bottom keep bits, bit-packed (not byte-aligned) across word boundaries.Where RZE can only exploit exact zeros, RAZE also captures small-magnitude nonzero values — e.g. a mostly-small-integer stream with a long tail of large outliers — without needing those small values to be literally zero to benefit.
Same chunk-size restriction and rationale as RZEStage (static __shared__ budget under the 48 KB cap). word_size selects the LC RAZE_1 / RAZE_2 / RAZE_4 / RAZE_8 variant.
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.
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.
Identical container to RZEStage — the per-chunk keep value RAZE adds lives inside each chunk's own compressed bytes (accounted for by that chunk's stored size), not in the container header, so the two stages share byte-for-byte the same host-side framing:
A chunk is stored verbatim (high-bit flag) when RAZE fails to shrink it. An all-zero chunk collapses to a 3-byte tag (RZE's equivalent fast path uses 2 bytes — RAZE's tag carries one extra byte because the general path's trailer is [keep][size_lo][size_hi] vs. RZE's [size_lo][size_hi]).
The GPU kernels in RAZEStage (modules/coders/lc_common/lc_chunk_components.cuh, d_PRencode/d_PRdecode<T, PartialReduceMode::ZERO>) are a faithful port of d_RAZE.h from the LC framework (Burtscher et al., Texas State University, BSD-3-Clause), sharing the histogram/keep-selection/bit-pack device code with RAREStage via a single template parameterized on the match predicate.
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.