|
FZGPUModules 2.0
GPU-accelerated modular compression pipelines
|
#include <zigzag_stage.h>
Inheritance diagram for fz::ZigzagStage< TIn, TOut >:Public Member Functions | |
| void | setInverse (bool inv) override |
| void | setByteTransparent (bool on) |
| FusionSpec | getFusionSpec () const override |
| FusedOpDecl | getFusedOp () 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< 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 |
| uint16_t | getStageTypeId () const override |
| uint8_t | getOutputDataType (size_t output_index) const override |
| uint8_t | getInputDataType (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 |
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 | 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 bool | isGraphCompatible () const |
| virtual void | setTerminalOutput (bool terminal) |
| virtual size_t | estimateScratchBytes (const std::vector< size_t > &input_sizes) 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 () |
Element-wise zigzag encode/decode stage. Forward: TIn[] → TOut[]; inverse: TOut[] → TIn[].
| TIn | Signed integer type. |
| TOut | Corresponding unsigned type (same width as TIn). |
|
inlineoverridevirtual |
Switch between forward (compression) and inverse (decompression) mode. Affects getNumInputs()/getNumOutputs() for stages with asymmetric port counts.
Reimplemented from fz::Stage.
|
inline |
Byte-transparent mode (default off). When enabled, the stage opts out of finalize() type checking by reporting DataType::UNKNOWN on both ports, while still applying the same sizeof(TIn)-word zigzag transform.
This is what lets ZigzagStage act as the LC framework's TCMS component (Two's-Complement → Magnitude-Sign) inside a byte-stream chain: LC's TCMS_N operates on a raw byte stream as N-byte words regardless of the upstream element type, so the typed default (which would reject, e.g., a uint16 codes port feeding Zigzag<int8_t>, or a UINT8 byte stream feeding Zigzag<int64_t>) must be relaxed for the chain to connect.
|
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.
|
inlineoverridevirtual |
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 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 |
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 |
|
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 |
|
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.