Feat/gpu wavelet codecs 64bit - #9
Open
ooreilly wants to merge 7 commits into
Open
Conversation
Introduces kernel 2 (coding) for the bitmap-significance-split encode: flat per-block fixed-width, per-line width, and two-level occupancy variants, plus validation/CR/throughput tests (test_bitmap_code_hip) and kernel-1 validation (test_bitmap_encode_hip, test_bitmap_vs_cpu_rle) with build targets. Adds waveletBitmapCodeTwoLevelOptKernel: a byte-for-byte-identical, faster two-level coder found by a polyopt LLM-agent campaign. Levers: 1024 threads (one z-line/thread), wave64 exclusive scans as DPP row shifts (no ds_bpermute), the two width-independent scans fused into one dual scan, a branch-free per-width group packer, values staged in an LDS image drained with 16-byte non-temporal stores, and single-ballot occupancy writes. Validated on gfx950 (MI355x): output is byte-exact vs the reference two-level kernel across sizes and quantization; the coding kernel is ~2.5-5.3x faster (3.7x at 256^3, 0.054 -> 0.014 ms). Uses ~132 KB LDS so it targets the larger LDS on gfx950; the original waveletBitmapCodeTwoLevelKernel remains the portable path for gfx90a. Build the gfx950 kernel with HIP_ARCH=gfx950.
…flag Fix the missing -mllvm -unroll-threshold=10000 on the test_bitmap_octree_hip and test_wavelet_buffer_hip build rules. Without it the ds79 Z-transform's tmp[32] + dynamic mirror-index loops spilled to scratch; adding it keeps them in registers and raises the forward/inverse transform from ~460 GB/s to ~2.3-2.4 TB/s (4x), which dominates full encode/decode. Add the octree significance coder (DFS reference, parallel level-major, fused kernel-2) and a full decode pipeline: stage-A (coded stream -> bitmap+int32, byte-exact inverse of the fused encoder) and stage-B (bitmap+values -> dequant + inverse wavelet). Round-trip is byte-exact vs kernel-1 and the reconstructed field is bit-identical to the RLE decode. Benchmark vs RLE on real panels (512^3, scale 8): octree CR 244x vs 36.8x (s1000) / 60x vs 26x (s3000) at identical distortion; full decode 1.6-1.9x faster; encode within ~5-9%.
Add HIP_COMPRESS_KERNEL_OCTREE as a first-class API kernel (3D only), end-to-end on the plan/scan/compact infrastructure: - Encode: k1 (bitmap+values) -> k2 (octree code) -> 4-align coded sizes -> exclusive scan -> woctCompactKernel writing a self-contained header [nb][nmf][offsets][sig_sizes][mulfac]. - Decode: header-addressed stage A (locates blocks via offsets, reads significance length from header, publishes 1/mulfac on device) -> dev-scale stage B; no host readback. - Refactor stage-A/stage-B into shared __device__ bodies so the validated prototype kernels keep byte-identical signatures. - 4-align each block in the packed stream so uint32 width-table / flat-mask reads stay aligned (the fixed-stride prototype masked this via 16B slots). - Plan gains octree scratch/sig/inv_scale buffers; MaxOutputSize and the header-size helper account for the octree layout. Add test_octree_round_trip and an octree case to bench_throughput. Validated on MI355x (gfx950): 38/38 API tests pass; at 512^3 octree CR 213.0 vs RLE 36.8, encode 0.91x RLE, decode 1.56x faster.
…ionality Replace the per-line width-table value coder in the octree (3D) and new quadtree (2D) significance codecs with a per-block PFOR (patched frame-of-reference) layout: a fixed W_lo base stream plus an exception bitmap and W_hi patch stream. All streams stay fixed-width and byte-aligned, so decode throughput is unchanged while compression ratio improves (~+5.6-10.4% at 512^3 for ~1-3% encode cost; near-0% decode). Add HIP_COMPRESS_KERNEL_AUTO and make it the plan default: resolves to QUADTREE for 2D (nz == 1) and OCTREE for 3D at plan creation, with ZLINE as the fallback for configurations the structured coders cannot handle. ZLINE remains the encode-throughput hedge for encode-bound callers. makefile: HIP_ARCH now accepts a space-separated list to build a fat binary (e.g. "gfx942 gfx950"). Validated: test_compress_api_hip 39/39 and test_compress_2d_hip 6/6 on gfx942 (MI300X) and gfx950 (MI355X).
Bump all AMD source-header and doc-footer copyrights in CvxCompress from 2025 to 2026. Add the missing MIT header to tests/test_compress_2d_hip.cpp. Refresh HIP_API.md and README.md for the current state: auto-select default (octree 3D / quadtree 2D), 2D support, gfx950/MI355X, and ROCm 7.2.1.
Replace the per-plane buffer-instruction loads/stores (32-bit intra-plane offset, 4 GB/plane cap) with global nontemporal load/store using size_t byte offsets, via new hip/hipPlaneIO.h. Applied across RLE/segrle forward, RLE inverse, bitmap, octree inverse, RLE2D, quadtree2D, and block-copy. - Add an alignment-tolerant float4 (aligned(4)) helper for the block-copy source-array fast paths, where arbitrary user ldimx/x0 can yield 4-byte-but-not-16-byte aligned offsets; codec/wavelet-side paths keep the 16-aligned helper (32-multiple dims guarantee alignment). - Relax the plan guard from nx*ny*4 <= 2^32 (4 GB) to nx*ny < 2^31 elements (~8 GB/plane); the remaining bound is the int plane stride. Update the CreatePlan/Copy* error-path tests to the new contract. - Remove the experimental saddr/saddr64 A-B kernels and scratch bench; the production kernels are now the single saddr64 path. Update the rle-fused test to the sole production launcher. Validated on gfx942 and gfx950: API, 2D, rle-fused, bitmap-encode, and bitmap-octree suites pass.
Opt-in, cached read of HIPCOMPRESS_DEBUG (0=silent). Level 1 prints the resolved scheme (with AUTO tag), dims, and block count at plan creation plus compressed bytes and CR at synchronize; level 2 adds device-memory detail. Zero output and a single cached branch when unset.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Improve CR via Octree/QuadTree for significance map