Optimized PFPL Performance for ROCm

CHANGELOG.md

@@ -9,7 +9,19 @@ Version numbers follow [Semantic Versioning](https://semver.org/).
 
 ## [Unreleased]
 
+### Added
+- `fz::pfpl` (`fused/chunk_fusion/pfpl_stream.h`): a float field compresses from one device buffer into another and back in one stream-ordered kernel each, with no host synchronization and the same archive bytes as the PFPL preset in ABS mode.
+- Wave64 AMD kernels for `fz::pfpl`, used by default on HIP; `FZ_PFPL_HARNESS=1` selects the generic harness kernels. See CN-PFPLSTREAM-1.
+- `fz::pfpl` takes double fields as well: a 64-bit zigzag code, difference and negabinary, a 64-plane bitshuffle, and the same RZE over 16 KiB chunks. Double streams carry their own magic and read back only as doubles. See CN-PFPLSTREAM-3.
+- `fz::pfpl::compressHost` and `decompressHost`: the same stream, byte for byte, in host memory, with no device, `Plan` or GPU. See CN-PFPLSTREAM-4.
+- `-DFZGMOD_BACKEND=HIP` configures again, with `hip` CMake presets. NVRTC specialization is not built on HIP, so every stage runs staged, and `test_fusion_f64`, `test_fusion_fixed`, `test_fusion_planner` and `test_lorenzo_quant_shared_init` are CUDA-only.
+
+### Fixed
+- On HIP, the ABS bound check could accept values that decode up to half an ulp past the bound, because clang fused its multiply and subtract into one FMA. `fz::backend::fmulRounded` keeps the product rounded in the staged quantizer and the chunk harness.
+- `fzgmod_modules` now links `fzgmod_mem`, which it calls into, so a consumer of the shared library links on toolchains that default to `--as-needed` (such as Ubuntu's GCC).
+
 ### Changed
+- The chunk harness decode copies payloads with `loadChunkBytes` instead of byte by byte, on CUDA as well.
 - CMake now locates the build-time Python interpreter with `find_package(Python3 3.7 REQUIRED COMPONENTS Interpreter)` and runs the header generators through `${Python3_EXECUTABLE}`, so a missing or too-old Python fails at configure time instead of mid-build. Override with `-DPython3_EXECUTABLE=...`.
 
 ## [2.0.1] — 2026-10-15

CMakeLists.txt

@@ -1,15 +1,26 @@
 cmake_minimum_required(VERSION 3.24 FATAL_ERROR)
 
 
-set(FZGMOD_BACKEND "CUDA" CACHE STRING "GPU backend (stable main supports CUDA)")
-set_property(CACHE FZGMOD_BACKEND PROPERTY STRINGS CUDA)
-if(NOT FZGMOD_BACKEND STREQUAL "CUDA")
-  message(FATAL_ERROR "Stable main supports CUDA only; backend development is on the development branch.")
-endif()
-if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
-  set(CMAKE_CUDA_ARCHITECTURES "86" CACHE STRING "CUDA architectures to compile for")
+set(FZGMOD_BACKEND "CUDA" CACHE STRING "GPU backend to build against: CUDA, HIP, or SYCL")
+set_property(CACHE FZGMOD_BACKEND PROPERTY STRINGS CUDA HIP SYCL)
+
+if(FZGMOD_BACKEND STREQUAL "CUDA")
+  set(FZGMOD_PROJECT_LANGUAGES CUDA CXX)
+  if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+    set(CMAKE_CUDA_ARCHITECTURES "86" CACHE STRING "CUDA architectures to compile for")
+  endif()
+elseif(FZGMOD_BACKEND STREQUAL "HIP")
+  set(FZGMOD_PROJECT_LANGUAGES CXX)
+  message(WARNING
+    "FZGMOD_BACKEND=HIP is experimental: backend coverage and install/consumer "
+    "compatibility are still being completed, and some stages are unavailable.")
+elseif(FZGMOD_BACKEND STREQUAL "SYCL")
+  message(FATAL_ERROR "FZGMOD_BACKEND=SYCL is not implemented yet.")
+else()
+  message(FATAL_ERROR "Unknown FZGMOD_BACKEND '${FZGMOD_BACKEND}' — must be CUDA, HIP, or SYCL.")
 endif()
-project(FZGPUModules LANGUAGES CUDA CXX VERSION 2.0.1)
+
+project(FZGPUModules LANGUAGES ${FZGMOD_PROJECT_LANGUAGES} VERSION 2.0.1)
 
 # General settings
 
@@ -39,18 +50,76 @@ endif()
 set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_STANDARD_REQUIRED True)
 
-# CUDA toolchain
-find_package(CUDAToolkit REQUIRED)
-if(CUDAToolkit_VERSION VERSION_LESS "11.2")
-  message(FATAL_ERROR "FZGPUModules requires CUDA Toolkit 11.2 or newer for stream-ordered allocation.")
+# GPU backend toolchain
+# Link the selected runtime through FZGMOD_GPU_RUNTIME.
+
+if(FZGMOD_BACKEND STREQUAL "CUDA")
+  find_package(CUDAToolkit REQUIRED)
+  if(CUDAToolkit_VERSION VERSION_LESS "11.2")
+    message(FATAL_ERROR
+      "FZGPUModules requires CUDA Toolkit 11.2 or newer.\n"
+      "  Reason: stream-ordered allocators (cudaMallocAsync / cudaMemPool_t) require 11.2+.\n"
+      "  Detected: ${CUDAToolkit_VERSION}")
+  endif()
+
+  # Set CUDA architectures before project() enables the language.
+  enable_language(CUDA)
+
+  set(FZGMOD_GPU_RUNTIME CUDA::cudart)
+
+else() # HIP — SYCL and unknown values already rejected above
+
+  # Set HIP architectures before enabling the language; default to MI100.
+  if(NOT DEFINED CMAKE_HIP_ARCHITECTURES)
+    set(CMAKE_HIP_ARCHITECTURES "gfx908" CACHE STRING
+        "AMD GPU architectures to compile for (gfx908=MI100, gfx90a=MI200, gfx942=MI300)")
+  endif()
+
+  enable_language(HIP)
+
+  # Pass explicit targets to hip-config, including on GPU-less build nodes.
+  set(GPU_TARGETS "${CMAKE_HIP_ARCHITECTURES}" CACHE STRING "AMD GPU targets for hip-config" FORCE)
+
+  # Put the ROCm root on CMAKE_PREFIX_PATH for CONFIG-mode discovery.
+  find_package(hip       CONFIG REQUIRED)
+  find_package(hipcub    CONFIG REQUIRED)  # cub facade — mirrors cub's API
+  find_package(rocthrust CONFIG REQUIRED)  # thrust facade, used by chunk_fusion/rze/golomb_rice
+
+  # Use hip::host: hip::device would apply -x hip to plain C++ sources.
+  # Details: docs/codebase_notes.md CN-BUILD-1.
+  set(FZGMOD_GPU_RUNTIME hip::host)
+
+  # Use header-only include paths to avoid transitively linking hip::device.
+  set(FZGMOD_HIP_HEADER_ONLY_INCLUDES "")
+  foreach(_hdr_tgt hip::hipcub roc::rocprim roc::rocthrust)
+    if(TARGET ${_hdr_tgt})
+      get_target_property(_inc ${_hdr_tgt} INTERFACE_INCLUDE_DIRECTORIES)
+      if(_inc)
+        list(APPEND FZGMOD_HIP_HEADER_ONLY_INCLUDES ${_inc})
+      endif()
+    endif()
+  endforeach()
+  list(REMOVE_DUPLICATES FZGMOD_HIP_HEADER_ONLY_INCLUDES)
+
 endif()
-set(FZGMOD_GPU_RUNTIME CUDA::cudart)
 
 # Build-time header generators (stage fingerprints, embedded NVRTC headers).
 find_package(Python3 3.7 REQUIRED COMPONENTS Interpreter)
 
-# Retain the module-local source hook; CUDA needs no retagging.
+# Retag .cu sources as HIP; define this before subdirectories use it.
+# Backend rationale: docs/codebase_notes.md CN-BUILD-1.
 function(fzgmod_set_gpu_source_language target)
+  if(NOT FZGMOD_BACKEND STREQUAL "HIP")
+    return()
+  endif()
+  get_target_property(_srcs ${target} SOURCES)
+  foreach(_src IN LISTS _srcs)
+    if(_src MATCHES "\\.cu$")
+      set_source_files_properties(${_src}
+        TARGET_DIRECTORY ${target}
+        PROPERTIES LANGUAGE HIP)
+    endif()
+  endforeach()
 endfunction()
 
 # WSL detection
@@ -135,7 +204,7 @@ elseif (USE_SANITIZER STREQUAL "Compute")
   if(NOT FZGMOD_BACKEND STREQUAL "CUDA")
     message(FATAL_ERROR
       "USE_SANITIZER=Compute requires the CUDA backend (compute-sanitizer is an "
-      "NVIDIA tool).")
+      "NVIDIA tool). Sanitizer parity for HIP is not implemented.")
   endif()
   message(STATUS "Configuring for Compute Sanitizer")
   set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -lineinfo -Xcompiler -rdynamic")
