[Seth-Wolfgang]

Restore the HIP backend on main; stream-ordered two-array PFPL with AMD kernels; fix FMA contraction in the ABS bound check

Base branch: main (e84726f)
Head: main of Seth-Wolfgang/FZGPUModules (74855a4)
Author: Seth Wolfgang (seawolfg@iu.edu)

Summary

This PR makes five changes:

• HIP on main: -DFZGMOD_BACKEND=HIP configures and builds again. This is CMake only, and it's a separate commit.
• fz::pfpl: a float or double field compresses from one device buffer straight into another and decompresses straight back, in one kernel each. Calls are ordered on a stream and need no host synchronization. A float stream's archive is byte-identical to the PFPL preset's in ABS mode. Host functions write and read the same stream with no GPU. On an A100 it runs about 2x the staged pipeline at 8 MiB and above.
• AMD kernels for it: on MI250X, at 32 MiB and above, they encode about 5x and decode about 3x faster than the staged pipeline. At 1 MiB the gap is about 8x and 7x.
• A correctness fix: on HIP, the ABS bound check could accept values that decode past the bound.
• A link fix: fzgmod_modules calls fz::MemoryPool but didn't link fzgmod_mem, so consumers failed to link on toolchains that default to --as-needed.

I'm opening this as a draft because fz::pfpl is a new API and because main currently rejects HIP on purpose ("Stable main supports CUDA only"), so I'd like your view on both. My fork's main carries all of them:

Commit	Contents
f826460	fz::pfpl, the AMD kernels and the bound fix
6d22386	HIP on main, CMake only
07919bf	The Plan race fix below, tests and notes
6ce75f5	The link fix
eef094a	Doubles
1811d45	Changelog and notes for doubles
c7b048a	The host coder
74855a4	Wording in docs/codebase_notes.md, no code

They stack in this order. The link fix also stands alone: fix-modules-mem-link has it on your main. The host coder and the wording commit after it are the last two, so they're easy to leave out.

Why

A staged pipeline is too heavy for a caller that compresses and decompresses many fields from several streams on each GPU, thousands of times per run: one pipeline per stream ran out of device memory or crashed. Fields also need to be ordered with the kernels around them on the same stream, without a host round trip.

The host coder is for reads that have no device to use: a process with no GPU, or a card too full to stage on. With 0.15 GiB of 64 free on a GCD, decodes staged through the device stalled inside ROCm, with stacks in libhsa-runtime64: 8 of 12 runs stuck.

The bound fix

__fmul_rn is a plain multiply on HIP, so clang's default -ffp-contract=fast-honor-pragmas fuses x - q * ebx2 into one v_fma_f32 on gfx90a. The check then tests the exact product instead of the rounded one the decoder stores, and passes values that miss the bound by up to half an ulp.

include/backend/fp.h adds fz::backend::fmulRounded(), which turns contraction off on HIP and stays __fmul_rn on CUDA (__dmul_rn for doubles). The staged quantizer (quantizer.cu) and the chunk harness (chunk_fusion.cuh) both use it now.

To reproduce (HIP backend, gfx90a, without this fix):

1. Compress a float field with the PFPL preset in ABS mode, with verify_reconstruction on.
2. Decompress it and compare each value with the input.
3. Some values exceed abs_eb by up to half an ulp. With the fix, the largest error is within abs_eb.

The link fix

fzgmod_modules calls fz::MemoryPool but never declared the dependency, so libfzgmod_modules.so had no DT_NEEDED entry for libfzgmod_mem.so. A consumer that lists fzgmod_mem before fzgmod_modules fails under --as-needed, which Ubuntu's GCC defaults to.

With GNU ld 2.45, the old library gives 6 undefined fz::MemoryPool references, and the fixed one links and runs. ROCm's lld links either way, but with the old library libfzgmod_mem.so is then never loaded.

The two-array API

#include "fused/chunk_fusion/pfpl_stream.h"

