Support of AMD GPUs

.github/workflows/build-hip.yml

@@ -0,0 +1,72 @@
+name: build-hip
+
+# Compile-only CI for the AMD GPU (HIP/ROCm) backend of HiPDLP.
+#
+# This job runs inside an official ROCm container and only *compiles and links*
+# the HIP backend (-DHIPDLP_HIP=ON). It does NOT execute any GPU code: the
+# GitHub-hosted runners have no AMD GPU. Its purpose is to catch HIP build
+# regressions (wrong hip* API names, desynchronised #ifdef guards, signature
+# mismatches) that are invisible to the CPU / CUDA builds.
+#
+# Running the GPU code (numerical validation) requires a self-hosted runner
+# with an AMD GPU and is intentionally out of scope here.
+
+on:
+  push:
+    paths:
+      - 'highs/pdlp/hipdlp/**'
+      - 'CMakeLists.txt'
+      - 'highs/CMakeLists.txt'
+      - '.github/workflows/build-hip.yml'
+  pull_request:
+    paths:
+      - 'highs/pdlp/hipdlp/**'
+      - 'CMakeLists.txt'
+      - 'highs/CMakeLists.txt'
+      - '.github/workflows/build-hip.yml'
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.ref }}
+  cancel-in-progress: true
+
+jobs:
+  hip-compile:
+    runs-on: ubuntu-latest
+    container:
+      # 24.04 + ROCm 10.x; the "-full" image bundles hipBLAS / hipSPARSE
+      image: rocm/dev-ubuntu-24.04:10.0.0-full
+
+    steps:
+      - uses: actions/checkout@v6
+
+      - name: Install build dependencies
+        shell: bash
+        run: |
+          apt-get update
+          apt-get install -y --no-install-recommends \
+            cmake make git ca-certificates zlib1g-dev
+          cmake --version
+          hipcc --version || /opt/rocm/bin/hipcc --version
+
+      - name: Configure CMake (HIP backend)
+        shell: bash
+        run: |
+          export PATH=/opt/rocm/bin:$PATH
+          export CMAKE_PREFIX_PATH=/opt/rocm
+          cmake -S . -B build_hip \
+            -DHIPDLP_HIP=ON \
+            -DBUILD_TESTING=OFF \
+            -DBUILD_EXAMPLES=ON \
+            -DCMAKE_C_COMPILER=amdclang \
+            -DCMAKE_CXX_COMPILER=amdclang++ \
+            -DCMAKE_BUILD_TYPE=Release
+
+      - name: Build (compile + link, no GPU execution)
+        shell: bash
+        run: |
+          export PATH=/opt/rocm/bin:$PATH
+          # Build the library and the HiPDLP example. The example only compiles
+          # and links against the HIP-enabled library; it is not executed here
+          # (the GitHub-hosted runners have no AMD GPU). BUILD_TESTING=OFF keeps
+          # it from being registered as a ctest.
+          cmake --build build_hip --target highs call_highs_hipdlp --parallel

.gitignore

@@ -220,6 +220,7 @@ $RECYCLE.BIN/
 dist/
 build/
 build_release*/
+install/
 eggs/
 parts/
 var/

CMakeLists.txt

@@ -172,11 +172,18 @@ message(STATUS "Use FindCUDAConf: ${CUPDLP_FIND_CUDA}")
 option(HIGHS_GPU_LIB "Build highs GPU as a single lib" OFF)
 message(STATUS "Build GPU with delayed loading: ${HIGHS_GPU_LIB}")
 
+option(HIPDLP_HIP "Build HiPDLP with AMD HIP (ROCm)" OFF)
+message(STATUS "Build HiPDLP with HIP: ${HIPDLP_HIP}")
+
 # For now
 if (HIGHS_GPU_LIB)
   set(CUPDLP_GPU ON)
 endif()
 
+if (HIPDLP_HIP)
+  set(HIPDLP_GPU ON)
+endif()
+
 if(CUPDLP_GPU AND CMAKE_VERSION VERSION_LESS "3.25.0")
     message("CUPDLP FindCUDAConf requires CMake version minumum 3.24. Please use a higher version of CMake.")
 endif()
@@ -208,30 +215,72 @@ if (HIPO AND NOT FAST_BUILD)
   message(ERROR "HIPO is only available with FAST_BUILD=ON.")
 endif()
 
-if (CUPDLP_GPU)
+if (CUPDLP_GPU OR HIPDLP_GPU)
   if (WIN32)
     set(BUILD_SHARED_LIBS ON)
   endif()
 
-  set (CUPDLP_CPU OFF)
-  message(NOTICE "Set build cuPDLP with CUDA")
+  set(CUPDLP_CPU OFF)
+
+  if (HIPDLP_HIP)
+    message(NOTICE "Set build HiPDLP with HIP (AMD ROCm)")
+
+    # ROCm/HIP requires CMake 3.21+
+    if(CMAKE_VERSION VERSION_LESS "3.21.0")
+      message(FATAL_ERROR "HIP support requires CMake >= 3.21")
+    endif()
+
+    # Allow the user to point to a non-standard ROCm install
+    if(DEFINED ENV{ROCM_PATH})
+      set(CMAKE_PREFIX_PATH "$ENV{ROCM_PATH}" ${CMAKE_PREFIX_PATH})
+    elseif(EXISTS "/opt/rocm")
+      set(CMAKE_PREFIX_PATH "/opt/rocm" ${CMAKE_PREFIX_PATH})
+    endif()
+
+    enable_language(HIP)
+    find_package(hipblas REQUIRED)
+    find_package(hipsparse REQUIRED)
+
+    # The default bfd linker is slow and memory-hungry with the ROCm toolchain
+    # and can stall on memory-constrained machines.
+    # The ROCm clang bundles lld and we use it if it is available.
+    find_program(HIGHS_LLD_LINKER NAMES ld.lld lld
+      HINTS ENV ROCM_PATH /opt/rocm
+      PATH_SUFFIXES lib/llvm/bin bin llvm/bin)
+    if (HIGHS_LLD_LINKER)
+      set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -fuse-ld=lld")
+      message(STATUS "HIP: linking with lld (${HIGHS_LLD_LINKER})")
+    else()
+      message(STATUS "HIP: lld not found; using the default linker")
+    endif()
+
+    set(GPU_LIBRARY roc::hipblas roc::hipsparse)
+    set(GPU_COMPILE_DEFINITIONS USE_HIP)
 
-  if (CUPDLP_FIND_CUDA)
-    # With FindCUDAConf.cmake
-    # Need to have the CUDA_HOME environment variable set.
-    include(FindCUDAConf)
   else()
-    # Without FindCUDAConf.cmake
-    enable_language(CUDA)
-    find_package(CUDAToolkit REQUIRED)
+    message(NOTICE "Set build cuPDLP with CUDA")
+
+    if (CUPDLP_FIND_CUDA)
+      # With FindCUDAConf.cmake
+      # Need to have the CUDA_HOME environment variable set.
+      include(FindCUDAConf)
+    else()
+      # Without FindCUDAConf.cmake
+      enable_language(CUDA)
+      find_package(CUDAToolkit REQUIRED)
+
+      set(CUDA_LIBRARY-NOTFOUND, OFF)
+      set(CUDA_LIBRARY CUDA::cudart CUDA::cublas CUDA::cusparse)
+    endif()
 
-    set(CUDA_LIBRARY-NOTFOUND, OFF)
-    set(CUDA_LIBRARY CUDA::cudart CUDA::cublas CUDA::cusparse)
+    set(GPU_LIBRARY ${CUDA_LIBRARY})
+    set(GPU_COMPILE_DEFINITIONS "")
   endif()
 
 else()
-  set (CUPDLP_CPU ON)
-    set(CUDA_LIBRARY-NOTFOUND true)
+  set(CUPDLP_CPU ON)
+  set(CUDA_LIBRARY-NOTFOUND true)
+  set(HIPDLP_GPU OFF)
 endif()
 
 if(BUILD_STATIC_EXE)
@@ -413,7 +462,7 @@ if (BUILD_CXX)
       message(STATUS "IPO / LTO: enabled")
     endif()
   endif()
-  if (CUPDLP_GPU AND CMAKE_INTERPROCEDURAL_OPTIMIZATION)
+  if ((CUPDLP_GPU OR HIPDLP_HIP) AND CMAKE_INTERPROCEDURAL_OPTIMIZATION)
       set(CMAKE_INTERPROCEDURAL_OPTIMIZATION FALSE)
       message(STATUS "IPO / LTO is not supported at the moment when PDLP is using GPU: LTO disabled.")
   endif()

cmake/cpp-highs.cmake

@@ -70,7 +70,7 @@ if (NOT HIGHS_COVERAGE)
     APPEND FILE "${HIGHS_BINARY_DIR}/highs-targets.cmake")
 endif()
 
-if (CUPDLP_GPU AND NOT HIGHS_GPU_LIB)
+if (CUPDLP_GPU AND NOT HIGHS_GPU_LIB AND NOT HIPDLP_HIP)
   install(TARGETS cudalin
       EXPORT ${lower}-targets
       INCLUDES DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}

docs/src/guide/gpu.md

@@ -6,6 +6,15 @@ NVIDIA GPU under Linux and Windows. However, to achieve this, CUDA
 utilities must be installed and HiGHS must be built locally using
 CMake, as described below.
 
+The native HiPDLP solver additionally supports AMD GPUs through
+[ROCm](https://rocm.docs.amd.com/) / HIP.
+
+Whether HiPDLP (and cuPDLP-C) runs on the CPU or on a GPU is fixed at
+build time: the GPU backend is only compiled with `-DHIPDLP_HIP=ON`
+(AMD) or `-DCUPDLP_GPU=ON` (NVIDIA). The runtime [__solver__](@ref
+option-solver) option selects the *solver*, not the *device*: on a build
+without GPU support, `solver = "hipdlp"` still runs, but on the CPU.
+
 ### PDLP: A health warning
 
 First order solvers for LP are still very much "work in
@@ -38,21 +47,28 @@ instance.
 
 ### Requirements
 
-CUDA Toolkit and CMake. 
+CMake, plus a CUDA Toolkit (for NVIDIA GPUs) or a ROCm installation
+(for AMD GPUs). HiGHS must be built locally with CMake.
 
-A [CUDA Toolkit](https://developer.nvidia.com/cuda-toolkit)
+For NVIDIA GPUs, a [CUDA Toolkit](https://developer.nvidia.com/cuda-toolkit)
 installation is required, along with the matching NVIDIA
 driver. Please install both following the instructions on NVIDIA's
-website.
-
-HiGHS must be build locally with CMake. 
-
-Make sure the CUDA compiler `nvcc` is installed by running 
+website. Make sure the CUDA compiler `nvcc` is installed by running
 
 ```
 nvcc --version
 ```
 
+For AMD GPUs, a [ROCm](https://rocm.docs.amd.com/) installation providing
+the HIP compiler and the hipBLAS / hipSPARSE libraries is required instead;
+see [Building HiGHS with AMD GPU support](@ref gpu-build-amd) for details.
+
 ### Build HiGHS with GPU support
 
-See [Building HiGHS with NVidia GPU support](@ref gpu-build).
+For NVIDIA GPUs, see [Building HiGHS with NVidia GPU support](@ref
+gpu-build).
+
+For AMD GPUs, see [Building HiGHS with AMD GPU support](@ref
+gpu-build-amd).
+This uses ROCm / HIP and its hipBLAS and hipSPARSE libraries
+instead of CUDA.

docs/src/installation.md

@@ -117,7 +117,7 @@ cmake --build build --parallel
 
 to build HiGHS.
 
-### Bazel build with Cuda
+### Bazel build with CUDA
 
 Alternatively, for Bazel run
 
@@ -130,3 +130,73 @@ It may be necessary to also specify the architecture, e.g.
 ```
 bazel build //... --//:cupdlp_gpu --@rules_cuda//cuda:archs=sm_89
 ```
+
+## [Building HiGHS with AMD GPU support](@id gpu-build-amd)
+
+The native HiPDLP solver can also run on an AMD GPU using
+[ROCm](https://rocm.docs.amd.com/) / HIP. This requires a ROCm
+installation providing the HIP compiler and the hipBLAS and hipSPARSE
+libraries. Make sure the HIP compiler is available by running
+
+```
+hipcc --version
+```
+
+ROCm 10.0 or newer is recommended (this is the version the HIP backend
+is tested against). On such a ROCm, the supported GPU architectures are
+`gfx908` and newer (e.g. `gfx908`/MI100, `gfx90a`/MI200,
+`gfx942`/MI300, and recent RDNA cards).
+
+Then build HiGHS, from the root directory, with
+
+```
+cmake -S. -Bbuild -DHIPDLP_HIP=ON
+cmake --build build --parallel
+```
+
+CMake must be able to find ROCm. If it is not installed in the default
+location, point it there, for example
+
+```
+export PATH=/opt/rocm/bin:$PATH
+export CMAKE_PREFIX_PATH=/opt/rocm
+```
+
+By default the HIP device code is compiled for a generic set of GPU
+architectures. To target the specific GPU on the build machine (which
+also speeds up compilation and linking), set `CMAKE_HIP_ARCHITECTURES`
+to its `gfx` target, for example
+
+```
+cmake -S. -Bbuild -DHIPDLP_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a
+```
+
+You can find the `gfx` identifier of the installed GPU with `rocminfo`
+(look for the `gfx` name, e.g. `gfx90a` for MI200-class cards or
+`gfx942` for MI300). Multiple architectures may be given as a
+semicolon-separated list, e.g. `-DCMAKE_HIP_ARCHITECTURES="gfx90a;gfx942"`.
+
+By default the host C/C++ sources are compiled with the system compiler
+(e.g. GCC) and only the HIP device code with ROCm's compiler. To build
+the whole of HiGHS with the ROCm toolchain instead, for a uniform
+Clang-based build, point CMake at `amdclang` / `amdclang++`
+
+```
+cmake -S. -Bbuild -DHIPDLP_HIP=ON \
+  -DCMAKE_C_COMPILER=amdclang -DCMAKE_CXX_COMPILER=amdclang++
+```
+
+The HIP backend compiles the same HiPDLP source as the CUDA backend,
+selecting the AMD implementation at build time. Once built, the solver
+is selected at run time by setting the [__solver__](@ref
+option-solver) option to "hipdlp".
+
+To check the ROCm / HIP backend on the local machine, run the example
+`call_highs_hipdlp` (also registered as the ctest
+`cxx_examples_call_highs_hipdlp`), which solves a small LP with
+`solver = "hipdlp"` and verifies the result. A successful run is a
+quick end-to-end sanity check of the GPU backend.
+
+To confirm the work is actually running on the GPU, watch `rocm-smi`
+(for example `watch -n 0.1 rocm-smi`) while the solve runs and check
+that GPU utilisation and memory usage rise.

docs/src/options/definitions.md

@@ -6,7 +6,7 @@
 - Default: "choose"
 
 ## [solver](@id option-solver)
-- LP/QP solver: "choose", "simplex", "ipm", "ipx", "hipo", "pdlp", "qpasm" or "hipdlp", 
+- LP/QP solver: "choose", "simplex", "ipm", "ipx", "hipo", "pdlp", "hipdlp" or "qpasm"
 - Type: string
 - Default: "choose"
 

docs/src/solvers.md

@@ -74,13 +74,13 @@ HiGHS has two interior point (IPM) solvers:
   The option [__hipo\_ordering__](@ref option-hipo-ordering) can be used
   to select the fill-reducing heuristic to use during the
   factorisation:
-  
+
   * Nested dissection, obtained setting the option
     [__hipo\_ordering__](@ref option-hipo-ordering) to "metis".
-  
+
   * Approximate minimum degree, obtained setting the option
     [__hipo\_ordering__](@ref option-hipo-ordering) to "amd".
-  
+
   * Reverse Cuthill-McKee, obtained setting the option
     [__hipo\_ordering__](@ref option-hipo-ordering) to "rcm".
 
@@ -94,8 +94,9 @@ HiGHS includes the [
 cuPDLP-C](https://github.com/COPT-Public/cuPDLP-C) primal-dual hybrid
 gradient method for LP (PDLP), and also has a native PDLP solver,
 HiPDLP. On Linux and Windows, these solvers can be run on an NVIDIA
-[GPU](@ref gpu). On a CPU, they are unlikely to be competitive with
-the HiGHS interior point or simplex solvers.
+[GPU](@ref gpu), and HiPDLP can additionally be run on an AMD GPU via
+ROCm / HIP. On a CPU, they are unlikely to be competitive with the HiGHS
+interior point or simplex solvers.
 
   * Setting the option [__solver__](@ref option-solver) to "pdlp" forces the cuPDLP-C solver to be used
   * Setting the option [__solver__](@ref option-solver) to "hipdlp" forces the HiPDLP solver to be used
@@ -191,4 +192,3 @@ The option [__solver__](@ref option-solver) is ignored and the default solver is
 * The problem is an LP and __solver__ is set to "qpasm".
 * The problem is a QP and __solver__ is set to "simplex", "ipx", "pdlp" or "hipdlp".
 * The problem is a MIP and __solver__ is not set to "choose".
-