cuda/apps/seeding_example_cuda.cpp

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2021-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/cuda/finding/combinatorial_kalman_filter_algorithm.hpp"
0010 #include "traccc/cuda/fitting/kalman_fitting_algorithm.hpp"
0011 #include "traccc/cuda/seeding/seed_parameter_estimation_algorithm.hpp"
0012 #include "traccc/cuda/seeding/triplet_seeding_algorithm.hpp"
0013 #include "traccc/cuda/utils/make_magnetic_field.hpp"
0014 #include "traccc/definitions/common.hpp"
0015 #include "traccc/device/container_d2h_copy_alg.hpp"
0016 #include "traccc/device/container_h2d_copy_alg.hpp"
0017 #include "traccc/efficiency/finding_performance_writer.hpp"
0018 #include "traccc/efficiency/nseed_performance_writer.hpp"
0019 #include "traccc/efficiency/seeding_performance_writer.hpp"
0020 #include "traccc/efficiency/track_filter.hpp"
0021 #include "traccc/examples/make_magnetic_field.hpp"
0022 #include "traccc/examples/print_fitted_tracks_statistics.hpp"
0023 #include "traccc/finding/combinatorial_kalman_filter_algorithm.hpp"
0024 #include "traccc/fitting/kalman_fitting_algorithm.hpp"
0025 #include "traccc/geometry/detector.hpp"
0026 #include "traccc/geometry/host_detector.hpp"
0027 #include "traccc/io/read_detector.hpp"
0028 #include "traccc/io/read_detector_description.hpp"
0029 #include "traccc/io/read_measurements.hpp"
0030 #include "traccc/io/read_spacepoints.hpp"
0031 #include "traccc/io/utils.hpp"
0032 #include "traccc/options/accelerator.hpp"
0033 #include "traccc/options/detector.hpp"
0034 #include "traccc/options/input_data.hpp"
0035 #include "traccc/options/magnetic_field.hpp"
0036 #include "traccc/options/performance.hpp"
0037 #include "traccc/options/program_options.hpp"
0038 #include "traccc/options/seed_matching.hpp"
0039 #include "traccc/options/track_finding.hpp"
0040 #include "traccc/options/track_fitting.hpp"
0041 #include "traccc/options/track_matching.hpp"
0042 #include "traccc/options/track_propagation.hpp"
0043 #include "traccc/options/track_seeding.hpp"
0044 #include "traccc/options/truth_finding.hpp"
0045 #include "traccc/performance/collection_comparator.hpp"
0046 #include "traccc/performance/soa_comparator.hpp"
0047 #include "traccc/performance/timer.hpp"
0048 #include "traccc/resolution/fitting_performance_writer.hpp"
0049 #include "traccc/seeding/detail/track_params_estimation_config.hpp"
0050 #include "traccc/seeding/seeding_algorithm.hpp"
0051 #include "traccc/seeding/track_params_estimation.hpp"
0052 #include "traccc/utils/propagation.hpp"
0053
0054 // VecMem include(s).
0055 #include <vecmem/memory/cuda/device_memory_resource.hpp>
0056 #include <vecmem/memory/cuda/host_memory_resource.hpp>
0057 #include <vecmem/memory/cuda/managed_memory_resource.hpp>
0058 #include <vecmem/memory/host_memory_resource.hpp>
0059 #include <vecmem/utils/cuda/async_copy.hpp>
0060 #include <vecmem/utils/cuda/copy.hpp>
0061
0062 #include "traccc/cuda/gbts_seeding/gbts_seeding_algorithm.hpp"
0063 #include "traccc/gbts_seeding/gbts_seeding_config.hpp"
0064 #include "traccc/options/track_gbts_seeding.hpp"
0065
0066 // System include(s).
0067 #include <exception>
0068 #include <iomanip>
0069 #include <iostream>
0070
0071 using namespace traccc;
0072
0073 int seq_run(const traccc::opts::track_seeding& seeding_opts,
0074             const traccc::opts::track_gbts_seeding& seeding_gbts_opts,
0075             const traccc::opts::track_finding& finding_opts,
0076             const traccc::opts::track_propagation& propagation_opts,
0077             [[maybe_unused]] const traccc::opts::track_fitting& fitting_opts,
0078             const traccc::opts::input_data& input_opts,
0079             const traccc::opts::detector& detector_opts,
0080             const traccc::opts::magnetic_field& bfield_opts,
0081             const traccc::opts::performance& performance_opts,
0082             const traccc::opts::accelerator& accelerator_opts,
0083             const traccc::opts::truth_finding& truth_finding_opts,
0084             const traccc::opts::seed_matching& seed_matching_opts,
0085             const traccc::opts::track_matching& track_matching_opts,
0086             std::unique_ptr<const traccc::Logger> ilogger, bool usingGBTS) {
0087   TRACCC_LOCAL_LOGGER(std::move(ilogger));
0088
0089   // Memory resources used by the application.
0090   vecmem::host_memory_resource host_mr;
0091   vecmem::cuda::host_memory_resource cuda_host_mr;
0092   vecmem::cuda::managed_memory_resource mng_mr;
0093   vecmem::cuda::device_memory_resource device_mr;
0094   traccc::memory_resource mr{device_mr, &cuda_host_mr};
0095
0096   // Performance writer
0097   traccc::seeding_performance_writer sd_performance_writer(
0098       traccc::seeding_performance_writer::config{
0099           .truth_config = truth_finding_opts,
0100           .seed_truth_config = seed_matching_opts},
0101       logger().clone("SeedingPerformanceWriter"));
0102   traccc::finding_performance_writer find_performance_writer(
0103       traccc::finding_performance_writer::config{
0104           .truth_config = truth_finding_opts,
0105           .track_truth_config = track_matching_opts,
0106           .require_fit = (traccc::finding_config(finding_opts).run_smoother !=
0107                           smoother_type::e_none)},
0108       logger().clone("FindingPerformanceWriter"));
0109   traccc::fitting_performance_writer fit_performance_writer(
0110       traccc::fitting_performance_writer::config{},
0111       logger().clone("FittingPerformanceWriter"));
0112
0113   traccc::nseed_performance_writer nsd_performance_writer(
0114       "nseed_performance_",
0115       std::make_unique<traccc::simple_charged_eta_pt_cut>(
0116           2.7f, 1.f * traccc::unit<traccc::scalar>::GeV),
0117       std::make_unique<traccc::stepped_percentage>(0.6f));
0118
0119   if (performance_opts.run) {
0120     nsd_performance_writer.initialize();
0121   }
0122
0123   // Output stats
0124   uint64_t n_spacepoints = 0;
0125   uint64_t n_seeds = 0;
0126   uint64_t n_seeds_cuda = 0;
0127   uint64_t n_found_tracks = 0;
0128   uint64_t n_found_tracks_cuda = 0;
0129   uint64_t n_fitted_tracks = 0;
0130   uint64_t n_fitted_tracks_cuda = 0;
0131
0132   /*****************************
0133    * Build a geometry
0134    *****************************/
0135
0136   traccc::detector_design_description::host host_det_descr{host_mr};
0137   traccc::detector_conditions_description::host host_det_cond{host_mr};
0138   traccc::io::read_detector_description(
0139       host_det_descr, host_det_cond, detector_opts.detector_file,
0140       detector_opts.digitization_file, detector_opts.conditions_file,
0141       traccc::data_format::json);
0142
0143   // B field value
0144   const traccc::vector3 field_vec(seeding_opts);
0145   const auto host_field = traccc::details::make_magnetic_field(bfield_opts);
0146   const auto device_field = traccc::cuda::make_magnetic_field(
0147       host_field, (accelerator_opts.use_gpu_texture_memory
0148                        ? traccc::cuda::magnetic_field_storage::texture_memory
0149                        : traccc::cuda::magnetic_field_storage::global_memory));
0150
0151   // Construct a Detray detector object, if supported by the configuration.
0152   traccc::host_detector host_det;
0153   traccc::io::read_detector(host_det, mng_mr, detector_opts.detector_file,
0154                             detector_opts.material_file,
0155                             detector_opts.grid_file);
0156
0157   // Copy objects
0158   vecmem::copy host_copy;
0159   vecmem::cuda::copy copy;
0160
0161   const traccc::detector_buffer detector_buffer =
0162       traccc::buffer_from_host_detector(host_det, mng_mr, host_copy);
0163
0164   // GBTS seeding configuration
0165   const traccc::gbts_seedfinder_config gbts_config(seeding_gbts_opts);
0166   // Seeding algorithm
0167   const traccc::seedfinder_config seedfinder_config(seeding_opts);
0168   const traccc::seedfilter_config seedfilter_config(seeding_opts);
0169   const traccc::spacepoint_grid_config spacepoint_grid_config(seeding_opts);
0170   traccc::host::seeding_algorithm sa(seedfinder_config, spacepoint_grid_config,
0171                                      seedfilter_config, host_mr,
0172                                      logger().clone("HostSeedingAlg"));
0173   const traccc::track_params_estimation_config track_params_estimation_config;
0174   traccc::host::track_params_estimation tp(
0175       track_params_estimation_config, host_mr,
0176       logger().clone("HostTrackParEstAlg"));
0177
0178   vecmem::cuda::stream_wrapper vecmem_stream;
0179   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0180
0181   vecmem::cuda::async_copy async_copy{stream.cudaStream()};
0182
0183   traccc::cuda::triplet_seeding_algorithm sa_cuda{
0184       seedfinder_config,
0185       spacepoint_grid_config,
0186       seedfilter_config,
0187       mr,
0188       async_copy,
0189       stream,
0190       logger().clone("CudaSeedingAlg")};
0191   traccc::cuda::gbts_seeding_algorithm gbts_sa_cuda(
0192       gbts_config, mr, copy, stream, logger().clone("CudaGbtsSeedingAlg"));
0193   traccc::cuda::seed_parameter_estimation_algorithm tp_cuda{
0194       track_params_estimation_config, mr, async_copy, stream,
0195       logger().clone("CudaTrackParEstAlg")};
0196
0197   // Propagation configuration
0198   detray::propagation::config propagation_config(propagation_opts);
0199
0200   // Finding algorithm configuration
0201   traccc::finding_config cfg(finding_opts);
0202   cfg.propagation = propagation_config;
0203
0204   // Finding algorithm object
0205   traccc::host::combinatorial_kalman_filter_algorithm host_finding(
0206       cfg, host_mr, logger().clone("HostFindingAlg"));
0207
0208   auto device_fitting =
0209       std::make_unique<traccc::cuda::kalman_fitting_algorithm>(
0210           cfg.kalman_smoother, mr, async_copy, stream,
0211           logger().clone("CudaFittingAlg"));
0212   traccc::cuda::combinatorial_kalman_filter_algorithm device_finding(
0213       cfg, mr, async_copy, stream, logger().clone("CudaFindingAlg"),
0214       std::move(device_fitting));
0215
0216   traccc::performance::timing_info elapsedTimes;
0217
0218   // Loop over events
0219   for (std::size_t event = input_opts.skip;
0220        event < input_opts.events + input_opts.skip; ++event) {
0221     // Instantiate host containers/collections
0222     traccc::edm::spacepoint_collection::host spacepoints_per_event{host_mr};
0223     traccc::edm::measurement_collection::host measurements_per_event{host_mr};
0224     traccc::host::seeding_algorithm::output_type seeds{host_mr};
0225     traccc::host::track_params_estimation::output_type params;
0226     traccc::edm::track_container<traccc::default_algebra>::host
0227         track_candidates{host_mr};
0228
0229     traccc::edm::seed_collection::buffer seeds_cuda_buffer;
0230     traccc::bound_track_parameters_collection_types::buffer params_cuda_buffer(
0231         0, *mr.host);
0232
0233     traccc::edm::track_container<traccc::default_algebra>::buffer
0234         track_candidates_cuda_buffer;
0235
0236     {  // Start measuring wall time
0237       traccc::performance::timer wall_t("Wall time", elapsedTimes);
0238
0239       /*-----------------
0240       hit file reading
0241       -----------------*/
0242       {
0243         traccc::performance::timer t("Hit reading  (cpu)", elapsedTimes);
0244         // Read the hits and measurements from the relevant event files
0245         traccc::io::read_spacepoints(
0246             spacepoints_per_event, measurements_per_event, event,
0247             input_opts.directory,
0248             (input_opts.use_acts_geom_source ? &host_det : nullptr),
0249             &host_det_descr, &host_det_cond, input_opts.format);
0250
0251       }  // stop measuring hit reading timer
0252
0253       /*----------------------------
0254           Seeding algorithm
0255       ----------------------------*/
0256
0257       /// CUDA
0258
0259       // Copy the spacepoint and module data to the device.
0260       traccc::edm::spacepoint_collection::buffer spacepoints_cuda_buffer(
0261           static_cast<unsigned int>(spacepoints_per_event.size()), mr.main);
0262       async_copy.setup(spacepoints_cuda_buffer)->wait();
0263       async_copy(vecmem::get_data(spacepoints_per_event),
0264                  spacepoints_cuda_buffer)
0265           ->wait();
0266
0267       traccc::edm::measurement_collection::buffer measurements_cuda_buffer(
0268           static_cast<unsigned int>(measurements_per_event.size()), mr.main);
0269       async_copy.setup(measurements_cuda_buffer)->wait();
0270       async_copy(vecmem::get_data(measurements_per_event),
0271                  measurements_cuda_buffer)
0272           ->wait();
0273
0274       {
0275         traccc::performance::timer t("Seeding (cuda)", elapsedTimes);
0276         // Reconstruct the spacepoints into seeds.
0277         if (usingGBTS) {
0278           seeds_cuda_buffer =
0279               gbts_sa_cuda(spacepoints_cuda_buffer, measurements_cuda_buffer);
0280         } else {
0281           seeds_cuda_buffer = sa_cuda(spacepoints_cuda_buffer);
0282         }
0283         stream.synchronize();
0284       }  // stop measuring seeding cuda timer
0285
0286       // CPU
0287
0288       if (accelerator_opts.compare_with_cpu) {
0289         {
0290           traccc::performance::timer t("Seeding  (cpu)", elapsedTimes);
0291           seeds = sa(vecmem::get_data(spacepoints_per_event));
0292         }
0293       }  // stop measuring seeding cpu timer
0294
0295       /*----------------------------
0296          Track params estimation
0297       ----------------------------*/
0298
0299       // CUDA
0300       {
0301         traccc::performance::timer t("Track params (cuda)", elapsedTimes);
0302         params_cuda_buffer =
0303             tp_cuda(device_field, measurements_cuda_buffer,
0304                     spacepoints_cuda_buffer, seeds_cuda_buffer);
0305         stream.synchronize();
0306       }  // stop measuring track params cuda timer
0307
0308       // CPU
0309       if (accelerator_opts.compare_with_cpu) {
0310         traccc::performance::timer t("Track params  (cpu)", elapsedTimes);
0311         params = tp(vecmem::get_data(measurements_per_event),
0312                     vecmem::get_data(spacepoints_per_event),
0313                     vecmem::get_data(seeds), field_vec);
0314       }  // stop measuring track params cpu timer
0315
0316       /*------------------------
0317          Track Finding with CKF
0318         ------------------------*/
0319
0320       {
0321         traccc::performance::timer t("Track finding with CKF (cuda)",
0322                                      elapsedTimes);
0323         track_candidates_cuda_buffer =
0324             device_finding(detector_buffer, device_field,
0325                            measurements_cuda_buffer, params_cuda_buffer);
0326       }
0327
0328       if (accelerator_opts.compare_with_cpu) {
0329         traccc::performance::timer t("Track finding with CKF (cpu)",
0330                                      elapsedTimes);
0331         track_candidates = host_finding(
0332             host_det, host_field, vecmem::get_data(measurements_per_event),
0333             vecmem::get_data(params));
0334       }
0335     }  // Stop measuring wall time
0336
0337     /*----------------------------------
0338       compare seeds from cpu and cuda
0339       ----------------------------------*/
0340
0341     // Copy the seeds to the host for comparisons
0342     traccc::edm::seed_collection::host seeds_cuda{host_mr};
0343     traccc::bound_track_parameters_collection_types::host params_cuda;
0344     async_copy(seeds_cuda_buffer, seeds_cuda)->wait();
0345     async_copy(params_cuda_buffer, params_cuda)->wait();
0346
0347     // Copy track candidates from device to host
0348     traccc::edm::track_container<traccc::default_algebra>::host
0349         track_candidates_cuda{host_mr,
0350                               vecmem::get_data(measurements_per_event)};
0351     async_copy(track_candidates_cuda_buffer.tracks,
0352                track_candidates_cuda.tracks)
0353         ->wait();
0354     async_copy(track_candidates_cuda_buffer.states,
0355                track_candidates_cuda.states)
0356         ->wait();
0357
0358     if (accelerator_opts.compare_with_cpu) {
0359       // Show which event we are currently presenting the results for.
0360       std::cout << "===>>> Event " << event << " <<<===" << std::endl;
0361
0362       // Compare the seeds made on the host and on the device
0363       traccc::soa_comparator<traccc::edm::seed_collection> compare_seeds{
0364           "seeds",
0365           traccc::details::comparator_factory<
0366               traccc::edm::seed_collection::const_device::const_proxy_type>{
0367               vecmem::get_data(spacepoints_per_event),
0368               vecmem::get_data(spacepoints_per_event)}};
0369       compare_seeds(vecmem::get_data(seeds), vecmem::get_data(seeds_cuda));
0370
0371       // Compare the track parameters made on the host and on the device.
0372       traccc::collection_comparator<traccc::bound_track_parameters<>>
0373           compare_track_parameters{"track parameters"};
0374       compare_track_parameters(vecmem::get_data(params),
0375                                vecmem::get_data(params_cuda));
0376     }
0377
0378     /*----------------
0379          Statistics
0380       ---------------*/
0381
0382     details::print_fitted_tracks_statistics(track_candidates_cuda, logger());
0383     n_spacepoints += spacepoints_per_event.size();
0384     n_seeds_cuda += seeds_cuda.size();
0385     n_seeds += seeds.size();
0386     n_found_tracks_cuda += track_candidates_cuda.tracks.size();
0387     n_found_tracks += track_candidates.tracks.size();
0388     n_fitted_tracks_cuda += track_candidates_cuda.tracks.size();
0389     n_fitted_tracks += track_candidates.tracks.size();
0390
0391     /*------------
0392       Writer
0393       ------------*/
0394
0395     if (performance_opts.run) {
0396       traccc::event_data evt_data(input_opts.directory, event, host_mr,
0397                                   input_opts.use_acts_geom_source, &host_det,
0398                                   input_opts.format, false);
0399
0400       sd_performance_writer.write(
0401           vecmem::get_data(seeds_cuda), vecmem::get_data(spacepoints_per_event),
0402           vecmem::get_data(measurements_per_event), evt_data);
0403
0404       std::cout << track_candidates_cuda.tracks.size() << ", "
0405                 << track_candidates_cuda.states.size() << std::endl;
0406
0407       find_performance_writer.write(
0408           traccc::edm::track_container<traccc::default_algebra>::const_data(
0409               track_candidates_cuda),
0410           evt_data);
0411
0412       for (unsigned int i = 0; i < track_candidates_cuda.tracks.size(); i++) {
0413         host_detector_visitor<detector_type_list>(
0414             host_det, [&]<typename detector_traits_t>(
0415                           const typename detector_traits_t::host& det) {
0416               fit_performance_writer.write(track_candidates_cuda.tracks.at(i),
0417                                            track_candidates_cuda.states,
0418                                            measurements_per_event, det,
0419                                            evt_data);
0420             });
0421       }
0422     }
0423   }
0424
0425   if (performance_opts.run) {
0426     sd_performance_writer.finalize();
0427     nsd_performance_writer.finalize();
0428     find_performance_writer.finalize();
0429     fit_performance_writer.finalize();
0430     std::cout << nsd_performance_writer.generate_report_str();
0431   }
0432
0433   std::cout << "==> Statistics ... " << std::endl;
0434   std::cout << "- read    " << n_spacepoints << " spacepoints" << std::endl;
0435   std::cout << "- created  (cpu)  " << n_seeds << " seeds" << std::endl;
0436   std::cout << "- created (cuda)  " << n_seeds_cuda << " seeds" << std::endl;
0437   std::cout << "- created  (cpu) " << n_found_tracks << " found tracks"
0438             << std::endl;
0439   std::cout << "- created (cuda) " << n_found_tracks_cuda << " found tracks"
0440             << std::endl;
0441   std::cout << "- created  (cpu) " << n_fitted_tracks << " fitted tracks"
0442             << std::endl;
0443   std::cout << "- created (cuda) " << n_fitted_tracks_cuda << " fitted tracks"
0444             << std::endl;
0445   std::cout << "==>Elapsed times...\n" << elapsedTimes << std::endl;
0446
0447   return 0;
0448 }
0449
0450 // The main routine
0451 //
0452 int main(int argc, char* argv[]) {
0453   std::unique_ptr<const traccc::Logger> logger = traccc::getDefaultLogger(
0454       "TracccExampleSeedingCuda", traccc::Logging::Level::INFO);
0455
0456   // Program options.
0457   traccc::opts::detector detector_opts;
0458   traccc::opts::magnetic_field bfield_opts;
0459   traccc::opts::input_data input_opts;
0460   traccc::opts::track_seeding seeding_opts;
0461   traccc::opts::track_gbts_seeding seeding_gbts_opts;
0462   traccc::opts::track_finding finding_opts;
0463   traccc::opts::track_propagation propagation_opts;
0464   traccc::opts::track_fitting fitting_opts;
0465   traccc::opts::performance performance_opts;
0466   traccc::opts::accelerator accelerator_opts;
0467   traccc::opts::truth_finding truth_finding_opts;
0468   traccc::opts::seed_matching seed_matching_opts;
0469   traccc::opts::track_matching track_matching_opts;
0470   traccc::opts::program_options program_opts{
0471       "Full Tracking Chain Using CUDA (without clusterization)",
0472       {detector_opts, bfield_opts, input_opts, seeding_opts, seeding_gbts_opts,
0473        finding_opts, propagation_opts, fitting_opts, performance_opts,
0474        accelerator_opts, truth_finding_opts, seed_matching_opts,
0475        track_matching_opts},
0476       argc,
0477       argv,
0478       logger->cloneWithSuffix("Options")};
0479
0480   // Run the application.
0481   return seq_run(seeding_opts, seeding_gbts_opts, finding_opts,
0482                  propagation_opts, fitting_opts, input_opts, detector_opts,
0483                  bfield_opts, performance_opts, accelerator_opts,
0484                  truth_finding_opts, seed_matching_opts, track_matching_opts,
0485                  logger->clone(), seeding_gbts_opts.useGBTS);
0486 }