fz::pfpl::Plan plan;                 // one per concurrently used stream
fz::pfpl::compress(plan, d_in, n, abs_eb, d_out, capacity, d_stream_bytes, stream);
fz::pfpl::decompress(plan, d_stream, d_out, n, stream);
size_t bytes = fz::pfpl::streamBytes(h_stream);   // reads the header in host memory

• Stream format: a 16-byte header (magic, archive bytes, abs_eb, n) followed by the LC/RZE archive. The pipeline is Quantizer in-place zigzag, Difference negabinary, Bitshuffle, then RZE over 16 KiB chunks.
• Device memory: only the two arrays, plus 8 bytes of look-back state per chunk and one ticket counter in the Plan.
• Overflow: nothing is written at or past capacity. The header, when it fits, records the length the stream needed.
• Host synchronization: only Plan::reserve synchronizes the device, and only when it has to replace a smaller state. New memory is cleared on the calling stream (CN-PFPLSTREAM-2).
• Types: float and double. A double stream has its own magic and reads back only as doubles (CN-PFPLSTREAM-3).
• Host coder: compressHost and decompressHost write and read the same stream, byte for byte, with no device, Plan or GPU (CN-PFPLSTREAM-4).
• Limits: the stream must be 4-byte aligned and the field aligned for its type, and a field must code to under 4 GiB.
• Invalid input: decoding a buffer that isn't a PFPL stream of n elements of that type stops the kernel with a device trap, and makes decompressHost return false, instead of decoding garbage.

Changes by area

• Root CMakeLists.txt: fzgmod_modules links fzgmod_mem.
• New files: pfpl_stream.h, pfpl_stream.cu, pfpl_stream_amd.inl and pfpl_stream_host.inl in modules/fused/chunk_fusion/, and include/backend/fp.h.
• AMD kernels (HIP default): wave64 with 19 KiB of shared memory and 3 blocks per CU. They use 6-7 barriers per chunk, where the generic harness uses about 20. FZ_PFPL_HARNESS=1 selects the generic harness kernels for comparison. CUDA always uses the generic harness.
• include/backend/warp.h: whole-wavefront primitives (kWaveSize, wave_mask_t, ballotWave, ffsWave, waveSum) and activeMask32.
• include/backend/api.h: aliases for cooperative launch.
• chunk_fusion.cuh:
  • Raw __shfl* and ballot intrinsics now go through the backend/warp.h wrappers.
  • The decode body is split out as chunk_inverse_pfpl_at.
  • A new loadChunkBytes replaces the byte-wise payload copy, so the CUDA decode path changes here too. On an A100 the specialized decode gets 7-16% faster (table below).
• HIP build fix: a const void* cast for cudaFuncSetAttribute in Bitshuffle.
• tests/stages/test_pfpl_stream.cpp: compares fz::pfpl with the staged pipeline, and the host coder with the kernels. On HIP, CTest runs it twice, once with the AMD kernels and once with the generic harness kernels.
• Notes: CN-PFPLSTREAM-1 holds the measurements, CN-PFPLSTREAM-2 the first-kernel race below, CN-PFPLSTREAM-3 doubles and CN-PFPLSTREAM-4 the host coder.

Restoring HIP on main

This commit changes CMake only:

• Root CMakeLists.txt, CMakePresets.json, cmake/FZGPUModulesConfig.cmake.in: the HIP hunks from development, without its profiling/ target, which main doesn't have. Module and test CMake files already branch on the backend, and include/backend already supports HIP.
• modules/fused/common/CMakeLists.txt: the NVRTC fused sources aren't built on HIP. findFusedImpl returns nothing and planAndInstallFusion records unsupported_backend, so every stage runs staged.
• Four tests are CUDA-only: test_fusion_f64 and test_fusion_planner call the NVRTC generators, test_fusion_fixed needs an installed specialization, and test_lorenzo_quant_shared_init includes cuda_runtime.h directly.

A first-kernel race in Plan

