#include <huffman_stage.h>
Inheritance diagram for fz::HuffmanStage< T >:Public Member Functions | |
| void | setBklen (uint32_t bklen) |
| void | setEncodeMode (HuffmanEncodeMode mode) |
| void | setBookSource (HuffmanBookSource src) |
| void | setFixedBookFromFreq (const uint32_t *h_freq, uint32_t n) |
| void | setFixedBookFromModel (const HuffmanBookSpec &spec) |
| void | setValidateSymbolRange (bool on) |
| const std::vector< uint32_t > & | getFixedBookFreq () const |
| Frequency table backing the fixed codebook; empty when none has been set. | |
| void | setAdaptiveFloorShift (uint8_t shift) |
| uint8_t | getAdaptiveFloorShiftUsed () const |
| bool | getAdaptiveFallbackUsed () const |
| std::vector< std::string > | getRunNotes () const override |
| void | setRefitThreshold (float ratio) |
| void | setRefitInterval (uint32_t n) |
| uint32_t | getRefitCount () const |
| double | getFitBitsPerSymbol () const |
| bool | getLastUsedFineEncode () const |
| uint8_t | getLastMaxCodeLen () const |
| bool | hasBookSpec () const |
| void | setInverse (bool inv) override |
| bool | isGraphCompatible () const override |
| void | onFinalize (size_t estimated_inlen, MemoryPool *pool) override |
| size_t | estimateDeviceFootprintBytes (size_t inlen) const override |
| size_t | estimatePinnedFootprintBytes (size_t inlen) const override |
| void | execute (fz::stream_t stream, MemoryPool *pool, const std::vector< void * > &inputs, const std::vector< void * > &outputs, const std::vector< size_t > &sizes) override |
| std::string | getName () const override |
| std::vector< size_t > | estimateOutputSizes (const std::vector< size_t > &input_sizes) const override |
| std::unordered_map< std::string, size_t > | getActualOutputSizesByName () const override |
| size_t | getActualOutputSize (int index) const override |
| uint16_t | getStageTypeId () const override |
| uint8_t | getOutputDataType (size_t) const override |
| uint8_t | getInputDataType (size_t) const override |
| size_t | serializeHeader (size_t, uint8_t *buf, size_t max_size) const override |
| void | deserializeHeader (const uint8_t *buf, size_t size) override |
| size_t | getMaxHeaderSize (size_t) const override |
| void | saveState () override |
| Public Member Functions inherited from fz::Stage | |
| virtual size_t | getRequiredInputAlignment () const |
| virtual std::vector< std::string > | getOutputNames () const |
| int | getOutputIndex (const std::string &name) const |
| virtual void | setDims (const std::array< size_t, 3 > &dims) |
| virtual void | postStreamSync (fz::stream_t stream) |
| virtual size_t | estimateScratchBytes (const std::vector< size_t > &input_sizes) const |
Additional Inherited Members | |
| Static Public Member Functions inherited from fz::Stage | |
| static constexpr bool | isSupportedOnBackend () |
Detailed Description
class fz::HuffmanStage< T >
Huffman entropy coding stage.
Forward: T[] → uint8_t[] PHF-encoded bitstream with embedded phf_header. Inverse: uint8_t[] → T[] Decoded symbol stream.
- Note
- Prior work: PHF source files (
hf.h,hf_bk*.cc,hf_buf.cc,hf_canon.cc,hf_hl.cc,hf_kernels.cu,hf_impl.hh) are vendored and adapted from the cuSZ PHF codec (origin/v1.1.0_dev), by the cuSZ team (BSD-3-Clause). Changes are documented at the top of each file. SeeTHIRD_PARTY.md.
- Template Parameters
-
T Input element type: uint8_t,uint16_t, oruint32_t.
Member Function Documentation
◆ setBklen()
|
inline |
Set the Huffman codebook length (number of distinct symbols).
Must be ≤ 2^(8*sizeof(T)). Typical values: uint8_t : 256 (covers all possible byte values) uint16_t : 1024 (covers quantization codes in [-512, 511]) uint32_t : 1024 (must be set explicitly; 2^32 is too large for a codebook)
Set before the first compress() call. Changing bklen after the first execute() forces Buf reallocation on the next call (old buffers returned to pool, new ones allocated). Buf is also reallocated when inlen grows past the previously allocated capacity; shrinking inlen reuses the existing allocation. Default: 256 for uint8_t, 1024 for uint16_t/uint32_t. bklen is rounded up until sizeof(T) * bklen is a multiple of 4.
The reverse codebook occupies 4*(2*32) + sizeof(T)*bklen bytes — a 4-aligned 256-byte prefix plus the symbol table — and the bitstream begins immediately after it. The decode kernel reads that bitstream as uint32*, so the symbol table must not push it off a 4-byte boundary.
This used to round odd up to even, which is only sufficient for a 2-byte symbol. With uint8_t symbols the requirement is a multiple of 4, and a bklen of, say, 30 left the bitstream at offset 2 mod 4. That did not fault — it decoded to wrong symbols, silently: an all-uint8_t round-trip at bklen 30 returned symbol 0 where symbol 1 began and reported success. uint32_t symbols are aligned at any bklen. Rounding up is always safe — a larger codebook only adds unused symbol slots.
◆ setEncodeMode()
|
inline |
Select the encode algorithm for the forward path.
Must be called before the first compress() / execute() call (or before the next one if changing mode at runtime — triggers Buf reallocation). Default: HuffmanEncodeMode::Coarse.
- Warning
Fineis experimental and will not engage on realistic data. It requires every code in the book to fit in 8 bits (four codes per 32-bit shard) and silently falls back toCoarseotherwise. Use getLastUsedFineEncode() to check which path ran rather than assuming. The 8-bit ceiling is a structural barrier, not a tuning matter, and the fix is a 2x16-bit shard geometry rather than length-limiting: docs/codebase_notes.md CN-HF-2
◆ setBookSource()
|
inline |
Select where the forward path gets its codebook.
Switching to Fixed without first calling setFixedBookFromFreq() or setFixedBookFromModel() throws on the next forward execute(): there is no sensible default distribution to guess.
The encoded bitstream still carries its reverse codebook in every mode, so a Fixed-mode stream decompresses with a stock decoder and no extra configuration. Nothing about the file format changes.
◆ setFixedBookFromFreq()
| void fz::HuffmanStage< T >::setFixedBookFromFreq | ( | const uint32_t * | h_freq, |
| uint32_t | n | ||
| ) |
Supply the frequency table that defines the fixed codebook and switch to HuffmanBookSource::Fixed.
n becomes the codebook length (subject to the same odd-value rounding as setBklen(); the padding slot gets frequency 1). Every entry must be non-zero — a zero-frequency symbol receives no code, and encoding it would corrupt the stream rather than fail loudly.
The table is held in host memory only. It is not written to the stage header (the per-stage config slot is 128 bytes; a reverse codebook is far larger), so a pipeline reconstructed from a saved config gets PerBlock back unless the caller re-supplies the table. Decompression is unaffected — see setBookSource().
- Exceptions
-
std::invalid_argument if nis 0 or any frequency is 0.
◆ setFixedBookFromModel()
| void fz::HuffmanStage< T >::setFixedBookFromModel | ( | const HuffmanBookSpec & | spec | ) |
Synthesize the fixed codebook from an analytic distribution and switch to HuffmanBookSource::Fixed.
Call setBklen() first: the model is evaluated over [0, bklen). Every symbol gets a strictly positive frequency, so the resulting book can encode any in-range symbol.
◆ setValidateSymbolRange()
|
inline |
Verify on the GPU that every symbol is in [0, bklen) when a codebook is pinned. Default on.
PerBlock gets this check for free — the histogram kernel skips out-of-range symbols, so sum(h_freq) != inlen betrays them. Adaptive and Fixed skip the histogram once a book is pinned, and the encode kernel then indexes the codebook with the raw symbol: an out-of-range value reads past d_bk4 and produces a stream that cannot be decoded, silently.
The check costs one grid-stride pass over the input and one stream sync. The sync is unavoidable: the verdict has to be known before encode() launches, since an out-of-range index faults inside the encode kernel and cannot be reported after the fact. It is still much cheaper than the PerBlock path it replaces — a 4-byte device-to-host copy instead of bklen words, and no host-side tree build — but it is not free, and it makes the stage ineligible for CUDA Graph capture.
Turn it off when the symbol range is guaranteed by construction upstream (for example LorenzoQuantStage with setZigzagCodes(true) and bklen == 2 * quant_radius), or when capturing a graph. Doing so with genuinely out-of-range symbols is undefined behaviour — in practice it takes down the CUDA context.
◆ setAdaptiveFloorShift()
|
inline |
Dynamic-range clamp used by HuffmanBookSource::Adaptive, as a right shift of the largest sampled frequency: every symbol's frequency is floored at max(1, max_freq >> shift).
The floor does two jobs. It guarantees a code for every symbol in [0, bklen), including ones absent from the sampled block — without it, a symbol that shows up only in a later block would have no code at all. And it bounds the frequency dynamic range, which bounds Huffman depth: an unbounded range can demand codes wider than the 27 bits HuffmanWord<4> holds.
A larger shift tracks the sampled distribution more closely (better ratio); a smaller one is flatter and safer. The build halves the shift and retries when the resulting book does not fit, so this is a starting point, not a hard setting — shift 0 is a uniform book, which always fits. Default 24.
◆ getAdaptiveFloorShiftUsed()
|
inline |
Shift actually used by the last Adaptive build, after any retries. Equals getAdaptiveFloorShift() unless the book had to be flattened to fit.
◆ getAdaptiveFallbackUsed()
|
inline |
True once a PerBlock/Fixed build has fallen back to an Adaptive book because the histogram drove a symbol past the 27-bit code field.
The fallback is not a relaxation of the error bound — the bound belongs to the quantizer, and Adaptive only flattens the frequency floor until every code fits. It does change which book the stage used, so a caller comparing compression ratios across fields needs to know it happened; the configured getBookSource() is deliberately left untouched so it still reports what was asked for.
◆ getRunNotes()
|
inlineoverridevirtual |
Surfaces getAdaptiveFallbackUsed() to Pipeline::collectRunNotes(), so a benchmark row can tell a fallback-encoded field from a normally-encoded one. See Stage::getRunNotes().
Reimplemented from fz::Stage.
◆ setRefitThreshold()
|
inline |
Refit trigger for Adaptive: rebuild the codebook once the encoded bit rate degrades past ratio times the rate the book achieved when it was fitted.
A pinned book goes stale when the symbol distribution moves. Detecting that would normally cost the histogram this mode exists to avoid, but encode already reports total_nbit, so bits-per-symbol is free. When it degrades, the next call histograms and re-pins; the current call is already encoded and is not retroactively improved.
1.2 (the default) refits after a 20% bit-rate regression. Lower refits more eagerly, trading throughput for ratio; PerBlock is the limit of that trade. 0 disables refitting, pinning the first book forever.
Drift measurements, and the degenerate-first-block case this guards against: docs/codebase_notes.md CN-HF-3
◆ setRefitInterval()
|
inline |
Refit unconditionally every n compress calls (0 = never).
The bit-rate trigger above only fires when the rate gets worse, so it catches a distribution drifting toward less compressible — but not a block that is more compressible than the fitted one while still being badly served by its codebook. That case is real and measured (CN-HF-3); deciding it properly needs the fresh book's rate, which means the histogram this mode exists to avoid.
A periodic refit sidesteps the question: it costs one histogram every n calls, bounding how long a stale book can persist regardless of which direction the rate moved. n = 1 is exactly PerBlock.
Default 0. Set it when successive calls carry genuinely different data (separate variables rather than successive slabs of one field); leave it off when they are slabs of one field, where the bit-rate trigger suffices.
◆ getRefitCount()
|
inline |
Number of times Adaptive has re-fitted its codebook. Diagnostic: a steadily climbing count means the data is drifting faster than one book can track, and PerBlock may be the better choice.
◆ getFitBitsPerSymbol()
|
inline |
Bits per symbol the resident Adaptive book achieved on the block it was fitted to; 0 before the first fit. This is the baseline the refit trigger compares against.
◆ getLastUsedFineEncode()
|
inline |
Whether the last forward call actually ran the ReVISIT-lite fine kernel.
setEncodeMode(Fine) is a request, not a guarantee: the fine path packs four codes into a 32-bit shard and so requires every code in the book to fit in 8 bits. When the built book has a longer code, encode() silently falls back to the coarse path. A Fine pipeline can therefore run coarse for its whole life without any outward sign, which quietly invalidates any "fine vs coarse" measurement taken from it.
False before the first forward call. Always false in Coarse mode.
◆ getLastMaxCodeLen()
|
inline |
Longest Huffman code, in bits, in the book used by the last forward call; 0 before the first call. Reported in both encode modes, because in Coarse mode this is what says whether switching to Fine would take effect: Fine engages only when this is ≤ 8.
Under Fixed/Adaptive the book is resident across calls, so a single reading after the first call characterizes every later one.
◆ hasBookSpec()
|
inline |
True when the fixed book came from setFixedBookFromModel(), i.e. when it is described by a handful of numbers and so can be written to a TOML config. A book set from a raw frequency table is not reproducible this way.
◆ setInverse()
|
inlineoverridevirtual |
Switch between forward (compression) and inverse (decompression) mode. Affects getNumInputs()/getNumOutputs() for stages with asymmetric port counts.
Reimplemented from fz::Stage.
◆ isGraphCompatible()
|
inlineoverridevirtual |
Whether this stage is safe inside a CUDA Graph capture.
A stage is graph-compatible if execute() enqueues only device-side work (kernel launches, cudaMemcpyAsync D2D/H2D) and makes no host-synchronous calls. Override and return false if execute() contains D2H copies or dynamic decisions based on device data — the DAG will throw at setCaptureMode(true) time rather than producing a broken graph.
Default: true. Inverse-mode stages that do D2H reads (e.g. RZE inverse) must return false.
Reimplemented from fz::Stage.
◆ onFinalize()
|
overridevirtual |
Called by Pipeline::finalize() after buffer-size propagation.
Pre-allocates phf::Buf<T> from the pool using the estimated input size so PREALLOCATE mode commits all memory at finalize time. If estimated_inlen is 0 (no size hint available), allocation is deferred to the first execute() call.
Reimplemented from fz::Stage.
◆ estimateDeviceFootprintBytes()
|
overridevirtual |
Estimated persistent device memory this stage allocates outside the pool (via pool->allocatePersistentDevice). Used for total footprint reporting. Default: 0.
Reimplemented from fz::Stage.
◆ estimatePinnedFootprintBytes()
|
overridevirtual |
Estimated persistent pinned-host memory this stage allocates outside the pool (via pool->allocatePersistentPinned). Used for total footprint reporting. Default: 0.
Reimplemented from fz::Stage.
◆ execute()
|
overridevirtual |
Execute the stage. Inputs, outputs, and sizes are device pointers/bytes.
Stages may call cudaStreamSynchronize(stream) or issue blocking D2H copies when the algorithm requires it (e.g. Huffman histogram readback for codebook construction, ANS renormalization tables). Such stages must return false from isGraphCompatible() and must document the sync points.
Note: the DAG dispatches sibling nodes (same topological level) via a sequential CPU loop, each enqueuing to its own stream. A sync inside execute() blocks the CPU from dispatching subsequent siblings until the synced stream is idle — this delays parallel branches in wide DAGs. In a linear pipeline there are no siblings and no extra cost.
Implements fz::Stage.
◆ getName()
|
inlineoverridevirtual |
Human-readable name used in error messages and debug output.
Implements fz::Stage.
◆ estimateOutputSizes()
|
inlineoverridevirtual |
Estimate output buffer sizes given input sizes. Used for buffer allocation planning in PREALLOCATE mode — must be a safe upper bound; under-estimation causes buffer overruns.
Implements fz::Stage.
◆ getActualOutputSizesByName()
|
inlineoverridevirtual |
◆ getActualOutputSize()
|
inlineoverridevirtual |
Actual size of a single output by index after execute(). Avoids constructing the map for the common single-output case. Default delegates to getActualOutputSizesByName(); override to return directly from an internal field.
Reimplemented from fz::Stage.
◆ getStageTypeId()
|
inlineoverridevirtual |
◆ getOutputDataType()
|
inlineoverridevirtual |
DataType enum of the given output port.
Implements fz::Stage.
◆ getInputDataType()
|
inlineoverridevirtual |
Expected DataType of the given input port.
Used by Pipeline::finalize() to detect type mismatches between connected stages before any execution. Return DataType::UNKNOWN to opt out of checking — byte-transparent stages (Bitshuffle, RZE, RRE) and mock stages must return UNKNOWN; finalize() skips any connection where either side is UNKNOWN.
Reimplemented from fz::Stage.
◆ serializeHeader()
|
inlineoverridevirtual |
Serialize stage config into header_buffer (max 128 bytes) for the FZM file. Return the number of bytes written, or 0 if the stage has no config.
Reimplemented from fz::Stage.
◆ deserializeHeader()
|
inlineoverridevirtual |
Restore stage config from header_buffer during decompression.
Reimplemented from fz::Stage.
◆ getMaxHeaderSize()
|
inlineoverridevirtual |
Maximum bytes this stage writes into its per-output FZM header slot.
Reimplemented from fz::Stage.
◆ saveState()
|
inlineoverridevirtual |
Save/restore config state around a decompression pass. deserializeHeader() overwrites the stage's forward-pass config; saveState() is called before and restoreState() after so the stage returns to its original configuration.
Reimplemented from fz::Stage.
