#include <gpulz_stage.h>

Inheritance diagram for fz::GPULZStage:

Public Member Functions
void	setInverse (bool inv) override

bool	isGraphCompatible () const override

void	setMatchLevel (int level)

void	setSplitMode (bool on)

size_t	getRequiredInputAlignment () 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

void	postStreamSync (fz::stream_t stream) override

std::string	getName () const override

std::vector< std::string >	getOutputNames () 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

size_t	estimateScratchBytes (const std::vector< size_t > &input_sizes) 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 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

void	saveState () override

Public Member Functions inherited from fz::Stage
int	getOutputIndex (const std::string &name) const

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

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

Detailed Description

GPU LZSS (GPULZ) coder stage.

setChunkSize(bytes) — chunk size in bytes (default 2048; supported: 1024, 2048, 4096 — must yield chunk_size / word_size >= 128 and a power of two). setWordSize(bytes) — word granularity 1/2/4/8 (default 4, matching the upstream reference's uint32_t default).

Note: Prior work: GPU kernels are a direct port of the compression and decompression kernels in gpulz.cu from GPULZ (Zhang, Tian, Di, Yu, Swany, Tao, Cappello — ICS '23; upstream repository declares no explicit license, see THIRD_PARTY.md). Upstream: https://github.com/hpdps-group/ICS23-GPULZ. The per-chunk container/offset-scan plumbing (raw-fallback flag, CUB exclusive scan for packing offsets, deferred tail-size readback) is FZGM's own, following the same pattern as RREStage/RZEStage. The all-zero-chunk fast path (skip encode entirely for chunks that are entirely zero) is adapted from the "sparse" GPULZ variant in boyuanzhang62/AIZ_VLDB26 (test/gpulz.cuh's notEmptyFlagArr).; CUDA Graph capture is supported for compression only (the final output-size readback is deferred to postStreamSync()). The inverse path reads the stream header with blocking D2H copies before it can launch the decode kernel, so it is not graph-capturable.

Member Function Documentation

◆ setInverse()

void fz::GPULZStage::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.

◆ isGraphCompatible()

bool fz::GPULZStage::isGraphCompatible ( ) const

inlineoverridevirtual

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 from fz::Stage.

◆ setMatchLevel()

void fz::GPULZStage::setMatchLevel ( int level )

inline

Match-search effort, 0 or 1 (default 1). Encode-side only — the stream format is identical either way, so a stream produced at one level decodes the same as at the other and the level is not serialized.

0 — exact longest match over the 32-element near window only. 1 — additionally consults a hashed table of two-word keys for long-range candidates (offsets up to 255).

Level 1 trades throughput for ratio. Measurements: docs/codebase_notes.md CN-GPULZ-1

◆ setSplitMode()

void fz::GPULZStage::setSplitMode ( bool on )

inline

Split mode (default off) — emit the compressed stream as four separate output ports instead of one interleaved stream:

literals the literal words, back to back (raw-fallback chunks land here too, since such a chunk is by definition all literal) lengths one match-length byte per match token offsets one match-offset byte per match token meta stream header + per-chunk size table + the flag bitmaps

This is the Zstandard split (literals separate from sequences), for the same reason: the parts have very different symbol distributions, and interleaving them into one byte stream raises the entropy a downstream coder sees.

The literals port keeps the data's natural word alphabet, so it should be fed to a symbol-width-matched coder (HuffmanStage<uint16_t> for uint16 quant codes) rather than a byte coder.

Every port must be entropy coded and all four re-merged: unlike the single-stream form, which codes the whole payload by construction, a split leaks any byte left out.

Measurements and why each port is folded in: docs/codebase_notes.md CN-GPULZ-1

◆ getRequiredInputAlignment()

size_t fz::GPULZStage::getRequiredInputAlignment ( ) const

inlineoverridevirtual

Deliberately 1, not chunk_size: this stage zero-pads its own tail chunk.

Requesting pipeline-level alignment does not actually work for a coder sitting behind a width-changing stage. Pipeline::finalize() pads the pipeline input to the LCM of every stage's alignment, but LorenzoQuant turns float32 into uint16 codes, so a 2048-aligned input arrives here as half as many bytes and need not be aligned at all. Worse, forcing the pipeline input up to a chunk multiple grows the upstream stage's output past its own estimate and trips the buffer-overwrite check. Padding internally is both simpler and correct for any upstream wiring.

Reimplemented from fz::Stage.

◆ execute()

void fz::GPULZStage::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.

◆ postStreamSync()

void fz::GPULZStage::postStreamSync ( fz::stream_t stream )

overridevirtual

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 from fz::Stage.

◆ getName()

std::string fz::GPULZStage::getName ( ) const

inlineoverridevirtual

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

Implements fz::Stage.

◆ getOutputNames()

std::vector< std::string > fz::GPULZStage::getOutputNames ( ) const

inlineoverridevirtual

Output port names in order. Default: single port named "output". Multi-output stages (e.g. Lorenzo: "codes", "outliers") override this.

Reimplemented from fz::Stage.

◆ estimateOutputSizes()

std::vector< size_t > fz::GPULZStage::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()

std::unordered_map< std::string, size_t > fz::GPULZStage::getActualOutputSizesByName ( ) const

overridevirtual

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

Implements fz::Stage.

◆ getActualOutputSize()

size_t fz::GPULZStage::getActualOutputSize ( int index ) const

overridevirtual

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.

◆ estimateScratchBytes()

size_t fz::GPULZStage::estimateScratchBytes ( const std::vector< size_t > & input_sizes ) const

inlineoverridevirtual

Forward pass allocates six persistent pool arrays proportional to n_chunks = ceil(input_bytes / chunk_size_): d_data_scratch_ : n_chunks * chunk_size_ (per-chunk worst-case compressed data) d_flag_scratch_ : n_chunks * (chunk_size_/word_size_/8) (per-chunk worst-case flag bits) d_flag_size_ : n_chunks * 4 (actual flag-array bytes per chunk) d_data_size_ : n_chunks * 4 (actual compressed-data bytes per chunk) d_clean_dev_ : n_chunks * 4 (raw-fallback-adjusted total size, scan input) d_dst_off_dev_ : n_chunks * 4 (exclusive prefix-sum packing offsets)

Reimplemented from fz::Stage.

◆ getStageTypeId()

uint16_t fz::GPULZStage::getStageTypeId ( ) const

inlineoverridevirtual

Stage type identifier written into the FZM file header.

Implements fz::Stage.

◆ getOutputDataType()

uint8_t fz::GPULZStage::getOutputDataType ( size_t output_index ) const

inlineoverridevirtual

DataType enum of the given output port.

Implements fz::Stage.

◆ getInputDataType()

uint8_t fz::GPULZStage::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()

size_t fz::GPULZStage::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()

void fz::GPULZStage::deserializeHeader	(	const uint8_t *	header_buffer,
		size_t	size
	)

inlineoverridevirtual

Restore stage config from header_buffer during decompression.

Reimplemented from fz::Stage.

◆ getMaxHeaderSize()

size_t fz::GPULZStage::getMaxHeaderSize ( size_t output_index ) const

inlineoverridevirtual

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

Reimplemented from fz::Stage.

◆ saveState()

void fz::GPULZStage::saveState ( )

inlineoverridevirtual

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 from fz::Stage.

Public Member Functions

Additional Inherited Members

Detailed Description

Member Function Documentation

◆ setInverse()

◆ isGraphCompatible()

◆ setMatchLevel()

◆ setSplitMode()

◆ getRequiredInputAlignment()

◆ execute()

◆ postStreamSync()

◆ getName()

◆ getOutputNames()

◆ estimateOutputSizes()

◆ getActualOutputSizesByName()

◆ getActualOutputSize()

◆ estimateScratchBytes()

◆ getStageTypeId()

◆ getOutputDataType()

◆ getInputDataType()

◆ serializeHeader()

◆ deserializeHeader()

◆ getMaxHeaderSize()

◆ saveState()