From eb2dd4784754fcae41f79484b926c73fba7e4c72 Mon Sep 17 00:00:00 2001 From: Joana Niermann Date: Tue, 24 Sep 2024 03:49:54 +0200 Subject: [PATCH 1/2] Update config --- .../benchmarks/toy_detector_benchmark.hpp | 21 +++++++++++++++---- 1 file changed, 17 insertions(+), 4 deletions(-) diff --git a/benchmarks/common/benchmarks/toy_detector_benchmark.hpp b/benchmarks/common/benchmarks/toy_detector_benchmark.hpp index 4b2601a199..9cc42c878b 100644 --- a/benchmarks/common/benchmarks/toy_detector_benchmark.hpp +++ b/benchmarks/common/benchmarks/toy_detector_benchmark.hpp @@ -55,13 +55,12 @@ class ToyDetectorBenchmark : public benchmark::Fixture { traccc::seedfinder_config seeding_cfg; traccc::seedfilter_config filter_cfg; traccc::spacepoint_grid_config grid_cfg{seeding_cfg}; - traccc::finding_config finding_cfg; + traccc::finding_config finding_cfg = get_trk_finding_config(); traccc::fitting_config fitting_cfg; 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{-3, 3}; static constexpr std::array mom_range{ 10.f * traccc::unit::GeV, 100.f * traccc::unit::GeV}; @@ -101,7 +100,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); @@ -130,6 +129,8 @@ class ToyDetectorBenchmark : public benchmark::Fixture { // Set constrained step size to 1 mm sim.get_config().propagation.stepping.step_constraint = 1.f * detray::unit::mm; + // Otherwise same propagation configuration for sim and reco + sim.get_config().propagation = finding_cfg.propagation; sim.run(); @@ -155,6 +156,18 @@ class ToyDetectorBenchmark : public benchmark::Fixture { return toy_cfg; } + traccc::finding_config get_trk_finding_config() const { + + traccc::finding_config finding_cfg{}; + + // Configure the propagation for the toy detector + finding_cfg.propagation.navigation.search_window = {3, 3}; + finding_cfg.propagation.navigation.overstep_tolerance = + -300.f * detray::unit::um; + + return finding_cfg; + } + void SetUp(::benchmark::State& /*state*/) { // Read events From e791de91056342fc99e6afa7d6002963548f45df Mon Sep 17 00:00:00 2001 From: Joana Niermann Date: Tue, 24 Sep 2024 04:12:48 +0200 Subject: [PATCH 2/2] Remove step size constraints from performance measurement --- .../benchmarks/toy_detector_benchmark.hpp | 30 +++++++++++-------- benchmarks/cpu/toy_detector_cpu.cpp | 5 ++-- benchmarks/cuda/toy_detector_cuda.cpp | 24 ++++++++------- device/cuda/src/finding/finding_algorithm.cu | 19 ++++++++---- device/cuda/src/fitting/fitting_algorithm.cu | 22 ++++++++++---- io/include/traccc/io/read_detector.hpp | 5 ++++ io/src/read_detector.cpp | 15 ++++++++-- 7 files changed, 80 insertions(+), 40 deletions(-) diff --git a/benchmarks/common/benchmarks/toy_detector_benchmark.hpp b/benchmarks/common/benchmarks/toy_detector_benchmark.hpp index 9cc42c878b..29d8fa5d2e 100644 --- a/benchmarks/common/benchmarks/toy_detector_benchmark.hpp +++ b/benchmarks/common/benchmarks/toy_detector_benchmark.hpp @@ -55,12 +55,12 @@ class ToyDetectorBenchmark : public benchmark::Fixture { traccc::seedfinder_config seeding_cfg; traccc::seedfilter_config filter_cfg; traccc::spacepoint_grid_config grid_cfg{seeding_cfg}; - traccc::finding_config finding_cfg = get_trk_finding_config(); + traccc::finding_config finding_cfg; traccc::fitting_config fitting_cfg; static constexpr std::array phi_range{ -traccc::constant::pi, traccc::constant::pi}; - static constexpr std::array eta_range{-3, 3}; + static constexpr std::array eta_range{-4, 4}; static constexpr std::array mom_range{ 10.f * traccc::unit::GeV, 100.f * traccc::unit::GeV}; @@ -82,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>; @@ -126,11 +130,11 @@ 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; - // Otherwise same propagation configuration for sim and reco - sim.get_config().propagation = finding_cfg.propagation; sim.run(); @@ -156,16 +160,18 @@ class ToyDetectorBenchmark : public benchmark::Fixture { return toy_cfg; } - traccc::finding_config get_trk_finding_config() const { - - traccc::finding_config finding_cfg{}; + template + void apply_propagation_config(config_t& cfg) const { // Configure the propagation for the toy detector - finding_cfg.propagation.navigation.search_window = {3, 3}; - finding_cfg.propagation.navigation.overstep_tolerance = - -300.f * detray::unit::um; - - return finding_cfg; + 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*/) { 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