#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	setDimsOverride (size_t x, size_t y, size_t z)

void	setTileShape (uint32_t tx, uint32_t ty=1, uint32_t tz=1)

void	setPredict (bool p)

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

FusionSpec	getInverseFusionSpec () const override

FusedOpDecl	getInverseFusedOp () const override

size_t	getFusedInverseElementCount () 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 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 double	getFusedForwardQuantStep () const

virtual double	getFusedInverseDequantStep () 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	setFusedInverseResult (size_t)

virtual void	setFusedSideOutput (int, size_t)

Additional Inherited Members
Static Public Member Functions inherited from fz::Stage
static constexpr bool	isSupportedOnBackend ()

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.

Template Parameters

T	Signed 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

Pipeline-driven dims push (at addStage and again at finalize). Ignored once setDimsOverride() has pinned explicit dims — see that method for why.

Reimplemented from fz::Stage.

◆ setDimsOverride()

template<typename T >

void fz::TiledLorenzoStage< T >::setDimsOverride	(	size_t	x,
		size_t	y,
		size_t	z
	)

inline

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.

◆ 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).

◆ setPredict()

template<typename T >

void fz::TiledLorenzoStage< T >::setPredict ( bool p )

inline

Enable/disable the separable Lorenzo delta. predict = false keeps the tile-major reorder and zero-padding but stores the quantizer codes directly (no prediction) — this reproduces cuSZp3 fixed mode's per-tile fixed-rate layout (8x8 in 2-D, 4x4x4 in 3-D) when paired with AdaptiveBitpackStage(block_size = tile_elems). Default true (plain/outlier).

◆ execute()

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.

◆ 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.

◆ getFusionSpec()

template<typename T >

FusionSpec fz::TiledLorenzoStage< T >::getFusionSpec ( ) const

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 Elementwise only in linear mode) must reflect that here. Forward-mode only; an inverse stage should report Unfusable.

Reimplemented from fz::Stage.

◆ getFusedOp()

template<typename T >

FusedOpDecl fz::TiledLorenzoStage< T >::getFusedOp ( ) const

inlineoverridevirtual

Warp-register predictor op (cuSZp3): 2-D separable tiled Lorenzo, EPL=2 (tile elems == 64). Declares the device policy type + its packed geometry (dims, tile, tiles-along-x) and the padded tile-major count n_ab, so the generic NVRTC warp runner composes the kernel with no per-predictor code. inv2eb is left 0 — the runner fills it from the resolved quantizer bound (see warp_op_params.h).

Reimplemented from fz::Stage.

◆ getInverseFusionSpec()

template<typename T >

FusionSpec fz::TiledLorenzoStage< T >::getInverseFusionSpec ( ) const

inlineoverridevirtual

Inverse-mode fusion role + geometry (valid only when isInverse()). The decompress-DAG analogue of getFusionSpec(): a stage that can participate in a fused inverse kernel overrides this to return a non-Unfusable spec. Kept a SEPARATE surface from getFusionSpec() (rather than un-gating it) so the forward fusion planner can never see an inverse stage. Default {} = not inverse-fusable.

Reimplemented from fz::Stage.

◆ getInverseFusedOp()

template<typename T >

FusedOpDecl fz::TiledLorenzoStage< T >::getInverseFusedOp ( ) const

inlineoverridevirtual

Inverse-mode device-op identity (valid only when isInverse()). Mirrors getFusedOp() for the decompress DAG: op_name is the device policy the inverse harness composes — the SAME policy class as the forward op, which carries decode()/undelta()/invert() alongside cost()/pack()/delta()/apply(). The generic inverse matcher/runner read these declarations by role, so a new warp predictor/coder that declares forward+inverse ops fuses in BOTH directions with no matcher edits. Default {} = not an inverse fused op.

Reimplemented from fz::Stage.

◆ getFusedInverseElementCount()

template<typename T >

size_t fz::TiledLorenzoStage< T >::getFusedInverseElementCount ( ) const

inlineoverridevirtual

Natural (row-major, unpadded) output element count of the tiled inverse — what the fused warp decode writes after the tile→natural scatter. The generic runner passes this as n_out so the coder's padded tile-major count is not mistaken for the output size.

Reimplemented from 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.

Public Member Functions

Additional Inherited Members

Detailed Description

Member Function Documentation

◆ setInverse()

◆ setDims()

◆ setDimsOverride()

◆ setTileShape()

◆ setPredict()

◆ execute()

◆ getName()

◆ getFusionSpec()

◆ getFusedOp()

◆ getInverseFusionSpec()

◆ getInverseFusedOp()

◆ getFusedInverseElementCount()

◆ estimateOutputSizes()

◆ getActualOutputSizesByName()

◆ getActualOutputSize()

◆ getStageTypeId()

◆ getOutputDataType()

◆ getInputDataType()

◆ serializeHeader()

◆ deserializeHeader()

◆ getMaxHeaderSize()