EIC code displayed by LXR master/​acts/​Traccc/​device/​hip/​src/​sanity/​contiguous_on.hpp	Source navigation Diff markup Identifier search General search
	Version: master test Revert   Change

File indexing completed on 2026-08-16 08:20:29

 /**
  * traccc library, part of the ACTS project (R&D line)
  *
  * (c) 2024-2026 CERN for the benefit of the ACTS project
  *
  * Mozilla Public License Version 2.0
  */
 
 #pragma once
 
 // Project include(s).
 #include "../utils/global_index.hpp"
 #include "../utils/hip_error_handling.hpp"
 #include "../utils/utils.hpp"
 #include "traccc/hip/utils/stream_wrapper.hpp"
 
 // VecMem include(s).
 #include <vecmem/containers/data/vector_buffer.hpp>
 #include <vecmem/containers/device_vector.hpp>
 #include <vecmem/memory/memory_resource.hpp>
 #include <vecmem/memory/unique_ptr.hpp>
 #include <vecmem/utils/copy.hpp>
 
 // HIP include(s).
 #include <hip/hip_runtime.h>
 
 // System include
 #include <concepts>
 #include <utility>
 
 namespace traccc::hip {
 
 namespace kernels {
 
 /// Kernel used in implementing @c traccc::hip::is_contiguous_on
 template <typename CONTAINER, typename P, typename VIEW,
           std::equality_comparable S>
   requires std::regular_invocable<P,
                                   decltype(std::declval<CONTAINER>().at(0))> &&
            std::semiregular<P>
 __global__ void is_contiguous_on_compress_adjacent(
     P projection, VIEW _in, vecmem::data::vector_view<S> out_view) {
   const device::global_index_t tid = details::global_index1();
 
   const CONTAINER in(_in);
   vecmem::device_vector<S> out(out_view);
 
   if (tid > 0 && tid < in.size()) {
     S v1 = projection(in.at(tid - 1));
     S v2 = projection(in.at(tid));
 
     if (v1 != v2) {
       out.push_back(v2);
     }
   } else if (tid == 0) {
     out.push_back(projection(in.at(tid)));
   }
 }
 
 /// Kernel used in implementing @c traccc::hip::is_contiguous_on
 template <std::equality_comparable T>
 __global__ void is_contiguous_on_all_unique(
     vecmem::data::vector_view<T> in_view, bool* out) {
   const device::global_index_t tid_x = threadIdx.x + blockIdx.x * blockDim.x;
   const device::global_index_t tid_y = threadIdx.y + blockIdx.y * blockDim.y;
 
   const vecmem::device_vector<T> in(in_view);
 
   if (tid_x < in.size() && tid_y < in.size() && tid_x != tid_y &&
       in.at(tid_x) == in.at(tid_y)) {
     *out = false;
   }
 }
 
 }  // namespace kernels
 
 /**
  * @brief Sanity check that a given container is contiguous on a given
  *        projection.
  *
  * For a container $v$ to be contiguous on a projection $\pi$, it must be the
  * case that for all indices $i$ and $j$, if $v_i = v_j$, then all indices $k$
  * between $i$ and $j$, $v_i = v_j = v_k$.
  *
  * @note This function runs in O(n^2) time.
  *
  * @tparam CONTAINER The type of the (device) container.
  * @tparam P The type of projection $\pi$, a callable which returns some
  * comparable type.
  * @tparam VIEW The type of the view for the container.
  * @param projection A projection object of type `P`.
  * @param mr A memory resource used for allocating intermediate memory.
  * @param view The container which to check for contiguity.
  * @return true If the container is contiguous on `P`.
  * @return false Otherwise.
  */
 template <typename CONTAINER, typename P, typename VIEW>
   requires std::regular_invocable<P,
                                   decltype(std::declval<CONTAINER>().at(0))> &&
            std::semiregular<P>
 bool is_contiguous_on(P&& projection, vecmem::memory_resource& mr,
                       const vecmem::copy& copy, const stream_wrapper& stream,
                       const VIEW& view) {
   // This should never be a performance-critical step, so we can keep the
   // block size fixed.
   constexpr int block_size = 512;
   constexpr int block_size_2d = 32;
 
   hipStream_t hip_stream = details::get_stream(stream);
 
   // Grab the number of elements in our container.
   const typename VIEW::size_type n = copy.get_size(view);
 
   // Exit early for empty containers.
   if (n == 0) {
     return true;
   }
 
   // Get the output type of the projection.
   using projection_t =
       std::invoke_result_t<P, decltype(std::declval<CONTAINER>().at(0))>;
 
   // Allocate memory for intermediate values and outputs, then set them up.
   vecmem::data::vector_buffer<projection_t> iout(
       n, mr, vecmem::data::buffer_type::resizable);
   copy.setup(iout)->ignore();
   vecmem::unique_alloc_ptr<bool> out = vecmem::make_unique_alloc<bool>(mr);
 
   bool initial_out = true;
 
   TRACCC_HIP_ERROR_CHECK(hipMemcpyAsync(out.get(), &initial_out, sizeof(bool),
                                         hipMemcpyHostToDevice, hip_stream));
 
   // Launch the first kernel, which will squash consecutive equal elements
   // into one element.
   hipLaunchKernelGGL(kernels::is_contiguous_on_compress_adjacent<CONTAINER>,
                      (n + block_size - 1) / block_size, block_size, 0,
                      hip_stream, projection, view, iout);
   TRACCC_HIP_ERROR_CHECK(hipGetLastError());
 
   // Launch the second kernel, which will check if the values are unique.
   uint32_t grid_size_rd =
       (copy.get_size(iout) + block_size_2d - 1) / block_size_2d;
   dim3 all_unique_grid_size(grid_size_rd, grid_size_rd);
   dim3 all_unique_block_size(block_size_2d, block_size_2d);
 
   hipLaunchKernelGGL(kernels::is_contiguous_on_all_unique, all_unique_grid_size,
                      all_unique_block_size, 0, hip_stream, iout, out.get());
   TRACCC_HIP_ERROR_CHECK(hipGetLastError());
 
   // Get the result from the device and return it.
   bool host_out;
 
   TRACCC_HIP_ERROR_CHECK(hipMemcpyAsync(&host_out, out.get(), sizeof(bool),
                                         hipMemcpyDeviceToHost, hip_stream));
 
   stream.synchronize();
 
   return host_out;
 }
 
 }  // namespace traccc::hip

This page was automatically generated by the 2.3.7 LXR engine. The LXR team