tests/sycl/test_ckf_combinatorics_telescope.cpp

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2024-2026 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 // Test include(s).
0009 #include "tests/ckf_telescope_test.hpp"
0010 #include "traccc/utils/seed_generator.hpp"
0011
0012 // Project include(s).
0013 #include "traccc/bfield/construct_const_bfield.hpp"
0014 #include "traccc/bfield/magnetic_field_types.hpp"
0015 #include "traccc/io/read_detector.hpp"
0016 #include "traccc/io/read_measurements.hpp"
0017 #include "traccc/io/utils.hpp"
0018 #include "traccc/simulation/event_generators.hpp"
0019 #include "traccc/simulation/simulator.hpp"
0020 #include "traccc/sycl/finding/combinatorial_kalman_filter_algorithm.hpp"
0021 #include "traccc/utils/event_data.hpp"
0022 #include "traccc/utils/ranges.hpp"
0023
0024 // VecMem include(s).
0025 #include <vecmem/memory/host_memory_resource.hpp>
0026 #include <vecmem/memory/sycl/device_memory_resource.hpp>
0027 #include <vecmem/memory/sycl/shared_memory_resource.hpp>
0028 #include <vecmem/utils/sycl/async_copy.hpp>
0029
0030 // GTest include(s).
0031 #include <gtest/gtest.h>
0032
0033 // System include(s).
0034 #include <filesystem>
0035 #include <string>
0036
0037 namespace traccc {
0038
0039 TEST_P(CkfCombinatoricsTelescopeTests, Run) {
0040   // Get the parameters
0041   const std::string name = std::get<0>(GetParam());
0042   const std::array<scalar, 3u> origin = std::get<1>(GetParam());
0043   const std::array<scalar, 3u> origin_stddev = std::get<2>(GetParam());
0044   const std::array<scalar, 2u> mom_range = std::get<3>(GetParam());
0045   const std::array<scalar, 2u> eta_range = std::get<4>(GetParam());
0046   const std::array<scalar, 2u> theta_range = eta_to_theta_range(eta_range);
0047   const std::array<scalar, 2u> phi_range = std::get<5>(GetParam());
0048   const traccc::pdg_particle<scalar> ptc = std::get<6>(GetParam());
0049   const unsigned int n_truth_tracks = std::get<7>(GetParam());
0050   const unsigned int n_events = std::get<8>(GetParam());
0051   const bool random_charge = std::get<9>(GetParam());
0052
0053   /*****************************
0054    * Build a telescope geometry
0055    *****************************/
0056
0057   // SYCL queue.
0058   vecmem::sycl::queue_wrapper vecmem_queue;
0059   traccc::sycl::queue_wrapper traccc_queue{vecmem_queue.queue()};
0060
0061   // Only run this test on NVIDIA and AMD backends.
0062   if (!(vecmem_queue.is_cuda() || vecmem_queue.is_hip())) {
0063     GTEST_SKIP();
0064   }
0065
0066   // Memory resources used by the application.
0067   vecmem::host_memory_resource host_mr;
0068   vecmem::sycl::device_memory_resource device_mr{vecmem_queue};
0069   traccc::memory_resource mr{device_mr, &host_mr};
0070   vecmem::sycl::shared_memory_resource shared_mr{vecmem_queue};
0071
0072   // Copy objects
0073   vecmem::sycl::async_copy copy{vecmem_queue};
0074
0075   // Path to the working directory.
0076   const std::filesystem::path path = std::filesystem::current_path() / name;
0077
0078   // Read in the detector geometry that was generated by the test fixture.
0079   traccc::host_detector host_detector;
0080   traccc::io::read_detector(
0081       host_detector, host_mr,
0082       std::filesystem::path(path / "telescope_detector_geometry.json").native(),
0083       std::filesystem::absolute(path /
0084                                 "telescope_detector_homogeneous_material.json")
0085           .native());
0086
0087   const traccc::detector_buffer detector_buffer =
0088       traccc::buffer_from_host_detector(host_detector, device_mr, copy);
0089   vecmem_queue.synchronize();
0090
0091   auto field = traccc::construct_const_bfield(std::get<13>(GetParam()));
0092
0093   /***************************
0094    * Generate simulation data
0095    ***************************/
0096
0097   // Track generator
0098   using generator_type =
0099       detray::random_track_generator<traccc::free_track_parameters<>,
0100                                      uniform_gen_t>;
0101   generator_type::configuration gen_cfg{};
0102   gen_cfg.n_tracks(n_truth_tracks);
0103   gen_cfg.origin(origin);
0104   gen_cfg.origin_stddev(origin_stddev);
0105   gen_cfg.phi_range(phi_range[0], phi_range[1]);
0106   gen_cfg.theta_range(theta_range[0], theta_range[1]);
0107   gen_cfg.mom_range(mom_range[0], mom_range[1]);
0108   gen_cfg.randomize_charge(random_charge);
0109   generator_type generator(gen_cfg);
0110
0111   // Smearing value for measurements
0112   traccc::measurement_smearer<traccc::default_algebra> meas_smearer(
0113       smearing[0], smearing[1]);
0114
0115   using writer_type = traccc::smearing_writer<
0116       traccc::measurement_smearer<traccc::default_algebra>>;
0117
0118   typename writer_type::config smearer_writer_cfg{meas_smearer};
0119   traccc::seed_generator<host_detector_type>::config seed_cfg{};
0120   seed_cfg.initial_sigmas = stddevs;
0121
0122   // Run simulator
0123   auto sim = traccc::simulator<host_detector_type, b_field_t, generator_type,
0124                                writer_type>(
0125       ptc, n_events, host_detector.as<detector_traits>(),
0126       field.as_field<traccc::const_bfield_backend_t<traccc::scalar>>(),
0127       std::move(generator), std::move(smearer_writer_cfg), path.native());
0128   sim.run();
0129
0130   /*****************************
0131    * Do the reconstruction
0132    *****************************/
0133
0134   // Seed generator
0135   seed_generator<host_detector_type> sg(host_detector.as<detector_traits>(),
0136                                         seed_cfg);
0137
0138   // Finding algorithm configuration
0139   traccc::sycl::combinatorial_kalman_filter_algorithm::config_type cfg_no_limit;
0140   cfg_no_limit.ptc_hypothesis = ptc;
0141   cfg_no_limit.max_num_branches_per_seed = 100000;
0142   cfg_no_limit.chi2_max = 30.f;
0143   cfg_no_limit.max_num_branches_per_surface = 10;
0144   cfg_no_limit.duplicate_removal_minimum_length = 100u;
0145   cfg_no_limit.run_smoother = smoother_type::e_none;
0146
0147   traccc::sycl::combinatorial_kalman_filter_algorithm::config_type cfg_limit;
0148   cfg_limit.ptc_hypothesis = ptc;
0149   cfg_limit.max_num_branches_per_seed = 500;
0150   cfg_limit.chi2_max = 30.f;
0151   cfg_limit.max_num_branches_per_surface = 10;
0152   cfg_limit.duplicate_removal_minimum_length = 100u;
0153   cfg_limit.run_smoother = smoother_type::e_none;
0154
0155   // Finding algorithm object
0156   traccc::sycl::combinatorial_kalman_filter_algorithm device_finding{
0157       cfg_no_limit, mr, copy, traccc_queue};
0158   traccc::sycl::combinatorial_kalman_filter_algorithm device_finding_limit{
0159       cfg_limit, mr, copy, traccc_queue};
0160
0161   // Iterate over events
0162   for (std::size_t i_evt = 0; i_evt < n_events; i_evt++) {
0163     // Truth Track Candidates
0164     traccc::event_data evt_data(path, i_evt, host_mr);
0165
0166     traccc::edm::measurement_collection::host truth_measurements{host_mr};
0167     traccc::edm::track_container<traccc::default_algebra>::host
0168         truth_track_candidates{host_mr};
0169     evt_data.generate_truth_candidates(truth_track_candidates,
0170                                        truth_measurements, sg, host_mr);
0171     truth_track_candidates.measurements = vecmem::get_data(truth_measurements);
0172
0173     ASSERT_EQ(truth_track_candidates.tracks.size(), n_truth_tracks);
0174
0175     // Prepare truth seeds
0176     traccc::bound_track_parameters_collection_types::host seeds(&host_mr);
0177     for (unsigned int i_trk = 0; i_trk < n_truth_tracks; i_trk++) {
0178       seeds.push_back(truth_track_candidates.tracks.at(i_trk).params());
0179     }
0180     ASSERT_EQ(seeds.size(), n_truth_tracks);
0181
0182     traccc::bound_track_parameters_collection_types::buffer seeds_buffer{
0183         static_cast<unsigned int>(seeds.size()), mr.main};
0184     copy.setup(seeds_buffer)->wait();
0185     copy(vecmem::get_data(seeds), seeds_buffer,
0186          vecmem::copy::type::host_to_device)
0187         ->wait();
0188
0189     // Read measurements
0190     traccc::edm::measurement_collection::host measurements_per_event{host_mr};
0191     traccc::io::read_measurements(measurements_per_event, i_evt, path.native());
0192
0193     traccc::edm::measurement_collection::buffer measurements_buffer(
0194         static_cast<unsigned int>(measurements_per_event.size()), mr.main);
0195     copy.setup(measurements_buffer)->wait();
0196     copy(vecmem::get_data(measurements_per_event), measurements_buffer)->wait();
0197
0198     // Run device finding
0199     auto track_candidates_buffer = device_finding(
0200         detector_buffer, field, measurements_buffer, seeds_buffer);
0201
0202     // Run device finding (Limit)
0203     auto track_candidates_limit_buffer = device_finding_limit(
0204         detector_buffer, field, measurements_buffer, seeds_buffer);
0205
0206     traccc::edm::track_collection<traccc::default_algebra>::host
0207         track_candidates{host_mr},
0208         track_candidates_limit{host_mr};
0209     copy(track_candidates_buffer.tracks, track_candidates,
0210          vecmem::copy::type::device_to_host)
0211         ->wait();
0212     copy(track_candidates_limit_buffer.tracks, track_candidates_limit,
0213          vecmem::copy::type::device_to_host)
0214         ->wait();
0215
0216     // Make sure that the number of found tracks = n_track ^ (n_planes + 1)
0217     EXPECT_GT(track_candidates.size(), track_candidates_limit.size());
0218     EXPECT_EQ(track_candidates.size(),
0219               std::pow(n_truth_tracks, std::get<11>(GetParam()) + 1));
0220     EXPECT_EQ(track_candidates_limit.size(),
0221               n_truth_tracks * cfg_limit.max_num_branches_per_seed);
0222   }
0223 }
0224
0225 INSTANTIATE_TEST_SUITE_P(
0226     SYCLCkfCombinatoricsTelescopeValidation, CkfCombinatoricsTelescopeTests,
0227     ::testing::Values(
0228         std::make_tuple("telescope_combinatorics_twin",
0229                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0230                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0231                         std::array<scalar, 2u>{100.f, 100.f},
0232                         std::array<scalar, 2u>{0.f, 0.f},
0233                         std::array<scalar, 2u>{0.f, 0.f},
0234                         traccc::muon<scalar>(), 2, 1, false, 20.f, 9u, 20.f,
0235                         vector3{2 * traccc::unit<scalar>::T, 0, 0}),
0236         std::make_tuple("telescope_combinatorics_trio",
0237                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0238                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0239                         std::array<scalar, 2u>{100.f, 100.f},
0240                         std::array<scalar, 2u>{0.f, 0.f},
0241                         std::array<scalar, 2u>{0.f, 0.f},
0242                         traccc::muon<scalar>(), 3, 1, false, 20.f, 9u, 20.f,
0243                         vector3{2 * traccc::unit<scalar>::T, 0, 0})));
0244
0245 }  // namespace traccc