|
FZGPUModules 2.0
GPU-accelerated modular compression pipelines
|
#include <szx_stage.h>
Inheritance diagram for fz::SZxStage< T >:Public Member Functions | |
| void | setInverse (bool inv) override |
| bool | isGraphCompatible () const override |
| double | getConstantBlockFraction () const |
| 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 |
| size_t | estimateScratchBytes (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 |
| std::vector< std::string > | getRunNotes () 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 | onFinalize (size_t, MemoryPool *) |
| virtual size_t | estimateDeviceFootprintBytes (size_t) const |
| virtual size_t | estimatePinnedFootprintBytes (size_t) const |
| virtual void | setTerminalOutput (bool terminal) |
| virtual FusionSpec | getFusionSpec () const |
| virtual FusedOpDecl | getFusedOp () const |
| virtual EncodingOracleDecl | getEncodingOracle () const |
| virtual bool | bindDownstreamEncodingOracle (const EncodingOracleDecl &) |
| virtual std::vector< FusedAuxOutputDecl > | getFusedAuxOutputs () const |
| virtual void | primeFusedForwardState (const FusedPrimeContext &) |
| virtual void | setFusedArchiveResult (size_t, size_t) |
| virtual void | setFusedSideOutput (int, size_t) |
Additional Inherited Members | |
Static Public Member Functions inherited from fz::Stage | |
| static constexpr bool | isSupportedOnBackend () |
SZx ultrafast error-bounded compressor.
Per block of block_size elements the forward pass:
max - min <= 2*eb (the whole block is representable by one value within the bound) or non-constant otherwise — a 2-bit type code per block;[-2^b, 2^b) where b is the block's required bit width, and bit-pack them (no Huffman).The archive is: [meta region: 2-bit type codes + per-block bit widths] followed by [payload region: reference values + packed residuals]. Output size is data dependent (constant blocks are ~one value), so estimateOutputSizes() returns a safe upper bound and postStreamSync() reads the true size back.
| T | float or double. |
|
inlineoverridevirtual |
Switch between forward (compression) and inverse (decompression) mode. Affects getNumInputs()/getNumOutputs() for stages with asymmetric port counts.
Reimplemented from fz::Stage.
|
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.
|
inline |
Fraction of blocks classified constant on the last forward encode — a cheap compressibility probe, reported through getRunNotes().
|
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.
|
overridevirtual |
|
inlineoverridevirtual |
Human-readable name used in error messages and debug output.
Implements fz::Stage.
|
overridevirtual |
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.
|
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 |
|
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.
|
overridevirtual |
Notes about what this stage actually did on the last run, when that differs from what was configured in a way that affects comparability.
Motivating case: HuffmanStage silently falls back to an Adaptive book when a PerBlock/Fixed build drives a symbol past the 27-bit code field. The fallback is correct — it does not relax the error bound — but a field encoded with a different codebook is not compression-ratio comparable to one that was not, and getBookSource() deliberately keeps reporting what was asked for. Without a channel like this, a benchmark row records the two cases identically and the difference is unrecoverable after the fact.
Returns short stable machine-readable tokens (e.g. "adaptive_fallback"), not prose — these are meant to land in a benchmark row and be grouped on. Empty by default; a stage that never surprises its caller need not implement 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.
|
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.