[ROCm] Add optional HIP build for AMD GPUs

CMakeLists.txt

@@ -3,19 +3,47 @@
 # see also CMake configuration in /docs/installation.rst
 
 # CMake
-
-cmake_minimum_required( VERSION 3.11 )
+
+# USE_HIP routes the .cu device sources through the ROCm/HIP toolchain for AMD
+# GPUs instead of CUDA. enable_language(HIP) needs CMake >= 3.21; the CUDA
+# (FindCUDA) path is unaffected by the higher floor. Declared here so every
+# subdirectory sees the option.
+option( USE_HIP "Build the device code with HIP for AMD GPUs" OFF )
+if( USE_HIP )
+  cmake_minimum_required( VERSION 3.21 )
+else()
+  cmake_minimum_required( VERSION 3.11 )
+endif()
 set_property( GLOBAL PROPERTY USE_FOLDERS ON )
 
 if( NOT PROJECT_NAME )
-  project( Gpufit VERSION 1.2.0 )
+  if( USE_HIP )
+    project( Gpufit VERSION 1.2.0 LANGUAGES CXX HIP )
+  else()
+    project( Gpufit VERSION 1.2.0 )
+  endif()
   include( CTest )
 endif()
 
 if( NOT CMAKE_CXX_STANDARD )
 	set( CMAKE_CXX_STANDARD 14 )
 endif()
 
+# On Windows with clang-cl (MSVC frontend), the HIP platform CMake module
+# (Windows-MSVC.cmake / __windows_compiler_msvc) unconditionally sets the HIP
+# compile-object rule to use the MSVC /Fo<out> output flag.  With /Fo, the
+# clang-cl driver routes the host object through the same path it passes to
+# clang-offload-bundler as input, which on Windows causes the device fatbinary
+# to be silently dropped (ERROR_USER_MAPPED_FILE or empty .hip_fatbin section).
+# Fix: override the rule after CMakeHIPInformation.cmake has run, using GNU
+# -o <OBJECT> instead, which keeps the host object in a distinct temp file.
+# Only needed for the MSVC/clang-cl frontend; GCC-frontend clang++ uses the
+# correct default rule (includes -x hip).
+if( USE_HIP AND WIN32 AND CMAKE_HIP_COMPILER MATCHES "clang-cl" )
+  set( CMAKE_HIP_COMPILE_OBJECT
+    "<CMAKE_HIP_COMPILER> <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> -c -- <SOURCE>" )
+endif()
+
 if( MSVC ) # link runtime statically with MSVC
   foreach( type ${CMAKE_CONFIGURATION_TYPES} ${CMAKE_BUILD_TYPE} )
     string( TOUPPER ${type} TYPE )
@@ -74,6 +102,11 @@ if( Boost_FOUND )
     )
     target_include_directories( ${target} PRIVATE ${PROJECT_SOURCE_DIR} )
     target_link_libraries( ${target} ${modules} Boost::boost )
+    if( USE_HIP )
+      # Lets tests select wave-size-tolerant bounds where wave32/wave64 change
+      # the FP reduction order of the GJ solver (e.g. Gauss_Fit_2D_Rotated).
+      target_compile_definitions( ${target} PRIVATE USE_HIP )
+    endif()
     set_property( TARGET ${target}
       PROPERTY RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}" )
     set_property( TARGET ${target} PROPERTY FOLDER Tests )

Gpufit/CMakeLists.txt

@@ -34,6 +34,11 @@
 #      Set CUDA_BIN_PATH before running CMake or CUDA_TOOLKIT_ROOT_DIR after first configuration
 #      to installation folder of desired CUDA version
 
