adds FAIR cards, fixes NOA eb issue missd in quatnizer stage fixed in lorenzo-quant, adds new profiling example

Author: skyler-ruiter

CHANGELOG.md

@@ -10,6 +10,8 @@ Version numbers follow [Semantic Versioning](https://semver.org/).
 ## [Unreleased] — 2.0.0
 
 ### Added
+- Added FAIR JSON cards for all 19 stages under `docs/cards/` (`fzgm-*.json`) and an `index.json` manifest; all validated against the Z-Hub card schema with 0 errors. New `isNewGeneralModule: true` cards: `fzgm-negabinary`, `fzgm-adm` (joins existing `fzgm-merge`). Provenance, algorithm slugs, config surfaces, and io dtype ranges are sourced from stage headers and `THIRD_PARTY.md`.
+- Replaced `basic_profiling_example.cpp` / `fzgmod-profile` target with `cuszp2_profile.cpp` / `fzgmod-profile-cuszp2`: profiles the cuSZp2 pipeline (`Quantizer(linear,ABS) → Lorenzo(32) → AdaptiveBitpack(32,outlier)`) matching `examples/presets/cuszp2.toml`; covers PREALLOCATE cold + steady-state compress/decompress and CUDA Graph compress replay with NVTX ranges; fixes the pool-ownership bug (pool-owned decompress outputs are no longer `cudaFree`'d).
 - `fzgmod-cli --graph` flag for the benchmark (`-b`) path: attempts CUDA Graph capture of the compress pipeline (forces `PREALLOCATE`, creates a real stream, calls `warmup(stream)` + `captureGraph(stream)`); silently falls back to the normal pipeline on incompatible stages (e.g. `HuffmanStage`, `BitplaneRLEStage`) and records the reason. `--report-json` schema bumped to `"1.1"` with a new `"graph"` block (`requested`, `active`, optional `incompatible_reason`).
 - CLI tests `BenchmarkGraphModeActivates` (graph capture succeeds for a `lorenzo->rle` pipeline, JSON reports `"active": true`) and `BenchmarkGraphModeFallsBackOnIncompatibleStage` (`huffman` terminal stage falls back silently, exit 0, `"active": false` with an `incompatible_reason`).
 
@@ -20,7 +22,7 @@ Version numbers follow [Semantic Versioning](https://semver.org/).
   3. `GInterpStage` outlier compaction scatters via `atomicAdd`, so the order of `outlier_vals`/`outlier_idxs` was non-deterministic (which thread claims which slot). Order doesn't affect correctness (the inverse scatters back by index) but changes how the downstream `Merge→Bitshuffle→RRE→RZE` chain compresses those arrays, so the blob **size** varied run-to-run — which undersized the reused inverse RZE output buffer (sized from a prior run's `cached_orig_bytes_`) and caused the E19 OOB. Now sorts `(idx, val)` by index after the forward kernel (`cub::DeviceRadixSort::SortPairs`); counts are ~0.001–0.004 % of N so the sort is negligible. Skipped under CUDA-graph capture (can't read the host count; capture only runs under PREALLOCATE, which sizes buffers to worst case once, so the size-variation crash can't occur there).
   4. Defensive: the GInterp codes buffer is sized to the padded element count `N` but the spline writes only the logical `dims` grid, and `Pipeline::concatOutputs()` pads each segment slot to 16 B (`ConcatLayout::slotSize`) — both left stale bytes the codes chain / final blob compressed. Now zero the codes buffer over `N` before the kernel and the concat buffer before the gather.
   Result: `cusz_hi_tp` and `cusz_hi_cr` are now bit-reproducible across runs, and `compute-sanitizer --tool memcheck` on `-b --runs 6` reports 0 errors (was a deterministic OOB). Round-trip fidelity/CR unchanged (CLDHGH 66.83 dB, 12.46×). New test `tests/pipeline/test_lc_presets.cpp` `LCPresets.DeterministicCompress` (outlier-heavy input, asserts 3 repeated compresses are byte-identical). Full suite green (ginterp 54, lc_presets 4, memory_strategies 21, concurrency 6, rre/rze 15/15).
-- `GInterpStage` / `LorenzoQuantStage` NOA/REL value-range scan included the input buffer's zero-padding tail, over-loosening the error bound on all-positive fields. `Pipeline::computeInputAlignment()` rounds the input buffer up to the LC chunk alignment (16384 B) and zero-pads the tail; the `computeValueBase()` min/max scan used the padded element count (`input_bytes / sizeof`), so a padding zero became the array minimum. For NOA (`abs_eb = eb·(max−min)`) on a strictly-positive field this collapsed the bound to `eb·max` — over-loosening it by `min/range` (e.g. CLDHGH ≈ 0.19%, NYX ≈ 0.05%), so the realized max error exceeded the intended bound by that ratio (the cross-tool "eb overshoot" of 1.00187 seen on CESM/CLDHGH). Fields straddling zero were unaffected (a padding zero sits inside their real range). Fix: scan only the logical grid (`dims` product), guarded to fall back to the full length when dims are unset. Overshoot drops to the float32 quantize/reconstruct floor (~1.00003, a handful of elements) — the same residual native cuSZ/fzgpu exhibit. Round-trip fidelity unchanged (CLDHGH 66.81→66.83 dB); CR unchanged (12.47→12.46×). `QuantizerStage` (flat, no `dims`) shares the latent pattern but does not manifest in the pfpl preset (its Difference predictor keeps residuals well under bound); left as-is.
+- `GInterpStage` / `LorenzoQuantStage` NOA/REL value-range scan included the input buffer's zero-padding tail, over-loosening the error bound on all-positive fields. `Pipeline::computeInputAlignment()` rounds the input buffer up to the LC chunk alignment (16384 B) and zero-pads the tail; the `computeValueBase()` min/max scan used the padded element count (`input_bytes / sizeof`), so a padding zero became the array minimum. For NOA (`abs_eb = eb·(max−min)`) on a strictly-positive field this collapsed the bound to `eb·max` — over-loosening it by `min/range` (e.g. CLDHGH ≈ 0.19%, NYX ≈ 0.05%), so the realized max error exceeded the intended bound by that ratio (the cross-tool "eb overshoot" of 1.00187 seen on CESM/CLDHGH). Fields straddling zero were unaffected (a padding zero sits inside their real range). Fix: scan only the logical grid (`dims` product), guarded to fall back to the full length when dims are unset. Overshoot drops to the float32 quantize/reconstruct floor (~1.00003, a handful of elements) — the same residual native cuSZ/fzgpu exhibit. Round-trip fidelity unchanged (CLDHGH 66.81→66.83 dB); CR unchanged (12.47→12.46×). Applied to all three predictors: `GInterpStage`, `LorenzoQuantStage`, and `QuantizerStage`. `QuantizerStage` had no `dims` member (flat stage), so a runtime-only `dims_` hint + `setDims()` override were added (not serialized — decode never scans); verified directly on a `Quantizer(NOA)→Bitshuffle→RZE` pipeline over all-positive CLDHGH, overshoot 1.00187 (11,666 violations) → 1.00007 (float floor). It didn't surface in pfpl only because its Difference predictor keeps residuals ≈0.1× the bound.
 - CLI test `BenchmarkGraphModeActivates` now skips (via a `cli_graph_supported()` mempool-fallback probe) instead of failing on vGPU / `cudaMalloc` fallback hosts where CUDA Graph capture is unsupported and the CLI silently falls back — matching the `GraphCapture` unit tests' `is_graph_supported()` guard.
 - **`fzgmod-cli -b --graph` aborted with a spurious `CUDA error … invalid argument` for any profiling-enabled pipeline that terminated in a deferred-size coder (e.g. the `cuszp`/`cuszp2` `AdaptiveBitpack` presets).** Graph capture succeeded, then compress() called `CompressionDAG::collectTimings()` after every `cudaGraphLaunch`; per-stage start/completion events are recorded by nodes baked into the captured graph, and `cudaEventElapsedTime()` across graph-recorded events returns `cudaErrorInvalidValue` on every node. That failure latched a sticky error into the CUDA context that then aborted the next unrelated kernel launch (surfacing as an "invalid argument" at the AdaptiveBitpack decode kernel; `CUDA_LAUNCH_BLOCKING=1` masked it). Fix: (1) skip the per-stage timing pass on a graph replay in `Pipeline::compress()` (the whole-pipeline `dag_elapsed_ms` from the outer `DagEventTimer`, recorded on `stream` outside the graph, is still valid); (2) `collectTimings()` now consumes the sticky error via `cudaGetLastError()` after a failed `cudaEventElapsedTime()` so it can never resurface downstream; (3) `CompressionDAG::execute()` no longer records the timing-only `start_event` during graph capture, avoiding dead nodes in the graph. `--graph` now round-trips `cuszp`/`cuszp2` at identical fidelity to non-graph (CLDHGH 70.83 dB). Per-stage `stages[]` for the graph-replayed compress is now empty by design (decompress, which runs the normal DAG, still reports it).
 - `examples/presets/cusz_hi_tp.toml` — set `code_type = "uint8"` + `quant_radius = 128` on the GInterp stage (was `uint16` / `32768`), matching cuSZ-Hi's actual spline ectrl (`ErrCtrlTrait<1> = u1`, `context.h` default `dict_size=256, radius=128`). Native's spline quant emits **1-byte** codes in `[0, 256)` and escapes any residual outside ±128 to the outlier compaction buffer (`spline3.inl:564`); our old preset used 2-byte codes with a wide radius that kept every residual inline, so the code stream was both twice as wide and higher-entropy. Halving the code width + tightening the radius roughly doubles TP CR at identical fidelity: **CLDHGH 7.68× → 12.47×** (native 13.33×), NYX 512³ 4.58× → 7.74×, CESM 4.11× → 7.65×; round-trip PSNR unchanged (CLDHGH 66.81 dB). This closes almost the entire TP-mode CR deficit vs native. Verified this was **not** an auto-tuning gap: enabling `auto_tuning = 3/4` on top adds < 0.5% CR (12.47× → 12.51×). Also confirmed not a BIT/RRE gap — our `RREStage` is the vendored LC `d_RRE`, and reimplementing bitshuffle as native's exact `d_BIT_1` gave zero change. (Supersedes the earlier `quant_radius = 32768` interim fix, which only reached 7.68×.)

docs/cards/fzgm-adaptive-bitpack.json

@@ -0,0 +1,56 @@
+{
+  "schemaVersion": "1.0",
+
+  "id": "fzgm-adaptive-bitpack",
+  "stage": "AdaptiveBitpackStage",
+  "name": "Per-Block Adaptive Bit-Plane Coder",
+
+  "zhubModule": "bit-packing",
+  "isNewGeneralModule": false,
+
+  "category": "coder",
+  "algorithm": "adaptive-bitpack",
+  "description": "Per-block adaptive fixed-rate bit-plane coder — the cuSZp lossless back-end. For each block of block_size signed elements emits one rate byte (bit-width of the largest magnitude), a per-block sign bitmap, and that many bit-planes. Per-block byte offsets resolved with a CUB exclusive scan. Pair with QuantizerStage(linearMode) → LorenzoStage(setBlockSize=32) for the full cuSZp-style pipeline. Optional outlierSelection (cuSZp2) stores element 0 as a raw outlier if smaller. Forward (compress) is CUDA Graph-capturable; inverse is not.",
+
+  "lossy": false,
+  "hardware": "gpu",
+
+  "io": {
+    "inputDtypes": ["i16", "i32"],
+    "outputDtypes": ["u8"]
+  },
+
+  "config": [
+    {
+      "name": "blockSize",
+      "type": "int",
+      "default": 32,
+      "description": "Elements per logical block (the fixed-rate granularity). Must be in [1, 1024]. cuSZp uses 32.",
+      "min": 1,
+      "max": 1024
+    },
+    {
+      "name": "outlierSelection",
+      "type": "bool",
+      "default": false,
+      "description": "cuSZp2 per-block plain/outlier selection: each block may store element 0 as a raw outlier and pack only the rest, whichever is smaller."
+    }
+  ],
+
+  "provenance": {
+    "sourceProject": "cuSZp",
+    "sourceUrl": "https://github.com/szcompressor/cuSZp",
+    "relationship": "direct-port",
+    "license": "BSD-3-Clause",
+    "paper": {
+      "title": "cuSZp: An Ultra-fast GPU Error-bounded Lossy Compression Framework with Optimized End-to-End Performance",
+      "doi": "10.1145/3581784.3607048",
+      "authors": "Yafan Huang, Sheng Di, Xiaodong Yu, Guanpeng Li, Franck Cappello",
+      "year": 2023
+    }
+  },
+
+  "docs": "stage_adaptive_bitpack.html",
+  "sourcePath": "modules/coders/adaptive_bitpack",
+  "versionAdded": "1.0"
+}

docs/cards/fzgm-adm.json

@@ -0,0 +1,43 @@
+{
+  "schemaVersion": "1.0",
+
+  "id": "fzgm-adm",
+  "stage": "ADMStage",
+  "name": "Adaptive Data Mapping",
+
+  "zhubModule": "adm",
+  "isNewGeneralModule": true,
+
+  "category": "transform",
+  "algorithm": "adm",
+  "description": "Adaptive Data Mapping (ADM): remaps u16/u32 integer streams into a compact 8-bit symbol domain before entropy coding. Partitions the input into 512-element warp blocks, computes a per-warp center (mean), and encodes each element as a compact (code, unary-signal) pair. The resulting byte codes are compressible by ANSStage with a significantly better ratio than the raw integer stream. D2H copies in both directions; not CUDA Graph-capturable. Lossless. GPU kernels adapted from the MANS project.",
+
+  "lossy": false,
+  "hardware": "gpu",
+
+  "io": {
+    "inputDtypes": ["u16", "u32"],
+    "outputDtypes": ["u8"]
+  },
+
+  "config": [
+    {
+      "name": "dtype",
+      "type": "enum",
+      "values": ["U16", "U32"],
+      "default": "U16",
+      "description": "Input element type. Must match the upstream stage's output. Set via setDtype(ADMDtype::U16) or setDtype(ADMDtype::U32)."
+    }
+  ],
+
+  "provenance": {
+    "sourceProject": "MANS",
+    "sourceUrl": "https://github.com/hpdps-group/MANS",
+    "relationship": "direct-port",
+    "license": "BSD-3-Clause"
+  },
+
+  "docs": "stage_adm.html",
+  "sourcePath": "modules/transforms/adm",
+  "versionAdded": "1.0"
+}

docs/cards/fzgm-ans.json

@@ -0,0 +1,44 @@
+{
+  "schemaVersion": "1.0",
+
+  "id": "fzgm-ans",
+  "stage": "ANSStage",
+  "name": "rANS Entropy Coder",
+
+  "zhubModule": "ans-encoding",
+  "isNewGeneralModule": false,
+
+  "category": "coder",
+  "algorithm": "ans",
+  "description": "GPU rANS (range Asymmetric Numeral Systems) entropy coder via vendored dietGPU kernel templates. Forward: uint8_t byte stream → ANS compressed bitstream (ANSCoalescedHeader + rANS data). Inverse: ANS bitstream → uint8_t byte stream. D2H copies occur in both directions; not CUDA Graph-capturable. Input must be 4-byte aligned. Only prob_bits=10 (1024-bucket table) is supported in this build.",
+
+  "lossy": false,
+  "hardware": "gpu",
+
+  "io": {
+    "inputDtypes": ["u8"],
+    "outputDtypes": ["u8"]
+  },
+
+  "config": [
+    {
+      "name": "probBits",
+      "type": "int",
+      "default": 10,
+      "description": "Probability resolution in bits (1024 buckets at default). Must be in [1, 11]. Only 10 is supported in the current build.",
+      "min": 1,
+      "max": 11
+    }
+  ],
+
+  "provenance": {
+    "sourceProject": "dietGPU",
+    "sourceUrl": "https://github.com/facebookresearch/dietgpu",
+    "relationship": "direct-port",
+    "license": "MIT"
+  },
+
+  "docs": "stage_ans.html",
+  "sourcePath": "modules/coders/ans",
+  "versionAdded": "1.0"
+}

docs/cards/fzgm-bitpack.json

@@ -0,0 +1,47 @@
+{
+  "schemaVersion": "1.0",
+
+  "id": "fzgm-bitpack",
+  "stage": "BitpackStage",
+  "name": "Fixed-Width Bit Packer",
+
+  "zhubModule": "bit-packing",
+  "isNewGeneralModule": false,
+
+  "category": "coder",
+  "algorithm": "bitpack",
+  "description": "GPU fixed-width bit-packing stage. Forward: packs each element using only the low nbits bits into a dense byte stream. Inverse: unpacks, zero-extending to full width. nbits must be a power of two (allowed: u8→{1,2,4,8}; u16→{1,2,4,8,16}; u32→{1,2,4,8,16,32}). Auto-detect mode (setAutoDetect) scans the input for its maximum value and selects the smallest valid nbits; incompatible with CUDA Graph capture.",
+
+  "lossy": false,
+  "hardware": "gpu",
+
+  "io": {
+    "inputDtypes": ["u8", "u16", "u32"],
+    "outputDtypes": ["u8"]
+  },
+
+  "config": [
+    {
+      "name": "nbits",
+      "type": "int",
+      "default": "8*sizeof(T)",
+      "description": "Bits per element. Must be a power of two in [1, 8*sizeof(T)]. Identity (no compression) at the default.",
+      "min": 1
+    },
+    {
+      "name": "autoDetect",
+      "type": "bool",
+      "default": false,
+      "description": "Scan input for max value and pick the smallest valid power-of-two nbits. Requires a D2H sync; incompatible with CUDA Graph capture."
+    }
+  ],
+
+  "provenance": {
+    "relationship": "original",
+    "license": "BSD-3-Clause"
+  },
+
+  "docs": "stage_bitpack.html",
+  "sourcePath": "modules/coders/bitpack",
+  "versionAdded": "1.0"
+}

docs/cards/fzgm-bitplane-rze.json

@@ -0,0 +1,40 @@
+{
+  "schemaVersion": "1.0",
+
+  "id": "fzgm-bitplane-rze",
+  "stage": "BitplaneRZEStage",
+  "name": "Fused Bitplane Transpose + Zero-Group RZE",
+
+  "zhubModule": "bit-shuffle",
+  "zhubModuleSecondary": ["rze-reducer"],
+  "isNewGeneralModule": false,
+
+  "category": "fused",
+  "algorithm": "bitplane-rze",
+  "description": "Fused bitplane-transpose + zero-group RZE — the FZ-GPU lossless encoder. In one shared-memory pass per 4096-byte chunk: (1) bit-transposes a 32×32 matrix of uint32_t words (bitshuffle at 4-byte width), then (2) eliminates all-zero 4-byte groups via per-block atomicAdd. Input is uint16_t symbols (two per word). Emits a self-describing byte archive (128-byte header + bitflag bitmap + per-block start positions + compacted bitstream). Not CUDA Graph-capturable (blocking D2H in both directions). Replaces a BitshuffleStage(element_width=4) → RZEStage(levels=1) pair with less global memory traffic.",
+
+  "lossy": false,
+  "hardware": "gpu",
+
+  "io": {
+    "inputDtypes": ["u16"],
+    "outputDtypes": ["u8"]
+  },
+
+  "provenance": {
+    "sourceProject": "cuSZ",
+    "sourceUrl": "https://github.com/szcompressor/cuSZ",
+    "relationship": "direct-port",
+    "license": "BSD-3-Clause",
+    "paper": {
+      "title": "FZ-GPU: A Fast and High-Ratio Lossy Compressor for Scientific Computing Applications on GPUs",
+      "doi": "10.1145/3588195.3592994",
+      "authors": "Boyuan Zhang, Jiannan Tian, Sheng Di, Xiaodong Yu, Yunhe Feng, Xin Liang, Dingwen Tao, Franck Cappello",
+      "year": 2023
+    }
+  },
+
+  "docs": "stage_bitplane_rze.html",
+  "sourcePath": "modules/fused/bitplane_rze",
+  "versionAdded": "1.0"
+}

docs/cards/fzgm-bitshuffle.json

@@ -0,0 +1,50 @@
+{
+  "schemaVersion": "1.0",
+
+  "id": "fzgm-bitshuffle",
+  "stage": "BitshuffleStage",
+  "name": "GPU Bit-Matrix Transpose",
+
+  "zhubModule": "bit-shuffle",
+  "isNewGeneralModule": false,
+
+  "category": "shuffler",
+  "algorithm": "bitshuffle",
+  "description": "GPU W×N bit-matrix transpose over fixed-size chunks. Given a chunk of N elements each W bits wide, the forward pass produces W groups each containing the k-th bit of all N elements (transposing the bit-plane matrix). Output is the same byte size as input. MSB-first output layout: bit-plane W-1 at index 0. 4- and 8-byte butterfly kernels adapted from the LC framework (d_BIT_4/d_BIT_8); 1- and 2-byte paths use standard __ballot_sync. Size-preserving; lossless.",
+
+  "lossy": false,
+  "hardware": "gpu",
+
+  "io": {
+    "inputDtypes": ["u8"],
+    "outputDtypes": ["u8"]
+  },
+
+  "config": [
+    {
+      "name": "blockSize",
+      "type": "int",
+      "default": 16384,
+      "description": "Chunk size in bytes. Must be a positive multiple of 1024*elementWidth.",
+      "min": 1
+    },
+    {
+      "name": "elementWidth",
+      "type": "int",
+      "values": [1, 2, 4, 8],
+      "default": 4,
+      "description": "Element width in bytes for the bit-matrix transpose. Determines W (bit-planes per chunk)."
+    }
+  ],
+
+  "provenance": {
+    "sourceProject": "LC-framework",
+    "sourceUrl": "https://github.com/burtscher/LC-framework",
+    "relationship": "direct-port",
+    "license": "BSD-3-Clause"
+  },
+
+  "docs": "stage_bitshuffle.html",
+  "sourcePath": "modules/shufflers/bitshuffle",
+  "versionAdded": "1.0"
+}