|
FZGPUModules 2.0
GPU-accelerated modular compression pipelines
|
#include <stage.h>
Inheritance diagram for fz::Stage:Public Member Functions | |
| virtual void | execute (fz::stream_t stream, MemoryPool *pool, const std::vector< void * > &inputs, const std::vector< void * > &outputs, const std::vector< size_t > &sizes)=0 |
| virtual std::string | getName () const =0 |
| virtual size_t | getRequiredInputAlignment () const |
| virtual std::vector< std::string > | getOutputNames () const |
| int | getOutputIndex (const std::string &name) const |
| virtual std::vector< size_t > | estimateOutputSizes (const std::vector< size_t > &input_sizes) const =0 |
| virtual std::unordered_map< std::string, size_t > | getActualOutputSizesByName () const =0 |
| virtual size_t | getActualOutputSize (int index) const |
| virtual void | setInverse (bool inverse) |
| virtual uint16_t | getStageTypeId () const =0 |
| virtual uint8_t | getOutputDataType (size_t output_index) const =0 |
| virtual uint8_t | getInputDataType (size_t) const |
| virtual size_t | serializeHeader (size_t output_index, uint8_t *header_buffer, size_t max_size) const |
| virtual void | deserializeHeader (const uint8_t *header_buffer, size_t size) |
| virtual void | saveState () |
| 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 | getMaxHeaderSize (size_t output_index) const |
| virtual bool | isGraphCompatible () const |
| virtual void | setTerminalOutput (bool terminal) |
| virtual size_t | estimateScratchBytes (const std::vector< size_t > &input_sizes) const |
| 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) |
Static Public Member Functions | |
| static constexpr bool | isSupportedOnBackend () |
Base class for all compression/decompression stages.
A stage is a single transformation in the pipeline (e.g. Lorenzo predictor, RLE encoder, bitshuffle). The pipeline interacts with stages exclusively through this interface — no downcasting or type-name branching anywhere in the pipeline or DAG code.
|
inlinestaticconstexpr |
Whether this stage type is supported on the backend the library was built for. Default true (every stage supports every backend) — a stage that doesn't (e.g. ANSStage on HIP/SYCL, whose vendored PTX lanemask assembly has no translation) hides this with its own static constexpr bool isSupportedOnBackend().
Deliberately static constexpr, not virtual: Pipeline::addStage<T>() must be able to check this before T is ever constructed, via if constexpr, so that on an unsupported backend the new T() branch is never instantiated at all — not merely never executed. That matters because an unsupported stage's own translation unit may be excluded from the build entirely (see CMakeLists.txt's HEADER_FILE_ONLY handling for ans_stage.cu on HIP); a virtual/instance method can't be called without an object to call it on, which would require the very constructor this mechanism exists to avoid referencing.
|
pure virtual |
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.
Implemented in fz::CLOGStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::GPULZStage, fz::HuffmanStage< T >, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
pure virtual |
Human-readable name used in error messages and debug output.
Implemented in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
inlinevirtual |
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 in fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::DifferenceStage< T, TOut, Mode >, fz::BitshuffleStage, and fz::TUPLStage.
|
inlinevirtual |
Output port names in order. Default: single port named "output". Multi-output stages (e.g. Lorenzo: "codes", "outliers") override this.
Reimplemented in fz::GPULZStage, fz::AdaptiveLorenzoStage< T >, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::LorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::MergeStage, fz::ROIBinSplitStage< TData >, and fz::LogTransformStage< TInput >.
|
inline |
Returns the index of a named output port, or -1 if not found.
|
pure virtual |
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.
Implemented in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
pure virtual |
Actual output sizes after execute(), keyed by output port name.
Implemented in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
inlinevirtual |
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 in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
inlinevirtual |
Switch between forward (compression) and inverse (decompression) mode. Affects getNumInputs()/getNumOutputs() for stages with asymmetric port counts.
Reimplemented in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::NegabinaryStage< TIn, TOut >, fz::ZigzagStage< TIn, TOut >, fz::LorenzoQuantStage< TInput, TCode >, fz::RLEStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::QuantizerStage< TInput, TCode >, and fz::LogTransformStage< TInput >.
|
pure virtual |
Stage type identifier written into the FZM file header.
Implemented in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
pure virtual |
DataType enum of the given output port.
Implemented in fz::RLEStage< T >, fz::AdaptiveLorenzoStage< T >, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::DifferenceStage< T, TOut, Mode >, fz::QuantizerStage< TInput, TCode >, fz::ROIBinSplitStage< TData >, fz::LogTransformStage< TInput >, fz::ZigzagStage< TIn, TOut >, fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::BitplaneRZEStage, fz::SZpStage< T >, fz::SZxStage< T >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::Cdf97Stage< TInput >, and fz::NegabinaryStage< TIn, TOut >.
|
inlinevirtual |
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 in fz::AdaptiveLorenzoStage< T >, fz::ROIBinSplitStage< TData >, fz::LogTransformStage< TInput >, fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::GPULZStage, fz::HuffmanStage< T >, fz::RLEStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::MergeStage, fz::Cdf97Stage< TInput >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
inlinevirtual |
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 in fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::GInterpStage< TInput, TCode >, fz::DifferenceStage< T, TOut, Mode >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::LogTransformStage< TInput >, fz::ZigzagStage< TIn, TOut >, fz::RLEStage< T >, fz::LorenzoQuantStage< TInput, TCode >, fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::HuffmanStage< T >, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::SZpStage< T >, fz::SZxStage< T >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, and fz::NegabinaryStage< TIn, TOut >.
|
inlinevirtual |
Restore stage config from header_buffer during decompression.
Reimplemented in fz::NegabinaryStage< TIn, TOut >, fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::ZigzagStage< TIn, TOut >, fz::RLEStage< T >, and fz::LorenzoQuantStage< TInput, TCode >.
|
inlinevirtual |
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 in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, and fz::LogTransformStage< TInput >.
|
inlinevirtual |
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 in fz::HuffmanStage< T >, fz::SZxStage< T >, and fz::ROIBinSplitStage< TData >.
|
inlinevirtual |
Called once by Pipeline::finalize() so stages can react to the dataset dimensions set via Pipeline::setDims() after construction.
| dims | {x, y, z} extents (z==1 → 2-D; y==z==1 → 1-D) |
Reimplemented in fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::ROIBinSplitStage< TData >, and fz::Cdf97Stage< TInput >.
|
inlinevirtual |
Called once by Pipeline::finalize() after buffer-size propagation, with this stage's estimated input size (bytes) and the pipeline pool.
Implement this to pre-allocate persistent stage-internal scratch (e.g. Huffman codebook/histogram buffers) via pool->allocatePersistentDevice and pool->allocatePersistentPinned rather than via cudaMalloc directly. Pre-allocating here makes PREALLOCATE mode semantically correct (all memory committed at finalize time) and makes the stage footprint visible via pool->getPersistentDeviceBytes() / getPersistentPinnedBytes().
Stages that also allow lazy allocation (e.g. for capacity-growth realloc in execute()) should check whether pool was already used to allocate here and skip the lazy path if so.
Default: no-op.
Reimplemented in fz::HuffmanStage< T >, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::QuantizerStage< TInput, TCode >, fz::ROIBinSplitStage< TData >, and fz::LogTransformStage< TInput >.
|
inlinevirtual |
Estimated persistent device memory this stage allocates outside the pool (via pool->allocatePersistentDevice). Used for total footprint reporting. Default: 0.
Reimplemented in fz::HuffmanStage< T >, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::QuantizerStage< TInput, TCode >, fz::ROIBinSplitStage< TData >, and fz::LogTransformStage< TInput >.
|
inlinevirtual |
Estimated persistent pinned-host memory this stage allocates outside the pool (via pool->allocatePersistentPinned). Used for total footprint reporting. Default: 0.
Reimplemented in fz::HuffmanStage< T >, and fz::GInterpStage< TInput, TCode >.
|
inlinevirtual |
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 in fz::CLOGStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::AdaptiveBitpackStage< T >, fz::GPULZStage, fz::HuffmanStage< T >, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::GInterpStage< TInput, TCode >, fz::LorenzoQuantStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::QuantizerStage< TInput, TCode >, and fz::LogTransformStage< TInput >.
|
inlinevirtual |
Maximum bytes this stage writes into its per-output FZM header slot.
Reimplemented in fz::RLEStage< T >, fz::LorenzoQuantStage< TInput, TCode >, fz::DifferenceStage< T, TOut, Mode >, fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, fz::TUPLStage, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, fz::LogTransformStage< TInput >, fz::NegabinaryStage< TIn, TOut >, and fz::ZigzagStage< TIn, TOut >.
|
inlinevirtual |
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 in fz::AdaptiveBitpackStage< T >, fz::BitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::HuffmanStage< T >, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::BitplaneRZEStage, fz::GInterpStage< TInput, TCode >, fz::SZpStage< T >, fz::SZxStage< T >, fz::MergeStage, fz::ROIBinSplitStage< TData >, fz::Cdf97Stage< TInput >, and fz::LogTransformStage< TInput >.
|
inlinevirtual |
Inform the stage whether all of its forward outputs are pipeline outputs.
A stage with deferred exact sizing may keep readback in postStreamSync when terminal, but must publish an exact size before returning from execute when a downstream stage will consume its output. Default: no placement-sensitive behavior.
Reimplemented in fz::HuffmanStage< T >.
|
inlinevirtual |
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 in fz::AdaptiveBitpackStage< T >, fz::CLOGStage, fz::GPULZStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::RLEStage< T >, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::SZpStage< T >, and fz::SZxStage< T >.
|
inlinevirtual |
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 in fz::AdaptiveBitpackStage< T >, fz::CLOGStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::AdaptiveLorenzoStage< T >, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, and fz::ZigzagStage< TIn, TOut >.
|
inlinevirtual |
Fused-kernel identity: the device-op this stage maps to, plus its runtime parameter bytes. The generic fused runner collects these across a fused group (after priming) and hands the ordered op list to the codegen, so no per-pipeline shape is hard-coded. Default (empty op_name) = not a fused op. Must agree with getFusionSpec() (a stage returning a fusable spec but no op cannot actually be composed). See include/stage/fusion.h and docs/codebase_notes.md CN-NVRTC-FUSE.
Reimplemented in fz::AdaptiveBitpackStage< T >, fz::CLOGStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::RREStage, fz::RZEStage, fz::DifferenceStage< T, TOut, Mode >, fz::LorenzoStage< T >, fz::TiledLorenzoStage< T >, fz::QuantizerStage< TInput, TCode >, fz::BitshuffleStage, and fz::ZigzagStage< TIn, TOut >.
|
inlinevirtual |
Exact local encoded-size policy exposed to a directly connected adaptive producer. This is an algorithmic semantic contract, not a fusion-cost estimate: staged and fused paths must make the same decision from it. Default = no oracle.
Reimplemented in fz::AdaptiveBitpackStage< T >.
|
inlinevirtual |
Bind an immediate downstream encoder's exact oracle during finalize(). Adaptive stages override this and return true only when type, unit size, exactness, and additivity are compatible. Default rejects the contract.
Reimplemented in fz::AdaptiveLorenzoStage< T >.
|
inlinevirtual |
Escaping outputs a generated fused op can produce beside its main port.
Reimplemented in fz::AdaptiveLorenzoStage< T >, and fz::QuantizerStage< TInput, TCode >.
|
inlinevirtual |
Establish forward-computed state this stage's own INVERSE will read, for a fused runner that bypasses forward execute(). Default no-op; a quantizer overrides it to run its value-range scan. Called once per fused group member before the fused kernel is generated (op params are read afterwards, so any param derived from primed state is valid).
Reimplemented in fz::QuantizerStage< TInput, TCode >.
|
inlinevirtual |
Tail-coder hook: a fused runner that produced this stage's archive without calling execute() reports the archive size and the ORIGINAL (uncompressed) input size, so the stage's inverse can size its output buffer. Without it a variable-length coder's inverse falls back to the compressed size and its decode overruns (see CN-CHUNK-WIRE). Distinct from any setFusedResult overload to avoid colliding with unrelated (size_t,size_t) signatures.
Reimplemented in fz::AdaptiveBitpackStage< T >, fz::CLOGStage, fz::HCLOGStage, fz::RAREStage, fz::RAZEStage, fz::RREStage, and fz::RZEStage.
|
inlinevirtual |
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 in fz::AdaptiveLorenzoStage< T >, and fz::QuantizerStage< TInput, TCode >.