|
| void | setInverse (bool inv) override |
| |
| void | setDims (const std::array< size_t, 3 > &dims) override |
| |
| void | execute (cudaStream_t stream, MemoryPool *pool, const std::vector< void * > &inputs, const std::vector< void * > &outputs, const std::vector< size_t > &sizes) override |
| |
|
bool | isGraphCompatible () const override |
| | No mid-execute sync or D2H: every kernel is enqueued on stream.
|
| |
| 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 |
| |
| void | saveState () override |
| |
| virtual size_t | getRequiredInputAlignment () const |
| |
| virtual std::vector< std::string > | getOutputNames () const |
| |
| int | getOutputIndex (const std::string &name) const |
| |
| 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 void | setTerminalOutput (bool terminal) |
| |
| virtual size_t | estimateScratchBytes (const std::vector< size_t > &input_sizes) const |
| |
| virtual FusionSpec | getFusionSpec () const |
| |
| virtual FusedOpDecl | getFusedOp () 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 TInput = double>
class fz::Cdf97Stage< TInput >
CDF 9/7 wavelet transform stage. Lossless, invertible, size-preserving.
- Template Parameters
-
| TInput | Floating-point element type (float or double). double is bit-exact with SPERR; float is a faster non-bit-exact variant. |
template<typename TInput = double>
| void fz::Cdf97Stage< 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.
template<typename TInput = double>
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.