Skip to content

Add AMD HIP/ROCm backend - #45

Open
KerelosGayed wants to merge 1 commit into
UWHPC:mainfrom
KerelosGayed:hip-backend
Open

KerelosGayed wants to merge 1 commit into
UWHPC:mainfrom
KerelosGayed:hip-backend

Conversation

@KerelosGayed

@KerelosGayed KerelosGayed commented Oct 1, 2026 •

Copy link
Copy Markdown
Contributor

Adds an AMD HIP/ROCm backend so applications can use xpu's allocation, buffers, structure-of-arrays storage, parallel algorithms, random generators, and optional linear algebra on AMD GPUs. The build selects one backend: CUDA takes precedence when available, HIP is selected next, and systems without either toolchain use the CPU implementation.

This is one PR because the runtime implementation, compiler configuration, and shared GPU tests together provide a usable backend.

Build configuration and requirements

  • Add XPU_ENABLE_HIP, enabled by default when CUDA is not selected. HIP targets AMD GPUs on Linux; CMake rejects a non-AMD HIP platform.
  • Require ROCm 7.0 or newer and the hip, rocprim, hipcub, and hiprand CMake packages. Optional linear algebra additionally requires hipsolver.
  • Discover ROCm packages through CMake's detected ROCm root, ROCM_PATH, CMAKE_PREFIX_PATH, and /opt/rocm. Reuse CMAKE_HIP_ARCHITECTURES for package configuration.
  • Check that the HIP compiler's C++23 standard library provides std::extents in <mdspan>, which hipCUB needs. Configuration errors explain the missing dependency and how to select a CPU build.
  • Link the HIP runtime and library dependencies through xpu::xpu and xpu::linalg. Import hipCUB's include directories without propagating its device compiler flags to every C++ source in a consuming project.

For example, build and test the HIP backend for an MI300-class gfx942 target with:

cmake -S . -B build-hip -G Ninja \
  -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \
  -DCMAKE_HIP_ARCHITECTURES=gfx942 \
  -DXPU_ENABLE_CUDA=OFF \
  -DXPU_ENABLE_HIP=ON \
  -DXPU_ENABLE_LINALG=ON
cmake --build build-hip
ctest --test-dir build-hip --output-on-failure

Use the architecture appropriate for the target GPU. A CPU-only build now explicitly sets both XPU_ENABLE_CUDA=OFF and XPU_ENABLE_HIP=OFF.

Backend configuration, memory, and math

  • Introduce XPU_HIP, the common XPU_GPU and XPU_DEVICE_COMPILE macros, and the xpu::xpu_hip / xpu::xpu_gpu constants. Reject simultaneously enabled CUDA/HIP macros and HIP headers compiled outside HIP mode.
  • Extend the existing callable/device annotations and cu_check helper to HIP, including backend-specific error messages and device-side traps for checked arithmetic failures.
  • Implement allocation, release, memory initialization, and inferred-direction copies through the HIP runtime. HIP uses the GPU alignment/padding policy, including the existing 128-byte alignment constant.
  • Select HIP device implementations for inverse square root, vector norms, sine/cosine, and atomic addition.

Parallel algorithms and random generation

  • Use hipcub::DeviceFor::Bulk for parallel_for and fill_n, with a fill functor because hipCUB does not provide the CUDA DeviceTransform::Fill operation. Empty fills and ranges return without launching work.
  • Use rocPRIM counting/transform iterators and hipCUB sum reduction while keeping the current three-argument parallel_reduce_sum API. The HIP implementation queries its temporary-storage requirement, allocates and releases scratch in order on the default stream with hipMallocAsync / hipFreeAsync, and retains the default pool for reuse. Empty reductions overwrite the output with zero.
  • Move the contribution constraint into a constexpr helper so HIP-Clang can check device-only contribution functions in the appropriate context.
  • Implement HIP random generators with hipRAND Philox, including seeding, stream IDs, offsets, bounded samples, and unbiased index generation. Float uniforms use the top 24 random bits to maintain the half-open [0, 1) interval. Sequences are backend-specific.

Optional linear algebra

Share the GPU implementation between cuSOLVER and hipSOLVER's compatible dense APIs. This covers single- and double-precision LU factorization, solves, inversion, Cholesky factorization/solves, and the supporting identity/transpose operations. Matrix storage remains row-major with strides measured in elements. Cholesky scratch is passed as a buffer_view, and diagnostics identify the selected solver.

Tests, CI, and documentation

  • Compile the existing GPU .cu test entry points as HIP, so both GPU backends exercise the same component cases and standalone-header checks.
  • Add backend-aware launch-error and synchronization helpers and use them across GPU buffer, memory/SoA, launch, math, and random tests. Extend configuration checks and host-access guards to recognize both GPU backends.
  • Add ./scripts/test.sh --hip. The default invocation attempts CPU, CUDA, and HIP; missing GPU compilers are skipped for automatic selection and cause an explicit error when their backend was requested. Compute Sanitizer remains specific to the CUDA suite.
  • Add an opt-in HIP Actions job requiring XPU_HIP_CI=true and a self-hosted runner labeled linux, x64, and rocm. Like the CUDA job, it is excluded from pull-request events and runs through eligible pushes or manual dispatches.
  • Document toolchain setup, backend precedence, source-language requirements, memory transfers, GPU padding, and the one-backend-per-program restriction.

Validation

Check Result
CPU Release build with Clang 21.1.7 on macOS arm64 Passed: 11 test executables and 9 standalone-header targets built.
Default CMake toolchain discovery Passed: missing CUDA and HIP compilers correctly select the CPU path.
bash -n scripts/test.sh Passed.
Explicit missing CUDA/HIP compiler, unavailable sanitizer, and unknown-option checks Passed: each exits with status 2 and the expected diagnostic.
git diff --check origin/main...HEAD Passed.
Local CPU runtime tests Attempted; macOS killed all generated executables, including the configuration-only test, both inside and outside the sandbox. No local runtime pass is claimed.
CPU linear algebra Not run locally: LAPACKE is unavailable.
CUDA/HIP compilation, GPU runtime tests, and Compute Sanitizer Not run locally: this machine has neither GPU toolchain nor GPU hardware for these backends.
Linux CPU Actions suite Passed: GCC 15.2.0 built all targets, including the linear-algebra header check, and all 12 runtime tests passed, including LAPACKE linear algebra.
CUDA/HIP Actions jobs Skipped as designed for pull-request events; GPU validation remains outstanding.

The local CPU build used XPU_ENABLE_CUDA=OFF, XPU_ENABLE_HIP=OFF, XPU_ENABLE_LINALG=OFF, and XPU_ARCH_NATIVE=OFF. It also used -D_LIBCPP_ENABLE_EXPERIMENTAL because the installed libc++ hides std::jthread behind that setting. This was a local CMake flag only. GPU execution and the hipSOLVER path still need validation on the corresponding hardware.

Add HIP build selection, runtime operations, hipCUB algorithms, hipRAND generators, and hipSOLVER linear algebra. Share the GPU test suite across CUDA and HIP and document the toolchain requirements and opt-in ROCm CI.
@KerelosGayed

Copy link
Copy Markdown
Contributor Author

@karl-kes review pls

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant