src/utils/barrier.hpp

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2023 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 #pragma once
0009
0010 // Project include(s).
0011 #include "traccc/device/concepts/barrier.hpp"
0012
0013 // SYCL include(s).
0014 #include <sycl/sycl.hpp>
0015
0016 namespace traccc::sycl::details {
0017
0018 /// A barrier type for SYCL kernels.
0019 template <std::size_t DIMENSIONS>
0020   requires(DIMENSIONS >= 1 && DIMENSIONS <= 3)
0021 struct barrier {
0022   /// Dimensions of the barrier
0023   static constexpr std::size_t dimensions = DIMENSIONS;
0024
0025   /// Constructor
0026   ///
0027   /// @param item The @c ::sycl::nd_item<1> to use for the barrier.
0028   ///
0029   explicit barrier(const ::sycl::nd_item<dimensions>& item) : m_item(item) {}
0030
0031   /// @name Function(s) implementing @c traccc::device::concepts::barrier
0032   /// @{
0033
0034   inline void blockBarrier() const {
0035     ::sycl::group_barrier(m_item.get_group());
0036   }
0037
0038   inline bool blockAnd(bool predicate) const {
0039     blockBarrier();
0040     return ::sycl::all_of_group(m_item.get_group(), predicate);
0041   }
0042
0043   inline bool blockOr(bool predicate) const {
0044     blockBarrier();
0045     return ::sycl::any_of_group(m_item.get_group(), predicate);
0046   }
0047
0048   inline unsigned int blockCount(bool predicate) const {
0049     blockBarrier();
0050     return ::sycl::reduce_over_group(m_item.get_group(), predicate ? 1u : 0u,
0051                                      ::sycl::plus<>());
0052   }
0053
0054   /// @}
0055
0056  private:
0057   /// Item object coming from the SYCL kernel
0058   const ::sycl::nd_item<dimensions>& m_item;
0059 };
0060
0061 /// Template deduction guide for the barrier type.
0062 template <int N>
0063 barrier(::sycl::nd_item<N>) -> barrier<N>;
0064
0065 /// Verify that @c traccc::sycl::details::barrier fulfills the
0066 /// @c traccc::device::concepts::barrier concept.
0067 static_assert(traccc::device::concepts::barrier<barrier<1>>);
0068 static_assert(traccc::device::concepts::barrier<barrier<2>>);
0069 static_assert(traccc::device::concepts::barrier<barrier<3>>);
0070
0071 }  // namespace traccc::sycl::details