diff --git a/benchmarks/common/benchmarks/toy_detector_benchmark.hpp b/benchmarks/common/benchmarks/toy_detector_benchmark.hpp index 4b2601a199..29d8fa5d2e 100644 --- a/benchmarks/common/benchmarks/toy_detector_benchmark.hpp +++ b/benchmarks/common/benchmarks/toy_detector_benchmark.hpp @@ -60,8 +60,7 @@ class ToyDetectorBenchmark : public benchmark::Fixture { static constexpr std::array phi_range{ -traccc::constant::pi, traccc::constant::pi}; - static constexpr std::array theta_range{ - 0.f, traccc::constant::pi}; + static constexpr std::array eta_range{-4, 4}; static constexpr std::array mom_range{ 10.f * traccc::unit::GeV, 100.f * traccc::unit::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>; @@ -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(), 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::mm; @@ -155,6 +160,20 @@ class ToyDetectorBenchmark : public benchmark::Fixture { return toy_cfg; } + template + 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::um; + cfg.propagation.navigation.min_mask_tolerance = + 1e-5f * detray::unit::mm; + cfg.propagation.navigation.max_mask_tolerance = + 3.f * detray::unit::mm; + cfg.propagation.navigation.mask_tolerance_scalor = 0.05f; + } + void SetUp(::benchmark::State& /*state*/) { // Read events diff --git a/benchmarks/cpu/toy_detector_cpu.cpp b/benchmarks/cpu/toy_detector_cpu.cpp index 4594c0e426..b7a66a7183 100644 --- a/benchmarks/cpu/toy_detector_cpu.cpp +++ b/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>; - 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; using host_fitter_type = traccc::kalman_fitter; diff --git a/benchmarks/cuda/toy_detector_cuda.cpp b/benchmarks/cuda/toy_detector_cuda.cpp index 47d321ae0f..f16d2c70ec 100644 --- a/benchmarks/cuda/toy_detector_cuda.cpp +++ b/benchmarks/cuda/toy_detector_cuda.cpp @@ -31,7 +31,6 @@ // VecMem include(s). #include #include -#include #include #include #include @@ -44,10 +43,9 @@ BENCHMARK_F(ToyDetectorBenchmark, CUDA)(benchmark::State& state) { // Type declarations using rk_stepper_type = detray::rk_stepper>; - 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; using device_fitter_type = traccc::kalman_fitter; @@ -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_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 @@ -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 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(); } } diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu index bc387a3bea..fe6ea7bcf1 100644 --- a/device/cuda/src/finding/finding_algorithm.cu +++ b/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::operator()( } // Explicit template instantiation -using default_detector_type = - detray::detector; -using default_stepper_type = +using debug_stepper_type = detray::rk_stepper::view_t, traccc::default_algebra, detray::constrained_step<>>; -using default_navigator_type = detray::navigator; +using default_stepper_type = + detray::rk_stepper::view_t, + traccc::default_algebra>; + +using default_navigator_type = + detray::navigator; +using toy_navigator_type = + detray::navigator; + +template class finding_algorithm; template class finding_algorithm; +template class finding_algorithm; } // namespace traccc::cuda diff --git a/device/cuda/src/fitting/fitting_algorithm.cu b/device/cuda/src/fitting/fitting_algorithm.cu index a8cf428d21..2809fe0fe4 100644 --- a/device/cuda/src/fitting/fitting_algorithm.cu +++ b/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::operator()( } // Explicit template instantiation -using default_detector_type = - detray::detector; -using default_stepper_type = +using debug_stepper_type = detray::rk_stepper::view_t, default_algebra, detray::constrained_step<>>; -using default_navigator_type = detray::navigator; +using default_stepper_type = + detray::rk_stepper::view_t, + traccc::default_algebra>; +using default_navigator_type = + detray::navigator; +using toy_navigator_type = + detray::navigator; + +using debug_fitter_type = + kalman_fitter; using default_fitter_type = kalman_fitter; +using toy_det_fitter_type = + kalman_fitter; + +template class fitting_algorithm; template class fitting_algorithm; +template class fitting_algorithm; } // namespace traccc::cuda diff --git a/io/include/traccc/io/read_detector.hpp b/io/include/traccc/io/read_detector.hpp index 40fe089901..966ca4e6e8 100644 --- a/io/include/traccc/io/read_detector.hpp +++ b/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 diff --git a/io/src/read_detector.cpp b/io/src/read_detector.cpp index 14702494dd..543dc696f3 100644 --- a/io/src/read_detector.cpp +++ b/io/src/read_detector.cpp @@ -20,7 +20,8 @@ namespace { /// Common implementation for constructing a detector from a set of input files -template +/// @tparam CAP Grid bin capacity 0: dynamic bin capacity +template 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(mr, cfg); + auto det = detray::io::read_detector(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(detector, mr, geometry_file, + material_file, grid_file); +} + } // namespace traccc::io