+# The CUDA (FindCUDA) toolchain detection and NVCC arch-flag logic is for the
+# NVIDIA build only. The HIP build (USE_HIP) drives the device language through
+# enable_language(HIP) further down and does not use FindCUDA.
+if( NOT USE_HIP )
+
 find_package( CUDA 6.5 REQUIRED )
 
 set( CUDA_ARCHITECTURES ${DEFAULT_CUDA_ARCH} CACHE STRING
@@ -128,6 +133,8 @@ if( NOT WIN32 )
 	list( APPEND CUDA_NVCC_FLAGS --std=c++11)
 endif()
 
+endif() # NOT USE_HIP
+
 # Gpufit
 
 set( GpuHeaders
@@ -177,6 +184,42 @@ source_group("CUDA Source Files" FILES ${GpuCudaSources})
 source_group("CUDA Model Files" FILES ${GpuCudaModels})
 source_group("CUDA Estimator Files" FILES ${GpuCudaEstimators})
 
+if( USE_HIP )
+	# Compile the .cu translation units as HIP. The compat header cuda_to_hip.h
+	# (force-included below) aliases the CUDA runtime symbols the project uses to
+	# their HIP equivalents, so the sources keep their plain CUDA spelling.
+	set_source_files_properties( ${GpuCudaSources} PROPERTIES LANGUAGE HIP )
+	add_library( Gpufit SHARED
+		${GpuHeaders}
+		${GpuSources}
+		${GpuCudaHeaders}
+		${GpuCudaSources}
+		${GpuCudaModels}
+		${GpuCudaEstimators}
+	)
+	# project(... LANGUAGES CXX HIP) in the top-level CMakeLists already honors an
+	# explicit -DCMAKE_HIP_ARCHITECTURES, otherwise auto-detects the host GPU(s)
+	# and errors on a no-GPU host, so no arch default is pinned here.
+	set_target_properties( Gpufit PROPERTIES HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}" )
+	# USE_HIP selects the HIP branch of cuda_to_hip.h. __HIP_PLATFORM_AMD__ is
+	# only defined once the HIP runtime header is included, so it cannot gate the
+	# include of that very header; the explicit define breaks the cycle.
+	target_compile_definitions( Gpufit PRIVATE USE_HIP )
+	# Force-include the compat header into every HIP TU so no per-file #include
+	# edits are needed (minimal footprint).
+	# clang-cl (MSVC frontend) ignores GNU-style -include; use /FI instead.
+	target_compile_options( Gpufit PRIVATE
+		$<$<COMPILE_LANGUAGE:HIP>:$<IF:$<STREQUAL:${CMAKE_HIP_COMPILER_FRONTEND_VARIANT},MSVC>,/FI${CMAKE_CURRENT_SOURCE_DIR}/cuda_to_hip.h,-include${CMAKE_CURRENT_SOURCE_DIR}/cuda_to_hip.h>> )
+	# hip_compat/ provides cuda_runtime.h and device_launch_parameters.h shims
+	# for the sources that include them directly; HIP build only.
+	target_include_directories( Gpufit PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/hip_compat )
+	# Host-language .cpp TUs (gpufit.cpp, lm_fit.cpp, ...) include gpu_data.cuh,
+	# which pulls in the cuda_to_hip.h shim and thus <hip/hip_runtime.h>. Link
+	# hip::host so those host TUs (compiled as CXX) get the ROCm include dirs and
+	# the __HIP_PLATFORM_AMD__ define.
+	find_package( hip REQUIRED )
+	target_link_libraries( Gpufit hip::host )
+else()
 cuda_add_library( Gpufit SHARED
 	${GpuHeaders}
 	${GpuSources}
@@ -185,6 +228,7 @@ cuda_add_library( Gpufit SHARED
 	${GpuCudaModels}
 	${GpuCudaEstimators}
 )
+endif()
 
 set_target_properties( Gpufit
 	PROPERTIES
@@ -193,7 +237,10 @@ set_target_properties( Gpufit
 )
 
 # USE_CUBLAS
-if( CMAKE_SIZEOF_VOID_P EQUAL 8 AND CUDA_VERSION VERSION_GREATER "6.5")
+# cuBLAS is a CUDA-only acceleration of the linear solver. The HIP build uses
+# the self-contained Gauss-Jordan solver (the upstream default, USE_CUBLAS OFF);
+# this block also depends on CUDA_VERSION, which is unset on the HIP path.
+if( NOT USE_HIP AND CMAKE_SIZEOF_VOID_P EQUAL 8 AND CUDA_VERSION VERSION_GREATER "6.5")
 	set( USE_CUBLAS ${DEFAULT_USE_CUBLAS} CACHE BOOL "ON | OFF")
 	if( USE_CUBLAS )
         if ( WIN32 )
@@ -230,6 +277,20 @@ elseif( CUDA_VERSION VERSION_LESS "7.0" )
 	message( STATUS "CUBLAS: CUDA Version < 7.0 detected; USE_CUBLAS flag ignored." )
 endif()
 
+# USE_CUBLAS on ROCm: the optional batched-LU solver maps onto hipBLAS. Off by
+# default (DEFAULT_USE_CUBLAS). When enabled, cuda_to_hip.h aliases the cublas*
+# batched symbols to hipblas* and the hip_compat/cublas_v2.h shim resolves the
+# project's #include "cublas_v2.h"; here we define USE_CUBLAS and link hipBLAS.
+if( USE_HIP )
+	set( USE_CUBLAS ${DEFAULT_USE_CUBLAS} CACHE BOOL "ON | OFF")
+	if( USE_CUBLAS )
+		find_package( hipblas REQUIRED )
+		target_compile_definitions( Gpufit PRIVATE USE_CUBLAS )
+		target_link_libraries( Gpufit roc::hipblas )
+		message( STATUS "CUBLAS: enabled on ROCm via hipBLAS (roc::hipblas)." )
+	endif()
+endif()
+
 #install( TARGETS Gpufit RUNTIME DESTINATION bin )
 
 # Examples using only Gpufit

Gpufit/cuda_to_hip.h

@@ -0,0 +1,82 @@
+//********************************************************//
+// CUDA-to-HIP compatibility shim for Gpufit              //
+//                                                        //
+// Minimal-footprint port: every other source file keeps  //
+// its plain CUDA spelling. On AMD this header includes    //
+// the HIP runtime and #defines the CUDA symbols the       //
+// project uses to their HIP equivalents. On NVIDIA it is  //
+// a no-op that pulls in <cuda_runtime.h>.                 //
+//                                                        //
+// Symbol names follow PyTorch's authoritative hipify map: //
+// torch/utils/hipify/cuda_to_hip_mappings.py             //
+//                                                        //
+// Scope: Gpufit's default solver (USE_CUBLAS=OFF) uses    //
+// the self-contained Gauss-Jordan kernel and no GPU math  //
+// library, so only the CUDA runtime API needs aliasing.   //
+// There are no warp intrinsics, textures, surfaces, or    //
+// constant memory in the project, hence none are mapped.  //
+//********************************************************//
+
+#ifndef GPUFIT_CUDA_TO_HIP_H
+#define GPUFIT_CUDA_TO_HIP_H
+
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+
+#include <hip/hip_runtime.h>
+
+// ---- Error handling ----
+#define cudaError              hipError_t
+#define cudaError_t            hipError_t
+#define cudaSuccess            hipSuccess
+#define cudaGetErrorString     hipGetErrorString
+#define cudaGetLastError       hipGetLastError
+#define cudaDeviceSynchronize  hipDeviceSynchronize
+
+// ---- Version queries ----
+#define cudaDriverGetVersion   hipDriverGetVersion
+#define cudaRuntimeGetVersion  hipRuntimeGetVersion
+
+// ---- Device management ----
+#define cudaGetDeviceCount       hipGetDeviceCount
+#define cudaGetDeviceProperties  hipGetDeviceProperties
+#define cudaSetDevice            hipSetDevice
+#define cudaDeviceProp           hipDeviceProp_t
+#define cudaMemGetInfo           hipMemGetInfo
+
+// ---- Linear memory ----
+#define cudaMalloc  hipMalloc
+#define cudaFree    hipFree
+#define cudaMemcpy  hipMemcpy
+#define cudaMemset  hipMemset
+
+// ---- memcpy kinds ----
+#define cudaMemcpyHostToDevice    hipMemcpyHostToDevice
+#define cudaMemcpyDeviceToHost    hipMemcpyDeviceToHost
+#define cudaMemcpyDeviceToDevice  hipMemcpyDeviceToDevice
+
+// ---- cuBLAS -> hipBLAS (optional batched-LU solver, USE_CUBLAS=ON) ----
+// The default Gauss-Jordan solver uses no GPU BLAS, so none of this is pulled
+// in unless USE_CUBLAS is requested. hipBLAS mirrors the cuBLAS batched-LU API,
+// so the cublas* calls map 1:1. One signature difference is handled at the call
+// site: getrsBatched takes float* const A[] on hipBLAS vs const float* const A[]
+// on cuBLAS (see lm_fit_cuda.cu).
+#if defined(USE_CUBLAS)
+#include <hipblas/hipblas.h>
+#define cublasHandle_t       hipblasHandle_t
+#define cublasStatus_t       hipblasStatus_t
+#define cublasCreate         hipblasCreate
+#define cublasDestroy        hipblasDestroy
+#define cublasSgetrfBatched  hipblasSgetrfBatched
+#define cublasDgetrfBatched  hipblasDgetrfBatched
+#define cublasSgetrsBatched  hipblasSgetrsBatched
+#define cublasDgetrsBatched  hipblasDgetrsBatched
+#define CUBLAS_OP_N          HIPBLAS_OP_N
+#endif // USE_CUBLAS
+
+#else // NVIDIA / CUDA
+
+#include <cuda_runtime.h>
+
+#endif
+
+#endif // GPUFIT_CUDA_TO_HIP_H

Gpufit/hip_compat/cublas_v2.h

@@ -0,0 +1,8 @@
+// HIP-only shim so the USE_CUBLAS path's `#include "cublas_v2.h"` resolves on
+// ROCm (where no CUDA headers exist). This directory is added to the include
+// path ONLY for the HIP build; the NVIDIA build uses the real cuBLAS header.
+// Routed through cuda_to_hip.h, which (under USE_CUBLAS) includes hipBLAS and
+// aliases the cublas* batched-LU symbols Gpufit uses to their hipblas*
+// equivalents.
+#pragma once
+#include "../cuda_to_hip.h"

Gpufit/hip_compat/cuda_runtime.h

@@ -0,0 +1,5 @@
+// HIP-only shim so source that does `#include <cuda_runtime.h>` resolves on
+// ROCm (where no CUDA headers exist). This directory is added to the include
+// path ONLY for the HIP build; the NVIDIA build uses the real CUDA header.
+#pragma once
+#include "../cuda_to_hip.h"

Gpufit/hip_compat/device_launch_parameters.h

@@ -0,0 +1,6 @@
+// HIP-only shim so source that does `#include <device_launch_parameters.h>`
+// resolves on ROCm. HIP provides threadIdx/blockIdx/blockDim/gridDim and the
+// kernel launch builtins via <hip/hip_runtime.h>, pulled in by cuda_to_hip.h;
+// there is no separate launch-parameters header on ROCm. HIP build only.
+#pragma once
+#include "../cuda_to_hip.h"

Gpufit/lm_fit_cuda.cu

@@ -35,7 +35,13 @@ void LMFitCUDA::solve_equation_systems_lup()
         CUBLAS_OP_N,
         info_.n_parameters_to_fit_,
         1,
+#ifdef USE_HIP
+        // hipBLAS getrsBatched takes float* const A[] (non-const float);
+        // cuBLAS takes const float* const A[]. Drop the leading const on HIP.
+        (REAL **)(gpu_data_.pointer_decomposed_hessians_.data()),
+#else
         (REAL const **)(gpu_data_.pointer_decomposed_hessians_.data()),
+#endif
         info_.n_parameters_to_fit_,
         gpu_data_.pivot_vectors_,
         gpu_data_.pointer_deltas_,

Gpufit/tests/Gauss_Fit_2D_Rotated.cpp

@@ -94,5 +94,13 @@ BOOST_AUTO_TEST_CASE( Gauss_Fit_2D_Rotated )
     BOOST_CHECK(std::abs(output_parameters[3] - true_parameters[3]) < 1e-6f);
     BOOST_CHECK(std::abs(output_parameters[4] - true_parameters[4]) < 1e-6f);
     BOOST_CHECK(std::abs(output_parameters[5] - true_parameters[5]) < 1e-6f);
+#if defined(USE_HIP)
+    // Rotation angle r: wave32 (gfx1100) vs wave64 (gfx90a) changes the GJ
+    // solver's block packing and thus the FP reduction order, shifting the
+    // converged r by ~1.1e-6 (chi^2 still ~1e-12). 3e-6f keeps margin over the
+    // observed error while staying tight; the CUDA path keeps strict 1e-6f.
+    BOOST_CHECK(std::abs(output_parameters[6] - true_parameters[6]) < 3e-6f);
+#else
     BOOST_CHECK(std::abs(output_parameters[6] - true_parameters[6]) < 1e-6f);
+#endif
 }

examples/c++/CMakeLists.txt

@@ -8,7 +8,24 @@ function( add_example module name )
 endfunction()
 
 function( add_cuda_example module name )
-	cuda_add_executable( ${name} ${name}.cu )
+	if( USE_HIP )
+		# Compile the .cu device-API example as a HIP TU. The compat header lives
+		# in the Gpufit module dir; force-include it and add the hip_compat shims.
+		set( gpufit_src_dir ${CMAKE_CURRENT_SOURCE_DIR}/../../Gpufit )
+		set_source_files_properties( ${name}.cu PROPERTIES LANGUAGE HIP )
+		add_executable( ${name} ${name}.cu )
+		# project(... LANGUAGES CXX HIP) in the top-level CMakeLists already honors an
+		# explicit -DCMAKE_HIP_ARCHITECTURES, otherwise auto-detects the host GPU(s)
+		# and errors on a no-GPU host, so no arch default is pinned here.
+		set_target_properties( ${name} PROPERTIES HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}" )
+		target_compile_definitions( ${name} PRIVATE USE_HIP )
+		# clang-cl (MSVC frontend) ignores GNU-style -include; use /FI instead.
+		target_compile_options( ${name} PRIVATE
+			$<$<COMPILE_LANGUAGE:HIP>:$<IF:$<STREQUAL:${CMAKE_HIP_COMPILER_FRONTEND_VARIANT},MSVC>,/FI${gpufit_src_dir}/cuda_to_hip.h,-include${gpufit_src_dir}/cuda_to_hip.h>> )
+		target_include_directories( ${name} PRIVATE ${gpufit_src_dir}/hip_compat )
+	else()
+		cuda_add_executable( ${name} ${name}.cu )
+	endif()
 	target_link_libraries( ${name} ${module} )
 	set_property( TARGET ${name}
 		PROPERTY RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}" )