src/sanity/ordered_on.hpp

0001 /**
0002  * traccc library, part of the ACTS project (R&D line)
0003  *
0004  * (c) 2024-2026 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 "../utils/get_queue.hpp"
0013 #include "traccc/sycl/utils/queue_wrapper.hpp"
0014
0015 // VecMem include(s).
0016 #include <vecmem/memory/memory_resource.hpp>
0017 #include <vecmem/memory/unique_ptr.hpp>
0018 #include <vecmem/utils/copy.hpp>
0019
0020 // SYCL include
0021 #include <sycl/sycl.hpp>
0022
0023 // System include
0024 #include <concepts>
0025
0026 namespace traccc::sycl {
0027
0028 namespace kernels {
0029
0030 /// Kernel used in implementing @c traccc::sycl::is_ordered_on.
0031 ///
0032 /// It has to be implemented as a functor, as oneAPI 2024.2.1 gets confused when
0033 /// trying to set up this type of code as a lambda.
0034 ///
0035 template <typename CONTAINER, typename R, typename VIEW>
0036 struct is_ordered_on {
0037   /// Constructor
0038   is_ordered_on(R relation, VIEW view, bool* out)
0039       : m_relation(relation), m_view(view), m_out(out) {}
0040
0041   /// Execution operator for the kernel
0042   void operator()(::sycl::nd_item<1> item) const {
0043     std::size_t tid = item.get_global_linear_id();
0044
0045     const CONTAINER in(m_view);
0046
0047     if (tid > 0 && tid < in.size()) {
0048       if (!m_relation(in.at(static_cast<CONTAINER::size_type>(tid - 1)),
0049                       in.at(static_cast<CONTAINER::size_type>(tid)))) {
0050         *m_out = false;
0051       }
0052     }
0053   }
0054
0055   /// The relation object to use
0056   R m_relation;
0057   /// View to the input container
0058   VIEW m_view;
0059   /// Output boolean
0060   bool* m_out;
0061 };
0062
0063 }  // namespace kernels
0064
0065 /**
0066  * @brief Sanity check that a given container is ordered on a given relation.
0067  *
0068  * For a container $v$ to be ordered on a relation $R$, it must be the case that
0069  * for all indices $i$ and $j$, if $i < j$, then $R(i, j)$.
0070  *
0071  * @note This function runs in O(n) time.
0072  *
0073  * @note Although functions like `std::sort` requires the relation to be strict
0074  * weak order, this function is more lax in its requirements. Rather, the
0075  * relation should be a total preorder, i.e. a non-strict weak order.
0076  *
0077  * @note For any strict weak order $R$, `is_ordered_on(sort(R, v))` is true.
0078  *
0079  * @tparam CONTAINER The type of the (device) container.
0080  * @tparam R The type of relation $R$, a callable which returns a bool if the
0081  * first argument can be immediately before the second type.
0082  * @tparam VIEW The type of the view for the container.
0083  * @param relation A relation object of type `R`.
0084  * @param mr A memory resource used for allocating intermediate memory.
0085  * @param view The container which to check for ordering.
0086  * @return true If the container is ordered on `R`.
0087  * @return false Otherwise.
0088  */
0089 template <typename CONTAINER, typename R, typename VIEW>
0090   requires std::regular_invocable<R, decltype(std::declval<CONTAINER>().at(0)),
0091                                   decltype(std::declval<CONTAINER>().at(0))> &&
0092            std::semiregular<R>
0093 bool is_ordered_on(R&& relation, vecmem::memory_resource& mr,
0094                    const vecmem::copy& copy, ::sycl::queue& queue,
0095                    const VIEW& view) {
0096   // This should never be a performance-critical step, so we can keep the
0097   // block size fixed.
0098   constexpr int block_size = 512;
0099
0100   // Grab the number of elements in our container.
0101   const typename VIEW::size_type n = copy.get_size(view);
0102
0103   // Exit early for empty containers.
0104   if (n == 0) {
0105     return true;
0106   }
0107
0108   // Initialize the output boolean.
0109   vecmem::unique_alloc_ptr<bool> out = vecmem::make_unique_alloc<bool>(mr);
0110   bool initial_out = true;
0111
0112   ::sycl::event kernel1_memcpy1 =
0113       queue.memcpy(out.get(), &initial_out, sizeof(bool));
0114
0115   ::sycl::nd_range<1> kernel_range{
0116       ::sycl::range<1>(((n + block_size - 1) / block_size) * block_size),
0117       ::sycl::range<1>(block_size)};
0118
0119   ::sycl::event kernel1 = queue.submit([&](::sycl::handler& h) {
0120     h.depends_on(kernel1_memcpy1);
0121     h.parallel_for<kernels::is_ordered_on<CONTAINER, R, VIEW>>(
0122         kernel_range,
0123         kernels::is_ordered_on<CONTAINER, R, VIEW>(relation, view, out.get()));
0124   });
0125
0126   // Copy the output to host, then return it.
0127   bool host_out;
0128
0129   queue.memcpy(&host_out, out.get(), sizeof(bool), {kernel1}).wait_and_throw();
0130
0131   return host_out;
0132 }
0133 }  // namespace traccc::sycl