FZGPUModules 2.0
GPU-accelerated modular compression pipelines
Loading...
Searching...
No Matches
fz::TiledLorenzoStage< T > Class Template Reference

#include <tiled_lorenzo_stage.h>

+ Inheritance diagram for fz::TiledLorenzoStage< T >:

Public Member Functions

void setInverse (bool inv) override
 
void setDims (const std::array< size_t, 3 > &dims) override
 
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 (cudaStream_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) 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
 
- 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 void onFinalize (size_t, MemoryPool *)
 
virtual size_t estimateDeviceFootprintBytes (size_t) const
 
virtual size_t estimatePinnedFootprintBytes (size_t) const
 
virtual void postStreamSync (cudaStream_t stream)
 
virtual bool isGraphCompatible () const
 
virtual size_t estimateScratchBytes (const std::vector< size_t > &input_sizes) const
 

Detailed Description

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. See also memory/cuszp_stages.md (Part 8).
Template Parameters
TSigned element type: int16_t or int32_t.

Member Function Documentation

◆ setInverse()

template<typename T >
void fz::TiledLorenzoStage< T >::setInverse ( bool  inverse)
inlineoverridevirtual

Switch between forward (compression) and inverse (decompression) mode. Affects getNumInputs()/getNumOutputs() for stages with asymmetric port counts.

Reimplemented from fz::Stage.

◆ setDims()

template<typename T >
void fz::TiledLorenzoStage< T >::setDims ( const std::array< size_t, 3 > &  dims)
inlineoverridevirtual

Called once by Pipeline::finalize() so stages can react to the dataset dimensions set via Pipeline::setDims() after construction.

Parameters
dims{x, y, z} extents (z==1 → 2-D; y==z==1 → 1-D)

Reimplemented from fz::Stage.

◆ setTileShape()

template<typename T >
void fz::TiledLorenzoStage< T >::setTileShape ( uint32_t  tx,
uint32_t  ty = 1,
uint32_t  tz = 1 
)
inline

Set the tile shape. tz == 1 ⇒ 2-D tile; ty == tz == 1 ⇒ 1-D tile. The product tx*ty*tz (= tile_elems, the natural AdaptiveBitpack block size) must be in [1, 1024]. Each extent must be in [1, 255]. Pass 0 for any extent to keep the ndim-derived default (2-D 8x8, 3-D 4x4x4, 1-D 64).

◆ execute()

template<typename T >
void fz::TiledLorenzoStage< T >::execute ( cudaStream_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.

◆ getName()

template<typename T >
std::string fz::TiledLorenzoStage< T >::getName ( ) const
inlineoverridevirtual

Human-readable name used in error messages and debug output.

Implements fz::Stage.

◆ estimateOutputSizes()

template<typename T >
std::vector< size_t > fz::TiledLorenzoStage< T >::estimateOutputSizes ( const std::vector< size_t > &  input_sizes) const
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.

◆ getActualOutputSizesByName()

template<typename T >
std::unordered_map< std::string, size_t > fz::TiledLorenzoStage< T >::getActualOutputSizesByName ( ) const
inlineoverridevirtual

Actual output sizes after execute(), keyed by output port name.

Implements fz::Stage.

◆ getActualOutputSize()

template<typename T >
size_t fz::TiledLorenzoStage< T >::getActualOutputSize ( int  index) const
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.

◆ getStageTypeId()

template<typename T >
uint16_t fz::TiledLorenzoStage< T >::getStageTypeId ( ) const
inlineoverridevirtual

Stage type identifier written into the FZM file header.

Implements fz::Stage.

◆ getOutputDataType()

template<typename T >
uint8_t fz::TiledLorenzoStage< T >::getOutputDataType ( size_t  output_index) const
inlineoverridevirtual

DataType enum of the given output port.

Implements fz::Stage.

◆ getInputDataType()

template<typename T >
uint8_t fz::TiledLorenzoStage< T >::getInputDataType ( size_t  ) const
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.

◆ serializeHeader()

template<typename T >
size_t fz::TiledLorenzoStage< T >::serializeHeader ( size_t  output_index,
uint8_t *  header_buffer,
size_t  max_size 
) const
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.

◆ deserializeHeader()

template<typename T >
void fz::TiledLorenzoStage< T >::deserializeHeader ( const uint8_t *  header_buffer,
size_t  size 
)
inlineoverridevirtual

Restore stage config from header_buffer during decompression.

Reimplemented from fz::Stage.

◆ getMaxHeaderSize()

template<typename T >
size_t fz::TiledLorenzoStage< T >::getMaxHeaderSize ( size_t  output_index) const
inlineoverridevirtual

Maximum bytes this stage writes into its per-output FZM header slot.

Reimplemented from fz::Stage.