diff --git a/CMakeLists.txt b/CMakeLists.txt index e2a06d0c0d..352ea98d3d 100644 --- a/CMakeLists.txt +++ b/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() diff --git a/CMakePresets.json b/CMakePresets.json index 71a285889a..5bb8a0a85a 100644 --- a/CMakePresets.json +++ b/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"] } ] } diff --git a/cmake/traccc-compiler-options-hip.cmake b/cmake/traccc-compiler-options-hip.cmake new file mode 100644 index 0000000000..b3ee8213f2 --- /dev/null +++ b/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() diff --git a/device/hip/CMakeLists.txt b/device/hip/CMakeLists.txt new file mode 100644 index 0000000000..30a807cd0b --- /dev/null +++ b/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 ) diff --git a/device/hip/include/traccc/hip/clusterization/clusterization_algorithm.hpp b/device/hip/include/traccc/hip/clusterization/clusterization_algorithm.hpp new file mode 100644 index 0000000000..13b2c26b07 --- /dev/null +++ b/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 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& disjoint_set, + edm::silicon_cluster_collection::view& cluster_data) const override; + + /// @} + +}; // class clusterization_algorithm + +} // namespace traccc::hip diff --git a/device/hip/include/traccc/hip/utils/algorithm_base.hpp b/device/hip/include/traccc/hip/utils/algorithm_base.hpp new file mode 100644 index 0000000000..da3207fb25 --- /dev/null +++ b/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 + +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 m_stream; + /// Warp size of the GPU being used + unsigned int m_warp_size; + +}; // class algorithm_base + +} // namespace traccc::hip diff --git a/device/hip/include/traccc/hip/utils/stream.hpp b/device/hip/include/traccc/hip/utils/stream.hpp new file mode 100644 index 0000000000..9cab8d8fd6 --- /dev/null +++ b/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 + +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 m_stream; + +}; // class stream + +} // namespace traccc::hip diff --git a/device/hip/src/clusterization/clusterization_algorithm.hip b/device/hip/src/clusterization/clusterization_algorithm.hip new file mode 100644 index 0000000000..ef806cb16a --- /dev/null +++ b/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 + +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 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& 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 diff --git a/device/hip/src/clusterization/kernels/ccl_kernel.hip b/device/hip/src/clusterization/kernels/ccl_kernel.hip new file mode 100644 index 0000000000..0a33d1f785 --- /dev/null +++ b/device/hip/src/clusterization/kernels/ccl_kernel.hip @@ -0,0 +1,64 @@ +/** 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 + */ + +// Local include(s). +#include "../../utils/barrier.hpp" +#include "../../utils/hip_error_handling.hpp" +#include "../../utils/thread_id.hpp" +#include "../../utils/utils.hpp" + +// Project include(s). +#include "traccc/clusterization/clustering_config.hpp" +#include "traccc/clusterization/device/ccl_kernel.hpp" +#include "traccc/clusterization/device/ccl_kernel_definitions.hpp" +#include "traccc/utils/projections.hpp" +#include "traccc/utils/relations.hpp" + +// Vecmem include(s). +#include +#include + +namespace traccc::hip::kernels { + +__global__ void ccl_kernel( + const clustering_config cfg, + const edm::silicon_cell_collection::const_view cells_view, + const silicon_detector_description::const_view det_descr_view, + edm::measurement_collection::view measurements_view, + vecmem::data::vector_view cell_links, + vecmem::data::vector_view f_backup_view, + vecmem::data::vector_view gf_backup_view, + vecmem::data::vector_view adjc_backup_view, + vecmem::data::vector_view adjv_backup_view, + unsigned int* backup_mutex_ptr, + vecmem::data::vector_view disjoint_set_view, + vecmem::data::vector_view cluster_size_view) { + + __shared__ std::size_t partition_start, partition_end; + __shared__ std::size_t outi; + extern __shared__ device::details::index_t shared_v[]; + vecmem::device_atomic_ref backup_mutex(*backup_mutex_ptr); + + using vector_size_t = + vecmem::data::vector_view::size_type; + + vecmem::data::vector_view f_view{ + static_cast(cfg.max_partition_size()), shared_v}; + vecmem::data::vector_view gf_view{ + static_cast(cfg.max_partition_size()), + shared_v + cfg.max_partition_size()}; + traccc::hip::barrier barry_r; + const details::thread_id1 thread_id; + + device::ccl_kernel( + cfg, thread_id, cells_view, det_descr_view, partition_start, + partition_end, outi, f_view, gf_view, f_backup_view, gf_backup_view, + adjc_backup_view, adjv_backup_view, backup_mutex, disjoint_set_view, + cluster_size_view, barry_r, measurements_view, cell_links); +} + +} // namespace traccc::hip::kernels diff --git a/device/hip/src/clusterization/kernels/ccl_kernel.hpp b/device/hip/src/clusterization/kernels/ccl_kernel.hpp new file mode 100644 index 0000000000..7c124db2a6 --- /dev/null +++ b/device/hip/src/clusterization/kernels/ccl_kernel.hpp @@ -0,0 +1,38 @@ +/** 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 + +// Project include(s). +#include "traccc/clusterization/clustering_config.hpp" +#include "traccc/clusterization/device/ccl_kernel_definitions.hpp" +#include "traccc/definitions/primitives.hpp" +#include "traccc/edm/measurement_collection.hpp" +#include "traccc/edm/silicon_cell_collection.hpp" +#include "traccc/geometry/silicon_detector_description.hpp" + +// VecMem include(s). +#include + +namespace traccc::hip::kernels { + +/// HIP kernel for running @c traccc::device::ccl_kernel +__global__ void ccl_kernel( + const clustering_config cfg, + const edm::silicon_cell_collection::const_view cells_view, + const silicon_detector_description::const_view det_descr_view, + edm::measurement_collection::view measurements_view, + vecmem::data::vector_view cell_links, + vecmem::data::vector_view f_backup_view, + vecmem::data::vector_view gf_backup_view, + vecmem::data::vector_view adjc_backup_view, + vecmem::data::vector_view adjv_backup_view, + unsigned int* backup_mutex_ptr, + vecmem::data::vector_view disjoint_set_view, + vecmem::data::vector_view cluster_size_view); + +} // namespace traccc::hip::kernels diff --git a/device/hip/src/clusterization/kernels/reify_cluster_data.hip b/device/hip/src/clusterization/kernels/reify_cluster_data.hip new file mode 100644 index 0000000000..070752865d --- /dev/null +++ b/device/hip/src/clusterization/kernels/reify_cluster_data.hip @@ -0,0 +1,25 @@ +/** 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 + */ + +// Local include(s). +#include "../../utils/thread_id.hpp" +#include "reify_cluster_data.hpp" + +// Project include(s). +#include "traccc/clusterization/device/reify_cluster_data.hpp" + +namespace traccc::hip::kernels { + +__global__ void reify_cluster_data( + vecmem::data::vector_view disjoint_set_view, + traccc::edm::silicon_cluster_collection::view cluster_view) { + + device::reify_cluster_data(details::thread_id1{}.getGlobalThreadId(), + disjoint_set_view, cluster_view); +} + +} // namespace traccc::hip::kernels diff --git a/device/hip/src/clusterization/kernels/reify_cluster_data.hpp b/device/hip/src/clusterization/kernels/reify_cluster_data.hpp new file mode 100644 index 0000000000..6d26960542 --- /dev/null +++ b/device/hip/src/clusterization/kernels/reify_cluster_data.hpp @@ -0,0 +1,27 @@ +/** 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 + */ + +#pragma once + +// Project include(s). +#include "traccc/edm/silicon_cluster_collection.hpp" + +// VecMem include(s). +#include + +namespace traccc::hip::kernels { + +/// Fill the cluster collection with the cell indices +/// +/// @param disjoint_set_view The cluster/measurement index of each cell +/// @param cluster_view The collection to fill +/// +__global__ void reify_cluster_data( + vecmem::data::vector_view disjoint_set_view, + traccc::edm::silicon_cluster_collection::view cluster_view); + +} // namespace traccc::hip::kernels diff --git a/device/hip/src/utils/algorithm_base.cpp b/device/hip/src/utils/algorithm_base.cpp new file mode 100644 index 0000000000..7f2d912308 --- /dev/null +++ b/device/hip/src/utils/algorithm_base.cpp @@ -0,0 +1,28 @@ +/** 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 + */ + +// Local include(s). +#include "traccc/hip/utils/algorithm_base.hpp" + +#include "../utils/utils.hpp" + +namespace traccc::hip { + +algorithm_base::algorithm_base(hip::stream& str) + : m_stream(str), m_warp_size(details::get_warp_size(str.device())) {} + +hip::stream& algorithm_base::stream() const { + + return m_stream.get(); +} + +unsigned int algorithm_base::warp_size() const { + + return m_warp_size; +} + +} // namespace traccc::hip diff --git a/device/hip/src/utils/barrier.hpp b/device/hip/src/utils/barrier.hpp new file mode 100644 index 0000000000..049d3d41b3 --- /dev/null +++ b/device/hip/src/utils/barrier.hpp @@ -0,0 +1,40 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/device/concepts/barrier.hpp" + +// HIP include(s). +#include + +namespace traccc::hip { + +/// A HIP barrier implementation +struct barrier { + + __device__ inline void blockBarrier() const { __syncthreads(); } + + __device__ inline bool blockAnd(bool predicate) const { + return __syncthreads_and(predicate); + } + + __device__ inline bool blockOr(bool predicate) const { + return __syncthreads_or(predicate); + } + + __device__ inline int blockCount(bool predicate) const { + return __syncthreads_count(predicate); + } + +}; // struct barrier + +// Ensure that the barrier concept is satisfied +static_assert(traccc::device::concepts::barrier); + +} // namespace traccc::hip diff --git a/device/hip/src/utils/hip_error_handling.cpp b/device/hip/src/utils/hip_error_handling.cpp new file mode 100644 index 0000000000..d360f482b5 --- /dev/null +++ b/device/hip/src/utils/hip_error_handling.cpp @@ -0,0 +1,30 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2024-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// Local include(s). +#include "hip_error_handling.hpp" + +// System include(s). +#include +#include +#include + +namespace traccc::hip::details { + +void throw_error(hipError_t errorCode, const char* expression, const char* file, + int line) { + + // Create a nice error message. + std::ostringstream errorMsg; + errorMsg << file << ":" << line << " Failed to execute: " << expression + << " (" << hipGetErrorString(errorCode) << ")"; + + // Now throw a runtime error with this message. + throw std::runtime_error(errorMsg.str()); +} + +} // namespace traccc::hip::details diff --git a/device/hip/src/utils/hip_error_handling.hpp b/device/hip/src/utils/hip_error_handling.hpp new file mode 100644 index 0000000000..5e947b91c5 --- /dev/null +++ b/device/hip/src/utils/hip_error_handling.hpp @@ -0,0 +1,29 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2021-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// HIP include(s). +#include + +/// Helper macro used for checking @c hipError_t type return values. +#define TRACCC_HIP_ERROR_CHECK(EXP) \ + do { \ + hipError_t errorCode = EXP; \ + if (errorCode != hipSuccess) { \ + traccc::hip::details::throw_error(errorCode, #EXP, __FILE__, \ + __LINE__); \ + } \ + } while (false) + +namespace traccc::hip::details { + +/// Function used to print and throw a user-readable error if something breaks +void throw_error(hipError_t errorCode, const char* expression, const char* file, + int line); + +} // namespace traccc::hip::details diff --git a/device/hip/src/utils/opaque_stream.cpp b/device/hip/src/utils/opaque_stream.cpp new file mode 100644 index 0000000000..c271aed143 --- /dev/null +++ b/device/hip/src/utils/opaque_stream.cpp @@ -0,0 +1,34 @@ +/** 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 + */ + +// Local include(s). +#include "opaque_stream.hpp" + +#include "hip_error_handling.hpp" + +namespace traccc::hip::details { + +opaque_stream::opaque_stream(int device) : m_device{device}, m_stream(nullptr) { + + TRACCC_HIP_ERROR_CHECK(hipStreamCreate(&m_stream)); +} + +opaque_stream::~opaque_stream() { + + // Don't check the return value of the stream destruction. This is because + // if the holder of this opaque stream is only destroyed during the + // termination of the application in which it was created, the HIP runtime + // may have already deleted all streams by the time that this function would + // try to delete it. + // + // This is not the most robust thing ever, but detecting reliably when this + // destructor is executed as part of the final operations of an application, + // would be too platform specific and fragile of an operation. + [[maybe_unused]] auto code = hipStreamDestroy(m_stream); +} + +} // namespace traccc::hip::details diff --git a/device/hip/src/utils/opaque_stream.hpp b/device/hip/src/utils/opaque_stream.hpp new file mode 100644 index 0000000000..b41217afe8 --- /dev/null +++ b/device/hip/src/utils/opaque_stream.hpp @@ -0,0 +1,34 @@ +/** 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 + +// HIP include(s). +#include + +namespace traccc::hip::details { + +/// RAII wrapper around @c hipStream_t +/// +/// It is used only internally by the HIP library, so it does not need to +/// provide any nice interface. +/// +struct opaque_stream { + + /// Default constructor + opaque_stream(int device); + /// Destructor + ~opaque_stream(); + + /// Device that the stream is associated to + int m_device; + /// Stream managed by the object + hipStream_t m_stream; + +}; // class opaque_stream + +} // namespace traccc::hip::details diff --git a/device/hip/src/utils/stream.cpp b/device/hip/src/utils/stream.cpp new file mode 100644 index 0000000000..52cdc33958 --- /dev/null +++ b/device/hip/src/utils/stream.cpp @@ -0,0 +1,51 @@ +/** 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 + */ + +// Local include(s). +#include "traccc/hip/utils/stream.hpp" + +#include "hip_error_handling.hpp" +#include "opaque_stream.hpp" +#include "utils.hpp" + +// HIP include(s). +#include + +namespace traccc::hip { + +stream::stream(int device) { + + // Make sure that the stream is constructed on the correct device. + details::select_device dev_selector{ + device == INVALID_DEVICE ? details::get_device() : device}; + + // Construct the stream. + m_stream = std::make_unique(dev_selector.device()); +} + +stream::stream(stream&& parent) noexcept = default; + +stream::~stream() = default; + +stream& stream::operator=(stream&& rhs) noexcept = default; + +int stream::device() const { + + return m_stream->m_device; +} + +void* stream::hipStream() const { + + return m_stream->m_stream; +} + +void stream::synchronize() const { + + TRACCC_HIP_ERROR_CHECK(hipStreamSynchronize(m_stream->m_stream)); +} + +} // namespace traccc::hip diff --git a/device/hip/src/utils/thread_id.hpp b/device/hip/src/utils/thread_id.hpp new file mode 100644 index 0000000000..59c785d66a --- /dev/null +++ b/device/hip/src/utils/thread_id.hpp @@ -0,0 +1,49 @@ +/** + * traccc library, part of the ACTS project (R&D line) + * + * (c) 2024-2025 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/device/concepts/thread_id.hpp" + +// HIP include(s). +#include + +namespace traccc::hip::details { + +/// A HIP thread identifier type +struct thread_id1 { + + __device__ thread_id1() {} + + inline unsigned int __device__ getLocalThreadId() const { + return threadIdx.x; + } + + inline unsigned int __device__ getLocalThreadIdX() const { + return threadIdx.x; + } + + inline unsigned int __device__ getGlobalThreadId() const { + return threadIdx.x + blockIdx.x * blockDim.x; + } + + inline unsigned int __device__ getGlobalThreadIdX() const { + return threadIdx.x + blockIdx.x * blockDim.x; + } + + inline unsigned int __device__ getBlockIdX() const { return blockIdx.x; } + + inline unsigned int __device__ getBlockDimX() const { return blockDim.x; } + + inline unsigned int __device__ getGridDimX() const { return gridDim.x; } +}; + +static_assert(traccc::device::concepts::thread_id1); + +} // namespace traccc::hip::details diff --git a/device/hip/src/utils/utils.cpp b/device/hip/src/utils/utils.cpp new file mode 100644 index 0000000000..353e749a05 --- /dev/null +++ b/device/hip/src/utils/utils.cpp @@ -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 + */ + +// Local include(s). +#include "utils.hpp" + +#include "hip_error_handling.hpp" + +namespace traccc::hip::details { + +int get_device() { + + int d = -1; + [[maybe_unused]] auto code = hipGetDevice(&d); + return d; +} + +unsigned int get_warp_size(int device) { + + int warp_size = 0; + TRACCC_HIP_ERROR_CHECK( + hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, device)); + return static_cast(warp_size); +} + +hipStream_t get_stream(const stream& stream) { + + return reinterpret_cast(stream.hipStream()); +} + +select_device::select_device(int device) { + /* + * When the object is constructed, grab the current device number and + * store it as a member variable. Then set the device to whatever was + * specified. + */ + TRACCC_HIP_ERROR_CHECK(hipGetDevice(&m_device)); + TRACCC_HIP_ERROR_CHECK(hipSetDevice(device)); +} + +select_device::~select_device() { + /* + * On destruction, reset the device number to whatever it was before the + * object was constructed. + */ + TRACCC_HIP_ERROR_CHECK(hipSetDevice(m_device)); +} + +int select_device::device() const { + + return m_device; +} + +} // namespace traccc::hip::details diff --git a/device/hip/src/utils/utils.hpp b/device/hip/src/utils/utils.hpp new file mode 100644 index 0000000000..c90944c520 --- /dev/null +++ b/device/hip/src/utils/utils.hpp @@ -0,0 +1,75 @@ +/** 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 + +// Local include(s). +#include "traccc/hip/utils/stream.hpp" + +// HIP include(s). +#include + +namespace traccc::hip::details { + +/// Get current HIP device number. +/// +/// This function wraps the hipGetDevice function in a way that returns the +/// device number rather than use a reference argument to write to. +/// +/// Note that calling the function on a machine with no HIP device does not +/// result in an error, the function just returns -1 in that case. +/// +int get_device(); + +/// Get the warp size for a given device. +/// +/// @param device The device to query. +/// +/// @return The warp size for the device. +/// +unsigned int get_warp_size(int device); + +/// Get concrete @c hipStream_t object out of our wrapper +hipStream_t get_stream(const stream& str); + +/// Class with RAII mechanism for selecting a HIP device. +/// +/// This class can be used to select HIP devices in a modern C++ way, with +/// scope safety. When an object of this class is constructed, it will switch +/// the thread-local device selector to the device number specified in the +/// constructor argument. When this object goes out of scope or gets +/// destructed in any other way, it will restore the device that was set +/// before the object was constructed. This allows us to easily write methods +/// with few side-effects. +/// +/// @warning The behaviour of this class is not well-defined if you construct +/// more than one in the same scope. +/// +class select_device { + + public: + /// Constructs the object, switching the current HIP device + /// to the requested number. + /// + /// @param device The HIP device number to switch to. + /// + select_device(int device); + + /// Deconstructs the object, returning to the device that was + /// selected before constructing this object. + ~select_device(); + + /// Return the identifier for the device being seleced + int device() const; + + private: + /// The old device number, this is what we restore when the + /// object goes out of scope. + int m_device; +}; + +} // namespace traccc::hip::details diff --git a/extern/rocThrust/CMakeLists.txt b/extern/rocThrust/CMakeLists.txt index f4e37213d1..ab879904dd 100644 --- a/extern/rocThrust/CMakeLists.txt +++ b/extern/rocThrust/CMakeLists.txt @@ -1,6 +1,6 @@ # TRACCC library, part of the ACTS project (R&D line) # -# (c) 2024 CERN for the benefit of the ACTS project +# (c) 2024-2025 CERN for the benefit of the ACTS project # # Mozilla Public License Version 2.0 @@ -27,6 +27,15 @@ FetchContent_Declare( rocThrust SYSTEM set( BUILD_TESTING FALSE ) set( HIP_COMPILER "clang" ) set( HIP_CXX_COMPILER "hipcc" ) +# The following is pretty horrendous. :-( If we let HIP_PLATFORM="amd" propagate +# to the ROCm CMake files, at least with ROCm 6.0, the hip::device target gets +# set up such that linking against it automatically adds "-x hip" to the +# compiler flags of the client. Whether or not the source file is HIP or not. +# I worked around this by telling rocThrust to "always use the NVIDIA platform". +# In which case it doesn't try to be cute about such compiler flags. And then I +# let the higher level configs (from vecmem/hip or Alpaka) sort out how +# they would configure the build of various source files. +set( HIP_PLATFORM "nvidia" ) # Get it into the current directory. FetchContent_MakeAvailable( rocThrust ) diff --git a/tests/hip/CMakeLists.txt b/tests/hip/CMakeLists.txt index 143be38c3f..ed22775314 100644 --- a/tests/hip/CMakeLists.txt +++ b/tests/hip/CMakeLists.txt @@ -1,19 +1,30 @@ # TRACCC library, part of the ACTS project (R&D line) # -# (c) 2024 CERN for the benefit of the ACTS project +# (c) 2024-2025 CERN for the benefit of the ACTS project # # Mozilla Public License Version 2.0 +# Enable the HIP language. enable_language(HIP) + +# Use the HIP runtime library. +find_package(HIPToolkit REQUIRED) + traccc_add_test( hip # Define the sources for the test. test_thrust.hip + test_cca.cpp LINK_LIBRARIES + HIP::hiprt rocthrust GTest::gtest_main vecmem::core vecmem::hip + traccc::core + traccc::device_common + traccc::hip + traccc_tests_common ) set_target_properties( traccc_test_hip PROPERTIES diff --git a/tests/hip/test_cca.cpp b/tests/hip/test_cca.cpp new file mode 100644 index 0000000000..d0b0d63edb --- /dev/null +++ b/tests/hip/test_cca.cpp @@ -0,0 +1,105 @@ +/** 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 + */ + +#include + +#include +#include +#include +#include + +#include "tests/cca_test.hpp" +#include "traccc/clusterization/clustering_config.hpp" +#include "traccc/clusterization/device/tags.hpp" +#include "traccc/geometry/silicon_detector_description.hpp" +#include "traccc/hip/clusterization/clusterization_algorithm.hpp" +#include "traccc/hip/utils/stream.hpp" + +namespace { +vecmem::host_memory_resource host_mr; + +cca_function_t get_f_with(traccc::clustering_config cfg) { + return [cfg](const traccc::edm::silicon_cell_collection::host& cells, + const traccc::silicon_detector_description::host& dd) + -> std::pair::host>, + traccc::edm::silicon_cluster_collection::host> { + std::map::host> + geom_to_meas_map; + + traccc::hip::stream stream; + vecmem::hip::device_memory_resource device_mr; + vecmem::hip::async_copy copy{stream.hipStream()}; + + traccc::hip::clusterization_algorithm cc({device_mr, &host_mr}, copy, + stream, cfg); + + traccc::silicon_detector_description::buffer dd_buffer{ + static_cast< + traccc::silicon_detector_description::buffer::size_type>( + dd.size()), + device_mr}; + copy.setup(dd_buffer)->ignore(); + copy(vecmem::get_data(dd), dd_buffer, + vecmem::copy::type::host_to_device) + ->ignore(); + + traccc::edm::silicon_cell_collection::buffer cells_buffer{ + static_cast< + traccc::edm::silicon_cell_collection::buffer::size_type>( + cells.size()), + device_mr}; + copy.setup(cells_buffer)->wait(); + copy(vecmem::get_data(cells), cells_buffer)->wait(); + + auto [measurements_buffer, cluster_buffer] = + cc(cells_buffer, dd_buffer, + traccc::device::clustering_keep_disjoint_set{}); + traccc::edm::measurement_collection::host + measurements{host_mr}; + copy(measurements_buffer, measurements)->wait(); + + traccc::edm::silicon_cluster_collection::host clusters{host_mr}; + copy(cluster_buffer, clusters)->wait(); + + for (std::size_t i = 0; i < measurements.size(); i++) { + if (geom_to_meas_map.contains( + measurements.at(i).surface_link().value()) == false) { + geom_to_meas_map.insert( + {measurements.at(i).surface_link().value(), + traccc::edm::measurement_collection< + traccc::default_algebra>::host{host_mr}}); + } + geom_to_meas_map.at(measurements.at(i).surface_link().value()) + .push_back(measurements.at(i)); + } + + return {std::move(geom_to_meas_map), std::move(clusters)}; + }; +} +} // namespace + +TEST_P(ConnectedComponentAnalysisTests, Run) { + test_connected_component_analysis(GetParam()); +} + +INSTANTIATE_TEST_SUITE_P( + HIPFastSvAlgorithm, ConnectedComponentAnalysisTests, + ::testing::Combine( + ::testing::Values(get_f_with(default_ccl_test_config())), + ::testing::ValuesIn(ConnectedComponentAnalysisTests::get_test_files())), + ConnectedComponentAnalysisTests::get_test_name); + +INSTANTIATE_TEST_SUITE_P( + HIPFastSvAlgorithmWithScratch, ConnectedComponentAnalysisTests, + ::testing::Combine( + ::testing::Values(get_f_with(tiny_ccl_test_config())), + ::testing::ValuesIn( + ConnectedComponentAnalysisTests::get_test_files_short())), + ConnectedComponentAnalysisTests::get_test_name);