[ROCm] Add a HIP backend (PSZ_BACKEND=HIP) targeting ROCm

CMakeLists.txt

@@ -1,9 +1,12 @@
 # select backend
-set(PSZ_BACKEND "CUDA" CACHE STRING "Select the desired backend: CUDA/cuda and ONEAPI/1api.")
+set(PSZ_BACKEND "CUDA" CACHE STRING "Select the desired backend: CUDA/cuda, HIP/hip, ONEAPI/1api.")
 
 if(PSZ_BACKEND STREQUAL "CUDA" OR PSZ_BACKEND STREQUAL "cuda")
   cmake_minimum_required(VERSION 3.18)
   project(CUSZ LANGUAGES CUDA CXX C VERSION 0.10.0)
+elseif(PSZ_BACKEND STREQUAL "HIP" OR PSZ_BACKEND STREQUAL "hip")
+  cmake_minimum_required(VERSION 3.21)
+  project(HIPSZ LANGUAGES HIP CXX C VERSION 0.10.0)
 elseif(PSZ_BACKEND STREQUAL "ONEAPI" OR PSZ_BACKEND STREQUAL "oneapi" OR PSZ_BACKEND STREQUAL "1api")
   cmake_minimum_required(VERSION 3.23)
   project(DPSZ LANGUAGES CXX C VERSION 0.6.0)
@@ -14,7 +17,15 @@ set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
 option(BUILD_SHARED_LIBS "prefer shared libaries" ON)
 
 # conditionally enable functionality
-option(PSZ_ACTIVATE_LC "build LC as a submodule" ON)
+# The bundled LC-framework (third_party/lc) is an external CUDA codebase that is
+# not yet ported to HIP; default it OFF on the HIP backend so the core
+# Lorenzo/spline + Huffman/HFR + FZG pipeline builds. It can be force-enabled
+# once LC-framework gains HIP support.
+if(PSZ_BACKEND STREQUAL "HIP" OR PSZ_BACKEND STREQUAL "hip")
+  option(PSZ_ACTIVATE_LC "build LC as a submodule" OFF)
+else()
+  option(PSZ_ACTIVATE_LC "build LC as a submodule" ON)
+endif()
 
 # example and testing
 option(PSZ_BUILD_EXAMPLES "build example codes" ON)
@@ -25,6 +36,9 @@ option(PSZ_BUILD_PYBINDING "enable Python binding for specific backend" OFF)
 if(PSZ_BACKEND STREQUAL "CUDA" OR PSZ_BACKEND STREQUAL "cuda")
   message("[psz::info] CUDA backend has been selected.")
   include(cmake/cuda.cmake)
+elseif(PSZ_BACKEND STREQUAL "HIP" OR PSZ_BACKEND STREQUAL "hip")
+  message("[psz::info] HIP (ROCm) backend has been selected.")
+  include(cmake/hip.cmake)
 elseif(PSZ_BACKEND STREQUAL "ONEAPI" OR PSZ_BACKEND STREQUAL "oneapi" OR PSZ_BACKEND STREQUAL "1api")
   message("[psz::info] oneAPI backend has been selected.")
   include(cmake/sycl.cmake)

README.md

@@ -8,7 +8,7 @@ pSZ/cuSZ: A GPU-Based Error-Bounded Lossy Compressor for Scientific Data
 <a href="./LICENSE"><img src="https://img.shields.io/badge/License-BSD%203--Clause-blue.svg"></a>
 </p>
 
