src/utils/thread_id.hpp

0001 /**
0002  * traccc library, part of the ACTS project (R&D line)
0003  *
0004  * (c) 2024-2025 CERN for the benefit of the ACTS project
0005  *
0006  * Mozilla Public License Version 2.0
0007  */
0008
0009 #pragma once
0010
0011 // Project include(s).
0012 #include "traccc/device/concepts/thread_id.hpp"
0013
0014 // SYCL include(s).
0015 #include <sycl/sycl.hpp>
0016
0017 namespace traccc::sycl::details {
0018
0019 /// A SYCL thread identifier type
0020 template <std::size_t DIMENSIONS>
0021   requires(DIMENSIONS >= 1 && DIMENSIONS <= 3)
0022 struct thread_id {
0023   /// Dimensions of the thread identifier
0024   static constexpr std::size_t dimensions = DIMENSIONS;
0025
0026   /// Constructor
0027   ///
0028   /// @param item The @c ::sycl::nd_item<1> to use for the thread ID.
0029   ///
0030   explicit thread_id(const ::sycl::nd_item<dimensions>& item) : m_item(item) {}
0031
0032   /// @name Function(s) implementing @c traccc::device::concepts::thread_id1
0033   /// @{
0034
0035   inline unsigned int getLocalThreadId() const {
0036     return static_cast<unsigned int>(m_item.get_local_linear_id());
0037   }
0038
0039   inline unsigned int getLocalThreadIdX() const {
0040     return static_cast<unsigned int>(m_item.get_local_id(0));
0041   }
0042
0043   inline unsigned int getGlobalThreadId() const {
0044     return static_cast<unsigned int>(m_item.get_global_linear_id());
0045   }
0046
0047   inline unsigned int getGlobalThreadIdX() const {
0048     return static_cast<unsigned int>(m_item.get_global_id(0));
0049   }
0050
0051   inline unsigned int getBlockIdX() const {
0052     return static_cast<unsigned int>(m_item.get_group(0));
0053   }
0054
0055   inline unsigned int getBlockDimX() const {
0056     return static_cast<unsigned int>(m_item.get_local_range(0));
0057   }
0058
0059   inline unsigned int getGridDimX() const {
0060     return static_cast<unsigned int>(m_item.get_global_range(0));
0061   }
0062
0063   /// @}
0064
0065  private:
0066   /// Item object coming from the SYCL kernel
0067   const ::sycl::nd_item<dimensions>& m_item;
0068
0069 };  // struct thread_id
0070
0071 /// Template deduction guide for the thread identifier type.
0072 template <int N>
0073 thread_id(::sycl::nd_item<N>) -> thread_id<N>;
0074
0075 /// Verify that @c traccc::sycl::details::thread_id fulfills the
0076 /// @c traccc::device::concepts::thread_id1 concept.
0077 static_assert(traccc::device::concepts::thread_id1<thread_id<1>>);
0078 static_assert(traccc::device::concepts::thread_id1<thread_id<2>>);
0079 static_assert(traccc::device::concepts::thread_id1<thread_id<3>>);
0080
0081 }  // namespace traccc::sycl::details