FZGPUModules 2.0
GPU-accelerated modular compression pipelines
Loading...
Searching...
No Matches
fz::FusedQuantAdaptiveLorenzoStage< T > Class Template Reference

#include <quant_adaptive_lorenzo_stage.h>

+ Inheritance diagram for fz::FusedQuantAdaptiveLorenzoStage< T >:

Public Member Functions

void setInverse (bool inv) override
 
double getComputedAbsErrorBound () const
 
bool bindDownstreamEncodingOracle (const EncodingOracleDecl &decl) override
 
FusionSpec getFusionSpec () const override
 
std::vector< FusedAuxOutputDeclgetFusedAuxOutputs () 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< std::string > getOutputNames () const override
 
void postStreamSync (fz::stream_t stream) override
 
bool isGraphCompatible () const override
 
size_t estimateScratchBytes (const std::vector< size_t > &input_sizes) const override
 
std::vector< size_t > estimateOutputSizes (const std::vector< size_t > &input_sizes) const override
 
void saveState () override
 
std::unordered_map< std::string, size_t > getActualOutputSizesByName () const override
 
size_t getActualOutputSize (int index) const override
 
void setFusedSideOutput (int output_index, size_t bytes) override
 
uint16_t getStageTypeId () const override
 
uint8_t getOutputDataType (size_t output_index) const override
 
uint8_t getInputDataType (size_t input_index) 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
 
- Public Member Functions inherited from fz::Stage
virtual size_t getRequiredInputAlignment () 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 setTerminalOutput (bool terminal)
 
virtual FusedOpDecl getFusedOp () const
 
virtual FusionSpec getInverseFusionSpec () const
 
virtual FusedOpDecl getInverseFusedOp () const
 
virtual size_t getFusedInverseElementCount () const
 
virtual double getFusedForwardQuantStep () const
 
virtual double getFusedInverseDequantStep () const
 
virtual EncodingOracleDecl getEncodingOracle () const
 
virtual void primeFusedForwardState (const FusedPrimeContext &)
 
virtual void setFusedArchiveResult (size_t, size_t)
 
virtual void setFusedInverseResult (size_t)
 

Additional Inherited Members

- Static Public Member Functions inherited from fz::Stage
static constexpr bool isSupportedOnBackend ()
 

Detailed Description

template<typename T = int32_t>
class fz::FusedQuantAdaptiveLorenzoStage< T >

