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

#include <log_transform_stage.h>

+ Inheritance diagram for fz::LogTransformStage< TInput >:

Classes

struct  Config
 

Public Member Functions

void execute (cudaStream_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
 
void onFinalize (size_t estimated_inlen, MemoryPool *pool) override
 
size_t estimateDeviceFootprintBytes (size_t) const override
 
bool isGraphCompatible () const override
 No mid-execute sync: the outlier-count D2H happens in postStreamSync().
 
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
 
void setInverse (bool inverse) override
 
uint16_t getStageTypeId () const override
 
uint8_t getOutputDataType (size_t output_index) const override
 
uint8_t getInputDataType (size_t input_index) const override
 
size_t serializeHeader (size_t output_index, uint8_t *buf, size_t max_size) const override
 
size_t getMaxHeaderSize (size_t) const override
 
void deserializeHeader (const uint8_t *buf, size_t size) override
 
void saveState () override
 
void setErrorBound (float delta)
 Set delta, the target point-wise relative error bound.
 
void setThreshold (float t)
 |x| < threshold => lossless outlier; 0 disables (specials only).
 
float quantizerErrorBound () const
 
- Public Member Functions inherited from fz::Stage
virtual size_t getRequiredInputAlignment () const
 
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 size_t estimatePinnedFootprintBytes (size_t) const
 
virtual size_t estimateScratchBytes (const std::vector< size_t > &input_sizes) const
 

Static Public Member Functions

static float quantizerErrorBoundFor (float delta)
 
static float minimumErrorBound ()
 
- Static Public Member Functions inherited from fz::Stage
static constexpr bool isSupportedOnBackend ()
 

Static Public Attributes

static constexpr float kLogRoundTripSlack = 1e-6f
 
static constexpr size_t kMinOutlierSlots = 8
 

Detailed Description

template<typename TInput = float>
class fz::LogTransformStage< TInput >

Log-space transform: turns a point-wise relative bound into a plain absolute bound, so an ordinary ABS quantizer downstream delivers the relative guarantee.

Note
Prior work: the transformation scheme, the IEEE-754 exponent shortcut for log2, and the near-zero threshold follow Liang et al., IEEE CLUSTER 2018. See THIRD_PARTY.md.

The identity

|x - x_hat| / |x| <= delta
<=> x_hat / x in [1-delta, 1+delta]
<=> log2|x_hat| - log2|x| in [log2(1-delta), log2(1+delta)]

A multiplicative bound on x is an additive bound on log2|x|. The interval is asymmetric and |log2(1-delta)| > log2(1+delta), so the binding side is the upper one and the absolute bound to use downstream is

e = log2(1 + delta) - kLogRoundTripSlack // quantizerErrorBound()
static constexpr float kLogRoundTripSlack
Definition log_transform_stage.h:172

(the slack pays for the transform's own float32 rounding — see below)

Symmetric bins of half-width e leave the negative side tighter than required, wasting roughly delta^2 / 2 of achievable relative error. Negligible for small delta.

Why this stage exists

QuantizerStage with ErrorBoundMode::REL already does log-space quantization and already gives an exact per-element bound — but it quantizes raw values, with no predictor in front, so its codes still carry the field's full spatial redundancy and barely compress. LorenzoQuantStage and GInterpStage compress well but cannot honour a per-element relative bound at all (see ErrorBoundMode::PREL).

Putting the log transform upstream of the predictor is what gets both:

\-> signs ------------------------------------------\
\-> outlier_vals/outlier_idxs ----------------------> MergeStage
Definition log_transform_stage.h:151
Definition lorenzo_stage.h:54
Definition merge_stage.h:68
Definition quantizer.h:129
@ ABS
Absolute error bound.

Downstream error bound is the caller's job

The quantizer's eb is a function of this stage's eb, and a stage cannot reach across the DAG to set it. Do it explicitly:

auto* lg = p.addStage<LogTransformStage<float>>();
lg->setErrorBound(1e-3f);
...
quant->setErrorBound(lg->quantizerErrorBound()); // ~log2(1 + 1e-3)
quant->setErrorBoundMode(ErrorBoundMode::ABS);
void setErrorBound(float delta)
Set delta, the target point-wise relative error bound.
Definition log_transform_stage.h:292

Passing the raw delta to the quantizer instead produces a far looser relative bound (by a factor of 1/log2(1+delta), ~693x at delta=1e-3) with no error and no warning. There is no cross-stage check for this — see the "Limitations" section of docs/stages/log_transform.md.

Preconditions worth knowing before you reach for this

  • Sign changes hurt. The sign is stripped into a separate bit-plane, so a sign flip between neighbours is invisible to the downstream predictor and costs a raw bit per element. Single-signed fields (density, pressure, magnitude) are the good case; fields oscillating about zero are not.
  • Near-zero values are outliers. log2|x| -> -inf, so zeros, denormals, inf/NaN and anything below threshold are stored losslessly. Fields with a lot of near-zero mass pay for it in the outlier list.
  • float32 log/exp2 round-trip costs ~1 ULP. quantizerErrorBound() already subtracts kLogRoundTripSlack to pay for it, and the forward pass escalates any element whose actual round-trip is worse than that. Below minimumErrorBound() (~1.4e-6) the slack would consume the whole budget and execute() throws, which is the honest answer: float32 log space cannot deliver that bound.

Ports

Forward (4 outputs):

[0] output - TInput[n] log2(|x|), or log_floor at outliers
[1] signs - uint8[ceil(n/8)] bit i set => element i is negative
[2] outlier_vals - TInput[k] original values at outlier positions
[3] outlier_idxs - uint32[k] indices of outlier positions

Inverse: those same 4 buffers -> TInput[n].

The outlier count is not a port — it lives in a stage-private 4-byte device scratch (pool->allocatePersistentDevice in onFinalize()), is D2H'd in postStreamSync(), and is serialized into the FZM stage header. The inverse path receives it as a kernel argument read from the deserialized header. This mirrors QuantizerStage exactly.

Template Parameters
TInputFloating-point input type. Only float is instantiated; float64 support is deliberately deferred (see the stage docs).

Member Function Documentation

◆ execute()

template<typename TInput = float>
void fz::LogTransformStage< TInput >::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.

◆ postStreamSync()

template<typename TInput = float>
void fz::LogTransformStage< TInput >::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.

◆ onFinalize()

template<typename TInput = float>
void fz::LogTransformStage< TInput >::onFinalize ( size_t  ,
MemoryPool  
)
overridevirtual

Called once by Pipeline::finalize() after buffer-size propagation, with this stage's estimated input size (bytes) and the pipeline pool.

Implement this to pre-allocate persistent stage-internal scratch (e.g. Huffman codebook/histogram buffers) via pool->allocatePersistentDevice and pool->allocatePersistentPinned rather than via cudaMalloc directly. Pre-allocating here makes PREALLOCATE mode semantically correct (all memory committed at finalize time) and makes the stage footprint visible via pool->getPersistentDeviceBytes() / getPersistentPinnedBytes().

Stages that also allow lazy allocation (e.g. for capacity-growth realloc in execute()) should check whether pool was already used to allocate here and skip the lazy path if so.

Default: no-op.

Reimplemented from fz::Stage.

◆ estimateDeviceFootprintBytes()

template<typename TInput = float>
size_t fz::LogTransformStage< TInput >::estimateDeviceFootprintBytes ( size_t  ) const
inlineoverridevirtual

Estimated persistent device memory this stage allocates outside the pool (via pool->allocatePersistentDevice). Used for total footprint reporting. Default: 0.

Reimplemented from fz::Stage.

◆ getName()

template<typename TInput = float>
std::string fz::LogTransformStage< TInput >::getName ( ) const
inlineoverridevirtual

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

Implements fz::Stage.

◆ getOutputNames()

template<typename TInput = float>
std::vector< std::string > fz::LogTransformStage< TInput >::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()

template<typename TInput = float>
std::vector< size_t > fz::LogTransformStage< TInput >::estimateOutputSizes ( const std::vector< size_t > &  input_sizes) const
overridevirtual

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 TInput = float>
std::unordered_map< std::string, size_t > fz::LogTransformStage< TInput >::getActualOutputSizesByName ( ) const
inlineoverridevirtual

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

Implements fz::Stage.

◆ getActualOutputSize()

template<typename TInput = float>
size_t fz::LogTransformStage< TInput >::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.

◆ setInverse()

template<typename TInput = float>
void fz::LogTransformStage< TInput >::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.

◆ getStageTypeId()

template<typename TInput = float>
uint16_t fz::LogTransformStage< TInput >::getStageTypeId ( ) const
inlineoverridevirtual

Stage type identifier written into the FZM file header.

Implements fz::Stage.

◆ getOutputDataType()

template<typename TInput = float>
uint8_t fz::LogTransformStage< TInput >::getOutputDataType ( size_t  output_index) const
inlineoverridevirtual

DataType enum of the given output port.

Implements fz::Stage.

◆ getInputDataType()

template<typename TInput = float>
uint8_t fz::LogTransformStage< TInput >::getInputDataType ( size_t  input_index) const
inlineoverridevirtual

Index-aware: the inverse path consumes the same four heterogeneous buffers the forward path produced, so reporting FLOAT32 for all of them mis-describes the sign bitmap and the outlier indices to the DAG's type check and buffer sizing.

Reimplemented from fz::Stage.

◆ serializeHeader()

template<typename TInput = float>
size_t fz::LogTransformStage< TInput >::serializeHeader ( size_t  output_index,
uint8_t *  header_buffer,
size_t  max_size 
) const
overridevirtual

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.

◆ getMaxHeaderSize()

template<typename TInput = float>
size_t fz::LogTransformStage< TInput >::getMaxHeaderSize ( size_t  output_index) const
inlineoverridevirtual

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

Reimplemented from fz::Stage.

◆ deserializeHeader()

template<typename TInput = float>
void fz::LogTransformStage< TInput >::deserializeHeader ( const uint8_t *  header_buffer,
size_t  size 
)
overridevirtual

Restore stage config from header_buffer during decompression.

Reimplemented from fz::Stage.

◆ saveState()

template<typename TInput = float>
void fz::LogTransformStage< TInput >::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.

◆ quantizerErrorBound()

template<typename TInput = float>
float fz::LogTransformStage< TInput >::quantizerErrorBound ( ) const
inline

The absolute error bound the downstream quantizer must use: log2(1 + delta) - kLogRoundTripSlack.

This is the whole point of the stage, and wiring it up is the caller's responsibility — see the class doc.

◆ quantizerErrorBoundFor()

template<typename TInput = float>
static float fz::LogTransformStage< TInput >::quantizerErrorBoundFor ( float  delta)
inlinestatic

quantizerErrorBound() as a free-standing calculation, for callers that want the number before building the stage.

◆ minimumErrorBound()

template<typename TInput = float>
static float fz::LogTransformStage< TInput >::minimumErrorBound ( )
inlinestatic

Smallest delta this stage can honour in float32.

Below this the round-trip slack swallows the whole budget and there is nothing left for the quantizer — the honest answer is that float32 log space cannot deliver the bound, so execute() throws rather than quietly returning a stream that violates it.

Member Data Documentation

◆ kLogRoundTripSlack

template<typename TInput = float>
constexpr float fz::LogTransformStage< TInput >::kLogRoundTripSlack = 1e-6f
staticconstexpr

Slack reserved for the transform's own float32 round-trip error, in log2 units.

The total log-space deviation is the quantizer's bound plus whatever exp2(log2(x)) loses to rounding. Handing the quantizer the full log2(1+delta) spends the entire budget before that second term exists, which puts a handful of elements marginally over the bound. So the quantizer gets log2(1+delta) - kLogRoundTripSlack and the transform keeps the remainder.

float32 log2/exp2 round-trips within ~1 ULP, i.e. ~1.2e-7 relative in value space, which is ~1.7e-7 in log2 units. 1e-6 is a comfortable several-x margin over that while costing <0.1% of the bound at delta = 1e-3 — an unmeasurable compression-ratio difference.

◆ kMinOutlierSlots

template<typename TInput = float>
constexpr size_t fz::LogTransformStage< TInput >::kMinOutlierSlots = 8
staticconstexpr

Minimum outlier slots reserved regardless of outlier_capacity, so a small input cannot round its reserve down to zero. See maxOutlierCount.