adds some library documentation

Author: skyler-ruiter

CHANGELOG.md

@@ -0,0 +1,79 @@
+# Changelog
+
+All notable changes to FZGPUModules are documented here.
+
+Format follows [Keep a Changelog](https://keepachangelog.com/en/1.1.0/).
+Version numbers follow [Semantic Versioning](https://semver.org/).
+
+---
+
+## [Unreleased] — 2.0.0
+
+### Added
+
+**Pipeline features**
+- Multi-source pipeline support: `InputSpec` API, `compress(std::vector<InputSpec>)`, `decompressMulti()`, `setInputSizeHint()` per source
+- `Pipeline::warmup(stream)` — forces PTX→SASS JIT compilation before timing-sensitive work
+- `Pipeline::enableBoundsCheck(bool)` — runtime toggle for buffer overwrite detection (always on in Debug builds)
+- `Pipeline::setCaptureMode(bool)` — CUDA Graph stream capture for steady-state compression
+- `Pipeline::setPoolManagedDecompOutput(bool)` — opt-in pool-owned decompression output (avoids D2D copy)
+- Cached inverse DAG: `buildInverseDAG()` result cached after first `decompress()` call; eliminates ~200–500 µs per-call DAG rebuild overhead
+- Logging system: `FZ_LOG(LEVEL, ...)` with compile-time gating; `FZ_LOG_MIN_LEVEL` CMake option (0=TRACE … 255=SILENT); `Logger::setLogCallback()` for custom sinks
+
+**Memory & DAG**
+- Buffer coloring: non-overlapping buffers in PREALLOCATE mode are aliased to reduce peak GPU memory
+- Pinned concat header buffer: reduces H2D API calls from `1+N` to 1 per `compress()` call
+- Custom gather kernel (`launch_gather_kernel`) for D2D segment copies: replaces N individual `cudaMemcpyAsync` calls with a single kernel dispatch
+- `getActualOutputSize(int index)` index-based accessor on `Stage` — eliminates per-call `unordered_map` allocations in the inner execute loop
+- Pool auto-sizing: `computeTopoPoolSize()` + `setReleaseThreshold()` for topology-aware pool configuration
+- `setExternalPointer()` zero-copy path: user-owned device buffer passed directly into DAG
+
+**File format**
+- FZM version bumped to 3 (`FZM_VERSION = 0x0300`); `FZMHeaderCore` extended to 80 bytes with `num_sources` and `source_uncompressed_sizes[4]` fields
+- CRC32 (IEEE 802.3) checksums on payload and header
+- `Pipeline::writeToFile()` / `Pipeline::decompressFromFile()` static utility
+
+**Stages**
+- `BitshuffleStage` — GPU bit-matrix transpose
+- `RZEStage` — recursive zero-byte elimination with optimized kernel
+- `ZigzagStage<TIn, TOut>` — zigzag encode/decode
+- `NegabinaryStage<TIn, TOut>` — negabinary encode/decode
+- `DifferenceStage<T, TOut>` — first-order difference / cumulative-sum coding
+- `QuantizerStage<TInput, TCode>` — direct-value quantizer with ABS/REL/NOA error modes
+- Multi-dimensional Lorenzo (2-D and 3-D predictor kernels)
+
+**Build & distribution**
+- `find_package(FZGPUModules REQUIRED)` support via `cmake/FZGPUModulesConfig.cmake.in`
+- Versioned shared library symlinks (`libfzgmod.so → .so.2 → .so.2.0.0`) via `VERSION`/`SOVERSION` target properties
+- Relocatable RPATH (`$ORIGIN/../lib` on Linux, `@loader_path/../lib` on macOS)
+- CUDA 11.2+ version floor check at CMake configure time
+- Little-endian host check at CMake configure time
+- `profiling/` directory for profiling programs (separated from `examples/`)
+- `scripts/new_stage.sh` scaffold script for adding new stages
+- CMakePresets with `asan` and `compute-san` presets for sanitizer builds
+- Doxygen CI GitHub Actions workflow publishing to GitHub Pages
+
+**Testing**
+- Comprehensive test suite: 20 test binaries covering pipeline, stages, file I/O, memory strategies, buffer coloring, CUDA Graphs, bounds checking, and error handling
+- All tests pass under CUDA Compute Sanitizer (memcheck, initcheck, racecheck, synccheck) and host ASan+UBSan
+
+### Fixed
+- Race condition in `CompressionDAG::execute()` for multi-source pipelines: internal per-branch streams now have a GPU-side happens-before edge into the caller stream via `cudaStreamWaitEvent`
+- `DifferenceStage` inverse: replaced `cub::DeviceScan` with a custom `cumsumChunkedKernel` that uses only shared memory (no device temp allocation, sanitizer-clean)
+- `decompressFromFile` cleanup frees use `stream=0` to avoid pool-destructor race with `cudaMemPoolDestroy`
+- Removed spurious `find_dependency(CCCL REQUIRED)` from installed CMake config file (would break downstream `find_package` for users without CCCL)
+- Removed C language from `project()` (only CUDA and CXX required)
+
+### Changed
+- Project version set to `2.0.0` in `CMakeLists.txt`
+- `buildInverseDAG()` return type changed from `{inv_dag, int}` to `{inv_dag, unordered_map<Stage*, int>}` for multi-source support (breaking for internal callers)
+- `Stage::saveState()` / `restoreState()` contract: any stage that modifies `actual_output_sizes_` in its inverse `execute()` must override both methods
+
+---
+
+## [1.0.0] — 2026-04-14
+
+Initial tagged release. cuSZ based compressor with modular design and experimental CUDASTF support.
+
+[Unreleased]: https://github.com/szcompressor/FZGPUModules/compare/v1.1.0...HEAD
+[1.1.0]: https://github.com/szcompressor/FZGPUModules/releases/tag/release

README.md

@@ -0,0 +1,243 @@
+# FZGPUModules
+
+GPU-accelerated modular lossy compression pipeline for scientific floating-point data.
+
+## Overview
+
+FZGPUModules is a CUDA library for building composable, high-throughput compression
+pipelines. Each pipeline is a directed acyclic graph (DAG) of stages — predictors,
+quantizers, encoders, and transforms — connected and executed entirely on the GPU
+with stream-ordered memory management.
+
+**Key properties:**
+- **Modular** — mix and match stages (Lorenzo, Quantizer, RLE, RZE, Bitshuffle, …)
+- **High throughput** — parallel level execution, persistent scratch, CUDA Graph support
+- **Memory-efficient** — MINIMAL and PREALLOCATE strategies; buffer coloring to alias non-overlapping allocations
+
+---
+
+## Requirements
+
+| Requirement | Minimum |
+|---|---|
+| CUDA Toolkit | 11.2+ (stream-ordered allocator) |
+| C++ Standard | C++17 |
+| CMake | 3.24+ |
+| Host byte order | Little-endian |
+
+---
+
+## Quick Start
+
+```cpp
+#include "fzgpumodules.h"
+
+// 1. Build a pipeline
+fz::Pipeline pipeline(fz::MemoryPoolConfig(input_bytes));
+
+auto* lrz = pipeline.addStage<fz::LorenzoStage<float, uint16_t>>(
+    fz::LorenzoStage<float, uint16_t>::Config{1e-4f});
+auto* rle = pipeline.addStage<fz::RLEStage<uint16_t>>();
+
+pipeline.connect(lrz, "codes", rle);
+pipeline.finalize();
+
+// 2. Compress
+void* d_compressed = nullptr;
+size_t compressed_size = 0;
+pipeline.compress(d_input, n * sizeof(float), &d_compressed, &compressed_size, stream);
+
+// 3. Decompress
+void* d_output = nullptr;
+size_t output_size = 0;
+pipeline.decompress(d_compressed, compressed_size, &d_output, &output_size, stream);
+cudaStreamSynchronize(stream);
+// d_output is caller-owned — call cudaFree() when done.
+cudaFree(d_output);
+```
+
+See `examples/` for more usage patterns including multi-branch pipelines, CUDA Graph
+capture, and the low-level DAG API.
+
+---
+
+## Caller-Allocated Output (With Size Query)
+
+If you want full memory control, use the `Into` APIs and ask the pipeline for
+max output sizes before allocating:
+
+```cpp
+// After finalize()
+size_t comp_capacity   = pipeline.getMaxCompressedOutputSize();
+size_t decomp_capacity = pipeline.getMaxDecompressedOutputSize();
+
+void* d_comp_user = nullptr;
+void* d_decomp_user = nullptr;
+cudaMalloc(&d_comp_user, comp_capacity);
+cudaMalloc(&d_decomp_user, decomp_capacity);
+
+size_t comp_size = 0;
+pipeline.compressInto(d_input, input_bytes,
+                      d_comp_user, comp_capacity,
+                      &comp_size, stream);
+
+size_t decomp_size = 0;
+pipeline.decompressInto(d_comp_user, comp_size,
+                        d_decomp_user, decomp_capacity,
+                        &decomp_size, stream);
+```
+
+For multi-source pipelines:
+- Use `compressInto(const std::vector<InputSpec>&, ...)`
+- Use `getMaxDecompressedOutputSizes()` to size one output buffer per source
+- Use `decompressMultiInto(...)` to decode directly into those buffers
+
+For `.fzm` files, you can query exact decompressed size(s) from header metadata:
+- `Pipeline::getDecompressedOutputSizeFromFile(path)`
+- `Pipeline::getDecompressedOutputSizesFromFile(path)`
+
+See `examples/caller_allocated_output.cpp` for a minimal end-to-end example.
+
+---
+
+## Building from Source
+
+```bash
+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
+cmake --build build -j
+```
+
+**CMake options:**
+
+| Option | Default | Description |
+|---|---|---|
+| `BUILD_SHARED_LIBS` | `ON` | Build shared libraries |
+| `BUILD_EXAMPLES` | `OFF` | Build example programs |
+| `BUILD_PROFILING` | `OFF` | Build profiling programs |
+| `BUILD_TESTING` | `OFF` | Build test suite |
+| `FZ_LOG_MIN_LEVEL` | `2` (INFO) | 0=TRACE 1=DEBUG 2=INFO 3=WARN 255=SILENT |
+
+---
+
+## Installing
+
+```bash
+cmake --install build --prefix /your/install/prefix
+```
+
+This installs headers, shared libraries with versioned symlinks
+(`libfzgmod.so → libfzgmod.so.2 → libfzgmod.so.2.0.0`), and CMake
+package config files.
+
+### Downstream CMake usage
+
+```cmake
+find_package(FZGPUModules REQUIRED)
+target_link_libraries(my_target PRIVATE FZGMOD::fzgmod)
+```
+
+All targets: `FZGMOD::fzgmod`, `FZGMOD::fzgmod_mem`, `FZGMOD::fzgmod_encoders`,
+`FZGMOD::fzgmod_predictors`, `FZGMOD::fzgmod_pipeline`.
+
+---
+
+## Available Stages
+
+| Stage | Description |
+|---|---|
+| `LorenzoStage<TInput, TCode>` | 1-D/2-D/3-D Lorenzo predictor |
+| `QuantizerStage<TInput, TCode>` | Direct-value quantizer (ABS/REL/NOA error modes) |
+| `RLEStage<T>` | Run-length encoding |
+| `DifferenceStage<T, TOut>` | First-order difference / cumulative-sum coding |
+| `BitshuffleStage` | GPU bit-matrix transpose |
+| `RZEStage` | Recursive zero-byte elimination |
+| `ZigzagStage<TIn, TOut>` | Zigzag encode/decode |
+| `NegabinaryStage<TIn, TOut>` | Negabinary encode/decode |
+
+---
+
+## Memory Strategies
+
+| Strategy | Description |
+|---|---|
+| `MINIMAL` | Allocate on demand, free at last consumer. Lowest peak GPU memory. |
+| `PREALLOCATE` | Allocate everything at `finalize()`. Required for CUDA Graph capture. Enables buffer coloring. |
+
+---
+
+## File I/O
+
+```cpp
+// Write to file after compressing
+pipeline.writeToFile("output.fzm", stream);
+
+// Decompress directly from file (no pipeline setup needed)
+void* d_out = nullptr;
+size_t out_size = 0;
+fz::Pipeline::decompressFromFile("output.fzm", &d_out, &out_size, stream);
+cudaStreamSynchronize(stream);
+cudaFree(d_out);
+```
+
+FZM files embed the full stage configuration and compressed payload with CRC32
+checksums. See `include/fzm_format.h` for the format specification.
+
+---
+
+## CUDA Graph Support
+
+For throughput-critical workloads, enable CUDA Graph capture to eliminate
+CPU-side kernel launch overhead on repeated compress calls:
+
+```cpp
+pipeline.setCaptureMode(true);           // requires PREALLOCATE strategy
+pipeline.finalize();
+pipeline.warmup(stream);                 // JIT-compiles all kernels once
+// subsequent compress() calls replay the captured graph
+```
+
+---
+
+## Key API Reference
+
+| Class | Header | Description |
+|---|---|---|
+| `fz::Pipeline` | `pipeline/compressor.h` | High-level builder and executor |
+| `fz::Stage` | `stage/stage.h` | Base class for all compression stages |
+| `fz::CompressionDAG` | `pipeline/dag.h` | Low-level DAG wiring and execution |
+| `fz::MemoryPool` | `mem/mempool.h` | Stream-ordered CUDA memory pool |
+| `fz::PipelinePerfResult` | `pipeline/perf.h` | Per-stage profiling results |
+
+Full API documentation is available on the [Doxygen site](https://szcompressor.github.io/FZGPUModules/).
+
+---
+
+## Citation
+
+If you reference this work, please cite the following publication.
+
+> Note: This citation corresponds to the v1.0.0 release of this codebase, not the current version.
+
+- **[DRBSD-11]** FZModules: A Heterogeneous Computing Framework for Customizable Scientific Data Compression Pipelines
+    ```bibtex
+    @inproceedings{ruiter2025fzmodules,
+        author = {Ruiter, Skyler and Tian, Jiannan and Song, Fengguang},
+        title = {FZModules: A Heterogeneous Computing Framework for Customizable Scientific Data Compression Pipelines},
+        year = {2025},
+        url = {https://doi.org/10.1145/3731599.3767376},
+        booktitle = {Proceedings of the SC '25 Workshops of the International Conference for High Performance Computing, Networking, Storage and Analysis},
+        pages = {332-338},
+        series = {SC Workshops '25}
+    }
+    ```
+
+---
+
+## Adding a Custom Stage
+
+See `memory/how_to_add_a_stage.md` for the step-by-step guide, or use the
+scaffold script:
+
+```bash
+scripts/new_stage.sh MyStageName transforms
+```