AdaptiveLorenzo with the upstream linear Quantizer folded into its forward kernel — the "M1" partial fusion (see FZGPUModules memory generic_fusion_plan.md's ARCHITECTURAL TENSION section, approach C, and quant_al_partial_fusion_probe.md for the validated probe this promotes to a real stage). Deletes the separate QuantizerStage kernel and its codes DRAM round-trip on compress only — decompress still runs two kernels (the existing cross-warp inverse scan, then a plain elementwise dequant), matching every other fusion result in this codebase (compress-side only).

Takes raw float input directly (no upstream Quantizer stage in the pipeline) and owns its own error-bound resolution — ABS/NOA/PREL, see ErrorBoundMode in fused/lorenzo_quant/lorenzo_quant.h, reusing the exact same computeValueBase() scan LorenzoQuantStage/QuantizerStage use. Quantization is the simple unbounded linear form (round(x * ebx2_r), matching QuantizerStage's linear_mode=true, linear_high_precision=false path bit-for-bit) — no overflow guard, unlike QuantizerStage: this is an explicit, narrower scope for an opt-in fused stage, not a regression in the staged path's safety net. Use AdaptiveLorenzoStage + a separate QuantizerStage instead if overflow protection matters for your data.

Forward outputs (compression):

  • [0] output — residuals for the selected variant (T, one per element)
  • [1] modes — one byte per tile: bit 0 = order 2, bit 1 = centering
  • [2] means — one T per tile (meaningful only where bit 1 is set)

Inverse (decompression): takes the three forward outputs, reconstructs the original float data directly (no downstream QuantizerStage needed).

Template Parameters
TSigned integer residual/code type: currently int32_t only.

Member Function Documentation

◆ setInverse()

template<typename T = int32_t>
void fz::FusedQuantAdaptiveLorenzoStage< 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.

◆ getComputedAbsErrorBound()

template<typename T = int32_t>
double fz::FusedQuantAdaptiveLorenzoStage< T >::getComputedAbsErrorBound ( ) const
inline

The absolute bound the forward pass actually quantized against, after NOA/PREL conversion. Valid only after compress(); 0 before that.

◆ bindDownstreamEncodingOracle()

template<typename T = int32_t>
bool fz::FusedQuantAdaptiveLorenzoStage< T >::bindDownstreamEncodingOracle ( const EncodingOracleDecl decl)
inlineoverridevirtual

Bind an exact additive fixed-rate policy exposed by a directly connected AdaptiveBitpack stage. See AdaptiveLorenzoStage's identical method.

Reimplemented from fz::Stage.

◆ getFusionSpec()

template<typename T = int32_t>
FusionSpec fz::FusedQuantAdaptiveLorenzoStage< T >::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 Elementwise only in linear mode) must reflect that here. Forward-mode only; an inverse stage should report Unfusable.

Reimplemented from fz::Stage.

◆ getFusedAuxOutputs()

template<typename T = int32_t>
std::vector< FusedAuxOutputDecl > fz::FusedQuantAdaptiveLorenzoStage< T >::getFusedAuxOutputs ( ) const
inlineoverridevirtual

Escaping outputs a generated fused op can produce beside its main port.

Reimplemented from fz::Stage.

◆ execute()

template<typename T = int32_t>
void fz::FusedQuantAdaptiveLorenzoStage< 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 = int32_t>
std::string fz::FusedQuantAdaptiveLorenzoStage< T >::getName ( ) const
inlineoverridevirtual

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

Implements fz::Stage.

◆ getOutputNames()

template<typename T = int32_t>
std::vector< std::string > fz::FusedQuantAdaptiveLorenzoStage< T >::getOutputNames ( ) const
inlineoverridevirtual

Output port names in order. Default: single port named "output". Multi-output stages (e.g. Lorenzo: "codes", "outliers") override this.

Reimplemented from fz::Stage.

◆ postStreamSync()

template<typename T = int32_t>
void fz::FusedQuantAdaptiveLorenzoStage< T >::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.

◆ isGraphCompatible()

template<typename T = int32_t>
bool fz::FusedQuantAdaptiveLorenzoStage< T >::isGraphCompatible ( ) const
inlineoverridevirtual

The means length is data-dependent (postStreamSync trim) and the forward pass does a D2H for NOA/PREL bound resolution — not graph-safe, same as AdaptiveLorenzoStage.

Reimplemented from fz::Stage.

◆ estimateScratchBytes()

template<typename T = int32_t>
size_t fz::FusedQuantAdaptiveLorenzoStage< T >::estimateScratchBytes ( const std::vector< size_t > &  input_sizes) const
overridevirtual

Peak persistent scratch bytes this stage holds in the MemoryPool.

Only count allocations that are drawn from the pool and kept alive across execute() calls. Transient scratch freed within execute() is already captured by the pool's high-water mark and must not be included. Used by CompressionDAG::computeTopoPoolSize() to size the release threshold.

Reimplemented from fz::Stage.

◆ estimateOutputSizes()

template<typename T = int32_t>
std::vector< size_t > fz::FusedQuantAdaptiveLorenzoStage< 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.

◆ saveState()

template<typename T = int32_t>
void fz::FusedQuantAdaptiveLorenzoStage< 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.

◆ getActualOutputSizesByName()

template<typename T = int32_t>
std::unordered_map< std::string, size_t > fz::FusedQuantAdaptiveLorenzoStage< T >::getActualOutputSizesByName ( ) const
inlineoverridevirtual

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

Implements fz::Stage.

◆ getActualOutputSize()

template<typename T = int32_t>
size_t fz::FusedQuantAdaptiveLorenzoStage< 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.

◆ setFusedSideOutput()

template<typename T = int32_t>
void fz::FusedQuantAdaptiveLorenzoStage< T >::setFusedSideOutput ( int  ,
size_t   
)
inlineoverridevirtual

Side-output hook: a fused runner that filled one of this stage's escaping output ports (a pipeline leaf, e.g. an outlier list) without calling execute() reports how many BYTES it wrote to port output_index. The stage updates whatever forward-computed state its serializeHeader depends on (e.g. the quantizer's outlier count) so the archive matches the fused result. Default no-op — single-output stages and stages the runner didn't feed ignore it.

Reimplemented from fz::Stage.

◆ getStageTypeId()

template<typename T = int32_t>
uint16_t fz::FusedQuantAdaptiveLorenzoStage< T >::getStageTypeId ( ) const
inlineoverridevirtual

Stage type identifier written into the FZM file header.

Implements fz::Stage.

◆ getOutputDataType()

template<typename T = int32_t>
uint8_t fz::FusedQuantAdaptiveLorenzoStage< T >::getOutputDataType ( size_t  output_index) const
inlineoverridevirtual

DataType enum of the given output port.

Implements fz::Stage.

◆ getInputDataType()

template<typename T = int32_t>
uint8_t fz::FusedQuantAdaptiveLorenzoStage< 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 = int32_t>
size_t fz::FusedQuantAdaptiveLorenzoStage< 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 = int32_t>
void fz::FusedQuantAdaptiveLorenzoStage< 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 = int32_t>
size_t fz::FusedQuantAdaptiveLorenzoStage< T >::getMaxHeaderSize ( size_t  output_index) const
inlineoverridevirtual

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

Reimplemented from fz::Stage.