tests/cuda/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 // Project include(s).
0009 #include "traccc/bfield/construct_const_bfield.hpp"
0010 #include "traccc/bfield/magnetic_field_types.hpp"
0011 #include "traccc/cuda/finding/combinatorial_kalman_filter_algorithm.hpp"
0012 #include "traccc/io/read_detector.hpp"
0013 #include "traccc/io/read_measurements.hpp"
0014 #include "traccc/io/utils.hpp"
0015 #include "traccc/simulation/event_generators.hpp"
0016 #include "traccc/simulation/simulator.hpp"
0017 #include "traccc/utils/event_data.hpp"
0018 #include "traccc/utils/ranges.hpp"
0019
0020 // Test include(s).
0021 #include "tests/ckf_telescope_test.hpp"
0022 #include "traccc/utils/seed_generator.hpp"
0023
0024 // VecMem include(s).
0025 #include <vecmem/memory/cuda/device_memory_resource.hpp>
0026 #include <vecmem/memory/cuda/managed_memory_resource.hpp>
0027 #include <vecmem/memory/host_memory_resource.hpp>
0028 #include <vecmem/utils/cuda/async_copy.hpp>
0029 #include <vecmem/utils/cuda/stream_wrapper.hpp>
0030
0031 // GTest include(s).
0032 #include <gtest/gtest.h>
0033
0034 // System include(s).
0035 #include <filesystem>
0036 #include <string>
0037
0038 using namespace traccc;
0039 // This defines the local frame test suite
0040 TEST_P(CudaCkfCombinatoricsTelescopeTests, Run) {
0041   // Get the parameters
0042   const std::string name = std::get<0>(GetParam());
0043   const std::array<scalar, 3u> origin = std::get<1>(GetParam());
0044   const std::array<scalar, 3u> origin_stddev = std::get<2>(GetParam());
0045   const std::array<scalar, 2u> mom_range = std::get<3>(GetParam());
0046   const std::array<scalar, 2u> eta_range = std::get<4>(GetParam());
0047   const std::array<scalar, 2u> theta_range = eta_to_theta_range(eta_range);
0048   const std::array<scalar, 2u> phi_range = std::get<5>(GetParam());
0049   const traccc::pdg_particle<scalar> ptc = std::get<6>(GetParam());
0050   const unsigned int n_truth_tracks = std::get<7>(GetParam());
0051   const unsigned int n_events = std::get<8>(GetParam());
0052   const bool random_charge = std::get<9>(GetParam());
0053
0054   /*****************************
0055    * Build a telescope geometry
0056    *****************************/
0057
0058   // Memory resources used by the application.
0059   vecmem::host_memory_resource host_mr;
0060   vecmem::cuda::device_memory_resource device_mr;
0061   traccc::memory_resource mr{device_mr, &host_mr};
0062   vecmem::cuda::managed_memory_resource mng_mr;
0063   vecmem::copy host_copy;
0064
0065   // Read back detector file
0066   const std::string path = name + "/";
0067   traccc::host_detector detector;
0068   traccc::io::read_detector(
0069       detector, mng_mr,
0070       std::filesystem::absolute(
0071           std::filesystem::path(path + "telescope_detector_geometry.json"))
0072           .native(),
0073       std::filesystem::absolute(
0074           std::filesystem::path(path +
0075                                 "telescope_detector_homogeneous_material.json"))
0076           .native());
0077
0078   const traccc::detector_buffer detector_buffer =
0079       traccc::buffer_from_host_detector(detector, mng_mr, host_copy);
0080
0081   const auto field = traccc::construct_const_bfield(std::get<13>(GetParam()));
0082
0083   /***************************
0084    * Generate simulation data
0085    ***************************/
0086
0087   // Track generator
0088   using generator_type =
0089       detray::random_track_generator<traccc::free_track_parameters<>,
0090                                      uniform_gen_t>;
0091   generator_type::configuration gen_cfg{};
0092   gen_cfg.n_tracks(n_truth_tracks);
0093   gen_cfg.origin(origin);
0094   gen_cfg.origin_stddev(origin_stddev);
0095   gen_cfg.phi_range(phi_range[0], phi_range[1]);
0096   gen_cfg.theta_range(theta_range[0], theta_range[1]);
0097   gen_cfg.mom_range(mom_range[0], mom_range[1]);
0098   gen_cfg.randomize_charge(random_charge);
0099   generator_type generator(gen_cfg);
0100
0101   // Smearing value for measurements
0102   traccc::measurement_smearer<traccc::default_algebra> meas_smearer(
0103       smearing[0], smearing[1]);
0104
0105   using writer_type = traccc::smearing_writer<
0106       traccc::measurement_smearer<traccc::default_algebra>>;
0107
0108   typename writer_type::config smearer_writer_cfg{meas_smearer};
0109   traccc::seed_generator<host_detector_type>::config seed_cfg{};
0110   seed_cfg.initial_sigmas = stddevs;
0111
0112   // Run simulator
0113   const std::string full_path = io::data_directory() + path;
0114   std::filesystem::create_directories(full_path);
0115   auto sim = traccc::simulator<host_detector_type, b_field_t, generator_type,
0116                                writer_type>(
0117       ptc, n_events, detector.as<detector_traits>(),
0118       field.as_field<traccc::const_bfield_backend_t<traccc::scalar>>(),
0119       std::move(generator), std::move(smearer_writer_cfg), full_path);
0120   sim.run();
0121
0122   /*****************************
0123    * Do the reconstruction
0124    *****************************/
0125
0126   // Stream object
0127   vecmem::cuda::stream_wrapper vecmem_stream;
0128   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0129
0130   // Copy objects
0131   vecmem::cuda::async_copy copy{stream.cudaStream()};
0132
0133   // Seed generator
0134   seed_generator<host_detector_type> sg(detector.as<detector_traits>(),
0135                                         seed_cfg);
0136
0137   // Finding algorithm configuration
0138   typename traccc::cuda::combinatorial_kalman_filter_algorithm::config_type
0139       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
0146   typename traccc::cuda::combinatorial_kalman_filter_algorithm::config_type
0147       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
0154   // Finding algorithm object
0155   traccc::cuda::combinatorial_kalman_filter_algorithm device_finding(
0156       cfg_no_limit, mr, copy, stream);
0157   traccc::cuda::combinatorial_kalman_filter_algorithm device_finding_limit(
0158       cfg_limit, mr, copy, stream);
0159
0160   // Iterate over events
0161   for (std::size_t i_evt = 0; i_evt < n_events; i_evt++) {
0162     // Truth Track Candidates
0163     traccc::event_data evt_data(path, i_evt, host_mr);
0164
0165     traccc::edm::measurement_collection::host truth_measurements{host_mr};
0166     traccc::edm::track_container<traccc::default_algebra>::host
0167         truth_track_candidates{host_mr};
0168     evt_data.generate_truth_candidates(truth_track_candidates,
0169                                        truth_measurements, sg, host_mr);
0170     truth_track_candidates.measurements = vecmem::get_data(truth_measurements);
0171
0172     ASSERT_EQ(truth_track_candidates.tracks.size(), n_truth_tracks);
0173
0174     // Prepare truth seeds
0175     traccc::bound_track_parameters_collection_types::host seeds(&host_mr);
0176     for (unsigned int i_trk = 0; i_trk < n_truth_tracks; i_trk++) {
0177       seeds.push_back(truth_track_candidates.tracks.at(i_trk).params());
0178     }
0179     ASSERT_EQ(seeds.size(), n_truth_tracks);
0180
0181     traccc::bound_track_parameters_collection_types::buffer seeds_buffer{
0182         static_cast<unsigned int>(seeds.size()), mr.main};
0183     copy.setup(seeds_buffer)->wait();
0184     copy(vecmem::get_data(seeds), seeds_buffer,
0185          vecmem::copy::type::host_to_device)
0186         ->wait();
0187
0188     // Read measurements
0189     traccc::edm::measurement_collection::host measurements_per_event{host_mr};
0190     traccc::io::read_measurements(measurements_per_event, i_evt, path);
0191
0192     traccc::edm::measurement_collection::buffer measurements_buffer(
0193         static_cast<unsigned int>(measurements_per_event.size()), mr.main);
0194     copy.setup(measurements_buffer)->wait();
0195     copy(vecmem::get_data(measurements_per_event), measurements_buffer)->wait();
0196
0197     // Run device finding
0198     traccc::edm::track_container<traccc::default_algebra>::buffer
0199         track_candidates_cuda_buffer = device_finding(
0200             detector_buffer, field, measurements_buffer, seeds_buffer);
0201
0202     // Run device finding (Limit)
0203     traccc::edm::track_container<traccc::default_algebra>::buffer
0204         track_candidates_limit_cuda_buffer = device_finding_limit(
0205             detector_buffer, field, measurements_buffer, seeds_buffer);
0206
0207     traccc::edm::track_collection<traccc::default_algebra>::host
0208         track_candidates_cuda{host_mr},
0209         track_candidates_limit_cuda{host_mr};
0210     copy(track_candidates_cuda_buffer.tracks, track_candidates_cuda,
0211          vecmem::copy::type::device_to_host)
0212         ->wait();
0213     copy(track_candidates_limit_cuda_buffer.tracks, track_candidates_limit_cuda,
0214          vecmem::copy::type::device_to_host)
0215         ->wait();
0216
0217     // Make sure that the number of found tracks = n_track ^ (n_planes +
0218     // 1)
0219     ASSERT_GT(track_candidates_cuda.size(), track_candidates_limit_cuda.size());
0220     ASSERT_EQ(track_candidates_cuda.size(),
0221               std::pow(n_truth_tracks, std::get<11>(GetParam()) + 1));
0222     ASSERT_EQ(track_candidates_limit_cuda.size(),
0223               n_truth_tracks * cfg_limit.max_num_branches_per_seed);
0224   }
0225 }
0226
0227 // Testing two identical tracks
0228 INSTANTIATE_TEST_SUITE_P(
0229     CUDACkfCombinatoricsTelescopeValidation, CudaCkfCombinatoricsTelescopeTests,
0230     ::testing::Values(
0231         std::make_tuple("telescope_combinatorics_twin",
0232                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0233                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0234                         std::array<scalar, 2u>{100.f, 100.f},
0235                         std::array<scalar, 2u>{0.f, 0.f},
0236                         std::array<scalar, 2u>{0.f, 0.f},
0237                         traccc::muon<scalar>(), 2, 1, false, 20.f, 9u, 20.f,
0238                         vector3{2 * traccc::unit<scalar>::T, 0, 0}),
0239         std::make_tuple("telescope_combinatorics_trio",
0240                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0241                         std::array<scalar, 3u>{0.f, 0.f, 0.f},
0242                         std::array<scalar, 2u>{100.f, 100.f},
0243                         std::array<scalar, 2u>{0.f, 0.f},
0244                         std::array<scalar, 2u>{0.f, 0.f},
0245                         traccc::muon<scalar>(), 3, 1, false, 20.f, 9u, 20.f,
0246                         vector3{2 * traccc::unit<scalar>::T, 0, 0})));