Clearing a new Plan with a stream-less cudaMemset isn't ordered with a non-blocking stream's kernels. The first encode could read the page's old contents as its ticket counter, compute a wild chunk id, and stop with HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION. The clear now runs on the calling stream. PfplStream.FreshPlanIgnoresStaleMemory faults in its first round without the fix, on both kernel sets, and passes with it. Details are in CN-PFPLSTREAM-2.

Results

MI250X

One GCD, ROCm 7.1.1, the polar field from test_pfpl_stream.cpp, abs_eb 1e-4, ratio about 1.7. In GB/s, measured on this branch:

Field	Encode: two-array	Encode: staged	Decode: two-array	Decode: staged
1 MiB	57.3	6.7	46.1	6.8
8 MiB	151.6	25.7	129.2	29.2
32 MiB	191.7	37.1	156.9	45.4
128 MiB	203.3	41.8	165.6	50.3
512 MiB	202.5	43.0	167.7	52.0

Doubles, on one MI250X GCD with a 128 MiB polar field, in GB/s. This is a separate run, so compare within this table:

Type, bound	Ratio	AMD encode	AMD decode	Generic encode	Generic decode
float, 1e-4	1.74	186.4	164.5	62.5	55.1
double, 1e-4	3.56	219.2	182.5	70.0	65.1
double, 1e-9	1.81	214.2	177.2
double, 1e-12	1.40	207.9	173.0	68.1	61.6

The host coder, on one Frontier login core with no GPU in the process and a 128 MiB polar field, encodes at 350-565 MB/s and decodes at 550-656 MB/s.

A100

A100-SXM4-40GB, CUDA 12.5, one stream. The field and bound are the MI250X table's, and the two-array columns are the generic harness, the only kernels CUDA has. Specialized is the same pipeline with SpecializationPolicy::Auto, which HIP doesn't have. In GB/s:

Field	Encode: two-array	Encode: staged	Encode: specialized	Decode: two-array	Decode: staged	Decode: specialized
1 MiB	44.4	11.2	12.9	39.2	11.6	21.6
8 MiB	111.7	53.5	65.6	132.8	55.6	100.9
32 MiB	164.4	76.2	125.5	171.7	76.5	158.7
128 MiB	193.8	88.6	165.5	187.5	85.5	188.2
512 MiB	201.1	93.7	176.6	191.7	93.3	197.9

Two-array is 2.1-2.4x staged at 8 MiB and above, and 3.4-4.0x at 1 MiB. Against specialized it encodes 1.1-3.4x faster, and decodes 1.8x faster at 1 MiB, 1.08x at 32 MiB and up to 3% slower at 128 MiB and above.

Each number is the median of 7 rounds of 11 to 41 calls, on fresh buffers each round. A second process agrees within 1% at 128 MiB and above, and within 9% below. This card rests at 1095 MHz between runs of work and takes a few tenths of a second of work to reach 1410 MHz, so every row warms up first; without that the large sizes read about 21% low.

Doubles, with the same 128 MiB field, in GB/s. The ratios match MI250X's, and every decoded value is within its bound:

Type, bound	Ratio	Encode	Decode
float, 1e-4	1.74	193.3	187.2
double, 1e-4	3.56	240.0	213.3
double, 1e-9	1.81	218.6	203.0
double, 1e-12	1.40	206.2	195.0

The host coder, on one EPYC 7713 core of the same VM and the same field, encodes at 204-380 MB/s and decodes at 308-333 MB/s (float at 1e-3, double at 1e-3 and 1e-9), and its streams equal the kernels' byte for byte.

The CUDA decode path against your main, specialized pipeline, in GB/s. Encode and the staged pipeline don't move beyond the noise above:

Field	Decode: e84726f	Decode: this branch
1 MiB	20.0	21.6
8 MiB	94.4	100.9
32 MiB	136.3	158.7
128 MiB	162.2	188.2
512 MiB	170.5	197.9

Testing

74855a4 changes only docs/codebase_notes.md, so the results at c7b048a below stand for it.

On c7b048a, HIP, ROCm 7.1.1, MI250X (gfx90a), from a clean clone:

• The library and all tests build, and cmake --preset hip configures.
• ctest: 50 of 53 pass, on one GCD.
• test_pfpl_stream passes all 488 of its stream tests with the AMD kernels and again with the generic harness kernels. They check:
  • archive and decode identity with the staged pipeline, and the bound (five field types, five lengths including ones that aren't whole chunks, three bounds);
  • doubles, with the expected codes computed on the host and the archive compared with the staged stages, since no staged pipeline matches the whole double chain;
  • the host coder against the kernels, byte for byte, for both types (six field kinds including one built to hit conversions out of range, lengths from 1 to 2^20, three bounds);
  • overflow containment;
  • 4 concurrent streams against serial encodes;
  • 300 repeated calls on one Plan;
  • a new Plan on dirty pages.
• The note and TOML documentation checks pass. check_doc_cpp.py skipped itself because this machine has no CUDA headers.

On c7b048a, CUDA 12.5 (nvcc from NVHPC 24.7, GCC 13.3), A100 (sm_80), from a clean clone:

• ctest: 57 of 57 pass. That includes the four CUDA-only tests, and test_huffman, test_log_transform and test_stage_registry, which fail on HIP for the reasons in the table below.
• test_pfpl_stream passes all 488, so doubles and the host coder hold on CUDA's 32-lane warps, including the host coder against the kernels byte for byte.
• ASan and UBSan (scripts/run_sanitizers.sh --mode asan, the asan preset's configuration at -DCMAKE_CUDA_ARCHITECTURES=native): all 57 test binaries pass with no reports, in 353 s.
• compute-sanitizer (scripts/run_sanitizers.sh --mode compute, memcheck, initcheck, racecheck and synccheck over all 57 test binaries): no errors and no hazards. Two things to know:
  • test_pfpl_stream is too slow under racecheck to finish whole, so racecheck covers 420 of its 488 tests. It leaves out the 2 concurrent-stream tests, 2 many-round tests, and the largest fields except one polar case per type. Memcheck, initcheck and synccheck run all 488.
  • DeviceBuffer.OwnedBufferFreesOnDestruction failed once under memcheck when another process started on the card during it, because it reads the whole card's free memory. It passes with the card to itself.
• As a control, memcheck does see these kernels. Decoding into a short buffer is reported in decodeChunk64 and chunk_inverse_pfpl_at, and encoding from a short one in encodeChunk64.
• check_doc_cpp.py, which skipped itself on Frontier, runs here with CUDA's headers on CPLUS_INCLUDE_PATH: no example references stale API (68 of 112 compile as they stand, 44 are fragments and 5 are skipped). The note and TOML checks pass too.

CUDA, compile only, nvcc 13.1 and GCC 12.2: at 07919bf, for sm_80 and sm_86, the library and every test build except test_jit_cache, which this branch doesn't change. Its cuCtxCreate(&ctx, 0, device) call has too few arguments for CUDA 13.1's cuCtxCreate_v4. At c7b048a, the library and test_pfpl_stream build for sm_80. backend/fp.h is in the embedded NVRTC headers.

The 3 ctest failures on HIP are outside PFPL, and I left them alone:

Test	Failing cases	Cause
test_huffman	3	phf_header has interior padding that {} leaves undefined on clang, so the host-coordinated path writes stale bytes into the archive and they differ from the device path's
test_log_transform	RelativeBoundHolds	ROCm's exp2f and log2f miss the round-trip slack on 1 value in 1536. The stage stores it as a lossless outlier, as designed, but the test requires 0 outliers
test_stage_registry	EveryShippedStageTypeHasAFactory	The test expects ANS, which isn't built on HIP

Not run:

• The invalid-input trap: nothing tests that decoding a foreign buffer stops the kernel.
• The installed-package path on HIP: FZGPUModulesConfig.cmake.in wasn't exercised, because the only consumer tested uses add_subdirectory.
• CUDA 13.1 at run time: the A100 runs used CUDA 12.5. CUDA 13.1 is compile-only, above.