ref: benchmark toy detector config

benchmarks/common/benchmarks/toy_detector_benchmark.hpp

@@ -60,8 +60,7 @@ class ToyDetectorBenchmark : public benchmark::Fixture {
 
     static constexpr std::array<float, 2> phi_range{
         -traccc::constant<float>::pi, traccc::constant<float>::pi};
-    static constexpr std::array<float, 2> theta_range{
-        0.f, traccc::constant<float>::pi};
+    static constexpr std::array<float, 2> eta_range{-4, 4};
     static constexpr std::array<float, 2> mom_range{
         10.f * traccc::unit<float>::GeV, 100.f * traccc::unit<float>::GeV};
 
@@ -83,6 +82,10 @@ class ToyDetectorBenchmark : public benchmark::Fixture {
                      "the simulation data."
                   << std::endl;
 
+        // Apply correct propagation config
+        apply_propagation_config(finding_cfg);
+        apply_propagation_config(fitting_cfg);
+
         // Use deterministic random number generator for testing
         using uniform_gen_t = detray::detail::random_numbers<
             traccc::scalar, std::uniform_real_distribution<traccc::scalar>>;
@@ -101,7 +104,7 @@ class ToyDetectorBenchmark : public benchmark::Fixture {
         generator_type::configuration gen_cfg{};
         gen_cfg.n_tracks(n_tracks);
         gen_cfg.phi_range(phi_range);
-        gen_cfg.theta_range(theta_range);
+        gen_cfg.eta_range(eta_range);
         gen_cfg.mom_range(mom_range);
         generator_type generator(gen_cfg);
 
@@ -127,6 +130,8 @@ class ToyDetectorBenchmark : public benchmark::Fixture {
             detray::muon<traccc::scalar>(), n_events, det, field,
             std::move(generator), std::move(smearer_writer_cfg), full_path);
 
+        // Same propagation configuration for sim and reco
+        apply_propagation_config(sim.get_config());
         // Set constrained step size to 1 mm
         sim.get_config().propagation.stepping.step_constraint =
             1.f * detray::unit<traccc::scalar>::mm;
@@ -155,6 +160,20 @@ class ToyDetectorBenchmark : public benchmark::Fixture {
         return toy_cfg;
     }
 
+    template <typename config_t>
+    void apply_propagation_config(config_t& cfg) const {
+
+        // Configure the propagation for the toy detector
+        cfg.propagation.navigation.search_window = {3, 3};
+        cfg.propagation.navigation.overstep_tolerance =
+            -300.f * detray::unit<float>::um;
+        cfg.propagation.navigation.min_mask_tolerance =
+            1e-5f * detray::unit<float>::mm;
+        cfg.propagation.navigation.max_mask_tolerance =
+            3.f * detray::unit<float>::mm;
+        cfg.propagation.navigation.mask_tolerance_scalor = 0.05f;
+    }
+
     void SetUp(::benchmark::State& /*state*/) {
 
         // Read events

benchmarks/cpu/toy_detector_cpu.cpp

@@ -42,9 +42,8 @@ BENCHMARK_F(ToyDetectorBenchmark, CPU)(benchmark::State& state) {
     // Type declarations
     using rk_stepper_type =
         detray::rk_stepper<b_field_t::view_t,
-                           typename detector_type::algebra_type,
-                           detray::constrained_step<>>;
-    using host_detector_type = traccc::default_detector::host;
+                           typename detector_type::algebra_type>;
+    using host_detector_type = traccc::toy_detector::host;
     using host_navigator_type = detray::navigator<const host_detector_type>;
     using host_fitter_type =
         traccc::kalman_fitter<rk_stepper_type, host_navigator_type>;

benchmarks/cuda/toy_detector_cuda.cpp

@@ -31,7 +31,6 @@
 // VecMem include(s).
 #include <vecmem/memory/cuda/device_memory_resource.hpp>
 #include <vecmem/memory/cuda/host_memory_resource.hpp>
-#include <vecmem/memory/cuda/managed_memory_resource.hpp>
 #include <vecmem/memory/host_memory_resource.hpp>
 #include <vecmem/utils/cuda/async_copy.hpp>
 #include <vecmem/utils/cuda/copy.hpp>
@@ -44,10 +43,9 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
     // Type declarations
     using rk_stepper_type =
         detray::rk_stepper<b_field_t::view_t,
-                           typename detector_type::algebra_type,
-                           detray::constrained_step<>>;
-    using host_detector_type = traccc::default_detector::host;
-    using device_detector_type = traccc::default_detector::device;
+                           typename detector_type::algebra_type>;
+    using host_detector_type = traccc::toy_detector::host;
+    using device_detector_type = traccc::toy_detector::device;
     using device_navigator_type = detray::navigator<const device_detector_type>;
     using device_fitter_type =
         traccc::kalman_fitter<rk_stepper_type, device_navigator_type>;
@@ -56,7 +54,6 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
     vecmem::cuda::host_memory_resource cuda_host_mr;
     vecmem::cuda::device_memory_resource device_mr;
     traccc::memory_resource mr{device_mr, &cuda_host_mr};
-    vecmem::cuda::managed_memory_resource mng_mr;
 
     // Copy and stream
     vecmem::copy host_copy;
@@ -65,9 +62,9 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
     vecmem::cuda::async_copy async_copy{stream.cudaStream()};
 
     // Read back detector file
-    host_detector_type det{mng_mr};
+    host_detector_type det{cuda_host_mr};
     traccc::io::read_detector(
-        det, mng_mr, sim_dir + "toy_detector_geometry.json",
+        det, cuda_host_mr, sim_dir + "toy_detector_geometry.json",
         sim_dir + "toy_detector_homogeneous_material.json",
         sim_dir + "toy_detector_surface_grids.json");
 
@@ -83,8 +80,10 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
     traccc::cuda::fitting_algorithm<device_fitter_type> device_fitting(
         fitting_cfg, mr, async_copy, stream);
 
+    // Copy detector to device
+    auto det_fixed_buff = detray::get_buffer(det, device_mr, copy);
     // Detector view object
-    auto det_view = detray::get_data(det);
+    auto det_view = detray::get_data(det_fixed_buff);
 
     // D2H copy object
     traccc::device::container_d2h_copy_alg<traccc::track_state_container_types>
@@ -109,7 +108,7 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
 
             auto& spacepoints_per_event = spacepoints[i_evt];
             auto& measurements_per_event = measurements[i_evt];
-
+            state.PauseTiming();
             // Copy the spacepoint and module data to the device.
             traccc::spacepoint_collection_types::buffer spacepoints_cuda_buffer(
                 spacepoints_per_event.size(), mr.main);
@@ -122,6 +121,7 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
             async_copy(vecmem::get_data(measurements_per_event),
                        measurements_cuda_buffer);
 
+            state.ResumeTiming();
             // Run seeding
             traccc::seed_collection_types::buffer seeds_cuda_buffer =
                 sa_cuda(spacepoints_cuda_buffer);
@@ -141,7 +141,7 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
             traccc::track_state_container_types::buffer
                 track_states_cuda_buffer = device_fitting(
                     det_view, field, track_candidates_cuda_buffer);
-
+            state.PauseTiming();
             // Create a temporary buffer that will receive the device memory.
             auto size = track_states_cuda_buffer.headers.size();
             std::vector<std::size_t> capacities(size, 0);
@@ -153,6 +153,8 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) {
             // Copy the track states back to the host.
             traccc::track_state_container_types::host track_states_host =
                 track_state_d2h(track_states_cuda_buffer);
+
+            state.ResumeTiming();
         }
     }
 

device/cuda/src/finding/finding_algorithm.cu

@@ -23,11 +23,10 @@
 #include "traccc/finding/device/make_barcode_sequence.hpp"
 #include "traccc/finding/device/propagate_to_next_surface.hpp"
 #include "traccc/finding/device/prune_tracks.hpp"
+#include "traccc/geometry/detector.hpp"
 #include "traccc/utils/projections.hpp"
 
 // detray include(s).
-#include "detray/core/detector.hpp"
-#include "detray/core/detector_metadata.hpp"
 #include "detray/detectors/bfield.hpp"
 #include "detray/navigation/navigator.hpp"
 #include "detray/propagator/rk_stepper.hpp"
@@ -534,12 +533,20 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
 }
 
 // Explicit template instantiation
-using default_detector_type =
-    detray::detector<detray::default_metadata, detray::device_container_types>;
-using default_stepper_type =
+using debug_stepper_type =
     detray::rk_stepper<covfie::field<detray::bfield::const_bknd_t>::view_t,
                        traccc::default_algebra, detray::constrained_step<>>;
-using default_navigator_type = detray::navigator<const default_detector_type>;
+using default_stepper_type =
+    detray::rk_stepper<covfie::field<detray::bfield::const_bknd_t>::view_t,
+                       traccc::default_algebra>;
+
+using default_navigator_type =
+    detray::navigator<const traccc::default_detector::device>;
+using toy_navigator_type =
+    detray::navigator<const traccc::toy_detector::device>;
+
+template class finding_algorithm<debug_stepper_type, default_navigator_type>;
 template class finding_algorithm<default_stepper_type, default_navigator_type>;
+template class finding_algorithm<default_stepper_type, toy_navigator_type>;
 
 }  // namespace traccc::cuda

device/cuda/src/fitting/fitting_algorithm.cu

@@ -12,9 +12,9 @@
 #include "traccc/fitting/device/fill_sort_keys.hpp"
 #include "traccc/fitting/device/fit.hpp"
 #include "traccc/fitting/kalman_filter/kalman_fitter.hpp"
+#include "traccc/geometry/detector.hpp"
 
 // detray include(s).
-#include "detray/core/detector_metadata.hpp"
 #include "detray/detectors/bfield.hpp"
 #include "detray/propagator/rk_stepper.hpp"
 
@@ -126,14 +126,26 @@ track_state_container_types::buffer fitting_algorithm<fitter_t>::operator()(
 }
 
 // Explicit template instantiation
-using default_detector_type =
-    detray::detector<detray::default_metadata, detray::device_container_types>;
-using default_stepper_type =
+using debug_stepper_type =
     detray::rk_stepper<covfie::field<detray::bfield::const_bknd_t>::view_t,
                        default_algebra, detray::constrained_step<>>;
-using default_navigator_type = detray::navigator<const default_detector_type>;
+using default_stepper_type =
+    detray::rk_stepper<covfie::field<detray::bfield::const_bknd_t>::view_t,
+                       traccc::default_algebra>;
+using default_navigator_type =
+    detray::navigator<const traccc::default_detector::device>;
+using toy_navigator_type =
+    detray::navigator<const traccc::toy_detector::device>;
+
+using debug_fitter_type =
+    kalman_fitter<debug_stepper_type, default_navigator_type>;
 using default_fitter_type =
     kalman_fitter<default_stepper_type, default_navigator_type>;
+using toy_det_fitter_type =
+    kalman_fitter<default_stepper_type, toy_navigator_type>;
+
+template class fitting_algorithm<debug_fitter_type>;
 template class fitting_algorithm<default_fitter_type>;
+template class fitting_algorithm<toy_det_fitter_type>;
 
 }  // namespace traccc::cuda

