HIP Library, main branch (2025.11.13.)

CMakeLists.txt

@@ -428,6 +428,9 @@ add_subdirectory( device/common )
 if( TRACCC_BUILD_CUDA )
    add_subdirectory( device/cuda )
 endif()
+if( TRACCC_BUILD_HIP )
+   add_subdirectory( device/hip )
+endif()
 if( TRACCC_BUILD_SYCL )
    add_subdirectory( device/sycl )
 endif()

CMakePresets.json

@@ -53,6 +53,26 @@
             "VECMEM_BUILD_CUDA_LIBRARY" : "TRUE"
          }
       },
+      {
+         "name" : "hip-fp32",
+         "displayName" : "HIP FP32 Code Development",
+         "inherits": ["base-fp32"],
+         "cacheVariables": {
+            "TRACCC_BUILD_HIP" : "TRUE",
+            "VECMEM_BUILD_HIP_LIBRARY" : "TRUE",
+            "TRACCC_SETUP_ROCTHRUST": "TRUE"
+         }
+      },
+      {
+         "name" : "hip-fp64",
+         "displayName" : "HIP FP64 Code Development",
+         "inherits": ["base-fp64"],
+         "cacheVariables": {
+            "TRACCC_BUILD_HIP" : "TRUE",
+            "VECMEM_BUILD_HIP_LIBRARY" : "TRUE",
+            "TRACCC_SETUP_ROCTHRUST": "TRUE"
+         }
+      },
       {
          "name" : "sycl-fp32",
          "displayName" : "SYCL FP32 Code Development",
@@ -146,12 +166,14 @@
       {
          "name" : "full-fp32",
          "displayName": "Full FP32 Code Development",
-         "inherits": ["host-fp32", "cuda-fp32", "sycl-fp32", "alpaka-fp32"]
+         "inherits": ["host-fp32", "cuda-fp32", "sycl-fp32", "alpaka-fp32",
+                      "hip-fp32"]
       },
       {
          "name" : "full-fp64",
          "displayName": "Full FP64 Code Development",
-         "inherits": ["host-fp64", "cuda-fp64", "sycl-fp64", "alpaka-fp64"]
+         "inherits": ["host-fp64", "cuda-fp64", "sycl-fp64", "alpaka-fp64",
+                      "hip-fp64"]
       }
    ]
 }

cmake/traccc-compiler-options-hip.cmake

@@ -0,0 +1,32 @@
+# TRACCC library, part of the ACTS project (R&D line)
+#
+# (c) 2025 CERN for the benefit of the ACTS project
+#
+# Mozilla Public License Version 2.0
+
+# Include the helper function(s).
+include( traccc-functions )
+
+# Warning flags for the AMD backend of the HIP compiler.
+if( "${CMAKE_HIP_PLATFORM}" STREQUAL "amd" )
+   traccc_add_flag( CMAKE_HIP_FLAGS "-Wall" )
+   traccc_add_flag( CMAKE_HIP_FLAGS "-Wextra" )
+   traccc_add_flag( CMAKE_HIP_FLAGS "-Wshadow" )
+   traccc_add_flag( CMAKE_HIP_FLAGS "-Wunused-local-typedefs" )
+   traccc_add_flag( CMAKE_HIP_FLAGS "-pedantic" )
+endif()
+
+# Specific flags for the NVIDIA backend of the HIP compiler.
+if( "${CMAKE_HIP_PLATFORM}" STREQUAL "nvidia" )
+   traccc_add_flag( CMAKE_HIP_FLAGS "--expt-relaxed-constexpr" )
+   traccc_add_flag( CMAKE_HIP_FLAGS "--use_fast_math" )
+endif()
+
+# Fail on warnings, if asked for that behaviour.
+if( TRACCC_FAIL_ON_WARNINGS )
+   if( "${CMAKE_HIP_PLATFORM}" STREQUAL "amd" )
+      traccc_add_flag( CMAKE_HIP_FLAGS "-Werror" )
+   elseif( "${CMAKE_HIP_PLATFORM}" STREQUAL "nvidia" )
+      traccc_add_flag( CMAKE_HIP_FLAGS "-Werror all-warnings" )
+   endif()
+endif()

device/hip/CMakeLists.txt