@@ -163,7 +232,20 @@ target_compile_options(
 )
 
 target_compile_features(fzgmod_compile_settings INTERFACE cxx_std_17)
-target_compile_features(fzgmod_compile_settings INTERFACE cuda_std_17)
+if(FZGMOD_BACKEND STREQUAL "CUDA")
+  target_compile_features(fzgmod_compile_settings INTERFACE cuda_std_17)
+else()
+  set(CMAKE_HIP_STANDARD 17)
+  set(CMAKE_HIP_STANDARD_REQUIRED True)
+
+  # HIP teardown paths intentionally ignore [[nodiscard]] runtime status.
+  target_compile_options(fzgmod_compile_settings INTERFACE
+    $<$<COMPILE_LANGUAGE:CXX,HIP>:-Wno-unused-result>)
+
+  # ROCm requires this flag for the sync warp intrinsics.
+  target_compile_definitions(fzgmod_compile_settings INTERFACE
+    HIP_ENABLE_WARP_SYNC_BUILTINS)
+endif()
 
 target_link_libraries(fzgmod_compile_settings INTERFACE ${FZGMOD_GPU_RUNTIME})
 
@@ -188,8 +270,18 @@ target_include_directories(fzgmod_compile_settings INTERFACE
   $<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/fzgmod>
 )
 
