#include <bitpack_stage.h>

Inheritance diagram for fz::BitpackStage< T >:

Public Member Functions
void	setInverse (bool inv) override

void	setNBits (uint8_t nbits)

void	setBase (T base)

void	setShift (uint8_t shift)

void	setAutoDetect (bool enable)

void	setAutoBase (bool enable)

void	setAutoShift (bool enable)

void	setAdaptive (bool enable)

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

bool	isGraphCompatible () const 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 std::vector< std::string >	getRunNotes () const

virtual void	setDims (const std::array< size_t, 3 > &dims)

virtual void	onFinalize (size_t, MemoryPool *)

virtual size_t	estimateDeviceFootprintBytes (size_t) const

virtual size_t	estimatePinnedFootprintBytes (size_t) const

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

template<typename T>
class fz::BitpackStage< T >

Bit-packing stage.

Forward: T[] → uint8_t[] Pack (v - base) >> shift using only the low nbits bits. Inverse: uint8_t[] → T[] Unpack, then restore (packed << shift) + base.

Template Parameters

T	Input element type: `uint8_t`, `uint16_t`, or `uint32_t`.

Member Function Documentation

◆ setInverse()

template<typename T >

void fz::BitpackStage< T >::setInverse ( bool inverse )

inlineoverridevirtual

Switch between forward (compression) and inverse (decompression) mode. Affects getNumInputs()/getNumOutputs() for stages with asymmetric port counts.

Reimplemented from fz::Stage.

◆ setNBits()

template<typename T >

void fz::BitpackStage< T >::setNBits ( uint8_t nbits )

inline

Set the number of bits per element.

Must be a power of two between 1 and 8*sizeof(T) inclusive. Allowed values: uint8_t : 1, 2, 4, 8 uint16_t : 1, 2, 4, 8, 16 uint32_t : 1, 2, 4, 8, 16, 32

Ignored during forward execute when setAutoDetect(true) is active.

◆ setBase()

template<typename T >

void fz::BitpackStage< T >::setBase ( T base )

inline

Frame-of-reference base: subtracted from every element before packing and added back on unpack. Removes dead high bits when the values are clustered away from zero. Always lossless.

Ignored during forward execute when setAutoBase(true) is active.

◆ setShift()

template<typename T >

void fz::BitpackStage< T >::setShift ( uint8_t shift )

inline

Right-shift applied after the base subtraction: (v - base) >> shift. Removes dead low bits. Must be in [0, 8*sizeof(T) - 1].

Warning: Lossy unless every (v - base) has shift trailing zeros. The inverse restores (packed << shift) + base, so any dropped low bits are gone. Use setAutoShift(true) to pick the largest shift that is provably lossless for the data at hand.

Ignored during forward execute when setAutoShift(true) is active.

◆ setAutoDetect()

template<typename T >

void fz::BitpackStage< T >::setAutoDetect ( bool enable )

inline

Enable automatic bit-width detection.

When true, forward execute scans the input for its maximum value and selects the smallest valid power-of-two nbits that covers the shifted, base-subtracted range. The chosen nbits is stored in the serialized header so the inverse pass can unpack correctly.

After compress(), getNBits() reflects the detected value.

Incompatible with CUDA Graph capture: isGraphCompatible() returns false while any auto-detect mode is enabled.

◆ setAutoBase()

template<typename T >

void fz::BitpackStage< T >::setAutoBase ( bool enable )

inline

Enable automatic frame-of-reference base selection: forward execute min-reduces the input and uses that minimum as base. Lossless. After compress(), getBase() reflects the detected value.

◆ setAutoShift()

template<typename T >

void fz::BitpackStage< T >::setAutoShift ( bool enable )

inline

Enable automatic shift selection: forward execute OR-reduces every (v - base) and uses the trailing-zero count of that OR as shift — i.e. the largest shift that drops no information. Always lossless, unlike a hand-set setShift(). After compress(), getShift() reflects the detected value.

◆ setAdaptive()

template<typename T >

void fz::BitpackStage< T >::setAdaptive ( bool enable )

inline

Convenience: enable auto base, auto shift, and auto nbits together — the fully adaptive lossless mode. Each element becomes (v - min) >> ctz(OR of (v - min)), packed at the tightest power-of-two width that fits.

◆ execute()

template<typename T >

void fz::BitpackStage< T >::execute	(	fz::stream_t	stream,
		MemoryPool *	pool,
		const std::vector< void * > &	inputs,
		const std::vector< void * > &	outputs,
		const std::vector< size_t > &	sizes
	)

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()

template<typename T >

std::string fz::BitpackStage< T >::getName ( ) const

inlineoverridevirtual

Human-readable name used in error messages and debug output.

Implements fz::Stage.

◆ estimateOutputSizes()

template<typename T >

std::vector< size_t > fz::BitpackStage< T >::estimateOutputSizes ( const std::vector< size_t > & input_sizes ) const

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()

template<typename T >

std::unordered_map< std::string, size_t > fz::BitpackStage< T >::getActualOutputSizesByName ( ) const

inlineoverridevirtual

Actual output sizes after execute(), keyed by output port name.

Implements fz::Stage.

◆ getActualOutputSize()

template<typename T >

size_t fz::BitpackStage< T >::getActualOutputSize ( int index ) const

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()

template<typename T >

uint16_t fz::BitpackStage< T >::getStageTypeId ( ) const

inlineoverridevirtual

Stage type identifier written into the FZM file header.

Implements fz::Stage.

◆ getOutputDataType()

template<typename T >

uint8_t fz::BitpackStage< T >::getOutputDataType ( size_t output_index ) const

inlineoverridevirtual

DataType enum of the given output port.

Implements fz::Stage.

◆ getInputDataType()

template<typename T >

uint8_t fz::BitpackStage< T >::getInputDataType ( size_t ) const

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()

template<typename T >

size_t fz::BitpackStage< T >::serializeHeader	(	size_t	output_index,
		uint8_t *	header_buffer,
		size_t	max_size
	)		const

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()

template<typename T >

void fz::BitpackStage< T >::deserializeHeader	(	const uint8_t *	header_buffer,
		size_t	size
	)

inlineoverridevirtual

Restore stage config from header_buffer during decompression.

Reimplemented from fz::Stage.

◆ getMaxHeaderSize()

template<typename T >

size_t fz::BitpackStage< T >::getMaxHeaderSize ( size_t output_index ) const

inlineoverridevirtual

Maximum bytes this stage writes into its per-output FZM header slot.

Reimplemented from fz::Stage.

◆ saveState()

template<typename T >

void fz::BitpackStage< T >::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.

◆ isGraphCompatible()

template<typename T >

bool fz::BitpackStage< T >::isGraphCompatible ( ) const

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.

Public Member Functions

Additional Inherited Members

Detailed Description

Member Function Documentation

◆ setInverse()

◆ setNBits()

◆ setBase()

◆ setShift()

◆ setAutoDetect()

◆ setAutoBase()

◆ setAutoShift()

◆ setAdaptive()

◆ execute()

◆ getName()

◆ estimateOutputSizes()

◆ getActualOutputSizesByName()

◆ getActualOutputSize()

◆ getStageTypeId()

◆ getOutputDataType()

◆ getInputDataType()

◆ serializeHeader()

◆ deserializeHeader()

◆ getMaxHeaderSize()

◆ saveState()

◆ isGraphCompatible()