@@ -0,0 +1,40 @@
+# TRACCC library, part of the ACTS project (R&D line)
+#
+# (c) 2025-2026 CERN for the benefit of the ACTS project
+#
+# Mozilla Public License Version 2.0
+
+# Enable HIP as a language.
+enable_language( HIP )
+
+# Project include(s).
+include( traccc-compiler-options-cpp )
+include( traccc-compiler-options-hip )
+
+# Find HIP/ROCm toolchain.
+find_package( HIPToolkit REQUIRED )
+
+# Set up the build of the traccc::hip library.
+traccc_add_library( traccc_hip hip TYPE SHARED
+   # Utility code.
+   "include/traccc/hip/utils/algorithm_base.hpp"
+   "src/utils/algorithm_base.cpp"
+   "include/traccc/hip/utils/stream.hpp"
+   "src/utils/stream.cpp"
+   "src/utils/opaque_stream.hpp"
+   "src/utils/opaque_stream.cpp"
+   "src/utils/hip_error_handling.hpp"
+   "src/utils/hip_error_handling.cpp"
+   "src/utils/utils.hpp"
+   "src/utils/utils.cpp"
+   # Clusterization code.
+   "include/traccc/hip/clusterization/clusterization_algorithm.hpp"
+   "src/clusterization/clusterization_algorithm.hip"
+   "src/clusterization/kernels/ccl_kernel.hpp"
+   "src/clusterization/kernels/ccl_kernel.hip"
+   "src/clusterization/kernels/reify_cluster_data.hpp"
+   "src/clusterization/kernels/reify_cluster_data.hip"
+)
+target_link_libraries( traccc_hip
+   PUBLIC traccc::core vecmem::core
+   PRIVATE HIP::hiprt traccc::device_common )

device/hip/include/traccc/hip/clusterization/clusterization_algorithm.hpp

@@ -0,0 +1,74 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2026 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/hip/utils/algorithm_base.hpp"
+
+// Project include(s).
+#include "traccc/clusterization/device/clusterization_algorithm.hpp"
+
+namespace traccc::hip {
+
+/// Algorithm performing hit clusterization
+///
+/// This algorithm implements hit clusterization in a massively-parallel
+/// approach. Each thread handles a pre-determined number of detector cells.
+///
+/// This algorithm returns a buffer which is not necessarily filled yet. A
+/// synchronisation statement is required before destroying the buffer.
+///
+class clusterization_algorithm : public device::clusterization_algorithm,
+                                 public algorithm_base {
+
+    public:
+    /// Constructor for clusterization algorithm
+    ///
+    /// @param mr The memory resource(s) to use in the algorithm
+    /// @param copy The copy object to use for copying data between device
+    ///             and host memory blocks
+    /// @param str The HIP stream to perform the operations in
+    /// @param config The clustering configuration partition
+    /// @param logger The logger instance to use for messaging
+    ///
+    clusterization_algorithm(
+        const traccc::memory_resource& mr, vecmem::copy& copy, hip::stream& str,
+        const config_type& config,
+        std::unique_ptr<const Logger> logger = getDummyLogger().clone());
+
+    private:
+    /// @name Function(s) inherited from the base class
+    /// @{
+
+    /// Function meant to perform sanity checks on the input data
+    ///
+    /// @param cells     All cells in an event
+    /// @return @c true if the input data is valid, @c false otherwise
+    ///
+    bool input_is_valid(
+        const edm::silicon_cell_collection::const_view& cells) const override;
+
+    /// Main CCL kernel launcher
+    void ccl_kernel(const ccl_kernel_payload& payload) const override;
+
+    /// Cluster data reification kernel launcher
+    ///
+    /// @param num_cells    Number of cells in the event
+    /// @param disjoint_set Buffer for the disjoint set data structure
+    /// @param cluster_data The cluster collection to fill
+    ///
+    void cluster_maker_kernel(
+        unsigned int num_cells,
+        const vecmem::data::vector_view<unsigned int>& disjoint_set,
+        edm::silicon_cluster_collection::view& cluster_data) const override;
+
+    /// @}
+
+};  // class clusterization_algorithm
+
+}  // namespace traccc::hip

device/hip/include/traccc/hip/utils/algorithm_base.hpp

@@ -0,0 +1,44 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2026 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// Local include(s).
+#include "traccc/hip/utils/stream.hpp"
+
+// System include(s).
+#include <functional>
+
+namespace traccc::hip {
+
+/// Base class for all HIP algorithms
+///
+/// Holding on to data that all HIP algorithms make use of.
+///
+class algorithm_base {
+
+    public:
+    /// Constructor for the algorithm base
+    ///
+    /// @param str The HIP stream to perform all operations on
+    ///
+    explicit algorithm_base(hip::stream& str);
+
+    /// Get the HIP stream of the algorithm
+    hip::stream& stream() const;
+    /// Get the warp size of the GPU being used
+    unsigned int warp_size() const;
+
+    private:
+    /// The HIP stream to use
+    std::reference_wrapper<hip::stream> m_stream;
+    /// Warp size of the GPU being used
+    unsigned int m_warp_size;
+
+};  // class algorithm_base
+
+}  // namespace traccc::hip

