tests/sycl/test_kalman_fitter_telescope.cpp

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2022-2026 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 // Project include(s).
0009 #include "traccc/bfield/construct_const_bfield.hpp"
0010 #include "traccc/bfield/magnetic_field_types.hpp"
0011 #include "traccc/device/container_d2h_copy_alg.hpp"
0012 #include "traccc/edm/track_container.hpp"
0013 #include "traccc/io/detector.hpp"
0014 #include "traccc/io/utils.hpp"
0015 #include "traccc/performance/details/is_same_object.hpp"
0016 #include "traccc/resolution/fitting_performance_writer.hpp"
0017 #include "traccc/simulation/event_generators.hpp"
0018 #include "traccc/simulation/simulator.hpp"
0019 #include "traccc/sycl/fitting/kalman_fitting_algorithm.hpp"
0020 #include "traccc/utils/memory_resource.hpp"
0021 #include "traccc/utils/ranges.hpp"
0022 #include "traccc/utils/seed_generator.hpp"
0023
0024 // Test include(s).
0025 #include "tests/kalman_fitting_telescope_test.hpp"
0026
0027 // VecMem include(s).
0028 #include <vecmem/memory/host_memory_resource.hpp>
0029 #include <vecmem/memory/sycl/device_memory_resource.hpp>
0030 #include <vecmem/memory/sycl/shared_memory_resource.hpp>
0031 #include <vecmem/utils/sycl/copy.hpp>
0032 #include <vecmem/utils/sycl/queue_wrapper.hpp>
0033
0034 // GTest include(s).
0035 #include <gtest/gtest.h>
0036
0037 // System include(s).
0038 #include <exception>
0039 #include <filesystem>
0040 #include <string>
0041
0042 using namespace traccc;
0043
0044 // This defines the local frame test suite
0045 TEST_P(KalmanFittingTelescopeTests, Run) {
0046   // Get the parameters
0047   const std::string name = std::get<0>(GetParam());
0048   const std::array<scalar, 3u> origin = std::get<1>(GetParam());
0049   const std::array<scalar, 3u> origin_stddev = std::get<2>(GetParam());
0050   const std::array<scalar, 2u> mom_range = std::get<3>(GetParam());
0051   const std::array<scalar, 2u> eta_range = std::get<4>(GetParam());
0052   const std::array<scalar, 2u> theta_range = eta_to_theta_range(eta_range);
0053   const std::array<scalar, 2u> phi_range = std::get<5>(GetParam());
0054   const traccc::pdg_particle<scalar> ptc = std::get<6>(GetParam());
0055   const unsigned int n_truth_tracks = std::get<7>(GetParam());
0056   const unsigned int n_events = std::get<8>(GetParam());
0057   const bool random_charge = std::get<9>(GetParam());
0058
0059   // Performance writer
0060   traccc::fitting_performance_writer::config fit_writer_cfg;
0061   fit_writer_cfg.file_path = "performance_track_fitting_" + name + ".root";
0062   traccc::fitting_performance_writer fit_performance_writer(
0063       fit_writer_cfg, traccc::getDefaultLogger("FittingPerformanceWriter",
0064                                                traccc::Logging::Level::INFO));
0065
0066   /*****************************
0067    * Build a telescope geometry
0068    *****************************/
0069
0070   // Creating SYCL queue object
0071   vecmem::sycl::queue_wrapper vecmem_queue;
0072   traccc::sycl::queue_wrapper traccc_queue{vecmem_queue.queue()};
0073
0074   // Only run this test on NVIDIA and AMD backends.
0075   if (!(vecmem_queue.is_cuda() || vecmem_queue.is_hip())) {
0076     GTEST_SKIP();
0077   }
0078
0079   std::cout << "Running on device: " << vecmem_queue.device_name() << "\n";
0080
0081   // Memory resources used by the application.
0082   vecmem::host_memory_resource host_mr;
0083   vecmem::sycl::device_memory_resource device_mr{vecmem_queue};
0084   traccc::memory_resource mr{device_mr, &host_mr};
0085   vecmem::sycl::shared_memory_resource shared_mr{vecmem_queue};
0086
0087   // Read back detector file
0088   const std::string path = name + "/";
0089   detray::io::detector_reader_config reader_cfg{};
0090   reader_cfg.add_file(path + "telescope_detector_geometry.json")
0091       .add_file(path + "telescope_detector_homogeneous_material.json");
0092
0093   auto [host_det, names] =
0094       detray::io::read_detector<host_detector_type>(shared_mr, reader_cfg);
0095
0096   traccc::host_detector polymorphic_detector;
0097   polymorphic_detector.set<detector_traits>(std::move(host_det));
0098
0099   const auto field = traccc::construct_const_bfield(std::get<13>(GetParam()));
0100
0101   /***************************
0102    * Generate simulation data
0103    ***************************/
0104
0105   // Track generator
0106   using generator_type =
0107       detray::random_track_generator<traccc::free_track_parameters<>,
0108                                      uniform_gen_t>;
0109   generator_type::configuration gen_cfg{};
0110   gen_cfg.n_tracks(n_truth_tracks);
0111   gen_cfg.origin(origin);
0112   gen_cfg.origin_stddev(origin_stddev);
0113   gen_cfg.phi_range(phi_range[0], phi_range[1]);
0114   gen_cfg.theta_range(theta_range[0], theta_range[1]);
0115   gen_cfg.mom_range(mom_range[0], mom_range[1]);
0116   gen_cfg.randomize_charge(random_charge);
0117   generator_type generator(gen_cfg);
0118
0119   // Smearing value for measurements
0120   traccc::measurement_smearer<traccc::default_algebra> meas_smearer(
0121       smearing[0], smearing[1]);
0122
0123   using writer_type = traccc::smearing_writer<
0124       traccc::measurement_smearer<traccc::default_algebra>>;
0125
0126   typename writer_type::config smearer_writer_cfg{meas_smearer};
0127   traccc::seed_generator<host_detector_type>::config seed_cfg{};
0128   seed_cfg.initial_sigmas = stddevs;
0129
0130   // Run simulator
0131   const std::string full_path = io::data_directory() + path;
0132   std::filesystem::create_directories(full_path);
0133   auto sim = traccc::simulator<host_detector_type, b_field_t, generator_type,
0134                                writer_type>(
0135       ptc, n_events, polymorphic_detector.as<detector_traits>(),
0136       field.as_field<traccc::const_bfield_backend_t<traccc::scalar>>(),
0137       std::move(generator), std::move(smearer_writer_cfg), full_path);
0138   sim.run();
0139
0140   /***************
0141    * Run fitting
0142    ***************/
0143
0144   vecmem::sycl::copy copy{vecmem_queue};
0145
0146   // Seed generator
0147   seed_generator<host_detector_type> sg(
0148       polymorphic_detector.as<detector_traits>(), seed_cfg);
0149
0150   // Fitting algorithm object
0151   typename traccc::sycl::kalman_fitting_algorithm::config_type fit_cfg;
0152   fit_cfg.ptc_hypothesis = ptc;
0153   fit_cfg.min_pT = 100.f * traccc::unit<float>::MeV;
0154   traccc::sycl::kalman_fitting_algorithm device_fitting(fit_cfg, mr, copy,
0155                                                         traccc_queue);
0156
0157   const traccc::detector_buffer detector_buffer =
0158       traccc::buffer_from_host_detector(polymorphic_detector, device_mr, copy);
0159
0160   // Iterate over events
0161   for (std::size_t i_evt = 0; i_evt < n_events; i_evt++) {
0162     // Event map
0163     traccc::event_data evt_data(path, i_evt, host_mr);
0164
0165     // Truth Track Candidates
0166     traccc::edm::measurement_collection::host measurements{host_mr};
0167     traccc::edm::track_container<traccc::default_algebra>::host
0168         track_candidates{host_mr};
0169     evt_data.generate_truth_candidates(track_candidates, measurements, sg,
0170                                        host_mr);
0171     track_candidates.measurements = vecmem::get_data(measurements);
0172
0173     // n_trakcs = 100
0174     ASSERT_EQ(track_candidates.tracks.size(), n_truth_tracks);
0175
0176     // track candidates buffer
0177     traccc::edm::measurement_collection::buffer measurements_buffer =
0178         copy.to(track_candidates.measurements, mr.main, mr.host,
0179                 vecmem::copy::type::host_to_device);
0180     traccc::edm::track_container<traccc::default_algebra>::buffer
0181         track_candidates_buffer{
0182             copy.to(vecmem::get_data(track_candidates.tracks), mr.main, mr.host,
0183                     vecmem::copy::type::host_to_device),
0184             {},
0185             measurements_buffer};
0186
0187     // Run fitting
0188     auto track_states_sycl_buffer =
0189         device_fitting(detector_buffer, field, track_candidates_buffer);
0190
0191     traccc::edm::track_container<traccc::default_algebra>::host
0192         track_states_sycl{host_mr};
0193     copy(track_states_sycl_buffer.tracks, track_states_sycl.tracks,
0194          vecmem::copy::type::device_to_host)
0195         ->wait();
0196     copy(track_states_sycl_buffer.states, track_states_sycl.states,
0197          vecmem::copy::type::device_to_host)
0198         ->wait();
0199
0200     const std::size_t n_fitted_tracks =
0201         count_successfully_fitted_tracks(track_states_sycl.tracks);
0202
0203     ASSERT_EQ(track_states_sycl.tracks.size(), n_truth_tracks);
0204     ASSERT_EQ(track_states_sycl.tracks.size(), n_fitted_tracks);
0205
0206     for (std::size_t i_trk = 0; i_trk < n_truth_tracks; i_trk++) {
0207       consistency_tests(track_states_sycl.tracks.at(i_trk),
0208                         track_states_sycl.states);
0209
0210       ndf_tests(track_states_sycl.tracks.at(i_trk), track_states_sycl.states,
0211                 measurements);
0212
0213       ASSERT_EQ(track_states_sycl.tracks.at(i_trk).nholes(), 0u);
0214
0215       fit_performance_writer.write(
0216           track_states_sycl.tracks.at(i_trk), track_states_sycl.states,
0217           measurements, polymorphic_detector.as<detector_traits>(), evt_data);
0218     }
0219   }
0220
0221   fit_performance_writer.finalize();
0222
0223   /********************
0224    * Pull value test
0225    ********************/
0226
0227   static const std::vector<std::string> pull_names{
0228       "pull_d0", "pull_z0", "pull_phi", "pull_theta", "pull_qop"};
0229   pull_value_tests(fit_writer_cfg.file_path, pull_names);
0230
0231   /********************
0232    * P-value test
0233    ********************/
0234
0235   p_value_tests(fit_writer_cfg.file_path);
0236
0237   /********************
0238    * Success rate test
0239    ********************/
0240
0241   float success_rate = static_cast<float>(n_success) /
0242                        static_cast<float>(n_truth_tracks * n_events);
0243
0244   ASSERT_FLOAT_EQ(success_rate, 1.00f);
0245 }
0246
0247 INSTANTIATE_TEST_SUITE_P(
0248     SYCLKalmanFitTelescopeValidation0, KalmanFittingTelescopeTests,
0249     ::testing::Values(std::make_tuple(
0250         "telescope_1_GeV_0_phi_muon", std::array<scalar, 3u>{0.f, 0.f, 0.f},
0251         std::array<scalar, 3u>{0.f, 0.f, 0.f}, std::array<scalar, 2u>{1.f, 1.f},
0252         std::array<scalar, 2u>{0.f, 0.f}, std::array<scalar, 2u>{0.f, 0.f},
0253         traccc::muon<scalar>(), 100, 100, false, 20.f, 20u, 20.f,
0254         vector3{0, 0, 2 * traccc::unit<scalar>::T})));
0255
0256 INSTANTIATE_TEST_SUITE_P(
0257     SYCLKalmanFitTelescopeValidation1, KalmanFittingTelescopeTests,
0258     ::testing::Values(std::make_tuple(
0259         "telescope_10_GeV_0_phi_muon", std::array<scalar, 3u>{0.f, 0.f, 0.f},
0260         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0261         std::array<scalar, 2u>{10.f, 10.f}, std::array<scalar, 2u>{0.f, 0.f},
0262         std::array<scalar, 2u>{0.f, 0.f}, traccc::muon<scalar>(), 100, 100,
0263         false, 20.f, 9u, 20.f, vector3{0, 0, 2 * traccc::unit<scalar>::T})));
0264
0265 INSTANTIATE_TEST_SUITE_P(
0266     SYCLKalmanFitTelescopeValidation2, KalmanFittingTelescopeTests,
0267     ::testing::Values(std::make_tuple(
0268         "telescope_100_GeV_0_phi_muon", std::array<scalar, 3u>{0.f, 0.f, 0.f},
0269         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0270         std::array<scalar, 2u>{100.f, 100.f}, std::array<scalar, 2u>{0.f, 0.f},
0271         std::array<scalar, 2u>{0.f, 0.f}, traccc::muon<scalar>(), 100, 100,
0272         false, 20.f, 9u, 20.f, vector3{0, 0, 2 * traccc::unit<scalar>::T})));
0273
0274 INSTANTIATE_TEST_SUITE_P(
0275     SYCLKalmanFitTelescopeValidation3, KalmanFittingTelescopeTests,
0276     ::testing::Values(std::make_tuple(
0277         "telescope_1_GeV_0_phi_anti_muon",
0278         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0279         std::array<scalar, 3u>{0.f, 0.f, 0.f}, std::array<scalar, 2u>{1.f, 1.f},
0280         std::array<scalar, 2u>{0.f, 0.f}, std::array<scalar, 2u>{0.f, 0.f},
0281         traccc::antimuon<scalar>(), 100, 100, false, 20.f, 9u, 20.f,
0282         vector3{0, 0, 2 * traccc::unit<scalar>::T})));
0283
0284 INSTANTIATE_TEST_SUITE_P(
0285     SYCLKalmanFitTelescopeValidation4, KalmanFittingTelescopeTests,
0286     ::testing::Values(std::make_tuple(
0287         "telescope_1_GeV_0_random_charge",
0288         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0289         std::array<scalar, 3u>{0.f, 0.f, 0.f}, std::array<scalar, 2u>{1.f, 1.f},
0290         std::array<scalar, 2u>{0.f, 0.f}, std::array<scalar, 2u>{0.f, 0.f},
0291         traccc::antimuon<scalar>(), 100, 100, true, 20.f, 9u, 20.f,
0292         vector3{0, 0, 2 * traccc::unit<scalar>::T})));