Skip to content

Prepare for MGARD 1.7.0 release - #259

Merged
JieyangChen7 merged 238 commits into
CODARcode:masterfrom
JieyangChen7:mgard-x
Sep 21, 2026
Merged

JieyangChen7 merged 238 commits into
CODARcode:masterfrom
JieyangChen7:mgard-x

Conversation

@JieyangChen7

Copy link
Copy Markdown
Collaborator

This PR bumps MGARD to 1.7.0 and rolls up development on mgard-x since the 1.6.0 release. The major additions are HP-MDR, a high-performance redesign of the MDR-X refactoring and reconstruction pipeline (batched bitplane encoders, adaptive/hybrid lossless compression, improved domain decomposition, and a QoI-guided progressive retrieval path), and BlockMGARD, a new block-based hybrid hierarchy compression pipeline with in-cache block refactoring and region-of-interest support. Alongside these, the legacy standalone MGARD-CUDA backend has been removed in favor of MGARD-X, which is now the sole GPU compression path going forward. See 1.7.0 release note for the full list of additions, changes, and bug fixes.

Highlights

  • HP-MDR: high-performance MDR-X refactoring/reconstruction pipeline
  • BlockMGARD: block-based hybrid hierarchy compression with RoI support
  • Deprecated and removed the legacy MGARD-CUDA backend (superseded by MGARD-X)
  • New lossless backends (rANS, BlockDelta) and a portable warp-cooperative LZ4 path
  • Blackwell (sm_120) GPU build support
  • Various performance improvements and bug fixes (see release notes for details)

Jieyang Chen and others added 30 commits March 23, 2025 19:52
JieyangChen7 and others added 29 commits September 10, 2026 21:54
nvcomp was only ever needed by the legacy mgard-cuda backend's
GPU_Huffman_LZ4 lossless option; mgard-x's CUDA backend never linked
against it. Every CUDA build script was building it from source
regardless, and a few CPU-only scripts referenced an
${nvcomp_install_dir} variable that was never even set.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
mgard-x has fully superseded it as the actively maintained, portable
GPU backend. Removes src/cuda, include/cuda, the mgard-gpu
executable, the MGARD_ENABLE_LEGACY_CUDA build option and everything
wired to it: CMake config propagation to downstream consumers
(mgard-config.cmake.in, MGARDXConfig.h.in), the legacy comparison
path in the BatchTests example, and the MGARD-GPU docs.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
New build_mgard_cuda_blackwell.sh, plus three CCCL-3.x compatibility
fixes in DeviceAdapterCuda.h needed to build against CUDA 13's
bundled CUB/Thrust: cub::Max() was removed (guarded fallback via
CUB_VERSION for older toolkits still on CCCL 2.x), cub::
TransformInputIterator was removed in favor of thrust::
make_transform_iterator, and thrust::binary_function was dropped
since neither user of it needed the inherited typedefs.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Remove the L2 projection from BlockMGARD's L-infinity path and use the hierarchical basis by default. Apply the change consistently to the 1D, 2D, and 3D local/global transforms and their fused and unfused variants.

Update the L-infinity error amplification constant to C=1, add an explicit projection-mode switch, and persist the resolved mode in metadata while retaining compatibility with legacy orthogonal streams.

Validate the new path on 294 single- and double-precision SDRBench cases across relative error bounds from 1e-1 through 1e-6; all cases satisfy the requested L-infinity error bound.
Document the 1D-3D hybrid hierarchy interface, projection-mode controls, and current high-level API support. Clarify that the hybrid hierarchy is not exposed through the public low-level API.

