src/sanity/ordered_on.cuh

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