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/global_index.hpp"
0013 #include "../utils/hip_error_handling.hpp"
0014 #include "../utils/utils.hpp"
0015 #include "traccc/hip/utils/stream_wrapper.hpp"
0016
0017 // VecMem include(s).
0018 #include <vecmem/memory/memory_resource.hpp>
0019 #include <vecmem/memory/unique_ptr.hpp>
0020 #include <vecmem/utils/copy.hpp>
0021
0022 // HIP include(s).
0023 #include <hip/hip_runtime.h>
0024
0025 // System include
0026 #include <concepts>
0027 #include <utility>
0028
0029 namespace traccc::hip {
0030
0031 namespace kernels {
0032
0033 template <typename CONTAINER, typename R, typename VIEW>
0034   requires std::regular_invocable<R, decltype(std::declval<CONTAINER>().at(0)),
0035                                   decltype(std::declval<CONTAINER>().at(0))> &&
0036            std::semiregular<R>
0037 __global__ void is_ordered_on_kernel(R relation, VIEW _in, bool* out) {
0038   const device::global_index_t tid = details::global_index1();
0039
0040   const CONTAINER in(_in);
0041
0042   if (tid > 0 && tid < in.size()) {
0043     if (!relation(in.at(tid - 1), in.at(tid))) {
0044       *out = false;
0045     }
0046   }
0047 }
0048
0049 }  // namespace kernels
0050
0051 /**
0052  * @brief Sanity check that a given container is ordered on a given relation.
0053  *
0054  * For a container $v$ to be ordered on a relation $R$, it must be the case that
0055  * for all indices $i$ and $j$, if $i < j$, then $R(i, j)$.
0056  *
0057  * @note This function runs in O(n) time.
0058  *
0059  * @note Although functions like `std::sort` requires the relation to be strict
0060  * weak order, this function is more lax in its requirements. Rather, the
0061  * relation should be a total preorder, i.e. a non-strict weak order.
0062  *
0063  * @note For any strict weak order $R$, `is_ordered_on(sort(R, v))` is true.
0064  *
0065  * @tparam CONTAINER The type of the (device) container.
0066  * @tparam R The type of relation $R$, a callable which returns a bool if the
0067  * first argument can be immediately before the second type.
0068  * @tparam VIEW The type of the view for the container.
0069  * @param relation A relation object of type `R`.
0070  * @param mr A memory resource used for allocating intermediate memory.
0071  * @param view The container which to check for ordering.
0072  * @return true If the container is ordered on `R`.
0073  * @return false Otherwise.
0074  */
0075 template <typename CONTAINER, typename R, typename VIEW>
0076   requires std::regular_invocable<R, decltype(std::declval<CONTAINER>().at(0)),
0077                                   decltype(std::declval<CONTAINER>().at(0))> &&
0078            std::semiregular<R>
0079 bool is_ordered_on(R&& relation, vecmem::memory_resource& mr,
0080                    const vecmem::copy& copy, const stream_wrapper& stream,
0081                    const VIEW& view) {
0082   // This should never be a performance-critical step, so we can keep the
0083   // block size fixed.
0084   constexpr int block_size = 512;
0085
0086   hipStream_t hip_stream = details::get_stream(stream);
0087
0088   // Grab the number of elements in our container.
0089   const typename VIEW::size_type n = copy.get_size(view);
0090
0091   // Exit early for empty containers.
0092   if (n == 0) {
0093     return true;
0094   }
0095
0096   // Initialize the output boolean.
0097   vecmem::unique_alloc_ptr<bool> out = vecmem::make_unique_alloc<bool>(mr);
0098   bool initial_out = true;
0099   TRACCC_HIP_ERROR_CHECK(hipMemcpyAsync(out.get(), &initial_out, sizeof(bool),
0100                                         hipMemcpyHostToDevice, hip_stream));
0101
0102   // Launch the kernel which will write its result to the `out` boolean.
0103   hipLaunchKernelGGL(kernels::is_ordered_on_kernel<CONTAINER>,
0104                      (n + block_size - 1) / block_size, block_size, 0,
0105                      hip_stream, relation, view, out.get());
0106   TRACCC_HIP_ERROR_CHECK(hipGetLastError());
0107
0108   // Copy the output to host, then return it.
0109   bool host_out;
0110
0111   TRACCC_HIP_ERROR_CHECK(hipMemcpyAsync(&host_out, out.get(), sizeof(bool),
0112                                         hipMemcpyDeviceToHost, hip_stream));
0113
0114   stream.synchronize();
0115
0116   return host_out;
0117 }
0118
0119 }  // namespace traccc::hip