|
FZGPUModules 2.0
GPU-accelerated modular compression pipelines
|
#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< FusedAuxOutputDecl > | getFusedAuxOutputs () 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 () |
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):
output — residuals for the selected variant (T, one per element)modes — one byte per tile: bit 0 = order 2, bit 1 = centeringmeans — 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).
| T | Signed integer residual/code type: currently int32_t only. |
|
inlineoverridevirtual |
Switch between forward (compression) and inverse (decompression) mode. Affects getNumInputs()/getNumOutputs() for stages with asymmetric port counts.
Reimplemented from fz::Stage.
|
inline |
The absolute bound the forward pass actually quantized against, after NOA/PREL conversion. Valid only after compress(); 0 before that.
|
inlineoverridevirtual |
Bind an exact additive fixed-rate policy exposed by a directly connected AdaptiveBitpack stage. See AdaptiveLorenzoStage's identical method.
Reimplemented from fz::Stage.
|
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.
|
inlineoverridevirtual |
Escaping outputs a generated fused op can produce beside its main port.
Reimplemented from fz::Stage.
|
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.
|
inlineoverridevirtual |
Human-readable name used in error messages and debug output.
Implements fz::Stage.
|
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.
|
overridevirtual |
|
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.
|
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.
|
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.
|
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.
|
inlineoverridevirtual |
|
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.
|
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.
|
inlineoverridevirtual |
|
inlineoverridevirtual |
DataType enum of the given output port.
Implements fz::Stage.
|
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.
|
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.
|
inlineoverridevirtual |
Restore stage config from header_buffer during decompression.
Reimplemented from fz::Stage.
|
inlineoverridevirtual |
Maximum bytes this stage writes into its per-output FZM header slot.
Reimplemented from fz::Stage.