Update the documented and built-in CLI lossless options to match the available string-based backends.
Document the 1D-3D hybrid ROI tolerance-map format and its CLI and high-level API usage. Add a zero-dependency generator for block-level maps, link the guide from the main README, include it in the documentation pages, and correct the ROI CLI option spelling.
Make the orthogonal basis the default hybrid projection policy and require an explicit hierarchical selection to remove the L2 projection. Remove the redundant Auto mode from the public enum and CLI, reject invalid projection modes, and update the user guides and tests accordingly.
8d7d33f added a required orthogonal_projection argument to
in_cache_block::decompose/recompose/decompose_quantize/recompose_dequantize
but missed the direct calls in test_DataRefactoring.cpp.in. Pass true to
keep testing the orthogonal basis path these cases were written for.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Run clang-format 14 (matching CI's run-clang-format check) over the
hybrid-projection-mode changes; only whitespace/line-wrapping is affected.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
HighLevelAPIsGPUBuffer/README.md was copy-pasted from HighLevelAPIs
(wrong directory name and source file extension); corrected.

LowLevelAPIs (SERIAL and CUDA) crashed with a heap corruption on
compress: lossless_type::Huffman_LZ4 hits a known buffer-sizing bug in
the library. Switch both examples to Huffman_Zstd until that is fixed
separately.

MDR-X (SERIAL/CUDA/HIP) no longer built: refactor.cpp/.cu still used
the old per-object MDR API (MGARDOrthoganalDecomposer, DirectInterleaver,
GroupedBPEncoder, DefaultLevelCompressor, MaxErrorCollector,
ConcatLevelFileWriter), which no longer exists in mgard_x::MDR. Rewrote
all three against the current ComposedRefactor/ComposedReconstructor +
MDRData/MDRMetadata API: Refactor() alone doesn't compress bitplanes or
populate metadata, so Compress()/StoreMetadata() must be called
explicitly; GenerateRequest() takes only the metadata now (tolerance/s
are set as fields on it first) and ProgressiveReconstruct() gained an
adaptive_resolution bool. refactor now does refactor+reconstruct in one
program (matching what CUDA's already-migrated code did), so the
separate reconstructor.cpp/.cu (still on the removed API) are dropped
and MDR-X/README.md documents the new single-executable interface.

Verified end-to-end on an H100 (CUDA sm_90, MDR enabled): all six
examples configure, build, and run; MDR-X SERIAL and CUDA agree on
MSE/PSNR/L2/L-inf error at every requested tolerance.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…and DeviceQueues

DeviceSpecification's per-device query loop and DeviceQueues::Initialize/
Destroy call cudaSetDevice(d)/hipSetDevice(d) for every device but never
restore the device that was active beforehand, leaving the real active
device at NumDevices-1 for the rest of the process. This silently
diverges from DeviceRuntime<CUDA/HIP>::cuda_dev_id/hip_dev_id (which stays
at its default of 0), so streams and allocations made before the next
explicit SelectDevice() end up tied to the wrong device's context -
surfacing much later as an unrelated illegal memory access.

Since these are eagerly-constructed static members, the bug fires before
main() even runs on any multi-GPU node. Fixed by saving/restoring the
originally-active device around each per-device loop.

Verified via full ctest suites: CUDA backend went from 238/265 to 262/265
passing (remaining 3 failures are an unrelated CUDA 13.3/CCCL 3.x
histogram bug on Blackwell). HIP backend: 265/265 passing on odyssey
(4x MI300A) with all GPUs visible, confirming the HIP_VISIBLE_DEVICES=0
workaround documented in build_mgard_hip_mi300.sh is no longer needed.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
- Bump the multi-core CPU (OpenMP) GCC requirement to 9+: its RLE and
  ZeroRLE lossless stages use std::inclusive_scan/exclusive_scan, which
  libstdc++ only provides from GCC 9 onward (compiles but throws a
  runtime error on older GCC).
- Note the CUDA 13.3/CCCL 3.x histogram bug on Blackwell (sm_120) GPUs,
  confirmed this session; CUDA 12.8 remains unaffected.
- Correct stale ZSTD/Protobuf versions and add the missing Catch2
  dependency to match what the build scripts actually pin.
- Update tested NVIDIA/AMD GPU lists.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Add unit tests for BPEncoderLocalityBlock and BPEncoderRegisterBlock
covering FP32/FP64 data, uint32_t/uint64_t bitplane words, binary vs.
negabinary encoding, and CONTROL_L2 (error collection) on/off. Each
configuration is generated into its own small translation unit
(tests/mgard-x/CMakeLists.txt), mirroring the library's own
per-type/dim source generation, since BPEncoderRegisterBlock's
progressive_decode() macro-unrolls 64 kernel variants per
configuration and templating across all of them in one file
overwhelms the CUDA/HIP device compiler.

Writing and running these tests (verified on SERIAL and on real
Blackwell GPUs) surfaced four real bugs, now fixed:

- BitplaneEncoder.hpp included 8 files that no longer exist in this
  branch, breaking anything that includes it.
- BPDecoderLocalityBlockFunctor::DecodeBinary wrote/read the `signs`
  array without the bounds guard `v` has, overflowing it whenever n
  isn't a multiple of the bitplane batch size -- exactly the case
  LocalityBlock exists to support.
- All four encode functors' CONTROL_L2 path accumulated into an
  uninitialized `errors[]` array, producing NaN error statistics.
- BPEncoderRegisterBlockFunctor::EncodeBinary packed sign bits by
  shifting a value cast to T_fp, but the shift amount runs up to
  BATCH_SIZE - 1 (63 when T_bitplane is uint64_t); with T_data=float
  (T_fp = uint32_t) this shifted a 32-bit value by up to 63 bits,
  undefined behavior that corrupted every sign bit for that
  combination.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Move the last BitplaneEncoder/ files not already under old/ (the
BPEncoderOpt* and BPEncoderRegister{Ballot,MatchAny,ReduceAll,Shift}
variants) alongside the other superseded encoders already archived
there, and drop the few whose old/ copy already existed. These are
the files that made BitplaneEncoder.hpp's includes stale.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…dLevelCompressor, ComposedRefactor/Reconstructor

Unit tests, verified on real GPU hardware (icarus, Blackwell):

- test_DirectInterleaver.cpp.in: interleave/reposition is a pure index
  permutation, so the round trip must be bit-exact -- 1D/2D/3D x
  float/double.
- test_MGARDDecomposer.cpp.in: decompose/recompose through the
  DecomposerInterface adapter ComposedRefactor actually drives (not just
  the underlying DataRefactor transform, already covered by
  test_DataRefactoring), across both bases and 1D/2D/3D x float/double.
- test_HybridLevelCompressor.cpp.in: compress_level/decompress_level must
  round-trip bitplane words exactly regardless of which of RLE, Huffman or
  the raw-copy fallback it picks; covers large-compressible (forces real
  compression), large-incompressible (forces the fallback after trying)
  and below-size-threshold (skips trying) for both bitplane word widths.
- test_ComposedRefactorReconstructor.cpp.in: the full refactor -> compress
  -> reconstruct pipeline across D x T_data x Basis x ControlL2 x
  NegaBinary, reconstructing progressively against a decreasing sequence
  of tolerances (single bound, progressive, and near-lossless in one
  pass) and verifying the achieved L-infinity or L2 error against each
  step. One file per configuration, same as the bitplane encoder tests,
  since ComposedRefactor/Reconstructor pull in BPEncoderRegisterBlock's
  64-way macro-unrolled decode kernels too.

Bugs found and fixed along the way:

- HybridLevelCompressor::Adapt took an unused max_level parameter; removed
  it and updated both call sites (ComposedRefactor/ComposedReconstructor).
- ComposedRefactor/ComposedReconstructor hardcoded ControlL2, Basis and
  NegaBinary as internal constants, so nothing outside the class could
  configure them. Moved all three to trailing template parameters
  (ControlL2, Basis, NegaBinary, in that order) defaulted to the previous
  hardcoded values, so every existing 3-argument call site keeps
  compiling unchanged.
- Renamed CONTROL_L2 to ControlL2 for consistency: it was spelled that way
  in ComposedRefactor/Reconstructor but as CollectError in the bitplane
  encoders it gets passed into: BitplaneEncoderInterface, BPEncoderBaseline,
  BPEncoderLocalityBlock, BPEncoderRegisterBlock.
- Writing test_ComposedRefactorReconstructor surfaced that
  examples/mgard-x/MDR-X/*/refactor.* (and the pattern they demonstrate)
  drive ProgressiveReconstruct without first calling LoadMetadata and
  Decompress, which MDRHighLevel/ReconstructPipeline.hpp does call. Without
  them the decoder reads a bitplane buffer that was never populated, so
  reconstruction silently returns near-zero data regardless of the
  requested tolerance. The test calls them in the correct order; the
  examples are unchanged.

Known, not fixed (flagged by the user in advance): Hierarchical basis +
NegaBinary misses only the tightest ("near lossless") requested tolerance,
narrowly (e.g. double: achieves 1.6e-9 against a requested 1e-9) -- every
looser step in the same progressive sequence passes. Reproduces
consistently across all three dimensions and both float and double.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
MaxErrorEstimatorOB (orthogonal basis) applies a 4x (c *= 4) safety
margin to account for negabinary needing 2 extra bits of range headroom
(BPEncoderLocalityBlock/RegisterBlock's EncodeNegaBinary does
`exp += 2`, so a given bitplane count buys less achievable precision
under negabinary than under binary). MaxErrorEstimatorHB (hierarchical
basis) never got the equivalent adjustment, so GenerateRequest's greedy
size interpreter under-requested bitplanes for Hierarchical + NegaBinary
reconstructions.

This is exactly what test_ComposedRefactorReconstructor's Hierarchical +
NegaBinary cases were catching: every tolerance in the progressive
sequence passed except the tightest ("near lossless") one, which missed
narrowly (e.g. double: achieved 1.6e-9 against a requested 1e-9).
Verified on real GPU hardware (icarus, Blackwell): the fix brings all 72
ComposedRefactor/Reconstructor test cases (216 assertions) to a clean
pass when run in isolation.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
level_data_array is allocated rounded up to the encoder's batch size,
but interleave() only writes the level's real elements, leaving the
padding as whatever cudaMalloc/pool memory previously held. AbsMax and
encode() both read the full padded length, so that padding could
pollute a level's abs_max scale factor whenever a prior, unrelated
allocation had left large leftover values in that memory -- causing
deterministic test failures when composed-refactor tests ran after the
larger pipeline tests, but not in isolation. Zero the padding once
after allocation; interleave() never touches it again for the object's
lifetime.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
MDRData::level_signs was allocated to a level's raw element count, but
BPDecoderRegisterBlockFunctor::DecodeBinary() (and the equivalent
LocalityBlock decoder) write signs across the full BATCH_SIZE-aligned
range, exactly like level_data_array does -- so any level whose element
count isn't already a multiple of the batch size overran the buffer.
The overflow corrupted adjacent heap metadata silently, surfacing later
as an unrelated `free(): invalid next size` abort whenever some other
allocation was freed. Pad level_signs the same way level_data_array is
padded, and relax VerifyLoadedBitplans' exact-size check to >=
accordingly.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Rename hybrid_projection_mode_type to compression_projection_mode_type
and add an Auto value (the new Config default): Hierarchical under an
L-infinity bound, Orthogonal otherwise. Config::projection_mode now
drives both the plain Compressor (which previously hardcoded this
choice from s alone, with no user control) and HybridHierarchyCompressor
through one shared resolve_projection_mode()/infer_orthogonal_projection()
in the new Utilities/ProjectionMode.h, instead of two separate
implementations.

Also simplifies HybridHierarchyCompressor's control flow: Quantize(),
Dequantize(), DecomposeQuantize(), and Decompress() each redundantly
recomputed and re-set orthogonal_projection even though they only ever
run after Compress() or DequantizeRecompose() already did so (confirmed
via the real call graph, including that the GPU/CPU domain-decomposition
pipelines call DequantizeRecompose() directly, bypassing Decompress()).
Reduced to exactly the two genuine resolution points.

Updates the CLI (-hp/--hybrid-projection -> -pm/--projection-mode, with
a new auto option), Metadata's type reference, and adds unit tests for
the Compressor across projection modes and error norms, mirroring the
existing hybrid coverage, plus rejection tests for hierarchical+L2 and
hierarchical above 3-D.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Rewraps template parameter lists and two long lines to match
run-clang-format's output exactly, per the CI check's reported diff.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…er config

Bump MGARD_VERSION and MGARD_FILE_VERSION to 1.7.0 and add version-history
release notes (0.0.0.2 through 1.7.0) linked from the README.

Remove Config::prefetch, which the 1.6.0 changes already made a no-op
(prefetching is now always enabled) but left as a dead struct field with
stale documentation. Remove Config::reorder and Metadata::reorder, which
only ever round-tripped a NOOP/SHUFFLE preprocessor flag through the file
header with no actual shuffle implementation behind it -- the header now
always records NOOP_PREPROCESSOR. Update the one test and doc references
that touched these fields.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Reflow the FillForCompression call sites/signature left mis-wrapped by
the previous commit's removal of the reorder parameter.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
… copyright

Bump copyright year to 2026 and author email to jieyang@uoregon.edu in
MGARD-X source file headers, and add the ORNL copyright line alongside
Brown University in LICENSE.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
@JieyangChen7
JieyangChen7 merged commit 6b43196 into CODARcode:master Sep 21, 2026
6 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants