[ROCm] Add optional HIP build for AMD GPUs - #149
Open
jeffdaily wants to merge 2 commits into
Open
Conversation
Adds an optional ROCm/HIP build (USE_HIP, default OFF) so Gpufit runs on AMD
GPUs, including the optional batched-LU solver (USE_CUBLAS) via hipBLAS.
Gpufit drives its device build through the legacy FindCUDA module; the HIP path
is a parallel modern-CMake build (enable_language(HIP), LANGUAGE HIP) gated on
USE_HIP, with every effort made to leave the existing CUDA path unchanged under
if(NOT USE_HIP). A single cuda_to_hip.h compat header (not included on the
NVIDIA path) aliases the cuda runtime symbols used, with small hip_compat shims
so the direct cuda_runtime.h / device_launch_parameters.h includes resolve on
ROCm. No solver-logic changes: the default Gauss-Jordan solver uses
shared-memory tree reductions (no warp intrinsics), so it is wave-size-agnostic.
HIP_ARCHITECTURES is configurable via -DCMAKE_HIP_ARCHITECTURES=<arch>,
defaulting to gfx90a when unset, so building for another target needs no
CMakeLists edit.
The optional batched-LU solver (USE_CUBLAS=ON, off by default) used cuBLAS and
was previously NVIDIA-only; the ROCm build fell back to Gauss-Jordan. The
cublas* batched symbols (create/destroy, S/D getrf/getrsBatched, handle, status,
OP_N) alias 1:1 to hipblas* in cuda_to_hip.h under USE_CUBLAS, a
hip_compat/cublas_v2.h shim resolves the project's #include "cublas_v2.h" on
ROCm, and the CMake HIP path links roc::hipblas when USE_CUBLAS is requested.
The one API difference -- hipBLAS getrsBatched takes float* const A[] where
cuBLAS takes const float* const A[] -- is handled by dropping the leading const
on the HIP call only; the CUDA path is unchanged.
One test tolerance change: in Gauss_Fit_2D_Rotated, on wave32 vs wave64 the
Gauss-Jordan solver's block-packing heuristic (lm_fit_cuda.cu, gated on the
runtime warpSize) changes the FP reduction order, shifting the converged
rotation angle by ~1.1e-6 while chi^2 stays ~1e-12. Under a USE_HIP guard on the
test TU, that single parameter bound is relaxed from 1e-6f to 3e-6f; the CUDA
assertion is untouched.
Windows/clang-cl: two build-system fixes let the HIP build link with clang-cl
(MSVC frontend). The CMake HIP platform module sets CMAKE_HIP_COMPILE_OBJECT to
use /Fo, which makes clang-cl drop the device fatbinary from the object; the
rule is overridden to GNU-style -o <OBJECT>, guarded to fire only for clang-cl.
clang-cl also ignores GNU-style -include<file>, so the force-include of
cuda_to_hip.h selects /FI<file> vs -include via
CMAKE_HIP_COMPILER_FRONTEND_VARIANT.
Test Plan:
cmake -S . -B build-hip -DUSE_HIP=ON -DUSE_CUBLAS=ON \
-DCMAKE_HIP_ARCHITECTURES=gfx90a \
-DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_BUILD_TYPE=Release
cmake --build build-hip -j
cd build-hip && HIP_VISIBLE_DEVICES=0 ctest --output-on-failure
Validated on AMD Instinct MI250X (gfx90a, wave64, ROCm 7.2) with USE_CUBLAS=ON:
all 11 C++ known-answer suites pass (including Cpufit/Gpufit consistency) and the
10000-fit Monte-Carlo example converges 99.99% deterministically. Re-validated
on AMD Radeon Pro W7800 (gfx1100, wave32): 10/10 ctest. On Windows with clang-cl
(TheRock ROCm), the RDNA3 and RDNA4 builds link device code into the DLL and
pass the known-answer suites. The default USE_HIP=OFF build is unchanged.
Authored with the assistance of Claude (AI).
The gfx90a defaults in Gpufit/CMakeLists.txt and examples/c++/CMakeLists.txt ran after project(... LANGUAGES CXX HIP) in the top-level CMakeLists, which already honors an explicit -DCMAKE_HIP_ARCHITECTURES, otherwise auto-detects the host GPU(s) and errors on a no-GPU host. The pins were unreachable dead code that duplicated CMake's own host-GPU detection; removing them leaves the HIP enable as the single source of the arch list. Test Plan: explicit-arch builds are byte-identical; passing -DCMAKE_HIP_ARCHITECTURES=gfx90a produces the same object code as before. cmake -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a .. Authored with the Claude AI assistant.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This adds an optional ROCm/HIP build so Gpufit runs on AMD GPUs, including the optional batched-LU solver (USE_CUBLAS) via hipBLAS.
Gpufit drives its device build through the legacy FindCUDA module; the HIP path is a parallel modern-CMake build (enable_language(HIP), LANGUAGE HIP) gated on USE_HIP (default OFF), so the CUDA path is left as-is under if(NOT USE_HIP). A single cuda_to_hip.h compat header (not included on the NVIDIA path) aliases the cuda runtime symbols used, with small hip_compat shims so the direct cuda_runtime.h / device_launch_parameters.h includes resolve on ROCm. No solver-logic changes: the default Gauss-Jordan solver uses shared-memory tree reductions (no warp intrinsics), so it is wave-size-agnostic. HIP_ARCHITECTURES is configurable via -DCMAKE_HIP_ARCHITECTURES=, defaulting to gfx90a when unset.
The optional batched-LU solver (USE_CUBLAS=ON, off by default) was previously NVIDIA-only; the ROCm build fell back to Gauss-Jordan. The cublas* batched symbols alias 1:1 to hipblas* under USE_CUBLAS, a cublas_v2.h shim resolves the include on ROCm, and CMake links roc::hipblas when USE_CUBLAS is requested. The one API difference (hipBLAS getrsBatched takes float* const A[] vs cuBLAS const float* const A[]) is handled on the HIP call only.
Notes for reviewers:
Validation (real GPU, all known-answer suites + Cpufit/Gpufit consistency):
Authored with the assistance of Claude (AI).