Add drift chamber tests and exmaples

README.md

@@ -249,6 +249,22 @@ cmake --build <build_directory> <options>
 <build_directory>/bin/traccc_throughput_mt_cuda --detector_file=tml_detector/trackml-detector.csv --digitization_config_file=tml_detector/default-geometric-config-generic.json --input_directory=tml_pixels/  --cold_run_events=100 --processed_events=1000 --threads=1
 ```
 
+```sh
+# Toy geometry examples
+<build_directory>/bin/traccc_simulate_toy_detector --gen-vertex-xyz-mm=0:0:0 --gen-vertex-xyz-std-mm=0:0:0 --gen-mom-gev=100:100 --gen-phi-degree=0:360 --events=10 --gen-nparticles=2000 --output_directory=detray_simulation/toy_detector/n_particles_2000/ --gen-eta=-3:3 --constraint-step-size-mm=1
+
+<build_directory>/bin/traccc_seeding_example_cuda --input_directory=detray_simulation/toy_detector/n_particles_2000/ --check_performance=true --detector_file=<detector_directory>/toy_detector_geometry.json --material_file=<detector_directory>/toy_detector_homogeneous_material.json --grid_file=<detector_directory>/toy_detector_surface_grids.json --event=1 --track_candidates_range=3:10 --constraint-step-size-mm=1000 --run_cpu=1
+
+<build_directory>/bin/traccc_truth_finding_example_cuda --input_directory=detray_simulation/toy_detector/n_particles_2000/ --check_performance=true --detector_file=<detector_directory>/toy_detector_geometry.json --material_file=<detector_directory>/toy_detector_homogeneous_material.json --grid_file=<detector_directory>/toy_detector_surface_grids.json --event=1 --track_candidates_range=3:10 --constraint-step-size-mm=1 --run_cpu=1
+```
+
+```sh
+# Wire chamber examples
+<build_directory>/bin/traccc_simulate_wire_chamber --gen-vertex-xyz-mm=0:0:0 --gen-vertex-xyz-std-mm=0:0:0 --gen-mom-gev=2:2 --gen-phi-degree=0:360 --events=10 --gen-nparticles=1000 --output_directory=detray_simulation/wire_chamber/n_particles_1000/ --gen-eta=-1:1 --constraint-step-size-mm=1
+
+<build_directory>/bin/traccc_truth_finding_example_cuda --input_directory=detray_simulation/wire_chamber/n_particles_1000/ --check_performance=true --detector_file=<detector_directory>/wire_chamber_geometry.json --material_file=<detector_directory>/wire_chamber_homogeneous_material.json --grid_file=<detector_directory>/wire_chamber_surface_grids.json  --event=10 --track_candidates_range=6:30 --constraint-step-size-mm=1
+```
+
 ### SYCL reconstruction chain
 
 - Users can generate SYCL examples by adding `-DTRACCC_BUILD_SYCL=ON` to cmake options
@@ -259,6 +275,49 @@ cmake --build <build_directory> <options>
 <build_directory>/bin/traccc_throughput_mt_sycl --detector_file=tml_detector/trackml-detector.csv --digitization_config_file=tml_detector/default-geometric-config-generic.json --input_directory=tml_pixels/  --cold_run_events=100 --processed_events=1000 --threads=1
 ```
 