-pSZ/cuSZ (cuSZ for short) is a GPU implementation of the seminal [SZ algorithm](https://github.com/szcompressor/SZ). It is the *first* GPU-practical framework of error-bounded lossy compression on GPU for scientific data (circa 2020), aiming to improve SZ's throughput on heterogeneous HPC systems. pSZ/cuSZ primarily focuses on CUDA backend support, with other GPU-parallel backends in development. pSZ/cuSZ is formerly known as cuSZ, which is also the short form of its current name. 
+pSZ/cuSZ (cuSZ for short) is a GPU implementation of the seminal [SZ algorithm](https://github.com/szcompressor/SZ). It is the *first* GPU-practical framework of error-bounded lossy compression on GPU for scientific data (circa 2020), aiming to improve SZ's throughput on heterogeneous HPC systems. pSZ/cuSZ primarily focuses on CUDA backend support, with other GPU-parallel backends in development, including an AMD GPU (ROCm/HIP) backend selected at build time with `-DPSZ_BACKEND=HIP`. pSZ/cuSZ is formerly known as cuSZ, which is also the short form of its current name. 
 
 (c) 2025 by Argonne National Laboratory and Oakland University. See [COPYRIGHT](https://github.com/szcompressor/cuSZ/blob/master/LICENSE) in top-level directory.
 

cmake/hip-compat/cooperative_groups.h

@@ -0,0 +1,8 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// Shim placed on the HIP include path so sources that #include
+// <cooperative_groups.h> resolve against HIP's cooperative-groups header.
+#ifndef PSZ_HIP_COMPAT_COOPERATIVE_GROUPS_H
+#define PSZ_HIP_COMPAT_COOPERATIVE_GROUPS_H
+#include <hip/hip_cooperative_groups.h>
+#endif

cmake/hip-compat/cooperative_groups/memcpy_async.h

@@ -0,0 +1,9 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// Shim for sources that #include <cooperative_groups/memcpy_async.h>. HIP's
+// cooperative-groups support lives in a single header and does not split out a
+// memcpy_async sub-header; redirect to it so the include resolves.
+#ifndef PSZ_HIP_COMPAT_CG_MEMCPY_ASYNC_H
+#define PSZ_HIP_COMPAT_CG_MEMCPY_ASYNC_H
+#include <hip/hip_cooperative_groups.h>
+#endif

cmake/hip-compat/cuda.h

@@ -0,0 +1,9 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// Shim placed on the HIP include path so sources that #include <cuda.h> resolve
+// against the HIP runtime; the CUDA->HIP translation macros are supplied by the
+// force-included psz_hip_compat.h prelude / the cuda_runtime.h shim.
+#ifndef PSZ_HIP_COMPAT_CUDA_H
+#define PSZ_HIP_COMPAT_CUDA_H
+#include <hip/hip_runtime.h>
+#endif

cmake/hip-compat/cuda_fp16.h

@@ -0,0 +1,7 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// Shim so sources that #include <cuda_fp16.h> resolve against HIP's half type.
+#ifndef PSZ_HIP_COMPAT_CUDA_FP16_H
+#define PSZ_HIP_COMPAT_CUDA_FP16_H
+#include <hip/hip_fp16.h>
+#endif

cmake/hip-compat/cuda_runtime.h

@@ -0,0 +1,17 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// Shim placed on the HIP include path so sources that #include <cuda_runtime.h>
+// (kernel .cu TUs and host .cc TUs alike) resolve against the HIP runtime and
+// pick up the CUDA->HIP translation macros. Device TUs additionally get the
+// force-included psz_hip_compat.h prelude, but a plain host .cc that only
+// includes <cuda_runtime.h> is fully covered by this shim alone.
+#ifndef PSZ_HIP_COMPAT_CUDA_RUNTIME_H
+#define PSZ_HIP_COMPAT_CUDA_RUNTIME_H
+
+#include <hip/hip_runtime.h>
+
+#include "macro/c_cu2hip_0_translation.h"
+#include "macro/c_cu2hip_1_fix_primitives.h"
+#include "macro/c_cu2hip_2_suppress_warning.h"
+
+#endif

cmake/hip-compat/cuda_runtime_api.h

@@ -0,0 +1,9 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// Shim so sources that #include <cuda_runtime_api.h> resolve against the HIP
+// runtime. The CUDA->HIP translation macros are supplied by the cuda_runtime.h
+// shim / the force-included psz_hip_compat.h prelude.
+#ifndef PSZ_HIP_COMPAT_CUDA_RUNTIME_API_H
+#define PSZ_HIP_COMPAT_CUDA_RUNTIME_API_H
+#include <hip/hip_runtime.h>
+#endif

cmake/hip-compat/curand.h

@@ -0,0 +1,18 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// Shim placed on the HIP include path so sources that #include <curand.h>
+// resolve against hipRAND. The cuRAND host-generator API used by cuSZ
+// (portable/src/utils/rand.cu.cc) maps 1:1 onto hipRAND.
+#ifndef PSZ_HIP_COMPAT_CURAND_H
+#define PSZ_HIP_COMPAT_CURAND_H
+
+#include <hiprand/hiprand.h>
+
+#define curandGenerator_t hiprandGenerator_t
+#define curandCreateGenerator hiprandCreateGenerator
+#define curandSetPseudoRandomGeneratorSeed hiprandSetPseudoRandomGeneratorSeed
+#define curandGenerateUniform hiprandGenerateUniform
+#define curandGenerateUniformDouble hiprandGenerateUniformDouble
+#define CURAND_RNG_PSEUDO_DEFAULT HIPRAND_RNG_PSEUDO_DEFAULT
+
+#endif

cmake/hip-compat/psz_hip_compat.h

@@ -0,0 +1,22 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+//
+// HIP/ROCm backend compatibility prelude for cuSZ.
+//
+// This header is force-included (compiler -include) into every translation
+// unit compiled for the HIP backend. The kernel and host sources are written
+// in CUDA spelling and reused as-is for HIP (single-source, no .hip mirrors);
+// this prelude pulls in the HIP runtime and the CUDA->HIP translation macros so
+// those sources compile unchanged against ROCm. It is never on the CUDA or SYCL
+// include path, so those backends are unaffected.
+
+#ifndef PSZ_HIP_COMPAT_H
+#define PSZ_HIP_COMPAT_H
+
+#include <hip/hip_cooperative_groups.h>
+#include <hip/hip_runtime.h>
+
+#include "macro/c_cu2hip_0_translation.h"
+#include "macro/c_cu2hip_1_fix_primitives.h"
+#include "macro/c_cu2hip_2_suppress_warning.h"
+
+#endif /* PSZ_HIP_COMPAT_H */