Add AMD GPU support via ROCm/HIP (#94)

* [ROCm] Add AMD GPU support via ROCm/HIP

This adds an optional AMD GPU build to cuPDLPx through ROCm/HIP, alongside
the existing CUDA path. The CUDA build is unchanged when USE_HIP is off.

To review: start with internal/cuda_to_hip.h, which routes the CUDA runtime,
cuBLAS, cuSPARSE, and CUB symbols used by the solver to their hipRT, hipBLAS,
hipSPARSE, and hipCUB equivalents on a HIP build, and includes the standard
CUDA headers otherwise. The device sources keep their CUDA spelling and are
compiled as HIP. internal/cusparse_compat.h selects the standard hipsparseSpMV
path on ROCm, since hipSPARSE does not provide the cusparseSpMVOp variant.

CMakeLists.txt gains a USE_HIP option (off by default). When enabled the
project is configured with the HIP language, the .cu sources are compiled as
HIP, and the targets link hipBLAS, hipSPARSE, and hipCUB instead of the CUDA
libraries. GPU architectures are chosen with CMAKE_HIP_ARCHITECTURES,
defaulting to gfx90a. On Windows the CLI-only mps_parser.c is excluded from
the core library because it relies on strtok_r. The interface test gains a
case that runs the GPU solver path with presolve disabled, exercising the
hipBLAS and hipSPARSE execution path end to end.

Test Plan:

Built and ran on an AMD Instinct MI200 (gfx90a) with ROCm 7.2.1:

```
cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_PREFIX_PATH=/opt/rocm \
  -DCUPDLPX_BUILD_CLI=ON -DCUPDLPX_BUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build -j$(nproc)
./build/tests/test_interface
```

The interface suite passes, including the GPU solver case (Status: OPTIMAL).
The same configuration builds cleanly for gfx1100 (RDNA3) and gfx1201
(RDNA4); the device code objects are identical across the documentation and
formatting commits that followed validation. The CUDA build path is
unaffected by these changes.

This work was authored with the assistance of Claude, an AI assistant by
Anthropic.

* [ROCm] Fix CUDA build: guard cub include for C translation units

The ROCm support commit routed all CUDA/HIP includes through
internal/cuda_to_hip.h and pulled it into utils.h and internal_types.h,
which are included by the C translation units (cli.c, cupdlpx.c,
mps_parser.c, presolve.c). On the CUDA path that header included
<cub/device/device_reduce.cuh> unconditionally; cub is C++ only, so the C
compiler failed with "unknown type name 'namespace'", breaking every
CUDA build job (all Linux and Windows toolchains, CUDA 12.4 through 13.1).
The HIP path was unaffected because its hipcub include was already guarded
with #ifdef __cplusplus.

The fix mirrors that guard on the CUDA branch: the cub header is only
included for C++ translation units (the .cu device sources that actually
use cub::DeviceReduce). The change is entirely within the #else CUDA
branch, so the HIP/ROCm device code is unchanged.

Authored with assistance from Claude.

Test Plan: reproduced and verified the CUDA path locally with the CUDA
12.8 toolkit (gcc 13, ninja), matching the upstream CI configure:

```
cmake -B build -G Ninja -DCMAKE_BUILD_TYPE=Release \
  -DCUPDLPX_BUILD_TESTS=OFF -DCMAKE_CUDA_ARCHITECTURES=80
cmake --build build --clean-first
```

Before: cc -std=gnu99 -c src/cupdlpx.c fails on cub/device/device_reduce.cuh.
After: clean build, links cupdlpx and libcupdlpx.so with 0 errors.

* [ROCm] Simplify HIP compat includes and strengthen interface test

Addresses review feedback on the ROCm support PR.

cuda_to_hip.h already includes the CUDA runtime/cuBLAS/cuSPARSE headers on
CUDA builds and their HIP equivalents on ROCm builds, so the per-file
`#if !defined(USE_HIP)` include blocks in utils.h and preconditioner.cu were
redundant. They are removed in favor of relying on cuda_to_hip.h alone.

The USE_HIP definition moves from directory-scoped add_compile_definitions to
target_compile_definitions on the cupdlpx_compile_flags interface target so it
travels reliably to every consumer, including the Python bindings.

test_interface Test 9 previously only checked for a non-NULL result, so it
could not catch a wrong answer. It now asserts TERMINATION_REASON_OPTIMAL and
the known optimum (objective 3.0 within 1e-4). The default 1e-4 relative
tolerance stops the solver around 3.0005, so the test tightens the convergence
tolerance to 1e-8, after which the objective reaches the true optimum.

On the device-link question: the HIP build does not use relocatable device
code (-fgpu-rdc is off), so each object is compiled whole-program and is
self-contained; no archive-boundary device link is needed for the static lib.
The CUDA path keeps CUDA_SEPARABLE_COMPILATION/CUDA_RESOLVE_DEVICE_SYMBOLS,
which it does require. The rationale is now recorded as comments in CMake.

Also drops the added per-file copyright/author lines from cuda_to_hip.h to
match the project's existing header convention.

This work was authored with the assistance of Claude, an AI assistant.

Test Plan:

Build and run the interface tests on gfx90a (MI250X, ROCm 7.2.1):

```
cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
  -DCUPDLPX_BUILD_CLI=ON -DCUPDLPX_BUILD_TESTS=ON -DCUPDLPX_BUILD_PYTHON=OFF \
  -DCMAKE_BUILD_TYPE=Release
cmake --build build -j$(nproc)
HIP_VISIBLE_DEVICES=0 ./build/tests/test_interface   # 9/9 pass, Test 9 obj=3.000000001
HIP_VISIBLE_DEVICES=0 ./build/cupdlpx 2club200v15p5scn.mps.gz .  # OPTIMAL, obj -121.2216698
```

Build and exercise the ROCm Python extension with a HIP-aware C++ compiler:

```
cmake -S . -B pybuild -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \
  -DCMAKE_CXX_COMPILER=amdclang++ -DCUPDLPX_BUILD_PYTHON=ON -DCMAKE_BUILD_TYPE=Release
cmake --build pybuild -j$(nproc)
# import + solve -> Status OPTIMAL, ObjVal 3.0, X [1, 2]
```

* apply clang format

* CI: add build-hip

* [ROCm] Drop remaining duplicate CUDA header includes

* Fix missing #endif dropped during moat-port merge conflict resolution

The merge commit's manual resolution of internal/cusparse_compat.h
dropped the closing #endif // USE_HIP, leaving the outer
#if/#else unterminated and breaking the CUDA build.

---------

Co-authored-by: ZedongPeng <peng_zedong@126.com>

Author: jeffdaily

.github/workflows/build.yml

@@ -8,7 +8,7 @@ on:
   workflow_dispatch:
 
 jobs:
-  build:
+  build-cuda:
     strategy:
       fail-fast: false
       matrix:
@@ -64,3 +64,51 @@ jobs:
         shell: pwsh
         run: |
           cmake --build build --clean-first --config Release
+
+  build-hip:
+    name: build-hip (rocm ${{ matrix.rocm }})
+    strategy:
+      fail-fast: false
+      matrix:
+        rocm: ["7.2", "7.2.4"]
+    runs-on: ubuntu-latest
+    container:
+      image: rocm/dev-ubuntu-22.04:${{ matrix.rocm }}
+    env:
+      DEBIAN_FRONTEND: noninteractive
+
+    steps:
+      # actions/checkout needs git, which the ROCm dev image does not ship with.
+      - name: Install checkout prerequisites
+        run: |
+          apt-get update
+          apt-get install -y --no-install-recommends git ca-certificates
+
+      - uses: actions/checkout@v4
+
+      - name: Install build deps
+        run: |
+          apt-get install -y --no-install-recommends \
+            cmake ninja-build zlib1g-dev \
+            hipblas-dev hipsparse-dev hipcub-dev rocprim-dev
+
+      # The ROCm apt packages don't add /opt/rocm/bin to PATH for non-interactive shells.
+      - name: Add ROCm to PATH
+        run: echo "/opt/rocm/bin" >> "$GITHUB_PATH"
+
+      - name: HIP info
+        run: |
+          hipcc --version
+
+      - name: Configure
+        run: |
+          cmake -B build -G Ninja \
+            -DCMAKE_BUILD_TYPE=Release \
+            -DCUPDLPX_BUILD_TESTS=OFF \
+            -DUSE_HIP=ON \
+            -DCMAKE_HIP_ARCHITECTURES=gfx90a \
+            -DCMAKE_PREFIX_PATH=/opt/rocm
+
+      - name: Build
+        run: |
+          cmake --build build --clean-first

CMakeLists.txt

@@ -3,26 +3,33 @@
 # -----------------------------------------------------------------------------
 cmake_minimum_required(VERSION 3.20)
 
+# HIP/ROCm support option (must be set before project() to influence language detection)
+option(USE_HIP "Build with HIP for AMD GPUs" OFF)
+
 # Project config
-project(cupdlpx LANGUAGES C CXX)
-
-# Default CUDA architectures: SASS for every current arch, PTX only for the newest one as a forward-compat fallback.
-find_package(CUDAToolkit REQUIRED)
-if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES AND NOT DEFINED ENV{CUDAARCHS})
-  if(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 13)
-    set(CMAKE_CUDA_ARCHITECTURES 75-real 80-real 86-real 89-real 90-real
-        100-real 120-real 120-virtual)
-  elseif(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 12.8)
-    set(CMAKE_CUDA_ARCHITECTURES 60-real 70-real 75-real 80-real 86-real
-        89-real 90-real 100-real 120-real 120-virtual)
-  else()
-    set(CMAKE_CUDA_ARCHITECTURES 60-real 70-real 75-real 80-real 86-real
-        89-real 90-real 90-virtual)
+if(USE_HIP)
+  project(cupdlpx LANGUAGES C CXX HIP)
+else()
+  project(cupdlpx LANGUAGES C CXX)
+
+  # Default CUDA architectures: SASS for every current arch, PTX only for the newest one as a forward-compat fallback.
+  find_package(CUDAToolkit REQUIRED)
+  if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES AND NOT DEFINED ENV{CUDAARCHS})
+    if(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 13)
+      set(CMAKE_CUDA_ARCHITECTURES 75-real 80-real 86-real 89-real 90-real
+          100-real 120-real 120-virtual)
+    elseif(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 12.8)
+      set(CMAKE_CUDA_ARCHITECTURES 60-real 70-real 75-real 80-real 86-real
+          89-real 90-real 100-real 120-real 120-virtual)
+    else()
+      set(CMAKE_CUDA_ARCHITECTURES 60-real 70-real 75-real 80-real 86-real
+          89-real 90-real 90-virtual)
+    endif()
   endif()