+### Running a partial chain with simplified simulation data
+
+Users can generate muon-like particle simulation data by running following example commands:
+
+```sh
+# Generate toy geometry data
+<build_directory>/bin/traccc_simulate_toy_detector --gen-vertex-xyz-mm=0:0:0 --gen-vertex-xyz-std-mm=0:0:0 --gen-mom-gev=100:100 --gen-phi-degree=0:360 --events=10 --gen-nparticles=2000 --output_directory=detray_simulation/toy_detector/n_particles_2000/ --gen-eta=-3:3 --constraint-step-size-mm=1
+
+# Generate drift chamber data
+<build_directory>/bin/traccc_simulate_wire_chamber --gen-vertex-xyz-mm=0:0:0 --gen-vertex-xyz-std-mm=0:0:0 --gen-mom-gev=2:2 --gen-phi-degree=0:360 --events=10 --gen-nparticles=100 --output_directory=detray_simulation/wire_chamber/n_particles_100/ --gen-eta=-1:1 --constraint-step-size-mm=1
+```
+
+The simulation will also generate the detector json files (geometry, material and surface_grid) in the current directory. It is user's responsibility to move them to an appropriate place (say, <detector_directory>) and match it to the input file arguments of the reconstruction examples.
+
+Currently, there are two types of partial reconstruction chain users can operate: seeding_example and truth_finding_example. seeding_example takes the truth measurement input w/o clusterization and it goes through seeding, finding and fitting, which generate performance root files, respectively. On the other hand, truth_finding_examples starts from the truth initial parameter of particles (no duplicate seeds; i.e. the number of seeds for CKF = the number of truth particles) and run track finding and track fitting.
+
+The dirft chamber will not produce meaningful results with seeding_examples as the current seeding algorithm is only designed for 2D measurement objects. Truth finding works OK in general but the combinatoric explosion can occur for a few unlucky events, leading to poor pull value distributions.
+
+```sh
+# Run cuda seeding example for toy geometry
+<build_directory>/bin/traccc_seeding_example_cuda --input_directory=detray_simulation/toy_detector/n_particles_2000/ --check_performance=true --detector_file=<detector_directory>/toy_detector_geometry.json --material_file=<detector_directory>/toy_detector_homogeneous_material.json --grid_file=<detector_directory>/toy_detector_surface_grids.json --event=1 --track_candidates_range=3:10 --constraint-step-size-mm=1000 --run_cpu=1
+```
+
+```sh
+# Run cuda truth finding example for toy geometry
+<build_directory>/bin/traccc_truth_finding_example_cuda --input_directory=detray_simulation/toy_detector/n_particles_2000/ --check_performance=true --detector_file=<detector_directory>/toy_detector_geometry.json --material_file=<detector_directory>/toy_detector_homogeneous_material.json --grid_file=<detector_directory>/toy_detector_surface_grids.json --event=1 --track_candidates_range=3:10 --constraint-step-size-mm=1000 --run_cpu=1
+```
+
+```sh
+# Run cuda truth finding example for drift chamber
+<build_directory>/bin/traccc_truth_finding_example_cuda --input_directory=detray_simulation/wire_chamber/n_particles_100/ --check_performance=true --detector_file=<detector_directory>/wire_chamber_geometry.json --material_file=<detector_directory>/wire_chamber_homogeneous_material.json --grid_file=<detector_directory>/wire_chamber_surface_grids.json  --event=10 --track_candidates_range=6:30 --constraint-step-size-mm=1 --run_cpu=1
+```
+
+Users can open the performance root files (with --check_performance=true) and draw the histograms. Following is the example of drawing track finding efficiency from track finding algorithm.
+
+```sh
+$ root -l performance_track_finding.root 
+root [0] 
+Attaching file performance_track_finding.root as _file0...
+(TFile *) 0x3871910
+root [1] finding_trackeff_vs_eta->Draw()
+```
+
 ## Troubleshooting
 
 The following are potentially useful instructions for troubleshooting various

core/include/traccc/finding/finding_config.hpp