+# Treat HIP dependency headers as system headers.
+if(FZGMOD_HIP_HEADER_ONLY_INCLUDES)
+  target_include_directories(fzgmod_compile_settings SYSTEM INTERFACE
+    ${FZGMOD_HIP_HEADER_ONLY_INCLUDES})
+endif()
+
+# Production annotations can be enabled without building research programs.
 option(ENABLE_NVTX "Enable NVTX annotations in production library code" OFF)
 if(ENABLE_NVTX)
+  if(NOT FZGMOD_BACKEND STREQUAL "CUDA")
+    message(FATAL_ERROR "ENABLE_NVTX requires the CUDA backend.")
+  endif()
   find_path(NVTX3_INCLUDE_DIR NAMES nvtx3/nvtx3.hpp
     HINTS ${CUDAToolkit_INCLUDE_DIRS} REQUIRED)
   target_compile_definitions(fzgmod_compile_settings INTERFACE
@@ -228,6 +320,7 @@ target_link_libraries(fzgmod_modules
     fzgmod_compile_settings
     ${FZGMOD_GPU_RUNTIME}
   PRIVATE
+    fzgmod_mem
     fzgmod_utils
 )
 add_subdirectory(modules)
@@ -456,6 +549,10 @@ endif()
 message(STATUS "FZGPUModules configuration complete")
 message(STATUS "  Backend: ${FZGMOD_BACKEND}")
 message(STATUS "  Build type: ${CMAKE_BUILD_TYPE}")
-message(STATUS "  CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}")
+if(FZGMOD_BACKEND STREQUAL "CUDA")
+  message(STATUS "  CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}")
+else()
+  message(STATUS "  HIP architectures: ${CMAKE_HIP_ARCHITECTURES}")
+endif()
 message(STATUS "  Install prefix: ${CMAKE_INSTALL_PREFIX}")
 message(STATUS "  Build examples: ${BUILD_EXAMPLES}")

CMakePresets.json

@@ -44,6 +44,16 @@
       }
     },
 
+    {
+      "name": "hip",
+      "inherits": "base",
+      "displayName": "HIP (AMD ROCm, experimental)",
+      "description": "Experimental release build against the HIP backend. Defaults to gfx908 (MI100); override with -DCMAKE_HIP_ARCHITECTURES=. Requires the ROCm root on CMAKE_PREFIX_PATH, e.g. export CMAKE_PREFIX_PATH=/opt/rocm:$CMAKE_PREFIX_PATH",
+      "cacheVariables": {
+        "CMAKE_BUILD_TYPE": "Release",
+        "FZGMOD_BACKEND":   "HIP"
+      }
+    },
 
     {
       "name": "cuda-h200",
@@ -87,6 +97,11 @@
       "configurePreset": "asan",
       "displayName": "Build ASan + UBSan"
     },