device/hip/include/traccc/hip/utils/stream.hpp

@@ -0,0 +1,58 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// System include(s).
+#include <memory>
+
+namespace traccc::hip {
+
+// Forward declaration(s).
+namespace details {
+struct opaque_stream;
+}
+
+/// Owning wrapper class around @c hipStream_t
+///
+/// It is necessary for passing around HIP stream objects in code that should
+/// not be directly exposed to the HIP header(s).
+///
+class stream {
+
+    public:
+    /// Invalid/default device identifier
+    static constexpr int INVALID_DEVICE = -1;
+
+    /// Construct a new stream (possibly for a specified device)
+    stream(int device = INVALID_DEVICE);
+
+    /// Move constructor
+    stream(stream&& parent) noexcept;
+
+    /// Destructor
+    ~stream();
+
+    /// Move assignment
+    stream& operator=(stream&& rhs) noexcept;
+
+    /// Device that the stream is associated to
+    int device() const;
+
+    /// Access a typeless pointer to the managed @c hipStream_t object
+    void* hipStream() const;
+
+    /// Wait for all queued tasks from the stream to complete
+    void synchronize() const;
+
+    private:
+    /// Smart pointer to the managed @c hipStream_t object
+    std::unique_ptr<details::opaque_stream> m_stream;
+
+};  // class stream
+
+}  // namespace traccc::hip

device/hip/src/clusterization/clusterization_algorithm.hip

@@ -0,0 +1,65 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2022-2026 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Local include(s).
+#include "../utils/hip_error_handling.hpp"
+#include "../utils/utils.hpp"
+#include "./kernels/ccl_kernel.hpp"
+#include "./kernels/reify_cluster_data.hpp"
+#include "traccc/hip/clusterization/clusterization_algorithm.hpp"
+
+// HIP include(s).
+#include <hip/hip_runtime.h>
+
+namespace traccc::hip {
+
+clusterization_algorithm::clusterization_algorithm(
+    const traccc::memory_resource& mr, vecmem::copy& copy, hip::stream& str,
+    const config_type& config, std::unique_ptr<const Logger> logger)
+    : device::clusterization_algorithm(mr, copy, config, std::move(logger)),
+      hip::algorithm_base(str) {}
+
+bool clusterization_algorithm::input_is_valid(
+    const edm::silicon_cell_collection::const_view& /*cells*/) const {
+
+    /// TODO: implement input checks
+    return true;
+}
+
+void clusterization_algorithm::ccl_kernel(
+    const ccl_kernel_payload& payload) const {
+
+    const unsigned int num_blocks =
+        (payload.n_cells + (payload.config.target_partition_size()) - 1) /
+        payload.config.target_partition_size();
+    hipLaunchKernelGGL(kernels::ccl_kernel, dim3{num_blocks},
+                       dim3{payload.config.threads_per_partition},
+                       2 * payload.config.max_partition_size() *
+                           sizeof(device::details::index_t),
+                       details::get_stream(stream()), payload.config,
+                       payload.cells, payload.det_descr, payload.measurements,
+                       payload.cell_links, payload.f_backup, payload.gf_backup,
+                       payload.adjc_backup, payload.adjv_backup,
+                       payload.backup_mutex, payload.disjoint_set,
+                       payload.cluster_sizes);
+    TRACCC_HIP_ERROR_CHECK(hipGetLastError());
+}
+
+void clusterization_algorithm::cluster_maker_kernel(
+    unsigned int num_cells,
+    const vecmem::data::vector_view<unsigned int>& disjoint_set,
+    edm::silicon_cluster_collection::view& cluster_data) const {
+
+    const unsigned int num_threads = warp_size() * 16u;
+    const unsigned int num_blocks = (num_cells + num_threads - 1) / num_threads;
+    hipLaunchKernelGGL(kernels::reify_cluster_data, num_blocks, num_threads, 0,
+                       details::get_stream(stream()), disjoint_set,
+                       cluster_data);
+    TRACCC_HIP_ERROR_CHECK(hipGetLastError());
+}
+
+}  // namespace traccc::hip