@@ -16,7 +16,7 @@ namespace traccc {
 template <typename scalar_t>
 struct finding_config {
     /// @NOTE: This paramter might be removed
-    unsigned int max_num_branches_per_seed = 10;
+    unsigned int max_num_branches_per_seed = 100;
 
     /// Maximum number of branches per surface
     unsigned int max_num_branches_per_surface = 10;

core/include/traccc/fitting/fitting_config.hpp

@@ -20,6 +20,7 @@ struct fitting_config {
     scalar_t pathlimit = std::numeric_limits<scalar_t>::max();
     scalar_t overstep_tolerance = -10 * detray::unit<scalar_t>::um;
     scalar_t step_constraint = std::numeric_limits<scalar_t>::max();
+    scalar_t mask_tolerance = 15.f * detray::unit<scalar_t>::um;
 };
 
 }  // namespace traccc

core/include/traccc/fitting/kalman_filter/gain_matrix_smoother.hpp

@@ -48,16 +48,18 @@ struct gain_matrix_smoother {
         track_state<algebra_t>& cur_state,
         const track_state<algebra_t>& next_state) {
 
+        using shape_type = typename mask_group_t::value_type::shape;
+
         const auto D = cur_state.get_measurement().meas_dim;
         assert(D == 1u || D == 2u);
         if (D == 1u) {
-            smoothe<1u>(cur_state, next_state);
+            smoothe<1u, shape_type>(cur_state, next_state);
         } else if (D == 2u) {
-            smoothe<2u>(cur_state, next_state);
+            smoothe<2u, shape_type>(cur_state, next_state);
         }
     }
 
-    template <size_type D>
+    template <size_type D, typename shape_t>
     TRACCC_HOST_DEVICE inline void smoothe(
         track_state<algebra_t>& cur_state,
         const track_state<algebra_t>& next_state) const {
@@ -110,8 +112,16 @@ struct gain_matrix_smoother {
         cur_state.smoothed().set_vector(smt_vec);
         cur_state.smoothed().set_covariance(smt_cov);
 
-        const matrix_type<D, e_bound_size> H =
-            meas.subs.template projector<D>();
+        matrix_type<D, e_bound_size> H = meas.subs.template projector<D>();
+
+        // Correct sign for line detector
+        if constexpr (std::is_same_v<shape_t, detray::line<true>> ||
+                      std::is_same_v<shape_t, detray::line<false>>) {
+
+            if (getter::element(smt_vec, e_bound_loc0, 0u) < 0) {
+                getter::element(H, 0u, e_bound_loc0) = -1;
+            }
+        }
 
         // Calculate smoothed chi square
         const matrix_type<D, 1>& meas_local =

core/include/traccc/fitting/kalman_filter/gain_matrix_updater.hpp

@@ -42,16 +42,18 @@ struct gain_matrix_updater {
         track_state<algebra_t>& trk_state,
         bound_track_parameters& bound_params) const {
 
+        using shape_type = typename mask_group_t::value_type::shape;
+
         const auto D = trk_state.get_measurement().meas_dim;
         assert(D == 1u || D == 2u);
         if (D == 1u) {
-            update<1u>(trk_state, bound_params);
+            update<1u, shape_type>(trk_state, bound_params);
         } else if (D == 2u) {
-            update<2u>(trk_state, bound_params);
+            update<2u, shape_type>(trk_state, bound_params);
         }
     }
 
-    template <size_type D>
+    template <size_type D, typename shape_t>
     TRACCC_HOST_DEVICE inline void update(
         track_state<algebra_t>& trk_state,
         bound_track_parameters& bound_params) const {
@@ -69,9 +71,9 @@ struct gain_matrix_updater {
         const matrix_type<D, D> I_m =
             matrix_operator().template identity<D, D>();
 
-        const matrix_type<D, e_bound_size> H =
-            meas.subs.template projector<D>();
+        matrix_type<D, e_bound_size> H = meas.subs.template projector<D>();
 
+        // Measurement data on surface
         const matrix_type<D, 1>& meas_local =
             trk_state.template measurement_local<D>();
 
@@ -87,6 +89,14 @@ struct gain_matrix_updater {
         trk_state.predicted().set_vector(predicted_vec);
         trk_state.predicted().set_covariance(predicted_cov);
 
+        if constexpr (std::is_same_v<shape_t, detray::line<true>> ||
+                      std::is_same_v<shape_t, detray::line<false>>) {
+
+            if (getter::element(predicted_vec, e_bound_loc0, 0u) < 0) {
+                getter::element(H, 0u, e_bound_loc0) = -1;
+            }
+        }
+
         // Spatial resolution (Measurement covariance)
         const matrix_type<D, D> V =
             trk_state.template measurement_covariance<D>();

core/include/traccc/fitting/kalman_filter/kalman_actor.hpp

@@ -93,7 +93,7 @@ struct kalman_actor : detray::actor {
 
             // Abort if the propagator fails to find the next measurement
             if (navigation.barcode() != trk_state.surface_link()) {
-                propagation._heartbeat &= navigation.abort();
+                // propagation._heartbeat &= navigation.abort();
                 return;
             }
 

core/include/traccc/fitting/kalman_filter/kalman_fitter.hpp

@@ -173,12 +173,13 @@ class kalman_fitter {
         // volume in the constructor
         propagation._navigation.set_volume(seed_params.surface_link().volume());
 
-        // Set overstep tolerance and stepper constraint
+        // Set overstep tolerance, stepper constraint and mask tolerance
         propagation._stepping().set_overstep_tolerance(
             m_cfg.overstep_tolerance);
         propagation._stepping
             .template set_constraint<detray::step::constraint::e_accuracy>(
                 m_cfg.step_constraint);
+        propagation.set_mask_tolerance(m_cfg.mask_tolerance);
 
         // Run forward filtering
         propagator.propagate(propagation, fitter_state());

device/common/include/traccc/finding/device/find_tracks.hpp

@@ -31,6 +31,7 @@ namespace traccc::device {
 /// @param[in] step               Step index
 /// @param[in] n_measurements_per_thread  Number of measurements per thread
 /// @param[in] n_total_threads    Number of total threads
+/// @param[in] n_max_candidates   Number of maximum candidates
 /// @param[out] out_params_view   Output parameters
 /// @param[out] links_view        link container for the current step
 /// @param[out] n_candidates      The number of candidates for the current step
@@ -45,7 +46,7 @@ TRACCC_DEVICE inline void find_tracks(
     bound_track_parameters_collection_types::const_view in_params_view,
     vecmem::data::vector_view<const unsigned int> n_threads_view,
     const unsigned int step, const unsigned int& n_measurements_per_thread,
-    const unsigned int& n_total_threads,
+    const unsigned int& n_total_threads, const unsigned int& n_max_candidates,
     bound_track_parameters_collection_types::view out_params_view,
     vecmem::data::vector_view<candidate_link> links_view,
     unsigned int& n_candidates);

device/common/include/traccc/finding/device/impl/find_tracks.ipp

@@ -25,7 +25,7 @@ TRACCC_DEVICE inline void find_tracks(
     bound_track_parameters_collection_types::const_view in_params_view,
     vecmem::data::vector_view<const unsigned int> n_threads_view,
     const unsigned int step, const unsigned int& n_measurements_per_thread,
-    const unsigned int& n_total_threads,
+    const unsigned int& n_total_threads, const unsigned int& n_max_candidates,
     bound_track_parameters_collection_types::view out_params_view,
     vecmem::data::vector_view<candidate_link> links_view,
     unsigned int& n_candidates) {
@@ -120,9 +120,14 @@ TRACCC_DEVICE inline void find_tracks(
             // Add measurement candidates to link
             vecmem::device_atomic_ref<unsigned int> num_candidates(
                 n_candidates);
+
             const unsigned int l_pos = num_candidates.fetch_add(1);
 
-            // @TODO; Consider max_num_branches_per_surface
+            if (l_pos >= n_max_candidates) {
+                n_candidates = n_max_candidates;
+                return;
+            }
+
             links[l_pos] = {{previous_step, in_param_id}, meas_idx};
 
             out_params[l_pos] = trk_state.filtered();

device/cuda/src/finding/finding_algorithm.cu

@@ -20,6 +20,7 @@
 
 // detray include(s).
 #include "detray/core/detector.hpp"
+#include "detray/core/detector_metadata.hpp"
 #include "detray/detectors/bfield.hpp"
 #include "detray/detectors/telescope_metadata.hpp"
 #include "detray/detectors/toy_metadata.hpp"
@@ -102,7 +103,7 @@ __global__ void find_tracks(
     bound_track_parameters_collection_types::const_view in_params_view,
     vecmem::data::vector_view<const unsigned int> n_threads_view,
     const unsigned int step, const unsigned int& n_measurements_per_thread,
-    const unsigned int& n_total_threads,
+    const unsigned int& n_total_threads, const unsigned int n_max_candidates,
     bound_track_parameters_collection_types::view out_params_view,
     vecmem::data::vector_view<candidate_link> links_view,
     unsigned int& n_candidates) {
@@ -112,7 +113,8 @@ __global__ void find_tracks(
     device::find_tracks<detector_t, config_t>(
         gid, cfg, det_data, measurements_view, barcodes_view, upper_bounds_view,
         in_params_view, n_threads_view, step, n_measurements_per_thread,
-        n_total_threads, out_params_view, links_view, n_candidates);
+        n_total_threads, n_max_candidates, out_params_view, links_view,
+        n_candidates);
 }
 
 /// CUDA kernel for running @c traccc::device::propagate_to_next_surface
@@ -341,6 +343,11 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
 
         // Buffer for kalman-updated parameters spawned by the measurement
         // candidates
+
+        const unsigned int n_max_candidates =
+            std::min(n_in_params * m_cfg.max_num_branches_per_surface,
+                     seeds.size() * m_cfg.max_num_branches_per_seed);
+
         bound_track_parameters_collection_types::buffer updated_params_buffer(
             n_in_params * m_cfg.max_num_branches_per_surface, m_mr.main);
 
@@ -357,7 +364,7 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
                     m_cfg, det_view, measurements, barcodes_buffer,
                     upper_bounds_buffer, in_params_buffer, n_threads_buffer,
                     step, (*global_counter_device).n_measurements_per_thread,
-                    (*global_counter_device).n_total_threads,
+                    (*global_counter_device).n_total_threads, n_max_candidates,
                     updated_params_buffer, link_map[step],
                     (*global_counter_device).n_candidates);
             CUDA_ERROR_CHECK(cudaGetLastError());
@@ -511,21 +518,12 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
 }
 
 // Explicit template instantiation
-using toy_detector_type =
-    detray::detector<detray::toy_metadata, detray::device_container_types>;
-using toy_stepper_type =
-    detray::rk_stepper<covfie::field_view<detray::bfield::const_bknd_t>,
-                       transform3, detray::constrained_step<>>;
-using toy_navigator_type = detray::navigator<const toy_detector_type>;
-template class finding_algorithm<toy_stepper_type, toy_navigator_type>;
-
-using device_detector_type =
-    detray::detector<detray::telescope_metadata<detray::rectangle2D<>>,
-                     detray::device_container_types>;
-using rk_stepper_type =
-    detray::rk_stepper<covfie::field_view<detray::bfield::const_bknd_t>,
+using default_detector_type =
+    detray::detector<detray::default_metadata, detray::device_container_types>;
+using default_stepper_type =
+    detray::rk_stepper<covfie::field<detray::bfield::const_bknd_t>::view_t,
                        transform3, detray::constrained_step<>>;
-using device_navigator_type = detray::navigator<const device_detector_type>;
-template class finding_algorithm<rk_stepper_type, device_navigator_type>;
+using default_navigator_type = detray::navigator<const default_detector_type>;
+template class finding_algorithm<default_stepper_type, default_navigator_type>;
 
 }  // namespace traccc::cuda

device/cuda/src/fitting/fitting_algorithm.cu

@@ -12,6 +12,7 @@
 #include "traccc/fitting/kalman_filter/kalman_fitter.hpp"
 
 // detray include(s).
+#include "detray/core/detector_metadata.hpp"
 #include "detray/detectors/bfield.hpp"
 #include "detray/detectors/telescope_metadata.hpp"
 #include "detray/detectors/toy_metadata.hpp"
@@ -99,24 +100,14 @@ track_state_container_types::buffer fitting_algorithm<fitter_t>::operator()(
 }
 
 // Explicit template instantiation
-using toy_detector_type =
-    detray::detector<detray::toy_metadata, detray::device_container_types>;
-using toy_stepper_type =
-    detray::rk_stepper<covfie::field_view<detray::bfield::const_bknd_t>,
+using default_detector_type =
+    detray::detector<detray::default_metadata, detray::device_container_types>;
+using default_stepper_type =
+    detray::rk_stepper<covfie::field<detray::bfield::const_bknd_t>::view_t,
                        transform3, detray::constrained_step<>>;
-using toy_navigator_type = detray::navigator<const toy_detector_type>;
-using toy_fitter_type = kalman_fitter<toy_stepper_type, toy_navigator_type>;
-template class fitting_algorithm<toy_fitter_type>;
-
-using device_detector_type =
-    detray::detector<detray::telescope_metadata<detray::rectangle2D<>>,
-                     detray::device_container_types>;
-using rk_stepper_type =
-    detray::rk_stepper<covfie::field_view<detray::bfield::const_bknd_t>,
-                       transform3, detray::constrained_step<>>;
-using device_navigator_type = detray::navigator<const device_detector_type>;
-using device_fitter_type =
-    kalman_fitter<rk_stepper_type, device_navigator_type>;
-template class fitting_algorithm<device_fitter_type>;
+using default_navigator_type = detray::navigator<const default_detector_type>;
+using default_fitter_type =
+    kalman_fitter<default_stepper_type, default_navigator_type>;
+template class fitting_algorithm<default_fitter_type>;
 
 }  // namespace traccc::cuda