-endif()
 
-enable_language(CUDA)
-message(STATUS "CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}")
+  enable_language(CUDA)
+  message(STATUS "CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}")
+endif()
 
 set(CUPDLPX_VERSION_MAJOR 0)
 set(CUPDLPX_VERSION_MINOR 2)
@@ -50,6 +57,14 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
     set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "Debug" "Release" "MinSizeRel" "RelWithDebInfo")
 endif()
 
+if(USE_HIP)
+  # HIP architecture configuration
+  # Default to gfx90a if not specified; can override with -DCMAKE_HIP_ARCHITECTURES=gfx1100, etc.
+  if(NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "")
+    set(CMAKE_HIP_ARCHITECTURES "gfx90a")
+  endif()
+endif()
+
 # -----------------------------------------------------------------------------
 # [ELEGANT DESIGN] Target-based Compile Flags
 # -----------------------------------------------------------------------------
@@ -75,9 +90,14 @@ else()
     endif()
 endif()
 
-# CUDA standards and RDC
-set(CMAKE_CUDA_STANDARD 17)
-set(CMAKE_CUDA_STANDARD_REQUIRED ON)
+# CUDA/HIP standards and RDC
+if(USE_HIP)
+  set(CMAKE_HIP_STANDARD 17)
+  set(CMAKE_HIP_STANDARD_REQUIRED ON)
+else()
+  set(CMAKE_CUDA_STANDARD 17)
+  set(CMAKE_CUDA_STANDARD_REQUIRED ON)
+endif()
 
 # -----------------------------------------------------------------------------
 # CONTROL OPTIONS