io/include/traccc/io/read_detector.hpp

@@ -35,4 +35,9 @@ void read_detector(default_detector::host& detector,
                    const std::string_view& material_file = "",
                    const std::string_view& grid_file = "");
 
+void read_detector(toy_detector::host& detector, vecmem::memory_resource& mr,
+                   const std::string_view& geometry_file,
+                   const std::string_view& material_file = "",
+                   const std::string_view& grid_file = "");
+
 }  // namespace traccc::io

io/src/read_detector.cpp

@@ -20,7 +20,8 @@
 namespace {
 
 /// Common implementation for constructing a detector from a set of input files
-template <typename detector_t>
+/// @tparam CAP Grid bin capacity 0: dynamic bin capacity
+template <typename detector_t, unsigned int CAP = 0u>
 void read_detector(detector_t& detector, vecmem::memory_resource& mr,
                    const std::string_view& geometry_file,
                    const std::string_view& material_file,
@@ -37,7 +38,7 @@ void read_detector(detector_t& detector, vecmem::memory_resource& mr,
     }
 
     // Read the detector.
-    auto det = detray::io::read_detector<detector_t>(mr, cfg);
+    auto det = detray::io::read_detector<detector_t, CAP>(mr, cfg);
     detector = std::move(det.first);
 }
 
@@ -50,8 +51,16 @@ void read_detector(default_detector::host& detector,
                    const std::string_view& geometry_file,
                    const std::string_view& material_file,
                    const std::string_view& grid_file) {
-
     ::read_detector(detector, mr, geometry_file, material_file, grid_file);
 }
 
+void read_detector(toy_detector::host& detector, vecmem::memory_resource& mr,
+                   const std::string_view& geometry_file,
+                   const std::string_view& material_file,
+                   const std::string_view& grid_file) {
+
+    ::read_detector<toy_detector::host, 1u>(detector, mr, geometry_file,
+                                            material_file, grid_file);
+}
+
 }  // namespace traccc::io