+    {
+      "name": "hip",
+      "configurePreset": "hip",
+      "displayName": "Build HIP (AMD ROCm, experimental)"
+    },
     {
       "name": "cuda-h200",
       "configurePreset": "cuda-h200",
@@ -129,6 +144,14 @@
       "filter": { "include": { "label": "pipeline" } }
     },
 
+    {
+      "name": "hip",
+      "configurePreset": "hip",
+      "displayName": "All tests (HIP, experimental)",
+      "description": "Run the full test suite against the experimental HIP build. ANSStage is excluded on HIP (see Stage::isSupportedOnBackend).",
+      "execution": { "noTestsAction": "error" },
+      "output": { "outputOnFailure": true }
+    },
 
     {
       "name": "cuda-h200",

cmake/FZGPUModulesConfig.cmake.in

@@ -4,7 +4,20 @@ include(CMakeFindDependencyMacro)
 
 set(FZGPUModules_BACKEND "@FZGMOD_BACKEND@")
 
-find_dependency(CUDAToolkit REQUIRED)
+if(FZGPUModules_BACKEND STREQUAL "CUDA")
+  find_dependency(CUDAToolkit REQUIRED)
+elseif(FZGPUModules_BACKEND STREQUAL "HIP")
+  # HIP support is experimental. These packages provide the runtime target and
+  # public headers referenced by the exported FZGPUModules targets.
+  find_dependency(hip CONFIG REQUIRED)
+  find_dependency(hipcub CONFIG REQUIRED)
+  find_dependency(rocthrust CONFIG REQUIRED)
+else()
+  set(FZGPUModules_FOUND FALSE)
+  set(FZGPUModules_NOT_FOUND_MESSAGE
+      "This FZGPUModules package was built with unknown backend '${FZGPUModules_BACKEND}'.")
+  return()
+endif()
 
 include("${CMAKE_CURRENT_LIST_DIR}/FZGPUModulesTargets.cmake")
 

docs/codebase_notes.md

@@ -1816,3 +1816,96 @@ it in `stage_registry.h`, and add FZGM integration tests. Use the existing SDK
 capability guards for optional families. Registration order retains the previous
 first-match behavior; individual handlers own their token aliases. The main
 plugin requires no stage-specific edits.
+
+## CN-PFPLSTREAM-1 — Why `fz::pfpl` has AMD kernels, and what they gain
+
+**Source:** `modules/fused/chunk_fusion/pfpl_stream_amd.inl`.
+
+The generic harness kernels keep a chunk in shared memory with block-wide scans
+and about 20 barriers per chunk. Profiled on MI250X, they wait about 65% of the
+time at 8 waves per CU. Their 36 KiB of shared memory also limits a CU to one
+block. The AMD kernels produce the same stream, byte for byte, with 6-7 barriers
+and 19 KiB of shared memory, which fits 3 blocks per CU.
+
+Measured on one MI250X GCD with ROCm 7.1.1, on a polar field with `abs_eb` 1e-4
+and a compression ratio of about 1.7. Throughput is in GB/s:
+
+| 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 |
+
+## CN-PFPLSTREAM-2 — A new `Plan` clears its memory on the calling stream
+
+**Source:** `Plan::reserve` and `Plan::nextEpoch` in
+`modules/fused/chunk_fusion/pfpl_stream.cu`.
+
+`Plan::reserve` used to clear its new ticket counter and look-back state with a
+stream-less `cudaMemset`. On MI250X with ROCm 7.1.1 that clear could land after
+the first encode kernel on a non-blocking stream had started. A read on the
+calling stream right after the call returned the page's previous contents, and
+the same read after a 50 ms wait, with no GPU synchronization, returned zero.
+
+A stale counter gives `atomicAdd(ticket, 1) - ticket_base` a wild chunk id, so the
+kernel reads from `in + cid * kElems`. Both the AMD and the harness encode kernels
+stopped with `HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`. Stale look-back words
+tagged with the same epoch would give a wrong offset with no fault.
+
+The clears now run in `Plan::nextEpoch` with `cudaMemsetAsync` on the calling
+stream, ahead of that call's kernel, so they need no host synchronization.
+`PfplStream.FreshPlanIgnoresStaleMemory` keeps the default stream busy so that a
+clear queued there always loses the race. Without this fix it faults in its first
+round on both kernel sets.
+
+## CN-PFPLSTREAM-3 — Doubles in `fz::pfpl`
+
+**Source:** `quantCode64` and `encodeChunk64` in `modules/fused/chunk_fusion/pfpl_stream.cu`, `Layout<Real>` in `pfpl_stream_amd.inl`.
+
+A double takes the float chain's steps in 64 bits. The chunk stays 16 KiB, so it holds 2048 elements in 64 planes of 256 bytes,
+and everything after the bitshuffle is unchanged. A 64-plane bitshuffle is two 32-plane transposes: a group's high words give
+planes 0..31 and its low words planes 32..63, with the groups of a plane in swapped pairs. That is the staged 8-byte
+`BitshuffleStage`'s order.
+
+The staged quantizer has no double form with in-place outliers, so no staged pipeline matches the whole chain.
+`PfplStream64Case.MatchesStagedStages` computes the expected codes on the host and compares the archive with staged
+`DifferenceStage<int64_t, uint64_t>`, `BitshuffleStage` and `RZEStage`, byte for byte.
+
+- **Quantization radius:** 2^51, the analog of the float preset's 2^22. A raw value's bits have to read as raw on decode, and
+  the bits of a normal double lie at or above 2 * radius. A subnormal always quantizes to index 0 at an allowed bound.
+- **Bound:** the 16-byte header is unchanged, so the stream records the bound as a float, rounded down, and that rounded value
+  is the bound the field is coded to. It has to be at least the smallest normal float. A power-of-two bound is exact.
+- **Rounding:** `rint`, with a range check before the conversion, which is undefined out of range.
+
+Measured on one MI250X GCD with ROCm 7.1.1, on a 128 MiB polar field (throughput in GB/s of raw data):
+
+| 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 |
+
+## CN-PFPLSTREAM-4 — The host coder
+
+**Source:** `modules/fused/chunk_fusion/pfpl_stream_host.inl`.
+
+A caller that has to read or write a stream off the GPU path used to stage it through the device, and that fails two ways.
+Threads reading fields back each allocate a field's worth of device memory beside whatever the card already holds, and double
+fields of 512 MiB stopped with out of memory. And a process that sees no GPU cannot code at all.
+
+`compressHost` and `decompressHost` write and read the same stream as the kernels, in plain C++, one thread per call, with stack
+buffers only. `PfplHostCase` holds them to the kernels byte for byte on both kernel sets, both types and every size class. Its
+`hostile` field exists for the one place they can differ, a float-to-int conversion out of range: the GPU saturates and the host
+conversion is undefined, so the host code tests the range first, and both then store the raw bits. `-ffp-contract` matters as
+it does in the kernels (see `fmulRounded`), so the rebuilt value goes through a `volatile`.
+
+Measured on one Frontier login core, with no GPU in the process, on a 128 MiB polar field:
+
+| Type, bound | Ratio | Encode | Decode |
+|---|---:|---:|---:|
+| float, 1e-3 | 2.43 | 350 MB/s | 550 MB/s |
+| double, 1e-3 | 4.87 | 565 MB/s | 656 MB/s |
+| double, 1e-9 | 1.92 | 460 MB/s | 638 MB/s |

include/backend/api.h

@@ -46,6 +46,7 @@
 #define cudaDevAttrMaxSharedMemoryPerBlock        hipDeviceAttributeMaxSharedMemoryPerBlock
 #define cudaDevAttrMaxSharedMemoryPerBlockOptin   hipDeviceAttributeSharedMemPerBlockOptin
 #define cudaDevAttrMultiProcessorCount            hipDeviceAttributeMultiprocessorCount
+#define cudaDevAttrCooperativeLaunch              hipDeviceAttributeCooperativeLaunch
 
 // ── Handle types
 #define cudaError_t                 hipError_t
@@ -75,6 +76,7 @@
 #define cudaDeviceGetAttribute                          hipDeviceGetAttribute
 #define cudaDeviceSynchronize                           hipDeviceSynchronize
 #define cudaFuncSetAttribute                            hipFuncSetAttribute
+#define cudaLaunchCooperativeKernel                     hipLaunchCooperativeKernel
 #define cudaFuncAttributeMaxDynamicSharedMemorySize     hipFuncAttributeMaxDynamicSharedMemorySize
 #define cudaOccupancyMaxActiveBlocksPerMultiprocessor   hipOccupancyMaxActiveBlocksPerMultiprocessor
 #define cudaProfilerStart                               hipProfilerStart