@@ -99,6 +119,14 @@ endif()
 # -----------------------------------------------------------------------------
 # FIND DEPENDENCIES
 # -----------------------------------------------------------------------------
+if(USE_HIP)
+  # Find ROCm/HIP libraries
+  find_package(hip REQUIRED)
+  find_package(hipblas REQUIRED)
+  find_package(hipsparse REQUIRED)
+  find_package(hipcub REQUIRED)
+  find_package(rocprim REQUIRED)
+endif()
 include(FetchContent)
 
 # 1. ZLIB Configuration
@@ -165,20 +193,41 @@ target_compile_definitions(cupdlpx_compile_flags INTERFACE PSLP_VERSION="${PSLP_
 file(GLOB C_SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/src/*.c")
 file(GLOB CU_SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/src/*.cu")
 list(REMOVE_ITEM C_SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/src/cli.c")
+if(WIN32)
+  # mps_parser.c is CLI-only; exclude it on Windows where strtok_r is unavailable
+  list(REMOVE_ITEM C_SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/src/mps_parser.c")
+endif()
 
 set(CORE_INCLUDE_DIRS
   PUBLIC  ${CMAKE_CURRENT_SOURCE_DIR}/include
   PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/internal
 )
 
-set(CORE_LINK_LIBS 
-  PUBLIC cupdlpx_compile_flags
-  PUBLIC CUDA::cudart
-  PUBLIC CUDA::cublas
-  PUBLIC CUDA::cusparse
-  PUBLIC ZLIB::ZLIB
-  PUBLIC PSLP
-)
+if(USE_HIP)
+  set(CORE_LINK_LIBS
+    PUBLIC cupdlpx_compile_flags
+    PUBLIC hip::device
+    PUBLIC roc::hipblas
+    PUBLIC roc::hipsparse
+    PUBLIC hip::hipcub
+    PUBLIC ZLIB::ZLIB
+    PUBLIC PSLP
+  )
+  # Mark .cu files as HIP language
+  set_source_files_properties(${CU_SOURCES} PROPERTIES LANGUAGE HIP)
+  # Define USE_HIP for the compat header; attach to the interface target so it
+  # travels to every consumer (core, shared, cli, tests, python bindings).
+  target_compile_definitions(cupdlpx_compile_flags INTERFACE USE_HIP)
+else()
+  set(CORE_LINK_LIBS
+    PUBLIC cupdlpx_compile_flags
+    PUBLIC CUDA::cudart
+    PUBLIC CUDA::cublas
+    PUBLIC CUDA::cusparse
+    PUBLIC ZLIB::ZLIB
+    PUBLIC PSLP
+  )
+endif()
 
 # 1. Core STATIC Library
 if(CUPDLPX_BUILD_STATIC_LIB)
@@ -187,9 +236,23 @@ if(CUPDLPX_BUILD_STATIC_LIB)
     target_link_libraries(cupdlpx_core ${CORE_LINK_LIBS})
     set_target_properties(cupdlpx_core PROPERTIES
       POSITION_INDEPENDENT_CODE ON
-      CUDA_SEPARABLE_COMPILATION ON
-      CUDA_RESOLVE_DEVICE_SYMBOLS ON
     )
+    if(USE_HIP)
+      # HIP compiles device code whole-program by default (-fgpu-rdc off), so
+      # each object is self-contained and the archive needs no separate device
+      # link. Host-only consumers (the pybind module) resolve everything at the
+      # normal link step. HIP_ARCHITECTURES is inherited from CMAKE_HIP_ARCHITECTURES.
+      set_target_properties(cupdlpx_core PROPERTIES
+        HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}"
+      )
+    else()
+      # CUDA uses relocatable device code here, so device symbols must be
+      # resolved at the archive boundary for host-only consumers to link.
+      set_target_properties(cupdlpx_core PROPERTIES
+        CUDA_SEPARABLE_COMPILATION ON
+        CUDA_RESOLVE_DEVICE_SYMBOLS ON
+      )
+    endif()
 endif()
 
 # 2. Shared Library
@@ -200,9 +263,17 @@ if(CUPDLPX_BUILD_SHARED_LIB)
     set_target_properties(cupdlpx_shared PROPERTIES
         OUTPUT_NAME "cupdlpx"
         RUNTIME_OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}"
+    )
+    if(USE_HIP)
+      set_target_properties(cupdlpx_shared PROPERTIES
+        HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}"
+      )
+    else()
+      set_target_properties(cupdlpx_shared PROPERTIES
         CUDA_SEPARABLE_COMPILATION ON
         CUDA_RESOLVE_DEVICE_SYMBOLS ON
-    )
+      )
+    endif()
 endif()
 
 # 3. CLI Executable
@@ -217,8 +288,12 @@ if(CUPDLPX_BUILD_CLI)
     set_target_properties(cupdlpx_cli PROPERTIES
         OUTPUT_NAME "cupdlpx"
         RUNTIME_OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}"
-        CUDA_RESOLVE_DEVICE_SYMBOLS ON
     )
+    if(NOT USE_HIP)
+      set_target_properties(cupdlpx_cli PROPERTIES
+        CUDA_RESOLVE_DEVICE_SYMBOLS ON
+      )
+    endif()
 endif()
 
 # 4. Tests
@@ -230,14 +305,18 @@ if(CUPDLPX_BUILD_TESTS)
     enable_testing()
     file(GLOB TEST_SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/test/*.c" "${CMAKE_CURRENT_SOURCE_DIR}/test/*.cu")
     foreach(TEST_SRC ${TEST_SOURCES})
-        get_filename_component(TEST_NAME ${TEST_SRC} NAME_WE) 
+        get_filename_component(TEST_NAME ${TEST_SRC} NAME_WE)
         add_executable(${TEST_NAME} ${TEST_SRC})
         target_link_libraries(${TEST_NAME} PRIVATE cupdlpx_core)
         target_include_directories(${TEST_NAME} PRIVATE include internal)
         set_target_properties(${TEST_NAME} PROPERTIES
             RUNTIME_OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/tests"
-            CUDA_RESOLVE_DEVICE_SYMBOLS ON
         )
+        if(NOT USE_HIP)
+          set_target_properties(${TEST_NAME} PROPERTIES
+            CUDA_RESOLVE_DEVICE_SYMBOLS ON
+          )
+        endif()
         add_test(NAME ${TEST_NAME} COMMAND ${TEST_NAME})
     endforeach()
 endif()

README.md

@@ -26,12 +26,14 @@ Our work is presented in two papers:
 ## Installation
 
 ### Requirements
-* **GPU:** NVIDIA GPU with CUDA 12.4+.
-* **Build Tools:** CMake (≥ 3.20), GCC, NVCC.
+* **GPU:** NVIDIA GPU with CUDA 12.4+, or AMD GPU with ROCm 7.2+.
+* **Build Tools:** CMake (≥ 3.20), GCC, and NVCC (CUDA) or hipcc (ROCm).
 
 > **SpMV backend** is selected automatically at compile time based on cuSPARSE version:
 > - `cusparseSpMV` — CUDA 12.4 – 13.2 (cuSPARSE < 12.8.1)
 > - `cusparseSpMVOp` — CUDA 13.3+ (cuSPARSE ≥ 12.8.1)
+>
+> On AMD GPUs the solver uses the `hipsparseSpMV` backend via hipSPARSE.
 
 ### Build from Source
 Clone the repository and compile the project using CMake.
@@ -43,6 +45,21 @@ cmake --build build --clean-first
 ```
 This will create the solver binary at `./build/cupdlpx`.
 
+#### Building for AMD GPUs (ROCm/HIP)
+To target AMD GPUs, configure with `-DUSE_HIP=ON` and select the GPU
+architecture with `-DCMAKE_HIP_ARCHITECTURES`. The CUDA sources are compiled
+as HIP and the cuBLAS/cuSPARSE/CUB calls are mapped to hipBLAS/hipSPARSE/hipCUB.
+```bash
+cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_PREFIX_PATH=/opt/rocm
+cmake --build build --clean-first
+```
+Set `CMAKE_HIP_ARCHITECTURES` to match your GPU (for example `gfx90a` for
+MI200, `gfx1100` for RDNA3 desktop, or `gfx1201` for RDNA4). If the ROCm
+install is not on CMake's default search path, point `-DCMAKE_PREFIX_PATH` at
+it (e.g. `/opt/rocm`) so `find_package` can locate hip, hipBLAS, hipSPARSE,
+and hipCUB. The resulting `./build/cupdlpx` binary is used exactly as in the
+CUDA build.
+
 #### Verifying the Installation
 Run a small test problem to confirm that the solver was built correctly.
 ```bash