FZGPUModules 2.0
GPU-accelerated modular compression pipelines
Loading...
Searching...
No Matches
compressor.h
Go to the documentation of this file.
1
5#pragma once
6
7#include "backend/types.h"
8#include "pipeline/dag.h"
9#include "pipeline/perf.h"
10#include "pipeline/config.h"
11#include "stage/stage.h"
12#include "stage/stage_factory.h"
13#include "mem/mempool.h"
14#include "fzm_format.h"
15
16#include <array>
17#include <memory>
18#include <stdexcept>
19#include <string>
20#include <unordered_map>
21#include <vector>
22
23namespace fz {
24
37class Pipeline {
38public:
44 explicit Pipeline(
45 size_t input_data_size = 0,
47 float pool_multiplier = 3.0f
48 );
49
57 explicit Pipeline(const std::string& config_path);
58
59 ~Pipeline();
60
61 // ── Configuration ─────────────────────────────────────────────────────────
62
65
67 void setNumStreams(int num_streams);
68
75 void setDims(size_t x, size_t y = 1, size_t z = 1) { dims_ = {x, y, z}; }
76 void setDims(std::array<size_t, 3> dims) { dims_ = dims; }
77 std::array<size_t, 3> getDims() const { return dims_; }
78
79 // ── Builder API ───────────────────────────────────────────────────────────
80
85 template<typename StageT, typename... Args>
86 StageT* addStage(Args&&... args);
87
95 int connect(Stage* dependent, Stage* producer, const std::string& output_name = "output");
96
98 int connect(Stage* dependent, const std::vector<Stage*>& producers);
99
105 void finalize();
106
112 void warmup(fz::stream_t stream = 0);
113
115 void setWarmupOnFinalize(bool enable) { warmup_on_finalize_ = enable; }
116 bool isWarmupOnFinalizeEnabled() const { return warmup_on_finalize_; }
117
124 void setPoolManagedDecompOutput(bool enable) { pool_managed_decomp_ = enable; }
125 bool isPoolManagedDecompOutput() const { return pool_managed_decomp_; }
126
140 size_t getMaxCompressedSize(size_t input_bytes) const;
141
161 size_t getLastUncompressedSize() const {
162 return original_input_size_ > 0 ? original_input_size_ : input_size_;
163 }
164
165 // ── Execution ─────────────────────────────────────────────────────────────
166
179 const void* d_input,
180 size_t input_size,
181 void** d_output,
182 size_t* output_size,
183 fz::stream_t stream = 0
184 );
185
213 const void* d_input,
214 size_t input_size,
215 void* d_output_buf,
216 size_t output_buf_capacity,
217 size_t* actual_output_size,
218 fz::stream_t stream = 0
219 );
220
236 const void* d_input,
237 size_t input_size,
238 void** d_output,
239 size_t* output_size,
240 fz::stream_t stream = 0
241 );
242
264 const void* d_input,
265 size_t input_size,
266 void* d_output_buf,
267 size_t output_buf_capacity,
268 size_t* actual_output_size,
269 fz::stream_t stream = 0
270 );
271
310 const void* d_input,
311 size_t input_size,
312 void* d_output_buf,
313 size_t output_buf_capacity,
314 size_t* actual_output_size,
315 fz::stream_t stream = 0
316 );
317
342 void prepareInverse(size_t uncompressed_size);
343
345 void reset(fz::stream_t stream = 0);
346
347 // ── Profiling ─────────────────────────────────────────────────────────────
348
353 void enableProfiling(bool enable);
354 bool isProfilingEnabled() const { return profiling_enabled_; }
355
357 const PipelinePerfResult& getLastPerfResult() const { return last_perf_result_; }
358
360 CompressionDAG* getDAG() { return dag_.get(); }
361
363 size_t getPoolThreshold() const;
364
375
381 void enableBoundsCheck(bool enable) { dag_->enableBoundsCheck(enable); }
382 bool isBoundsCheckEnabled() const { return dag_->isBoundsCheckEnabled(); }
383
389 void setColoringEnabled(bool enable) {
390 coloring_enabled_ = enable; // survives a setMemoryStrategy() DAG swap
391 dag_->setColoringEnabled(enable);
392 }
393 bool isColoringEnabled() const { return dag_->isColoringEnabled(); }
394 size_t getColorRegionCount() const { return dag_->getColorRegionCount(); }
395
403 bool isColoringRequested() const { return coloring_enabled_; }
404
414 std::unordered_map<std::string, std::vector<std::string>> collectRunNotes() const {
415 std::unordered_map<std::string, std::vector<std::string>> notes;
416 for (const auto& s : stages_) {
417 if (!s) continue;
418 auto n = s->getRunNotes();
419 if (!n.empty()) notes.emplace(s->getName(), std::move(n));
420 }
421 return notes;
422 }
423
424 // ── CUDA Graph Capture (compression-only) ─────────────────────────────────
425
433 void enableGraphMode(bool enable);
434 bool isGraphModeEnabled() const { return graph_mode_enabled_; }
435
446 void captureGraph(fz::stream_t stream = 0);
447 bool isGraphCaptured() const { return graph_captured_; }
448
449 size_t getPeakMemoryUsage() const;
450 size_t getCurrentMemoryUsage() const;
451 void printPipeline() const;
452
453 // ── File Serialization ────────────────────────────────────────────────────
454
457 FZMHeaderCore core;
458 std::vector<FZMStageInfo> stages;
459 std::vector<FZMBufferEntry> buffers;
460 };
461
463 void writeToFile(const std::string& filename, fz::stream_t stream = 0);
464
466 static FZMFileHeader readHeader(const std::string& filename);
467
470
471 // ── In-memory metadata header (decode without a prior compress) ────────────
472
487 std::vector<uint8_t> serializeHeaderToMemory() const;
488
508 void primeInverseFromHeader(const void* header_bytes, size_t header_size);
509
525 const std::string& filename,
526 void** d_output,
527 size_t* output_size,
528 fz::stream_t stream = 0,
529 PipelinePerfResult* perf_out = nullptr,
530 size_t pool_override_bytes = 0
531 );
532
551 const std::string& filename,
552 void** d_output,
553 size_t* output_size,
554 fz::stream_t stream = 0,
555 PipelinePerfResult* perf_out = nullptr
556 );
557
583 const void* header_bytes,
584 size_t header_size,
585 const void* d_blob,
586 size_t blob_size,
587 void** d_output,
588 size_t* output_size,
589 fz::stream_t stream = 0
590 );
591
592 // ── Config File ───────────────────────────────────────────────────────────
593
607 void loadConfig(const std::string& path);
608
617 void saveConfig(const std::string& path) const;
618
619private:
620 // ── RAII buffer wrappers (private implementation detail) ─────────────────
621
622 // Pool-allocated persistent device buffer.
623 struct PoolBuffer {
624 void* ptr = nullptr;
625 size_t capacity = 0;
626 MemoryPool* pool = nullptr;
627
628 ~PoolBuffer() { free(0); }
629 PoolBuffer() = default;
630 PoolBuffer(const PoolBuffer&) = delete;
631 PoolBuffer& operator=(const PoolBuffer&) = delete;
632
633 void free(fz::stream_t s) {
634 if (ptr && pool) { pool->free(ptr, s); ptr = nullptr; capacity = 0; }
635 }
636 bool allocate(MemoryPool* p, size_t bytes, fz::stream_t s,
637 const char* tag, bool persistent = false) {
638 free(s);
639 pool = p;
640 ptr = pool->allocate(bytes, s, tag, persistent);
641 if (ptr) capacity = bytes;
642 return ptr != nullptr;
643 }
644 };
645
646 // cudaHostAlloc pinned host buffer — grows on demand, never shrinks.
647 struct PinnedBuffer {
648 void* ptr = nullptr;
649 size_t capacity = 0;
650
651 ~PinnedBuffer() { if (ptr) cudaFreeHost(ptr); }
652 PinnedBuffer() = default;
653 PinnedBuffer(const PinnedBuffer&) = delete;
654 PinnedBuffer& operator=(const PinnedBuffer&) = delete;
655
656 // Returns false on CUDA allocation failure.
657 bool ensureCapacity(size_t bytes) {
658 if (capacity >= bytes) return true;
659 if (ptr) { cudaFreeHost(ptr); ptr = nullptr; capacity = 0; }
660 if (cudaHostAlloc(&ptr, bytes, cudaHostAllocDefault) != cudaSuccess) return false;
661 capacity = bytes;
662 return true;
663 }
664 };
665
666 // cudaMalloc device buffer — grows on demand, never shrinks.
667 struct DeviceBuffer {
668 void* ptr = nullptr;
669 size_t capacity = 0;
670
671 ~DeviceBuffer() { if (ptr) cudaFree(ptr); }
672 DeviceBuffer() = default;
673 DeviceBuffer(const DeviceBuffer&) = delete;
674 DeviceBuffer& operator=(const DeviceBuffer&) = delete;
675
676 // Returns false on CUDA allocation failure.
677 bool ensureCapacity(size_t bytes) {
678 if (capacity >= bytes) return true;
679 if (ptr) { cudaFree(ptr); ptr = nullptr; capacity = 0; }
680 if (cudaMalloc(&ptr, bytes) != cudaSuccess) return false;
681 capacity = bytes;
682 return true;
683 }
684 };
685
686 // ── Internal helpers ──────────────────────────────────────────────────────
687
688 Stage* addRawStage(Stage* stage);
689
690 struct OutputBuffer {
691 void* d_ptr;
692 size_t actual_size;
693 size_t allocated_size;
694 std::string name;
695 int buffer_id;
696 };
697 std::vector<OutputBuffer> getOutputBuffers() const;
698
699 static void* loadCompressedData(
700 const std::string& filename,
701 const FZMFileHeader& header,
702 fz::stream_t stream = 0,
703 MemoryPool* pool = nullptr
704 );
705
706 void validate();
707 std::pair<std::vector<Stage*>, std::vector<Stage*>> identifyTopology();
708 void setupInputBuffers(const std::vector<Stage*>& sources);
709 int autoDetectUnconnectedOutputs();
710 void detectMultiOutputScenario(int pipeline_outputs);
711 void configureStreamsIfNeeded();
712
713 // finalize() sub-steps
714 void typeCheckConnections();
715 void computeInputAlignment();
716 void notifyStagesFinalizeHooks();
717 void refinePoolSize();
718 void setupGraphModeInput();
719 void preallocatePadBuffer();
720 void preallocateConcatBuffers();
721
722 // compress() helper: handles graph-mode copy or alignment padding.
723 // Returns the effective source pointer and padded source size.
724 std::pair<const void*, size_t> prepareInputSource(
725 const void* d_input, size_t input_size, fz::stream_t stream);
726
732 void propagateBufferSizes(bool force_from_current_inputs = false);
733
734 std::vector<Stage*> getSourceStages() const;
735 std::vector<Stage*> getSinkStages() const;
736
737 // ── Inverse DAG helpers ───────────────────────────────────────────────────
738
740 struct FwdStageDesc {
741 Stage* stage;
742 std::vector<int> output_buf_ids;
743 std::vector<int> input_buf_ids;
744 };
745
747 using PipelineOutputMap = std::unordered_map<int, std::pair<void*, size_t>>;
748
749 // decompress() helper: builds or reuses the inverse DAG cache.
750 void buildOrReuseInvCache(
751 const PipelineOutputMap& po_map,
752 Stage* src_stage,
753 size_t src_sz,
754 fz::stream_t stream);
755
768 void decompressCore(
769 const void* d_input,
770 size_t input_size,
771 void* caller_output,
772 size_t caller_capacity,
773 bool synchronize,
774 void** d_output,
775 size_t* output_size,
776 fz::stream_t stream);
777
785 void buildStaticBufferMetadata();
786
792 std::vector<size_t> readConcatSegmentSizes(
793 const void* d_blob, size_t n, fz::stream_t stream) const;
794
795 // decompressFromFile() helpers.
797 static FZMFileHeader parseHeaderFromMemory(const void* data, size_t size);
798 static size_t computeFilePoolSize(const FZMFileHeader& fh, size_t pool_override_bytes);
799 static std::pair<std::vector<std::unique_ptr<Stage>>, std::vector<FwdStageDesc>>
800 reconstructForwardTopology(const FZMFileHeader& fh);
801 static std::unordered_map<Stage*, size_t> buildSourceSizesFromHeader(
802 const FZMFileHeader& fh, const std::vector<FwdStageDesc>& fwd_topology);
803
809 static std::pair<std::unique_ptr<CompressionDAG>,
810 std::unordered_map<Stage*, int>>
811 buildInverseDAG(
812 const std::vector<FwdStageDesc>& fwd_stages,
813 const PipelineOutputMap& pipeline_outputs,
814 MemoryPool* pool,
815 MemoryStrategy strategy,
816 const std::unordered_map<Stage*, size_t>& source_sizes,
817 bool enable_profiling
818 );
819
820 // ── Concat helpers ────────────────────────────────────────────────────────
821
822 struct OutputBufferInfo {
823 int buffer_id;
824 void* d_ptr;
825 size_t actual_size;
826 std::string stage_name;
827 std::string output_name;
828 };
829
830 std::vector<OutputBufferInfo> collectOutputBuffers() const;
831
833 size_t calculateConcatSize(const std::vector<OutputBufferInfo>& outputs) const;
834
835 size_t writeConcatBuffer(
836 const std::vector<OutputBufferInfo>& outputs,
837 uint8_t* d_concat_bytes,
838 fz::stream_t stream
839 ) const;
840
841 void concatOutputs(void** d_output, size_t* output_size, fz::stream_t stream);
842
843 // ── Member variables ──────────────────────────────────────────────────────
844
845 std::unique_ptr<MemoryPool> mem_pool_;
846 std::unique_ptr<CompressionDAG> dag_;
847 MemoryStrategy strategy_;
848
849 std::vector<std::unique_ptr<Stage>> stages_;
850 std::unordered_map<Stage*, DAGNode*> stage_to_node_;
851
852 struct ConnectionInfo {
853 Stage* dependent;
854 Stage* producer;
855 std::string output_name;
856 int output_index;
857 };
858 std::vector<ConnectionInfo> connections_;
859
860 int num_streams_;
861 bool is_finalized_;
862 bool warmup_on_finalize_;
863 bool pool_managed_decomp_;
864
865 // is_compressed_: true after the first successful compress() (gates writeToFile).
866 // was_compressed_: true between compress() and the next reset() (gates captureGraph).
867 bool is_compressed_;
868 bool was_compressed_;
869
870 bool profiling_enabled_;
873 bool coloring_enabled_ = true;
874 PipelinePerfResult last_perf_result_;
875
876 std::vector<DAGNode*> input_nodes_;
877 std::vector<DAGNode*> output_nodes_;
878 std::vector<int> input_buffer_ids_;
879 std::vector<int> output_buffer_ids_;
880
881 PoolBuffer d_concat_buffer_;
882 bool needs_concat_;
883
884 // Pool-persistent decompress output buffers (one per source stage).
885 // Only used when pool_managed_decomp_ == true.
886 std::vector<void*> d_decomp_outputs_;
887
888 // Pinned host buffer for concat header (one H2D copy instead of N).
889 PinnedBuffer h_concat_header_;
890 // Persistent pinned host + device descriptor buffers for the gather kernel.
891 PinnedBuffer h_copy_descs_;
892 DeviceBuffer d_copy_descs_;
893
894 size_t input_size_;
895
896 // Per-source input sizes from the most recent compress(), ordered to match
897 // input_nodes_. Used by decompress() to size each inverse result buffer.
898 std::vector<size_t> source_input_sizes_;
899
900 // Input alignment in bytes — LCM of all stage getRequiredInputAlignment() values.
901 // compress() zero-pads to this boundary transparently.
902 size_t input_alignment_bytes_;
903 PoolBuffer d_pad_buf_;
904
905 // Original (pre-padding) input size. decompress() uses this to trim the
906 // reported output back to what the caller provided. 0 when no padding.
907 size_t original_input_size_;
908
909 size_t input_size_hint_;
910 float pool_multiplier_;
911
912 // Dataset dimensions (x=fast, y, z). Pushed to each stage on addStage() and
913 // again at finalize(). Default {0,1,1} = 1-D, infer x from input size.
914 std::array<size_t, 3> dims_;
915
924 struct InvDAGCache {
925 std::unique_ptr<CompressionDAG> inv_dag;
926 std::unordered_map<Stage*, int> inv_result_map;
927 std::unordered_map<int, int> fwd_to_inv_ext_buf;
928 std::unordered_map<Stage*, size_t> source_sizes;
929 };
930 std::unique_ptr<InvDAGCache> inv_cache_;
931
932 struct BufferMetadata {
933 int buffer_id;
934 size_t actual_size;
935 size_t allocated_size;
936 std::string name;
937 DAGNode* producer;
938 int output_index;
939 };
940 std::vector<BufferMetadata> buffer_metadata_;
941
942 bool graph_mode_enabled_;
943 bool graph_captured_;
944
945 // Fixed device input buffer whose address is baked into the captured graph.
946 // compress() copies user input here before cudaGraphLaunch().
947 PoolBuffer d_graph_input_;
948 size_t d_graph_input_size_;
949
950 fz::graph_t captured_graph_;
951 fz::graph_exec_t graph_exec_;
952};
953
954// ── Template implementation ───────────────────────────────────────────────────
955
956template<typename StageT, typename... Args>
957StageT* Pipeline::addStage(Args&&... args) {
958 if (is_finalized_) {
959 throw std::runtime_error("Cannot add stages after finalization");
960 }
961
962 // if constexpr, not a runtime check: on an unsupported backend, StageT's
963 // constructor may not exist in the build at all (its .cu translation
964 // unit excluded — see Stage::isSupportedOnBackend()'s doc comment), so
965 // every line below that references `new StageT()` must never be
966 // instantiated at all, not merely never executed — hence the whole rest
967 // of the function lives in the `if constexpr` branch rather than after
968 // a standalone early-throw.
969 if constexpr (!StageT::isSupportedOnBackend()) {
970 throw std::runtime_error(
971 "addStage(): this stage type is not supported on the current "
972 "GPU backend (FZGMOD_BACKEND) this library was built for");
973 } else {
974 auto stage_ptr = std::make_unique<StageT>(std::forward<Args>(args)...);
975 StageT* stage = stage_ptr.get();
976
977 stage->setDims(dims_);
978
979 DAGNode* node = dag_->addStage(stage, stage->getName());
980 size_t num_outputs = stage->getNumOutputs();
981 auto output_names = stage->getOutputNames();
982
983 // Pre-allocate all output slots as unconnected (size=1 placeholder).
984 // connect() will promote any that get wired to downstream stages.
985 for (size_t i = 0; i < num_outputs; i++) {
986 std::string out_name = i < output_names.size() ? output_names[i] : std::to_string(i);
987 dag_->addUnconnectedOutput(node, 1, i, stage->getName() + "." + out_name + "_unconnected");
988 }
989
990 stage_to_node_[stage] = node;
991 stages_.push_back(std::move(stage_ptr));
992 return stage;
993 }
994}
995
996} // namespace fz
Definition dag.h:92
Definition mempool.h:82
void free(void *ptr, fz::stream_t stream)
void * allocate(size_t size, fz::stream_t stream, const std::string &tag="", bool persistent=false)
Definition compressor.h:37
void setDims(size_t x, size_t y=1, size_t z=1)
Definition compressor.h:75
void compress(const void *d_input, size_t input_size, void *d_output_buf, size_t output_buf_capacity, size_t *actual_output_size, fz::stream_t stream=0)
int connect(Stage *dependent, const std::vector< Stage * > &producers)
void setPoolManagedDecompOutput(bool enable)
Definition compressor.h:124
std::vector< uint8_t > serializeHeaderToMemory() const
void warmup(fz::stream_t stream=0)
size_t getPoolThreshold() const
void enableBoundsCheck(bool enable)
Definition compressor.h:381
bool isMemPoolFallbackMode() const
void reset(fz::stream_t stream=0)
void saveConfig(const std::string &path) const
std::unordered_map< std::string, std::vector< std::string > > collectRunNotes() const
Definition compressor.h:414
size_t getLastUncompressedSize() const
Definition compressor.h:161
void loadConfig(const std::string &path)
static void decompressFromFile(const std::string &filename, void **d_output, size_t *output_size, fz::stream_t stream=0, PipelinePerfResult *perf_out=nullptr, size_t pool_override_bytes=0)
void setWarmupOnFinalize(bool enable)
Definition compressor.h:115
CompressionDAG * getDAG()
Definition compressor.h:360
void captureGraph(fz::stream_t stream=0)
StageT * addStage(Args &&... args)
Definition compressor.h:957
const PipelinePerfResult & getLastPerfResult() const
Definition compressor.h:357
void decompressInto(const void *d_input, size_t input_size, void *d_output_buf, size_t output_buf_capacity, size_t *actual_output_size, fz::stream_t stream=0)
Pipeline(const std::string &config_path)
void decompress(const void *d_input, size_t input_size, void *d_output_buf, size_t output_buf_capacity, size_t *actual_output_size, fz::stream_t stream=0)
int connect(Stage *dependent, Stage *producer, const std::string &output_name="output")
Pipeline(size_t input_data_size=0, MemoryStrategy strategy=MemoryStrategy::MINIMAL, float pool_multiplier=3.0f)
void setMemoryStrategy(MemoryStrategy strategy)
void compress(const void *d_input, size_t input_size, void **d_output, size_t *output_size, fz::stream_t stream=0)
void prepareInverse(size_t uncompressed_size)
static FZMFileHeader readHeader(const std::string &filename)
void setColoringEnabled(bool enable)
Definition compressor.h:389
size_t getMaxCompressedSize(size_t input_bytes) const
void finalize()
void primeInverseFromHeader(const void *header_bytes, size_t header_size)
void enableGraphMode(bool enable)
void decompressFromFileInstance(const std::string &filename, void **d_output, size_t *output_size, fz::stream_t stream=0, PipelinePerfResult *perf_out=nullptr)
void enableProfiling(bool enable)
FZMFileHeader buildHeader() const
void decompress(const void *d_input, size_t input_size, void **d_output, size_t *output_size, fz::stream_t stream=0)
bool isColoringRequested() const
Definition compressor.h:403
void decompressFromMemory(const void *header_bytes, size_t header_size, const void *d_blob, size_t blob_size, void **d_output, size_t *output_size, fz::stream_t stream=0)
void setNumStreams(int num_streams)
void writeToFile(const std::string &filename, fz::stream_t stream=0)
Definition stage.h:30
TOML-based pipeline configuration file support.
Compression DAG wiring, execution, and memory strategy types.
FZM binary file format definitions — structs, enums, and helpers.
Stream-ordered CUDA memory pool for pipeline buffer management.
Definition algorithms.h:48
MemoryStrategy
Definition dag.h:23
@ MINIMAL
Allocate on-demand, free at last consumer. Lowest peak memory.
Pipeline and per-stage profiling result types.
Base class interface for all compression stages.
Factory function for reconstructing pipeline stages from serialized FZM headers.
Definition dag.h:52
Fixed-size FZM file header core (80 bytes).
Definition fzm_format.h:250
Definition perf.h:78
Definition compressor.h:456
Backend-neutral GPU type aliases.