fz::RREStage Class Reference

#include <rre_stage.h>

Inheritance diagram for fz::RREStage:

Public Member Functions
void	setInverse (bool inv) override
bool	isGraphCompatible () const override
size_t	getRequiredInputAlignment () const override
FusionSpec	getFusionSpec () const override
FusedOpDecl	getFusedOp () const override
	Chunk-cooperative coder op (the swappable variable-length sink). Stateless.
void	setFusedArchiveResult (size_t archive_bytes, size_t orig_bytes) override
	Base-class tail hook → the existing coder result setter (archive, orig).
void	setFusedResult (size_t archive_bytes, size_t orig_bytes)
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
void	postStreamSync (fz::stream_t stream) 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
size_t	estimateScratchBytes (const std::vector< size_t > &input_sizes) const override
uint16_t	getStageTypeId () const override
uint8_t	getOutputDataType (size_t) const override
size_t	serializeHeader (size_t output_index, 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 std::vector< std::string >	getOutputNames () const
int	getOutputIndex (const std::string &name) const
virtual uint8_t	getInputDataType (size_t) 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	setTerminalOutput (bool terminal)
virtual EncodingOracleDecl	getEncodingOracle () const
virtual bool	bindDownstreamEncodingOracle (const EncodingOracleDecl &)
virtual std::vector< FusedAuxOutputDecl >	getFusedAuxOutputs () const
virtual void	primeFusedForwardState (const FusedPrimeContext &)
virtual void	setFusedSideOutput (int, size_t)

Additional Inherited Members
Static Public Member Functions inherited from fz::Stage
static constexpr bool	isSupportedOnBackend ()

Detailed Description

LC RRE repeated-word bitmap reducer.

setChunkSize(bytes) — chunk size (default 16384; only 16384 is supported). setWordSize(bytes) — word granularity 1/2/4/8 (default 1).

Note

Prior work: GPU kernels are a faithful port of d_RRE.h and d_repetition_elimination.h from the LC framework (Burtscher et al., BSD-3-Clause). See THIRD_PARTY.md and modules/coders/lc_common/lc_chunk_components.cuh.

CUDA Graph capture is supported for compression only. The inverse path reads the stream header with blocking D2H copies before it can launch the decode kernel (same constraint as RZEStage).

Member Function Documentation

setInverse()

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

isGraphCompatible()

bool fz::RREStage::isGraphCompatible ( ) const

inlineoverridevirtual

Forward (compress) is graph-capturable; inverse reads the stream header with blocking D2H copies and therefore is not. See RZEStage for the rationale (decompress-only graph capture has no practical use case).

Reimplemented from fz::Stage.

getRequiredInputAlignment()

size_t fz::RREStage::getRequiredInputAlignment ( ) const

inlineoverridevirtual

Minimum input size alignment in bytes. Chunked stages return their chunk size; the pipeline uses the LCM of all stage alignments at finalize() to transparently zero-pad the input. Default: 1 (no alignment requirement).

Reimplemented from fz::Stage.

getFusionSpec()

FusionSpec fz::RREStage::getFusionSpec ( ) const

inlineoverridevirtual

Fusion contract: how this stage accesses its input, which decides whether the fusion planner may fold it into a single kernel with its neighbours (see include/stage/fusion.h and docs/codebase_notes.md CN-FUSE-PROOF).

Default is Unfusable — a stage is only ever fused if it opts in by overriding this. Stages whose fusability depends on configuration (e.g. a quantizer is a pure Map only in linear mode) must reflect that here. Forward-mode only; an inverse stage should report Unfusable.

Reimplemented from fz::Stage.

setFusedResult()

void fz::RREStage::setFusedResult ( size_t  archive_bytes, size_t  orig_bytes  )

inline

Set by a fused runner that produced this coder's archive without execute(). Also sets cached_orig_bytes_ (uncompressed input size) — the inverse output buffer is sized from it (else it defaults to the compressed size → OOB).

execute()

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

postStreamSync()

void fz::RREStage::postStreamSync ( fz::stream_t  stream )

overridevirtual

Called after dag->execute() and stream sync, before compress() returns. Use for D2H transfers that must not block mid-pipeline (e.g. Lorenzo's outlier count readback). The stream is already idle so a plain cudaMemcpy is safe here.

Reimplemented from fz::Stage.

getName()

std::string fz::RREStage::getName ( ) const

inlineoverridevirtual

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

Implements fz::Stage.

estimateOutputSizes()

std::vector< size_t > fz::RREStage::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()

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

overridevirtual

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

Implements fz::Stage.

getActualOutputSize()

size_t fz::RREStage::getActualOutputSize ( int  index ) const

overridevirtual

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.

estimateScratchBytes()

size_t fz::RREStage::estimateScratchBytes ( const std::vector< size_t > &  input_sizes ) const

inlineoverridevirtual

Forward pass allocates four persistent pool arrays proportional to n_chunks = ceil(input_bytes / chunk_size_): d_scratch_ : n_chunks * chunk_size_ (per-chunk worst-case output) d_sizes_dev_ : n_chunks * 4 (raw compressed sizes) d_clean_dev_ : n_chunks * 4 (flag-stripped sizes) d_dst_off_dev_: n_chunks * 4 (exclusive prefix-sum offsets)

Reimplemented from fz::Stage.

getStageTypeId()

uint16_t fz::RREStage::getStageTypeId ( ) const

inlineoverridevirtual

Stage type identifier written into the FZM file header.

Implements fz::Stage.

getOutputDataType()

uint8_t fz::RREStage::getOutputDataType ( size_t  output_index ) const

inlineoverridevirtual

DataType enum of the given output port.

Implements fz::Stage.

serializeHeader()

size_t fz::RREStage::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()

void fz::RREStage::deserializeHeader ( const uint8_t *  header_buffer, size_t  size  )

inlineoverridevirtual

Restore stage config from header_buffer during decompression.

Reimplemented from fz::Stage.

getMaxHeaderSize()

size_t fz::RREStage::getMaxHeaderSize ( size_t  output_index ) const

inlineoverridevirtual

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

Reimplemented from fz::Stage.

saveState()

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