New example for demonstrating how to run reconstruction on GPU with time-varying misalignments

Author: Vakho Tsulaia

examples/run/cuda/CMakeLists.txt

@@ -26,6 +26,10 @@ traccc_add_executable( truth_fitting_example_cuda "truth_fitting_example_cuda.cp
    LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance
                   traccc::core traccc::device_common traccc::cuda
                   traccc::options )
+traccc_add_executable( misaligned_truth_fitting_example_cuda "misaligned_truth_fitting_example_cuda.cpp"
+   LINK_LIBRARIES vecmem::core vecmem::cuda traccc::io traccc::performance
+                  traccc::core traccc::device_common traccc::cuda
+                  traccc::options )
 #
 # Set up the "throughput applications".
 #

examples/run/cuda/misaligned_truth_fitting_example_cuda.cpp

@@ -0,0 +1,303 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2023-2025 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+// Project include(s).
+#include "traccc/cuda/fitting/fitting_algorithm.hpp"
+#include "traccc/cuda/utils/stream.hpp"
+#include "traccc/definitions/common.hpp"
+#include "traccc/definitions/primitives.hpp"
+#include "traccc/device/container_d2h_copy_alg.hpp"
+#include "traccc/device/container_h2d_copy_alg.hpp"
+#include "traccc/fitting/kalman_filter/kalman_fitter.hpp"
+#include "traccc/fitting/kalman_fitting_algorithm.hpp"
+#include "traccc/geometry/detector.hpp"
+#include "traccc/io/read_geometry.hpp"
+#include "traccc/io/read_measurements.hpp"
+#include "traccc/io/utils.hpp"
+#include "traccc/options/accelerator.hpp"
+#include "traccc/options/detector.hpp"
+#include "traccc/options/input_data.hpp"
+#include "traccc/options/performance.hpp"
+#include "traccc/options/program_options.hpp"
+#include "traccc/options/track_fitting.hpp"
+#include "traccc/options/track_propagation.hpp"
+#include "traccc/performance/collection_comparator.hpp"
+#include "traccc/performance/container_comparator.hpp"
+#include "traccc/performance/timer.hpp"
+#include "traccc/resolution/fitting_performance_writer.hpp"
+#include "traccc/utils/bfield.hpp"
+#include "traccc/utils/propagation.hpp"
+#include "traccc/utils/seed_generator.hpp"
+
+// Detray include(s).
+#include <detray/core/detail/alignment.hpp>
+#include <detray/io/frontend/detector_reader.hpp>
+
+// VecMem include(s).
+#include <vecmem/memory/cuda/device_memory_resource.hpp>
+#include <vecmem/memory/cuda/host_memory_resource.hpp>
+#include <vecmem/memory/cuda/managed_memory_resource.hpp>
+#include <vecmem/memory/host_memory_resource.hpp>
+#include <vecmem/utils/cuda/async_copy.hpp>
+#include <vecmem/utils/cuda/copy.hpp>
+
+// System include(s).
+#include <cstdlib>
+#include <exception>
+#include <iomanip>
+#include <iostream>
+
+using namespace traccc;
+
+// The main routine
+//
+int main(int argc, char* argv[]) {
+    std::unique_ptr<const traccc::Logger> ilogger =
+        traccc::getDefaultLogger("TracccExampleMisalignedTruthFittingCuda",
+                                 traccc::Logging::Level::INFO);
+    TRACCC_LOCAL_LOGGER(std::move(ilogger));
+
+    // Program options.
+    traccc::opts::detector detector_opts;
+    traccc::opts::input_data input_opts;
+    traccc::opts::track_propagation propagation_opts;
+    traccc::opts::track_fitting fitting_opts;
+    traccc::opts::performance performance_opts;
+    traccc::opts::accelerator accelerator_opts;
+    traccc::opts::program_options program_opts{
+        "Misaligned Truth Track Fitting Using CUDA",
+        {detector_opts, input_opts, propagation_opts, performance_opts,
+         accelerator_opts},
+        argc,
+        argv,
+        logger().cloneWithSuffix("Options")};
+
+    /// Type declarations
+    using host_detector_type = traccc::default_detector::host;
+    using device_detector_type = traccc::default_detector::device;
+
+    using scalar_type = device_detector_type::scalar_type;
+    using b_field_t =
+        covfie::field<traccc::const_bfield_backend_t<scalar_type>>;
+    using rk_stepper_type =
+        detray::rk_stepper<b_field_t::view_t, traccc::default_algebra,
+                           detray::constrained_step<scalar_type>>;
+    using device_navigator_type = detray::navigator<const device_detector_type>;
+    using device_fitter_type =
+        traccc::kalman_fitter<rk_stepper_type, device_navigator_type>;
+
+    // Memory resources used by the application.
+    vecmem::host_memory_resource host_mr;
+    vecmem::cuda::host_memory_resource cuda_host_mr;
+    vecmem::cuda::managed_memory_resource mng_mr;
+    vecmem::cuda::device_memory_resource device_mr;
+    vecmem::cuda::copy cuda_cpy;
+    traccc::memory_resource mr{device_mr, &cuda_host_mr};
+
+    // Performance writer
+    traccc::fitting_performance_writer fit_performance_writer(
+        traccc::fitting_performance_writer::config{},
+        logger().clone("FittingPerformanceWriter"));
+
+    // Output Stats
+    std::size_t n_fitted_tracks = 0;
+    std::size_t n_fitted_tracks_cuda = 0;
+
+    /*****************************
+     * Build a geometry
+     *****************************/
+
+    // B field value and its type
+    // @TODO: Set B field as argument
+    const traccc::vector3 B{0, 0, 2 * traccc::unit<traccc::scalar>::T};
+    auto field = traccc::construct_const_bfield<traccc::scalar>(B);
+
+    // Read the detector
+    detray::io::detector_reader_config reader_cfg{};
+    reader_cfg.add_file(
+        traccc::io::get_absolute_path(detector_opts.detector_file));
+    if (!detector_opts.material_file.empty()) {
+        reader_cfg.add_file(
+            traccc::io::get_absolute_path(detector_opts.material_file));
+    }
+    if (!detector_opts.grid_file.empty()) {
+        reader_cfg.add_file(
+            traccc::io::get_absolute_path(detector_opts.grid_file));
+    }
+    auto [host_det, names] =
+        detray::io::read_detector<host_detector_type>(host_mr, reader_cfg);
+
+    // Copy detector to the device
+    auto det_buff_static = detray::get_buffer(host_det, device_mr, cuda_cpy);
+
+    // Detector view object
+    auto det_view_static = detray::get_data(det_buff_static);
+
+    /// Create a "misaligned" context
+    using xf_container = host_detector_type::transform_container;
+    xf_container tf_store_aligned_host;
+    tf_store_aligned_host.reserve(
+        host_det.transform_store().size(),
+        typename host_detector_type::transform_container::context_type{});
+    for (const auto& tf : host_det.transform_store()) {
+        tf_store_aligned_host.push_back(tf);
+    }
+
+    // Copy the vector of "misaligned" transforms to the device
+    auto tf_buff_aligned = detray::get_buffer(
+        tf_store_aligned_host, device_mr, cuda_cpy, detray::copy::sync,
+        vecmem::data::buffer_type::fixed_size);
+
+    // Get the view of the "misaligned" detector using the vector of
+    // "misaligned" transforms and the static part of the detector copied to the
+    // device earlier
+    auto det_view_aligned =
+        detray::detail::misaligned_detector_view<host_detector_type>(
+            det_buff_static, tf_buff_aligned);
+
+    /*****************************
+     * Do the reconstruction
+     *****************************/
+
+    // Stream object
+    traccc::cuda::stream stream;
+
+    // Copy object
+    vecmem::copy host_copy;
+    vecmem::cuda::async_copy async_copy{stream.cudaStream()};
+
+    traccc::device::container_d2h_copy_alg<traccc::track_state_container_types>
+        track_state_d2h{mr, async_copy, logger().clone("TrackStateD2HCopyAlg")};
+
+    /// Standard deviations for seed track parameters
+    static constexpr std::array<scalar, e_bound_size> stddevs = {
+        0.03f * traccc::unit<scalar>::mm,
+        0.03f * traccc::unit<scalar>::mm,
+        0.017f,
+        0.017f,
+        0.001f / traccc::unit<scalar>::GeV,
+        1.f * traccc::unit<scalar>::ns};
+
+    // Fitting algorithm object
+    traccc::fitting_config fit_cfg(fitting_opts);
+    fit_cfg.propagation = propagation_opts;
+
+    traccc::host::kalman_fitting_algorithm host_fitting(
+        fit_cfg, host_mr, host_copy, logger().clone("HostFittingAlg"));
+    traccc::cuda::fitting_algorithm<device_fitter_type> device_fitting(
+        fit_cfg, mr, async_copy, stream, logger().clone("CudaFittingAlg"));
+
+    // Seed generator
+    traccc::seed_generator<host_detector_type> sg(host_det, stddevs);
+
+    traccc::performance::timing_info elapsedTimes;
+
+    // Iterate over events
+    for (std::size_t event = input_opts.skip;
+         event < input_opts.events + input_opts.skip; ++event) {
+
+        // Truth Track Candidates
+        traccc::event_data evt_data(input_opts.directory, event, host_mr,
+                                    input_opts.use_acts_geom_source, &host_det,
+                                    input_opts.format, false);
+
+        traccc::edm::track_candidate_container<traccc::default_algebra>::host
+            truth_track_candidates{host_mr};
+        evt_data.generate_truth_candidates(truth_track_candidates, sg, host_mr);
+
+        // track candidates buffer
+        traccc::edm::track_candidate_container<traccc::default_algebra>::buffer
+            truth_track_candidates_buffer{
+                async_copy.to(vecmem::get_data(truth_track_candidates.tracks),
+                              mr.main, mr.host,
+                              vecmem::copy::type::host_to_device),
+                async_copy.to(
+                    vecmem::get_data(truth_track_candidates.measurements),
+                    mr.main, vecmem::copy::type::host_to_device)};
+
+        // Instantiate cuda containers/collections
+        traccc::track_state_container_types::buffer track_states_cuda_buffer{
+            {{}, *(mr.host)}, {{}, *(mr.host), mr.host}};
+
+        // Run fitting
+        {
+            traccc::performance::timer t("Track fitting  (cuda)", elapsedTimes);
+
+            // For the first half of events use the static detector view
+            // For the second half of events switch to the "misaligned" detector
+            // view
+            bool firstHalf =
+                ((event - input_opts.skip) / (input_opts.events / 2) == 0);
+            track_states_cuda_buffer = device_fitting(
+                (firstHalf ? det_view_static : det_view_aligned), field,
+                {truth_track_candidates_buffer.tracks,
+                 truth_track_candidates_buffer.measurements});
+        }
+
+        traccc::track_state_container_types::host track_states_cuda =
+            track_state_d2h(track_states_cuda_buffer);
+
+        // CPU container(s)
+        traccc::host::kalman_fitting_algorithm::output_type track_states;
+
+        if (accelerator_opts.compare_with_cpu) {
+
+            {
+                traccc::performance::timer t("Track fitting  (cpu)",
+                                             elapsedTimes);
+
+                // Run fitting
+                track_states = host_fitting(
+                    host_det, field,
+                    {vecmem::get_data(truth_track_candidates.tracks),
+                     vecmem::get_data(truth_track_candidates.measurements)});
+            }
+        }
+
+        if (accelerator_opts.compare_with_cpu) {
+            // Show which event we are currently presenting the results for.
+            std::cout << "===>>> Event " << event << " <<<===" << std::endl;
+
+            // Compare the track parameters made on the host and on the device.
+            traccc::collection_comparator<
+                traccc::fitting_result<traccc::default_algebra>>
+                compare_fitting_results{"fitted tracks"};
+            compare_fitting_results(
+                vecmem::get_data(track_states.get_headers()),
+                vecmem::get_data(track_states_cuda.get_headers()));
+        }
+
+        // Statistics
+        n_fitted_tracks += track_states.size();
+        n_fitted_tracks_cuda += track_states_cuda.size();
+
+        if (performance_opts.run) {
+            for (unsigned int i = 0; i < track_states_cuda.size(); i++) {
+                const auto& trk_states_per_track =
+                    track_states_cuda.at(i).items;
+
+                const auto& fit_res = track_states_cuda[i].header;
+
+                fit_performance_writer.write(trk_states_per_track, fit_res,
+                                             host_det, evt_data);
+            }
+        }
+    }
+
+    if (performance_opts.run) {
+        fit_performance_writer.finalize();
+    }
+
+    std::cout << "==> Statistics ... " << std::endl;
+    std::cout << "- created (cuda) " << n_fitted_tracks_cuda << " fitted tracks"
+              << std::endl;
+    std::cout << "- created  (cpu) " << n_fitted_tracks << " fitted tracks"
+              << std::endl;
+    std::cout << "==>Elapsed times...\n" << elapsedTimes << std::endl;
+
+    return EXIT_SUCCESS;
+}