|
| void | setInverse (bool inv) override |
| |
| void | setDims (const std::array< size_t, 3 > &dims) override |
| |
| void | setDimsOverride (size_t x, size_t y, size_t z) |
| |
| void | setTileShape (uint32_t tx, uint32_t ty=1, uint32_t tz=1) |
| |
|
uint32_t | getTileElems () const |
| | Elements per tile = the AdaptiveBitpack block_size that aligns blocks to tiles.
|
| |
| 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 |
| |
| FusionSpec | getFusionSpec () const override |
| |
| FusedOpDecl | getFusedOp () 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) 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 |
| |
| 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 | 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) |
| |
template<typename T>
class fz::TiledLorenzoStage< T >
Dimension-aware (tiled separable) Lorenzo predictor (cuSZp3). Lossless.
- Note
- Prior work: the dimension-aware separable delta is the cuSZp3 design (Yafan Huang et al., SC'25, BSD-3-Clause). This stage is a direct port of the cuSZp3 delta kernel logic (cuSZp_kernels_{2D,3D}_f32.cu); the tile-major modular decomposition, FZM header, and MemoryPool integration are FZGPUModules code. The cuSZp BSD-3-Clause copyright is reproduced in
THIRD_PARTY.md.
- Template Parameters
-
| T | Signed element type: int16_t or int32_t. |
Pin this stage's dimensions independently of the pipeline's.
A branch may carry data whose shape is not the pipeline's input shape — the binned background produced by ROIBinSplitStage is ceil(nx/b) x ceil(ny/b) x nz, not nx x ny x nz. Without pinning, Pipeline::finalize() re-pushes the global dims to every stage and silently overwrites anything set at construction, so the predictor would decorrelate the binned image using the full-resolution row stride and produce garbage residuals that still round-trip (the inverse makes the same mistake) — wrong ratio, no error. The pinned dims are serialized in this stage's own header, so the inverse pass recovers them from the archive.
template<typename T >
| void fz::TiledLorenzoStage< 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.