From e07cfb831a446fceee1688ec6fb7b58655f366bd Mon Sep 17 00:00:00 2001 From: Shima Shimizu Date: Thu, 16 Nov 2023 17:43:50 +0900 Subject: [PATCH 01/33] Parameter to limit the N track cands per initial seed. --- core/include/traccc/finding/candidate_link.hpp | 6 +++++- .../traccc/finding/finding_algorithm.ipp | 18 ++++++++++++++++-- core/include/traccc/finding/finding_config.hpp | 5 ++++- 3 files changed, 25 insertions(+), 4 deletions(-) diff --git a/core/include/traccc/finding/candidate_link.hpp b/core/include/traccc/finding/candidate_link.hpp index 00870c4b07..971ff4dca0 100644 --- a/core/include/traccc/finding/candidate_link.hpp +++ b/core/include/traccc/finding/candidate_link.hpp @@ -30,6 +30,10 @@ struct candidate_link { // Measurement index unsigned int meas_idx; + + // Index to the initial seed + unsigned int initial_seed; + }; -} // namespace traccc \ No newline at end of file +} // namespace traccc diff --git a/core/include/traccc/finding/finding_algorithm.ipp b/core/include/traccc/finding/finding_algorithm.ipp index e3a45ffcec..86b69171ba 100644 --- a/core/include/traccc/finding/finding_algorithm.ipp +++ b/core/include/traccc/finding/finding_algorithm.ipp @@ -75,9 +75,13 @@ finding_algorithm::operator()( // Copy seed to input parameters std::vector in_params; + std::vector n_trks_per_seed; + n_trks_per_seed.reserve(seeds.size()); + in_params.reserve(seeds.size()); for (const auto& seed : seeds) { in_params.push_back(seed); + n_trks_per_seed.push_back(0); } std::vector out_params; @@ -99,11 +103,16 @@ finding_algorithm::operator()( // Previous step ID const unsigned int previous_step = (step == 0) ? std::numeric_limits::max() : step - 1; + + for (unsigned int i_seed = 0; i_seed < seeds.size() ; i_seed ++) + n_trks_per_seed[i_seed] = 0; for (unsigned int in_param_id = 0; in_param_id < n_in_params; in_param_id++) { bound_track_parameters& in_param = in_params[in_param_id]; + unsigned int orig_param_id = (step ==0 ? in_param_id : + links[step-1][param_to_link[step-1][in_param_id]].initial_seed); /************************* * Material interaction @@ -165,6 +174,9 @@ finding_algorithm::operator()( if (n_branches > m_cfg.max_num_branches_per_surface) { break; } + if (n_trks_per_seed[orig_param_id] >= m_cfg.max_num_branches_per_initial_seed) { + break; + } bound_track_parameters bound_param(in_param.surface_link(), in_param.vector(), @@ -188,9 +200,11 @@ finding_algorithm::operator()( static_cast(links[step].size()); n_branches++; + n_trks_per_seed[orig_param_id]++; - links[step].push_back( - {{previous_step, in_param_id}, item_id}); + // links[step].push_back( + // {{previous_step, in_param_id}, item_id}); + links[step].push_back({{previous_step, in_param_id}, item_id, orig_param_id}); /********************************* * Propagate to the next surface diff --git a/core/include/traccc/finding/finding_config.hpp b/core/include/traccc/finding/finding_config.hpp index a484e98df3..fce1428182 100644 --- a/core/include/traccc/finding/finding_config.hpp +++ b/core/include/traccc/finding/finding_config.hpp @@ -25,6 +25,9 @@ struct finding_config { unsigned int min_track_candidates_per_track = 3; unsigned int max_track_candidates_per_track = 30; + /// Maximum number of branches per initial seed + unsigned int max_num_branches_per_initial_seed = 1000; + /// Minimum step length that track should make to reach the next surface. It /// should be set higher than the overstep tolerance not to make it stay on /// the same surface @@ -42,4 +45,4 @@ struct finding_config { unsigned int n_avg_threads_per_track = 4; }; -} // namespace traccc \ No newline at end of file +} // namespace traccc From 1e06426bbcef177557061a3576f9ebbab88799a6 Mon Sep 17 00:00:00 2001 From: Shima Shimizu Date: Thu, 16 Nov 2023 19:04:02 +0900 Subject: [PATCH 02/33] format corrected --- .../include/traccc/finding/candidate_link.hpp | 3 +-- .../traccc/finding/finding_algorithm.ipp | 27 +++++++++++-------- 2 files changed, 17 insertions(+), 13 deletions(-) diff --git a/core/include/traccc/finding/candidate_link.hpp b/core/include/traccc/finding/candidate_link.hpp index 971ff4dca0..cdfb02b725 100644 --- a/core/include/traccc/finding/candidate_link.hpp +++ b/core/include/traccc/finding/candidate_link.hpp @@ -32,8 +32,7 @@ struct candidate_link { unsigned int meas_idx; // Index to the initial seed - unsigned int initial_seed; - + unsigned int initial_seed; }; } // namespace traccc diff --git a/core/include/traccc/finding/finding_algorithm.ipp b/core/include/traccc/finding/finding_algorithm.ipp index 86b69171ba..32f987475c 100644 --- a/core/include/traccc/finding/finding_algorithm.ipp +++ b/core/include/traccc/finding/finding_algorithm.ipp @@ -81,7 +81,7 @@ finding_algorithm::operator()( in_params.reserve(seeds.size()); for (const auto& seed : seeds) { in_params.push_back(seed); - n_trks_per_seed.push_back(0); + n_trks_per_seed.push_back(0); } std::vector out_params; @@ -103,16 +103,19 @@ finding_algorithm::operator()( // Previous step ID const unsigned int previous_step = (step == 0) ? std::numeric_limits::max() : step - 1; - - for (unsigned int i_seed = 0; i_seed < seeds.size() ; i_seed ++) - n_trks_per_seed[i_seed] = 0; + + for (unsigned int i_seed = 0; i_seed < seeds.size(); i_seed++) + n_trks_per_seed[i_seed] = 0; for (unsigned int in_param_id = 0; in_param_id < n_in_params; in_param_id++) { bound_track_parameters& in_param = in_params[in_param_id]; - unsigned int orig_param_id = (step ==0 ? in_param_id : - links[step-1][param_to_link[step-1][in_param_id]].initial_seed); + unsigned int orig_param_id = + (step == 0 + ? in_param_id + : links[step - 1][param_to_link[step - 1][in_param_id]] + .initial_seed); /************************* * Material interaction @@ -174,9 +177,10 @@ finding_algorithm::operator()( if (n_branches > m_cfg.max_num_branches_per_surface) { break; } - if (n_trks_per_seed[orig_param_id] >= m_cfg.max_num_branches_per_initial_seed) { - break; - } + if (n_trks_per_seed[orig_param_id] >= + m_cfg.max_num_branches_per_initial_seed) { + break; + } bound_track_parameters bound_param(in_param.surface_link(), in_param.vector(), @@ -200,11 +204,12 @@ finding_algorithm::operator()( static_cast(links[step].size()); n_branches++; - n_trks_per_seed[orig_param_id]++; + n_trks_per_seed[orig_param_id]++; // links[step].push_back( // {{previous_step, in_param_id}, item_id}); - links[step].push_back({{previous_step, in_param_id}, item_id, orig_param_id}); + links[step].push_back( + {{previous_step, in_param_id}, item_id, orig_param_id}); /********************************* * Propagate to the next surface From 9fe7d1a27f376741d8eb41ab7f6c94c41a5eaff4 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Mon, 20 Nov 2023 13:11:48 +0100 Subject: [PATCH 03/33] Reduce the number of threads to 2 for gitlab CI build --- .gitlab-ci.yml | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index 7b1cf07d7b..7f97c0d2af 100644 --- a/.gitlab-ci.yml +++ b/.gitlab-ci.yml @@ -18,8 +18,7 @@ build_cuda: -DCMAKE_BUILD_TYPE=Release -DTRACCC_BUILD_TESTING=ON -DTRACCC_BUILD_CUDA=ON - - cmake --build build - + - cmake --build build --parallel 2 test_cuda: stage: test From d99f884eef54a34dddece68b7a19215c3f0d8d68 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Mon, 20 Nov 2023 13:11:53 +0100 Subject: [PATCH 04/33] Avoid racing condition in geometry read/write of Kalman filter tests --- .../tests/kalman_fitting_telescope_test.hpp | 6 ++++-- .../tests/kalman_fitting_wire_chamber_test.hpp | 5 +++-- tests/cpu/test_ckf_sparse_tracks_telescope.cpp | 8 ++++---- tests/cpu/test_kalman_fitter_telescope.cpp | 8 ++++---- tests/cpu/test_kalman_fitter_wire_chamber.cpp | 6 +++--- tests/cuda/test_ckf_sparse_tracks_telescope.cpp | 16 ++++++++-------- tests/cuda/test_kalman_fitter_telescope.cpp | 14 +++++++------- tests/sycl/test_kalman_fitter_telescope.sycl | 14 +++++++------- 8 files changed, 40 insertions(+), 37 deletions(-) diff --git a/tests/common/tests/kalman_fitting_telescope_test.hpp b/tests/common/tests/kalman_fitting_telescope_test.hpp index 71321ef1a3..dff2265771 100644 --- a/tests/common/tests/kalman_fitting_telescope_test.hpp +++ b/tests/common/tests/kalman_fitting_telescope_test.hpp @@ -62,7 +62,8 @@ class KalmanFittingTelescopeTests : public KalmanFittingTests { } protected: - static void SetUpTestCase() { + virtual void SetUp() override { + vecmem::host_memory_resource host_mr; detray::tel_det_config<> tel_cfg{rectangle}; @@ -77,7 +78,8 @@ class KalmanFittingTelescopeTests : public KalmanFittingTests { // Write detector file auto writer_cfg = detray::io::detector_writer_config{} .format(detray::io::format::json) - .replace_files(true); + .replace_files(true) + .path(std::get<0>(GetParam())); detray::io::write_detector(det, name_map, writer_cfg); } }; diff --git a/tests/common/tests/kalman_fitting_wire_chamber_test.hpp b/tests/common/tests/kalman_fitting_wire_chamber_test.hpp index bf262461f4..5b2b2961f5 100644 --- a/tests/common/tests/kalman_fitting_wire_chamber_test.hpp +++ b/tests/common/tests/kalman_fitting_wire_chamber_test.hpp @@ -63,7 +63,7 @@ class KalmanFittingWireChamberTests : public KalmanFittingTests { } protected: - static void SetUpTestCase() { + virtual void SetUp() override { vecmem::host_memory_resource host_mr; detray::wire_chamber_config wire_chamber_cfg; @@ -76,7 +76,8 @@ class KalmanFittingWireChamberTests : public KalmanFittingTests { // Write detector file auto writer_cfg = detray::io::detector_writer_config{} .format(detray::io::format::json) - .replace_files(true); + .replace_files(true) + .path(std::get<0>(GetParam())); detray::io::write_detector(det, name_map, writer_cfg); } }; diff --git a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp index fa1d3b0fd7..045af72d76 100644 --- a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp @@ -78,10 +78,11 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { detray::io::write_detector(det, name_map, writer_cfg); // Read back detector file + const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file("telescope_detector_geometry.json") - .add_file("telescope_detector_homogeneous_material.json") - .add_file("telescope_detector_surface_grids.json"); + reader_cfg.add_file(path + "telescope_detector_geometry.json") + .add_file(path + "telescope_detector_homogeneous_material.json") + .add_file(path + "telescope_detector_surface_grids.json"); const auto [host_det, names] = detray::io::read_detector(host_mr, reader_cfg); @@ -115,7 +116,6 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { typename writer_type::config smearer_writer_cfg{meas_smearer}; // Run simulator - const std::string path = name + "/"; const std::string full_path = io::data_directory() + path; std::filesystem::create_directories(full_path); auto sim = traccc::simulator(host_mr, reader_cfg); @@ -97,7 +98,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { typename writer_type::config smearer_writer_cfg{meas_smearer}; // Run simulator - const std::string path = name + "/"; const std::string full_path = io::data_directory() + path; std::filesystem::create_directories(full_path); auto sim = traccc::simulator(mng_mr, reader_cfg); @@ -127,7 +128,6 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { typename writer_type::config smearer_writer_cfg{meas_smearer}; // Run simulator - const std::string path = name + "/"; const std::string full_path = io::data_directory() + path; std::filesystem::create_directories(full_path); auto sim = traccc::simulator{0.f, 0.f, 0.f}, + "cuda_single_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 1, 5000))); @@ -295,7 +295,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation1, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "double_tracks", std::array{0.f, 0.f, 0.f}, + "cuda_double_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 2, 2500))); @@ -303,7 +303,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation2, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "quadra_tracks", std::array{0.f, 0.f, 0.f}, + "cuda_quadra_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 4, 1250))); @@ -311,7 +311,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation3, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "decade_tracks", std::array{0.f, 0.f, 0.f}, + "cuda_decade_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 10, 500))); \ No newline at end of file diff --git a/tests/cuda/test_kalman_fitter_telescope.cpp b/tests/cuda/test_kalman_fitter_telescope.cpp index c086b97f08..2f8d98ca64 100644 --- a/tests/cuda/test_kalman_fitter_telescope.cpp +++ b/tests/cuda/test_kalman_fitter_telescope.cpp @@ -74,10 +74,11 @@ TEST_P(KalmanFittingTelescopeTests, Run) { vecmem::cuda::managed_memory_resource mng_mr; // Read back detector file + const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file("telescope_detector_geometry.json") - .add_file("telescope_detector_homogeneous_material.json") - .add_file("telescope_detector_surface_grids.json"); + reader_cfg.add_file(path + "telescope_detector_geometry.json") + .add_file(path + "telescope_detector_homogeneous_material.json") + .add_file(path + "telescope_detector_surface_grids.json"); auto [host_det, names] = detray::io::read_detector(mng_mr, reader_cfg); @@ -114,7 +115,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { typename writer_type::config smearer_writer_cfg{meas_smearer}; // Run simulator - const std::string path = name + "/"; const std::string full_path = io::data_directory() + path; std::filesystem::create_directories(full_path); auto sim = traccc::simulator{0.f, 0.f, 0.f}, + "cuda_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -231,7 +231,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "cuda_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -239,7 +239,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "cuda_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); diff --git a/tests/sycl/test_kalman_fitter_telescope.sycl b/tests/sycl/test_kalman_fitter_telescope.sycl index ea8cc78e53..95dec38edb 100644 --- a/tests/sycl/test_kalman_fitter_telescope.sycl +++ b/tests/sycl/test_kalman_fitter_telescope.sycl @@ -95,10 +95,11 @@ TEST_P(KalmanFittingTelescopeTests, Run) { vecmem::sycl::shared_memory_resource shared_mr; // Read back detector file + const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; - reader_cfg.add_file("telescope_detector_geometry.json") - .add_file("telescope_detector_homogeneous_material.json") - .add_file("telescope_detector_surface_grids.json"); + reader_cfg.add_file(path + "telescope_detector_geometry.json") + .add_file(path + "telescope_detector_homogeneous_material.json") + .add_file(path + "telescope_detector_surface_grids.json"); auto [host_det, names] = detray::io::read_detector(shared_mr, reader_cfg); @@ -134,7 +135,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { typename writer_type::config smearer_writer_cfg{meas_smearer}; // Run simulator - const std::string path = name + "/"; const std::string full_path = io::data_directory() + path; std::filesystem::create_directories(full_path); auto sim = traccc::simulator{0.f, 0.f, 0.f}, + "sycl_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -238,7 +238,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "sycl_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -246,7 +246,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "sycl_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); From 03805eac38e6b2c5dca0d865218a62e83657a52d Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Wed, 15 Nov 2023 11:43:00 +0100 Subject: [PATCH 05/33] Fix the track candidate range in finding config --- examples/options/src/options/finding_input_options.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/options/src/options/finding_input_options.cpp b/examples/options/src/options/finding_input_options.cpp index 8f4dc18558..18ae557125 100644 --- a/examples/options/src/options/finding_input_options.cpp +++ b/examples/options/src/options/finding_input_options.cpp @@ -14,7 +14,7 @@ traccc::finding_input_config::finding_input_config( desc.add_options()("track_candidates_range", po::value>() ->value_name("MIN:MAX") - ->default_value({6, 30}), + ->default_value({3, 30}), "Range of track candidates number"); } From 8fac82c011e3feabb75d9e73547ecb5f12715033 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Tue, 21 Nov 2023 16:52:19 +0100 Subject: [PATCH 06/33] Better thread balance in CKF find_tracks_kernel --- .../include/traccc/finding/finding_config.hpp | 6 +- device/common/CMakeLists.txt | 4 +- .../edm/device/finding_global_counter.hpp | 7 +- .../finding/device/count_measurements.hpp | 42 ++++++++ .../traccc/finding/device/count_threads.hpp | 43 -------- .../traccc/finding/device/find_tracks.hpp | 16 ++- .../device/impl/count_measurements.ipp | 51 +++++++++ .../finding/device/impl/count_threads.ipp | 67 ------------ .../finding/device/impl/find_tracks.ipp | 90 ++++++---------- device/cuda/src/finding/finding_algorithm.cu | 102 ++++++++---------- 10 files changed, 189 insertions(+), 239 deletions(-) create mode 100644 device/common/include/traccc/finding/device/count_measurements.hpp delete mode 100644 device/common/include/traccc/finding/device/count_threads.hpp create mode 100644 device/common/include/traccc/finding/device/impl/count_measurements.ipp delete mode 100644 device/common/include/traccc/finding/device/impl/count_threads.ipp diff --git a/core/include/traccc/finding/finding_config.hpp b/core/include/traccc/finding/finding_config.hpp index a484e98df3..26067fa826 100644 --- a/core/include/traccc/finding/finding_config.hpp +++ b/core/include/traccc/finding/finding_config.hpp @@ -37,9 +37,9 @@ struct finding_config { /// @TODO: Make a separate file for propagation config? scalar_t constrained_step_size = std::numeric_limits::max(); - /// GPU-specific parameter to evaluate the number of measurements to be - /// iterated per track - unsigned int n_avg_threads_per_track = 4; + /// GPU-specific parameter for the number of measurements to be + /// iterated per thread + unsigned int n_measurements_per_thread = 8; }; } // namespace traccc \ No newline at end of file diff --git a/device/common/CMakeLists.txt b/device/common/CMakeLists.txt index f2b36d19da..f33903f976 100644 --- a/device/common/CMakeLists.txt +++ b/device/common/CMakeLists.txt @@ -61,13 +61,13 @@ traccc_add_library( traccc_device_common device_common TYPE SHARED # Track finding funtions(s). "include/traccc/finding/device/apply_interaction.hpp" "include/traccc/finding/device/build_tracks.hpp" - "include/traccc/finding/device/count_threads.hpp" + "include/traccc/finding/device/count_measurements.hpp" "include/traccc/finding/device/find_tracks.hpp" "include/traccc/finding/device/make_barcode_sequence.hpp" "include/traccc/finding/device/propagate_to_next_surface.hpp" "include/traccc/finding/device/impl/apply_interaction.ipp" "include/traccc/finding/device/impl/build_tracks.ipp" - "include/traccc/finding/device/impl/count_threads.ipp" + "include/traccc/finding/device/impl/count_measurements.ipp" "include/traccc/finding/device/impl/find_tracks.ipp" "include/traccc/finding/device/impl/make_barcode_sequence.ipp" "include/traccc/finding/device/impl/propagate_to_next_surface.ipp" diff --git a/device/common/include/traccc/edm/device/finding_global_counter.hpp b/device/common/include/traccc/edm/device/finding_global_counter.hpp index d72ac4e6d4..977052ad12 100644 --- a/device/common/include/traccc/edm/device/finding_global_counter.hpp +++ b/device/common/include/traccc/edm/device/finding_global_counter.hpp @@ -11,11 +11,8 @@ namespace traccc::device { struct finding_global_counter { - // Number of measurements per thread - unsigned int n_measurements_per_thread; - - // Number of threads for find_track kernel - unsigned int n_total_threads; + // Sum of the number of measurements for every parameter + unsigned int n_measurements_sum; // Number of found measurements for the current step unsigned int n_candidates; diff --git a/device/common/include/traccc/finding/device/count_measurements.hpp b/device/common/include/traccc/finding/device/count_measurements.hpp new file mode 100644 index 0000000000..eb3f603662 --- /dev/null +++ b/device/common/include/traccc/finding/device/count_measurements.hpp @@ -0,0 +1,42 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +// Project include(s). +#include "traccc/definitions/qualifiers.hpp" + +namespace traccc::device { + +/// Function evalulating the number of measurements to be iterated per thread +/// and the total number of measurements +/// +/// @param[in] globalIndex The index of the current thread +/// @param[in] params_view Input parameters view object +/// @param[in] barcodes_view Barcodes view object +/// @param[in] upper_bounds Upper bounds of measurements w.r.t geometry +/// ID +/// @param[out] n_measurements_view The number of measurements per parameter +/// @param[out] ref_meas_idx The first index of measurements per +/// parameter +/// @param[out] n_measurements_sum The sum of the number of measurements per +/// parameter +/// +TRACCC_DEVICE inline void count_measurements( + std::size_t globalIndex, + bound_track_parameters_collection_types::const_view params_view, + vecmem::data::vector_view barcodes_view, + vecmem::data::vector_view upper_bounds_view, + const unsigned int n_in_params, + vecmem::data::vector_view n_measurements_view, + vecmem::data::vector_view ref_meas_idx_view, + unsigned int& n_measurements_sum); + +} // namespace traccc::device + +// Include the implementation. +#include "traccc/finding/device/impl/count_measurements.ipp" \ No newline at end of file diff --git a/device/common/include/traccc/finding/device/count_threads.hpp b/device/common/include/traccc/finding/device/count_threads.hpp deleted file mode 100644 index 2fe43ed619..0000000000 --- a/device/common/include/traccc/finding/device/count_threads.hpp +++ /dev/null @@ -1,43 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2023 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -#pragma once - -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" - -namespace traccc::device { - -/// Function evalulating the number of measurements to be iterated per thread -/// and the total number of threads required -/// -/// @param[in] globalIndex The index of the current thread -/// @param[in] cfg Track finding config object -/// @param[in] params_view Input parameters view object -/// @param[in] barcodes_view Barcodes view object -/// @param[in] sizes_view The number of measurements of each module -/// @param[in] n_in_params The number of input track parameters -/// @param[in] n_total_measurements Total number of meausurments -/// @param[out] n_threads_view The number of threads per tracks -/// @param[out] n_measurements_per_thread The number of measurements to be -/// iterated per thread -/// @param[out] n_total_thread Total number of threads -/// -template -TRACCC_DEVICE inline void count_threads( - std::size_t globalIndex, const config_t cfg, - bound_track_parameters_collection_types::const_view params_view, - vecmem::data::vector_view barcodes_view, - vecmem::data::vector_view sizes_view, - const int n_in_params, const int n_total_measurements, - vecmem::data::vector_view n_threads_view, - unsigned int& n_measurements_per_thread, unsigned int& n_total_threads); - -} // namespace traccc::device - -// Include the implementation. -#include "traccc/finding/device/impl/count_threads.ipp" \ No newline at end of file diff --git a/device/common/include/traccc/finding/device/find_tracks.hpp b/device/common/include/traccc/finding/device/find_tracks.hpp index b38f187094..ea0a9f27f3 100644 --- a/device/common/include/traccc/finding/device/find_tracks.hpp +++ b/device/common/include/traccc/finding/device/find_tracks.hpp @@ -23,14 +23,13 @@ namespace traccc::device { /// @param[in] cfg Track finding config object /// @param[in] det_data Detector view object /// @param[in] measurements_view Measurements container view -/// @param[in] barcodes_view Barcode sequence view object /// @param[in] upper_bounds_view Upper bounds of measurements unique w.r.t /// barcode /// @param[in] in_params_view Input parameters -/// @param[in] n_threads_view The number of threads per tracks +/// @param[in] n_measurements_prefix_sum_view Prefix sum of the number of +/// measurements per parameter +/// @param[in] ref_meas_idx_view The first index of measurements per parameter /// @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 @@ -41,12 +40,11 @@ TRACCC_DEVICE inline void find_tracks( std::size_t globalIndex, const config_t cfg, typename detector_t::view_type det_data, measurement_collection_types::const_view measurements_view, - vecmem::data::vector_view barcodes_view, - vecmem::data::vector_view upper_bounds_view, bound_track_parameters_collection_types::const_view in_params_view, - vecmem::data::vector_view n_threads_view, - const unsigned int step, const unsigned int& n_measurements_per_thread, - const unsigned int& n_total_threads, const unsigned int& n_max_candidates, + vecmem::data::vector_view + n_measurements_prefix_sum_view, + vecmem::data::vector_view ref_meas_idx_view, + const unsigned int step, const unsigned int& n_max_candidates, bound_track_parameters_collection_types::view out_params_view, vecmem::data::vector_view links_view, unsigned int& n_candidates); diff --git a/device/common/include/traccc/finding/device/impl/count_measurements.ipp b/device/common/include/traccc/finding/device/impl/count_measurements.ipp new file mode 100644 index 0000000000..058b76445e --- /dev/null +++ b/device/common/include/traccc/finding/device/impl/count_measurements.ipp @@ -0,0 +1,51 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2023 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +#pragma once + +namespace traccc::device { + +TRACCC_DEVICE inline void count_measurements( + std::size_t globalIndex, + bound_track_parameters_collection_types::const_view params_view, + vecmem::data::vector_view barcodes_view, + vecmem::data::vector_view upper_bounds_view, + const unsigned int n_in_params, + vecmem::data::vector_view n_measurements_view, + vecmem::data::vector_view ref_meas_idx_view, + unsigned int& n_measurements_sum) { + + bound_track_parameters_collection_types::const_device params(params_view); + vecmem::device_vector barcodes( + barcodes_view); + vecmem::device_vector upper_bounds(upper_bounds_view); + vecmem::device_vector n_measurements(n_measurements_view); + vecmem::device_vector ref_meas_idx(ref_meas_idx_view); + + if (globalIndex >= n_in_params) { + return; + } + + // Get barcode + const auto bcd = params.at(globalIndex).surface_link(); + const auto lo = + thrust::lower_bound(thrust::seq, barcodes.begin(), barcodes.end(), bcd); + const auto bcd_id = std::distance(barcodes.begin(), lo); + + // Get the reference measurement index and the number of measurements per + // parameter + ref_meas_idx.at(globalIndex) = + lo == barcodes.begin() ? 0u : upper_bounds[bcd_id - 1]; + n_measurements.at(globalIndex) = + upper_bounds[bcd_id] - ref_meas_idx.at(globalIndex); + + // Increase the total number of measurements with atomic addition + vecmem::device_atomic_ref n_meas_sum(n_measurements_sum); + n_meas_sum.fetch_add(n_measurements.at(globalIndex)); +} + +} // namespace traccc::device diff --git a/device/common/include/traccc/finding/device/impl/count_threads.ipp b/device/common/include/traccc/finding/device/impl/count_threads.ipp deleted file mode 100644 index e232ecf4a3..0000000000 --- a/device/common/include/traccc/finding/device/impl/count_threads.ipp +++ /dev/null @@ -1,67 +0,0 @@ -/** TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2023 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -#pragma once - -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" - -namespace traccc::device { - -template -TRACCC_DEVICE inline void count_threads( - std::size_t globalIndex, const config_t cfg, - bound_track_parameters_collection_types::const_view params_view, - vecmem::data::vector_view barcodes_view, - vecmem::data::vector_view sizes_view, - const int n_in_params, const int n_total_measurements, - vecmem::data::vector_view n_threads_view, - unsigned int& n_measurements_per_thread, unsigned int& n_total_threads) { - - bound_track_parameters_collection_types::const_device params(params_view); - vecmem::device_vector barcodes( - barcodes_view); - vecmem::device_vector sizes(sizes_view); - vecmem::device_vector n_threads(n_threads_view); - - const unsigned int n_params = params.size(); - - if (globalIndex >= n_in_params) { - return; - } - - // Get barcode - const auto bcd = params.at(globalIndex).surface_link(); - - // Search for the corresponding index of unique vector - const auto lower = - thrust::lower_bound(thrust::seq, barcodes.begin(), barcodes.end(), bcd); - const auto idx = thrust::distance(barcodes.begin(), lower); - - // The averaged number of measurement per track - const unsigned int n_avg_meas_per_track = - (n_total_measurements + n_params - 1) / n_params; - - // Estimate how many measurements should be handled per thread in CKF - const unsigned int n_meas_per_thread = - (n_avg_meas_per_track + cfg.n_avg_threads_per_track - 1) / - cfg.n_avg_threads_per_track; - - if (globalIndex == 0) { - n_measurements_per_thread = n_meas_per_thread; - } - - // Set the number of threads assigned per track - n_threads.at(globalIndex) = - (sizes.at(idx) + n_meas_per_thread - 1) / n_meas_per_thread; - - // Estimate the total number of threads we need for CKF - vecmem::device_atomic_ref num_total_threads(n_total_threads); - num_total_threads.fetch_add(n_threads.at(globalIndex)); -} - -} // namespace traccc::device diff --git a/device/common/include/traccc/finding/device/impl/find_tracks.ipp b/device/common/include/traccc/finding/device/impl/find_tracks.ipp index c051bc7de1..20d4c1ef0f 100644 --- a/device/common/include/traccc/finding/device/impl/find_tracks.ipp +++ b/device/common/include/traccc/finding/device/impl/find_tracks.ipp @@ -20,28 +20,20 @@ TRACCC_DEVICE inline void find_tracks( std::size_t globalIndex, const config_t cfg, typename detector_t::view_type det_data, measurement_collection_types::const_view measurements_view, - vecmem::data::vector_view barcodes_view, - vecmem::data::vector_view upper_bounds_view, bound_track_parameters_collection_types::const_view in_params_view, - vecmem::data::vector_view n_threads_view, - const unsigned int step, const unsigned int& n_measurements_per_thread, - const unsigned int& n_total_threads, const unsigned int& n_max_candidates, + vecmem::data::vector_view + n_measurements_prefix_sum_view, + vecmem::data::vector_view ref_meas_idx_view, + const unsigned int step, const unsigned int& n_max_candidates, bound_track_parameters_collection_types::view out_params_view, vecmem::data::vector_view links_view, unsigned int& n_candidates) { - if (globalIndex >= n_total_threads) { - return; - } - // Detector detector_t det(det_data); // Measurement measurement_collection_types::const_device measurements(measurements_view); - vecmem::device_vector barcodes( - barcodes_view); - vecmem::device_vector upper_bounds(upper_bounds_view); // Input parameters bound_track_parameters_collection_types::const_device in_params( @@ -53,65 +45,53 @@ TRACCC_DEVICE inline void find_tracks( // Links vecmem::device_vector links(links_view); - // n threads - vecmem::device_vector n_threads(n_threads_view); - - // Search for out_param index - const auto lo1 = thrust::lower_bound(thrust::seq, n_threads.begin(), - n_threads.end(), globalIndex + 1); - const auto in_param_id = std::distance(n_threads.begin(), lo1); - - // Get measurements on surface - unsigned int ref; - if (lo1 == n_threads.begin()) { - ref = 0; - } else { - ref = *(lo1 - 1); - } - - // Get barcode - const auto bcd = in_params.at(in_param_id).surface_link(); - - const auto lo2 = - thrust::lower_bound(thrust::seq, barcodes.begin(), barcodes.end(), bcd); - const auto bcd_id = std::distance(barcodes.begin(), lo2); - - // Find the range for measurements - std::pair range; - if (lo2 == barcodes.begin()) { - range.first = 0u; - range.second = upper_bounds[bcd_id]; - } else { - range.first = upper_bounds[bcd_id - 1]; - range.second = upper_bounds[bcd_id]; - } - - const unsigned int offset = globalIndex - ref; - const unsigned int stride = offset * n_measurements_per_thread; - const unsigned int n_meas_on_surface = range.second - range.first; + // Prefix sum of the number of measurements per parameter + vecmem::device_vector n_measurements_prefix_sum( + n_measurements_prefix_sum_view); - // Iterate over the measurements - const detray::surface sf{det, bcd}; + // Reference (first) measurement index per parameter + vecmem::device_vector ref_meas_idx(ref_meas_idx_view); // Last step ID const unsigned int previous_step = (step == 0) ? std::numeric_limits::max() : step - 1; - for (unsigned int i = 0; i < n_measurements_per_thread; i++) { - if (i + stride >= n_meas_on_surface) { + const unsigned int n_measurements_sum = n_measurements_prefix_sum.back(); + const unsigned int stride = globalIndex * cfg.n_measurements_per_thread; + + vecmem::device_vector::iterator lo1; + + for (unsigned int i_meas = 0; i_meas < cfg.n_measurements_per_thread; + i_meas++) { + const unsigned int idx = stride + i_meas; + + if (idx >= n_measurements_sum) { break; } - const unsigned int meas_idx = i + stride + range.first; + if (i_meas == 0 || idx == *lo1) { + lo1 = thrust::lower_bound(thrust::seq, + n_measurements_prefix_sum.begin(), + n_measurements_prefix_sum.end(), idx + 1); + } + + const unsigned int in_param_id = + std::distance(n_measurements_prefix_sum.begin(), lo1); + const detray::geometry::barcode bcd = + in_params.at(in_param_id).surface_link(); + const unsigned int offset = + lo1 == n_measurements_prefix_sum.begin() ? idx : idx - *(lo1 - 1); + const unsigned int meas_idx = ref_meas_idx.at(in_param_id) + offset; bound_track_parameters in_par = in_params.at(in_param_id); - const auto meas = measurements.at(meas_idx); + + const auto& meas = measurements.at(meas_idx); track_state trk_state(meas); + const detray::surface sf{det, bcd}; // Run the Kalman update sf.template visit_mask< gain_matrix_updater>(trk_state, in_par); - // Get the chi-square const auto chi2 = trk_state.filtered_chi2(); diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu index ae51f2daf2..795b24ce6c 100644 --- a/device/cuda/src/finding/finding_algorithm.cu +++ b/device/cuda/src/finding/finding_algorithm.cu @@ -14,7 +14,7 @@ #include "traccc/finding/candidate_link.hpp" #include "traccc/finding/device/apply_interaction.hpp" #include "traccc/finding/device/build_tracks.hpp" -#include "traccc/finding/device/count_threads.hpp" +#include "traccc/finding/device/count_measurements.hpp" #include "traccc/finding/device/find_tracks.hpp" #include "traccc/finding/device/make_barcode_sequence.hpp" #include "traccc/finding/device/propagate_to_next_surface.hpp" @@ -75,23 +75,21 @@ __global__ void apply_interaction( n_params, params_view); } -/// CUDA kernel for running @c traccc::device::count_threads -template -__global__ void count_threads( - const config_t cfg, +/// CUDA kernel for running @c traccc::device::count_measurements +__global__ void count_measurements( bound_track_parameters_collection_types::const_view params_view, vecmem::data::vector_view barcodes_view, - vecmem::data::vector_view sizes_view, - const int n_in_params, const int n_total_measurements, - vecmem::data::vector_view n_threads_view, - unsigned int& n_measurements_per_thread, unsigned int& n_total_threads) { + vecmem::data::vector_view upper_bounds_view, + const unsigned int n_in_params, + vecmem::data::vector_view n_measurements_view, + vecmem::data::vector_view ref_meas_idx_view, + unsigned int& n_measurements_sum) { int gid = threadIdx.x + blockIdx.x * blockDim.x; - device::count_threads(gid, cfg, params_view, barcodes_view, - sizes_view, n_in_params, - n_total_measurements, n_threads_view, - n_measurements_per_thread, n_total_threads); + device::count_measurements( + gid, params_view, barcodes_view, upper_bounds_view, n_in_params, + n_measurements_view, ref_meas_idx_view, n_measurements_sum); } /// CUDA kernel for running @c traccc::device::find_tracks @@ -99,12 +97,11 @@ template __global__ void find_tracks( const config_t cfg, typename detector_t::view_type det_data, measurement_collection_types::const_view measurements_view, - vecmem::data::vector_view barcodes_view, - vecmem::data::vector_view upper_bounds_view, bound_track_parameters_collection_types::const_view in_params_view, - vecmem::data::vector_view n_threads_view, - const unsigned int step, const unsigned int& n_measurements_per_thread, - const unsigned int& n_total_threads, const unsigned int n_max_candidates, + vecmem::data::vector_view + n_measurements_prefix_sum_view, + vecmem::data::vector_view ref_meas_idx_view, + const unsigned int step, const unsigned int n_max_candidates, bound_track_parameters_collection_types::view out_params_view, vecmem::data::vector_view links_view, unsigned int& n_candidates) { @@ -112,10 +109,9 @@ __global__ void find_tracks( int gid = threadIdx.x + blockIdx.x * blockDim.x; device::find_tracks( - 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, n_max_candidates, out_params_view, links_view, - n_candidates); + gid, cfg, det_data, measurements_view, in_params_view, + n_measurements_prefix_sum_view, ref_meas_idx_view, step, + n_max_candidates, out_params_view, links_view, n_candidates); } /// CUDA kernel for running @c traccc::device::propagate_to_next_surface @@ -248,17 +244,6 @@ finding_algorithm::operator()( uniques.begin(), uniques.begin() + n_modules, upper_bounds.begin(), measurement_sort_comp()); - // Get the number of measurements of each module - vecmem::data::vector_buffer sizes_buffer{n_modules, - m_mr.main}; - vecmem::device_vector sizes(sizes_buffer); - thrust::adjacent_difference(thrust::cuda::par.on(stream), - upper_bounds.begin(), upper_bounds.end(), - sizes.begin()); - - // Number of total measurements - const unsigned int n_total_measurements = measurements_device.size(); - /***************************************************************** * Kernel1: Create barcode sequence *****************************************************************/ @@ -312,20 +297,22 @@ finding_algorithm::operator()( CUDA_ERROR_CHECK(cudaGetLastError()); /***************************************************************** - * Kernel3: Count the number of threads per parameter + * Kernel3: Count the number of measurements per parameter ****************************************************************/ - // Create a buffer for the number of threads per parameter - vecmem::data::vector_buffer n_threads_buffer(n_in_params, - m_mr.main); + vecmem::data::vector_buffer n_measurements_buffer( + n_in_params, m_mr.main); + + // Create a buffer for the first measurement index of parameter + vecmem::data::vector_buffer ref_meas_idx_buffer( + n_in_params, m_mr.main); + nThreads = WARP_SIZE * 2; nBlocks = (n_in_params + nThreads - 1) / nThreads; - - kernels::count_threads<<>>( - m_cfg, in_params_buffer, barcodes_buffer, sizes_buffer, n_in_params, - n_total_measurements, n_threads_buffer, - (*global_counter_device).n_measurements_per_thread, - (*global_counter_device).n_total_threads); + kernels::count_measurements<<>>( + in_params_buffer, barcodes_buffer, upper_bounds_buffer, n_in_params, + n_measurements_buffer, ref_meas_idx_buffer, + (*global_counter_device).n_measurements_sum); CUDA_ERROR_CHECK(cudaGetLastError()); // Global counter object: Device -> Host @@ -336,10 +323,17 @@ finding_algorithm::operator()( m_stream.synchronize(); - // Get Prefix Sum of n_thread vector - vecmem::device_vector n_threads(n_threads_buffer); - thrust::inclusive_scan(thrust::cuda::par.on(stream), n_threads.begin(), - n_threads.end(), n_threads.begin()); + // Create the buffer for the prefix sum of the number of measurements + // per parameter + vecmem::device_vector n_measurements( + n_measurements_buffer); + vecmem::data::vector_buffer + n_measurements_prefix_sum_buffer(n_in_params, m_mr.main); + vecmem::device_vector n_measurements_prefix_sum( + n_measurements_prefix_sum_buffer); + thrust::inclusive_scan(thrust::cuda::par.on(stream), + n_measurements.begin(), n_measurements.end(), + n_measurements_prefix_sum.begin()); /***************************************************************** * Kernel4: Find valid tracks @@ -347,7 +341,6 @@ finding_algorithm::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); @@ -359,17 +352,16 @@ finding_algorithm::operator()( link_map[step] = {n_in_params * m_cfg.max_num_branches_per_surface, m_mr.main}; m_copy.setup(link_map[step]); - nBlocks = - (global_counter_host.n_total_threads + nThreads - 1) / nThreads; + nBlocks = (global_counter_host.n_measurements_sum + + nThreads * m_cfg.n_measurements_per_thread - 1) / + (nThreads * m_cfg.n_measurements_per_thread); if (nBlocks > 0) { kernels::find_tracks <<>>( - 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, n_max_candidates, - updated_params_buffer, link_map[step], + m_cfg, det_view, measurements, in_params_buffer, + n_measurements_prefix_sum_buffer, ref_meas_idx_buffer, step, + n_max_candidates, updated_params_buffer, link_map[step], (*global_counter_device).n_candidates); CUDA_ERROR_CHECK(cudaGetLastError()); } From 276cc44d12ecc7cc876d93689be3f7281d5c24cb Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Wed, 22 Nov 2023 20:22:02 +0100 Subject: [PATCH 07/33] Update detray version to v49 --- .../seeding/experimental/spacepoint_formation.ipp | 2 +- .../device/experimental/impl/form_spacepoints.ipp | 2 +- extern/detray/CMakeLists.txt | 2 +- .../common/tests/kalman_fitting_wire_chamber_test.hpp | 5 +++++ tests/cpu/test_ckf_sparse_tracks_telescope.cpp | 11 +++++------ tests/cpu/test_kalman_fitter_telescope.cpp | 9 ++++----- tests/cpu/test_kalman_fitter_wire_chamber.cpp | 8 ++++---- tests/cuda/test_ckf_sparse_tracks_telescope.cpp | 11 +++++------ tests/cuda/test_kalman_fitter_telescope.cpp | 9 ++++----- tests/sycl/test_kalman_fitter_telescope.sycl | 9 ++++----- 10 files changed, 34 insertions(+), 34 deletions(-) diff --git a/core/include/traccc/seeding/experimental/spacepoint_formation.ipp b/core/include/traccc/seeding/experimental/spacepoint_formation.ipp index 1c703304e8..438354a8f0 100644 --- a/core/include/traccc/seeding/experimental/spacepoint_formation.ipp +++ b/core/include/traccc/seeding/experimental/spacepoint_formation.ipp @@ -34,7 +34,7 @@ spacepoint_collection_types::host spacepoint_formation::operator()( // This local to global transformation only works for 2D planar // measurement // (e.g. barrel pixel and endcap pixel detector) - const auto global = sf.local_to_global({}, ms.local, {}); + const auto global = sf.bound_to_global({}, ms.local, {}); // Fill result with this spacepoint result.push_back({global, ms}); diff --git a/device/common/include/traccc/seeding/device/experimental/impl/form_spacepoints.ipp b/device/common/include/traccc/seeding/device/experimental/impl/form_spacepoints.ipp index 5ff2134f18..42de4905a6 100644 --- a/device/common/include/traccc/seeding/device/experimental/impl/form_spacepoints.ipp +++ b/device/common/include/traccc/seeding/device/experimental/impl/form_spacepoints.ipp @@ -39,7 +39,7 @@ TRACCC_HOST_DEVICE inline void form_spacepoints( // This local to global transformation only works for 2D planar measurement // (e.g. barrel pixel and endcap pixel detector) - const auto global = sf.local_to_global({}, ms.local, {}); + const auto global = sf.bound_to_global({}, ms.local, {}); spacepoints.push_back({global, ms}); } diff --git a/extern/detray/CMakeLists.txt b/extern/detray/CMakeLists.txt index af094a3fbb..c07b48bd33 100644 --- a/extern/detray/CMakeLists.txt +++ b/extern/detray/CMakeLists.txt @@ -18,7 +18,7 @@ message( STATUS "Building Detray as part of the TRACCC project" ) # Declare where to get Detray from. set( TRACCC_DETRAY_SOURCE -"URL;https://github.com/acts-project/detray/archive/refs/tags/v0.47.0.tar.gz;URL_MD5;200f7fb00ac1d16b22e7330da62d03fb" +"URL;https://github.com/acts-project/detray/archive/refs/tags/v0.49.0.tar.gz;URL_MD5;d017a48de31e2ee116af870c917bce7f" CACHE STRING "Source for Detray, when built as part of this project" ) mark_as_advanced( TRACCC_DETRAY_SOURCE ) diff --git a/tests/common/tests/kalman_fitting_wire_chamber_test.hpp b/tests/common/tests/kalman_fitting_wire_chamber_test.hpp index 5b2b2961f5..d10bef8166 100644 --- a/tests/common/tests/kalman_fitting_wire_chamber_test.hpp +++ b/tests/common/tests/kalman_fitting_wire_chamber_test.hpp @@ -70,6 +70,11 @@ class KalmanFittingWireChamberTests : public KalmanFittingTests { wire_chamber_cfg.n_layers(n_wire_layers); wire_chamber_cfg.half_z(half_z); + wire_chamber_cfg.mapped_material(detray::vacuum()); + //@NOTE: 2 GeV test fails in pull check with the following setup + // wire_chamber_cfg.mapped_material(detray::beryllium()); + // wire_chamber_cfg.m_thickness = 100.f * detray::unit::um; + // Create telescope detector auto [det, name_map] = create_wire_chamber(host_mr, wire_chamber_cfg); diff --git a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp index 045af72d76..f476ff56b6 100644 --- a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp @@ -81,8 +81,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json") - .add_file(path + "telescope_detector_surface_grids.json"); + .add_file(path + "telescope_detector_homogeneous_material.json"); const auto [host_det, names] = detray::io::read_detector(host_mr, reader_cfg); @@ -221,7 +220,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation0, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "single_tracks", std::array{0.f, 0.f, 0.f}, + "telescope_single_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 1, 5000))); @@ -229,7 +228,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation1, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "double_tracks", std::array{0.f, 0.f, 0.f}, + "telescope_double_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 2, 2500))); @@ -237,7 +236,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation2, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "quadra_tracks", std::array{0.f, 0.f, 0.f}, + "telescope_quadra_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 4, 1250))); @@ -245,7 +244,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation3, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "decade_tracks", std::array{0.f, 0.f, 0.f}, + "telescope_decade_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 10, 500))); diff --git a/tests/cpu/test_kalman_fitter_telescope.cpp b/tests/cpu/test_kalman_fitter_telescope.cpp index 1925b7cb1d..e415b62759 100644 --- a/tests/cpu/test_kalman_fitter_telescope.cpp +++ b/tests/cpu/test_kalman_fitter_telescope.cpp @@ -64,8 +64,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json") - .add_file(path + "telescope_detector_surface_grids.json"); + .add_file(path + "telescope_detector_homogeneous_material.json"); const auto [host_det, names] = detray::io::read_detector(host_mr, reader_cfg); @@ -178,7 +177,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation0, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "telescope_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -186,7 +185,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "telescope_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -194,7 +193,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "telescope_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); \ No newline at end of file diff --git a/tests/cpu/test_kalman_fitter_wire_chamber.cpp b/tests/cpu/test_kalman_fitter_wire_chamber.cpp index 79aedfca32..37808a89a7 100644 --- a/tests/cpu/test_kalman_fitter_wire_chamber.cpp +++ b/tests/cpu/test_kalman_fitter_wire_chamber.cpp @@ -63,7 +63,7 @@ TEST_P(KalmanFittingWireChamberTests, Run) { detray::io::detector_reader_config reader_cfg{}; reader_cfg.add_file(path + "wire_chamber_geometry.json") .add_file(path + "wire_chamber_homogeneous_material.json") - //.add_file("wire_chamber_surface_grids.json") + //.add_file(path + "wire_chamber_surface_grids.json") .do_check(true); const auto [host_det, names] = @@ -185,7 +185,7 @@ TEST_P(KalmanFittingWireChamberTests, Run) { INSTANTIATE_TEST_SUITE_P( KalmanFitWireChamberValidation0, KalmanFittingWireChamberTests, ::testing::Values(std::make_tuple( - "2_GeV", std::array{0.f, 0.f, 0.f}, + "wire_2_GeV", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{2.f, 2.f}, std::array{-1.f, 1.f}, std::array{0.f, 2.0f * detray::constant::pi}, 100, @@ -194,7 +194,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitWireChamberValidation1, KalmanFittingWireChamberTests, ::testing::Values(std::make_tuple( - "10_GeV", std::array{0.f, 0.f, 0.f}, + "wire_10_GeV", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{-1.f, 1.f}, std::array{0.f, 2.0f * detray::constant::pi}, 100, @@ -203,7 +203,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitWireChamberValidation2, KalmanFittingWireChamberTests, ::testing::Values(std::make_tuple( - "100_GeV", std::array{0.f, 0.f, 0.f}, + "wire_100_GeV", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{-1.f, 1.f}, std::array{0.f, 2.0f * detray::constant::pi}, 100, diff --git a/tests/cuda/test_ckf_sparse_tracks_telescope.cpp b/tests/cuda/test_ckf_sparse_tracks_telescope.cpp index caec252ad1..cb5c271cb8 100644 --- a/tests/cuda/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cuda/test_ckf_sparse_tracks_telescope.cpp @@ -91,8 +91,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json") - .add_file(path + "telescope_detector_surface_grids.json"); + .add_file(path + "telescope_detector_homogeneous_material.json"); auto [host_det, names] = detray::io::read_detector(mng_mr, reader_cfg); @@ -287,7 +286,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation0, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "cuda_single_tracks", std::array{0.f, 0.f, 0.f}, + "cuda_telescope_single_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 1, 5000))); @@ -295,7 +294,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation1, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "cuda_double_tracks", std::array{0.f, 0.f, 0.f}, + "cuda_telescope_double_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 2, 2500))); @@ -303,7 +302,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation2, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "cuda_quadra_tracks", std::array{0.f, 0.f, 0.f}, + "cuda_telescope_quadra_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 4, 1250))); @@ -311,7 +310,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation3, CkfSparseTrackTelescopeTests, ::testing::Values(std::make_tuple( - "cuda_decade_tracks", std::array{0.f, 0.f, 0.f}, + "cuda_telescope_decade_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 10, 500))); \ No newline at end of file diff --git a/tests/cuda/test_kalman_fitter_telescope.cpp b/tests/cuda/test_kalman_fitter_telescope.cpp index 2f8d98ca64..1e60eccc66 100644 --- a/tests/cuda/test_kalman_fitter_telescope.cpp +++ b/tests/cuda/test_kalman_fitter_telescope.cpp @@ -77,8 +77,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json") - .add_file(path + "telescope_detector_surface_grids.json"); + .add_file(path + "telescope_detector_homogeneous_material.json"); auto [host_det, names] = detray::io::read_detector(mng_mr, reader_cfg); @@ -223,7 +222,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation0, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "cuda_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "cuda_telescope_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -231,7 +230,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "cuda_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "cuda_telescope_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -239,7 +238,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "cuda_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "cuda_telescope_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); diff --git a/tests/sycl/test_kalman_fitter_telescope.sycl b/tests/sycl/test_kalman_fitter_telescope.sycl index 95dec38edb..2ebde173aa 100644 --- a/tests/sycl/test_kalman_fitter_telescope.sycl +++ b/tests/sycl/test_kalman_fitter_telescope.sycl @@ -98,8 +98,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; reader_cfg.add_file(path + "telescope_detector_geometry.json") - .add_file(path + "telescope_detector_homogeneous_material.json") - .add_file(path + "telescope_detector_surface_grids.json"); + .add_file(path + "telescope_detector_homogeneous_material.json"); auto [host_det, names] = detray::io::read_detector(shared_mr, reader_cfg); @@ -230,7 +229,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation0, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "sycl_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "sycl_telescope_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -238,7 +237,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "sycl_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "sycl_telescope_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); @@ -246,7 +245,7 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "sycl_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "sycl_telescope_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, 100))); From 5110cd8d73e7159302faeb08b0e15a3c83d9e2e2 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Fri, 24 Nov 2023 01:26:16 +0100 Subject: [PATCH 08/33] Include simulation directory in check format --- .github/check_format.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/check_format.sh b/.github/check_format.sh index d03763e494..16f08942bf 100755 --- a/.github/check_format.sh +++ b/.github/check_format.sh @@ -7,7 +7,7 @@ set -e # abort on error -INCLUDE_DIRS=("core device examples io tests performance plugins") +INCLUDE_DIRS=("core device examples io tests performance plugins simulation") if [ $# -ne 1 ]; then echo "wrong number of arguments" From 0a9a63c79ee8f1691edf4e63e8baa78c9a8336c2 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Fri, 24 Nov 2023 20:48:39 +0100 Subject: [PATCH 09/33] Add a options for charge and corresponding tests --- .../traccc/options/particle_gen_options.hpp | 6 +++ examples/simulation/simulate.cpp | 1 + examples/simulation/simulate_telescope.cpp | 1 + examples/simulation/simulate_toy_detector.cpp | 1 + examples/simulation/simulate_wire_chamber.cpp | 1 + tests/common/tests/kalman_fitting_test.hpp | 7 ++-- .../cpu/test_ckf_sparse_tracks_telescope.cpp | 14 ++++--- tests/cpu/test_kalman_fitter_telescope.cpp | 30 ++++++++++---- tests/cpu/test_kalman_fitter_wire_chamber.cpp | 33 ++++++++++----- tests/cpu/test_simulation.cpp | 41 ++++++++++--------- .../cuda/test_ckf_sparse_tracks_telescope.cpp | 14 ++++--- tests/cuda/test_kalman_fitter_telescope.cpp | 14 ++++--- tests/sycl/test_kalman_fitter_telescope.sycl | 14 ++++--- 13 files changed, 111 insertions(+), 66 deletions(-) diff --git a/examples/options/include/traccc/options/particle_gen_options.hpp b/examples/options/include/traccc/options/particle_gen_options.hpp index 01c913a621..4be4d29b4e 100644 --- a/examples/options/include/traccc/options/particle_gen_options.hpp +++ b/examples/options/include/traccc/options/particle_gen_options.hpp @@ -29,6 +29,7 @@ struct particle_gen_options { Reals mom_range{1., 1.}; Reals phi_range{0., 0.}; Reals theta_range{0., 0.}; + scalar_t charge{-1.f}; particle_gen_options(po::options_description& desc) { desc.add_options()("gen-nparticles", @@ -59,6 +60,10 @@ struct particle_gen_options { ->value_name("MIN:MAX") ->default_value({0.f, 0.f}), "Range of eta"); + desc.add_options()( + "charge", + po::value()->value_name("charge")->default_value(-1.f), + "Charge of particles"); } void read(const po::variables_map& vm) { gen_nparticles = vm["gen-nparticles"].as(); @@ -71,6 +76,7 @@ struct particle_gen_options { theta_range = eta_to_theta_range(eta_range); phi_range = {phi_range_degree[0] * detray::unit::degree, phi_range_degree[1] * detray::unit::degree}; + charge = vm["charge"].as(); } }; diff --git a/examples/simulation/simulate.cpp b/examples/simulation/simulate.cpp index c6bcfe3177..c3d3a972dc 100644 --- a/examples/simulation/simulate.cpp +++ b/examples/simulation/simulate.cpp @@ -112,6 +112,7 @@ int main(int argc, char* argv[]) { gen_cfg.phi_range(pg_opts.phi_range[0], pg_opts.phi_range[1]); gen_cfg.theta_range(pg_opts.theta_range[0], pg_opts.theta_range[1]); gen_cfg.mom_range(pg_opts.mom_range[0], pg_opts.mom_range[1]); + gen_cfg.charge(pg_opts.charge); generator_type generator(gen_cfg); // Smearing value for measurements diff --git a/examples/simulation/simulate_telescope.cpp b/examples/simulation/simulate_telescope.cpp index 614d36051c..e264205fce 100644 --- a/examples/simulation/simulate_telescope.cpp +++ b/examples/simulation/simulate_telescope.cpp @@ -103,6 +103,7 @@ int simulate(std::string output_directory, unsigned int events, gen_cfg.phi_range(pg_opts.phi_range[0], pg_opts.phi_range[1]); gen_cfg.theta_range(pg_opts.theta_range[0], pg_opts.theta_range[1]); gen_cfg.mom_range(pg_opts.mom_range[0], pg_opts.mom_range[1]); + gen_cfg.charge(pg_opts.charge); generator_type generator(gen_cfg); // Smearing value for measurements diff --git a/examples/simulation/simulate_toy_detector.cpp b/examples/simulation/simulate_toy_detector.cpp index f97cac8e37..8f665a12b4 100644 --- a/examples/simulation/simulate_toy_detector.cpp +++ b/examples/simulation/simulate_toy_detector.cpp @@ -76,6 +76,7 @@ int simulate(std::string output_directory, unsigned int events, gen_cfg.phi_range(pg_opts.phi_range[0], pg_opts.phi_range[1]); gen_cfg.theta_range(pg_opts.theta_range[0], pg_opts.theta_range[1]); gen_cfg.mom_range(pg_opts.mom_range[0], pg_opts.mom_range[1]); + gen_cfg.charge(pg_opts.charge); generator_type generator(gen_cfg); // Smearing value for measurements diff --git a/examples/simulation/simulate_wire_chamber.cpp b/examples/simulation/simulate_wire_chamber.cpp index a5935a728e..99ce8175a0 100644 --- a/examples/simulation/simulate_wire_chamber.cpp +++ b/examples/simulation/simulate_wire_chamber.cpp @@ -81,6 +81,7 @@ int simulate(std::string output_directory, unsigned int events, gen_cfg.phi_range(pg_opts.phi_range[0], pg_opts.phi_range[1]); gen_cfg.theta_range(pg_opts.theta_range[0], pg_opts.theta_range[1]); gen_cfg.mom_range(pg_opts.mom_range[0], pg_opts.mom_range[1]); + gen_cfg.charge(pg_opts.charge); generator_type generator(gen_cfg); // Smearing value for measurements diff --git a/tests/common/tests/kalman_fitting_test.hpp b/tests/common/tests/kalman_fitting_test.hpp index e6776fb4ea..664daf4344 100644 --- a/tests/common/tests/kalman_fitting_test.hpp +++ b/tests/common/tests/kalman_fitting_test.hpp @@ -40,13 +40,14 @@ namespace traccc { /// (4) momentum range /// (5) eta range /// (6) phi range -/// (7) number of tracks per event -/// (8) number of events +/// (7) charge +/// (8) number of tracks per event +/// (9) number of events class KalmanFittingTests : public ::testing::TestWithParam, std::array, std::array, std::array, - std::array, unsigned int, unsigned int>> { + std::array, scalar, unsigned int, unsigned int>> { public: /// Type declarations diff --git a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp index f476ff56b6..cde4899e3e 100644 --- a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp @@ -47,8 +47,9 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { const std::array eta_range = std::get<4>(GetParam()); const std::array theta_range = eta_to_theta_range(eta_range); const std::array phi_range = std::get<5>(GetParam()); - const unsigned int n_truth_tracks = std::get<6>(GetParam()); - const unsigned int n_events = std::get<7>(GetParam()); + const scalar charge = std::get<6>(GetParam()); + const unsigned int n_truth_tracks = std::get<7>(GetParam()); + const unsigned int n_events = std::get<8>(GetParam()); // Performance writer traccc::fitting_performance_writer::config fit_writer_cfg; @@ -103,6 +104,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { gen_cfg.phi_range(phi_range[0], phi_range[1]); gen_cfg.theta_range(theta_range[0], theta_range[1]); gen_cfg.mom_range(mom_range[0], mom_range[1]); + gen_cfg.charge(charge); generator_type generator(gen_cfg); // Smearing value for measurements @@ -223,7 +225,7 @@ INSTANTIATE_TEST_SUITE_P( "telescope_single_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 1, 5000))); + std::array{0.f, 0.f}, -1.f, 1, 5000))); INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation1, CkfSparseTrackTelescopeTests, @@ -231,7 +233,7 @@ INSTANTIATE_TEST_SUITE_P( "telescope_double_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 2, 2500))); + std::array{0.f, 0.f}, -1.f, 2, 2500))); INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation2, CkfSparseTrackTelescopeTests, @@ -239,7 +241,7 @@ INSTANTIATE_TEST_SUITE_P( "telescope_quadra_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 4, 1250))); + std::array{0.f, 0.f}, -1.f, 4, 1250))); INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation3, CkfSparseTrackTelescopeTests, @@ -247,4 +249,4 @@ INSTANTIATE_TEST_SUITE_P( "telescope_decade_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 10, 500))); + std::array{0.f, 0.f}, -1.f, 10, 500))); diff --git a/tests/cpu/test_kalman_fitter_telescope.cpp b/tests/cpu/test_kalman_fitter_telescope.cpp index e415b62759..e6e75ea006 100644 --- a/tests/cpu/test_kalman_fitter_telescope.cpp +++ b/tests/cpu/test_kalman_fitter_telescope.cpp @@ -45,8 +45,9 @@ TEST_P(KalmanFittingTelescopeTests, Run) { const std::array eta_range = std::get<4>(GetParam()); const std::array theta_range = eta_to_theta_range(eta_range); const std::array phi_range = std::get<5>(GetParam()); - const unsigned int n_truth_tracks = std::get<6>(GetParam()); - const unsigned int n_events = std::get<7>(GetParam()); + const scalar charge = std::get<6>(GetParam()); + const unsigned int n_truth_tracks = std::get<7>(GetParam()); + const unsigned int n_events = std::get<8>(GetParam()); // Performance writer traccc::fitting_performance_writer::config fit_writer_cfg; @@ -85,6 +86,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { gen_cfg.phi_range(phi_range[0], phi_range[1]); gen_cfg.theta_range(theta_range[0], theta_range[1]); gen_cfg.mom_range(mom_range[0], mom_range[1]); + gen_cfg.charge(charge); generator_type generator(gen_cfg); // Smearing value for measurements @@ -177,23 +179,33 @@ TEST_P(KalmanFittingTelescopeTests, Run) { INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation0, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "telescope_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "telescope_1_GeV_0_phi_muon", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, - std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, - 100))); + std::array{0.f, 0.f}, std::array{0.f, 0.f}, + -1.f, 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "telescope_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "telescope_10_GeV_0_phi_muon", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 100, 100))); + std::array{0.f, 0.f}, -1.f, 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, ::testing::Values(std::make_tuple( - "telescope_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, + "telescope_100_GeV_0_phi_muon", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 100, 100))); \ No newline at end of file + std::array{0.f, 0.f}, -1.f, 100, 100))); + +INSTANTIATE_TEST_SUITE_P( + KalmanFitTelescopeValidation3, KalmanFittingTelescopeTests, + ::testing::Values(std::make_tuple("telescope_1_GeV_0_phi_anti_muon", + std::array{0.f, 0.f, 0.f}, + std::array{0.f, 0.f, 0.f}, + std::array{1.f, 1.f}, + std::array{0.f, 0.f}, + std::array{0.f, 0.f}, 1.f, + 100, 100))); diff --git a/tests/cpu/test_kalman_fitter_wire_chamber.cpp b/tests/cpu/test_kalman_fitter_wire_chamber.cpp index 37808a89a7..0b5f772cbd 100644 --- a/tests/cpu/test_kalman_fitter_wire_chamber.cpp +++ b/tests/cpu/test_kalman_fitter_wire_chamber.cpp @@ -44,8 +44,9 @@ TEST_P(KalmanFittingWireChamberTests, Run) { const std::array eta_range = std::get<4>(GetParam()); const std::array theta_range = eta_to_theta_range(eta_range); const std::array phi_range = std::get<5>(GetParam()); - const unsigned int n_truth_tracks = std::get<6>(GetParam()); - const unsigned int n_events = std::get<7>(GetParam()); + const scalar charge = std::get<6>(GetParam()); + const unsigned int n_truth_tracks = std::get<7>(GetParam()); + const unsigned int n_events = std::get<8>(GetParam()); // Performance writer traccc::fitting_performance_writer::config fit_writer_cfg; @@ -85,6 +86,7 @@ TEST_P(KalmanFittingWireChamberTests, Run) { gen_cfg.phi_range(phi_range[0], phi_range[1]); gen_cfg.theta_range(theta_range[0], theta_range[1]); gen_cfg.mom_range(mom_range[0], mom_range[1]); + gen_cfg.charge(charge); generator_type generator(gen_cfg); // Smearing value for measurements @@ -185,26 +187,35 @@ TEST_P(KalmanFittingWireChamberTests, Run) { INSTANTIATE_TEST_SUITE_P( KalmanFitWireChamberValidation0, KalmanFittingWireChamberTests, ::testing::Values(std::make_tuple( - "wire_2_GeV", std::array{0.f, 0.f, 0.f}, + "wire_2_GeV_muon", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{2.f, 2.f}, std::array{-1.f, 1.f}, - std::array{0.f, 2.0f * detray::constant::pi}, 100, - 100))); + std::array{0.f, 2.0f * detray::constant::pi}, -1.f, + 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitWireChamberValidation1, KalmanFittingWireChamberTests, ::testing::Values(std::make_tuple( - "wire_10_GeV", std::array{0.f, 0.f, 0.f}, + "wire_10_GeV_muon", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{-1.f, 1.f}, - std::array{0.f, 2.0f * detray::constant::pi}, 100, - 100))); + std::array{0.f, 2.0f * detray::constant::pi}, -1.f, + 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitWireChamberValidation2, KalmanFittingWireChamberTests, ::testing::Values(std::make_tuple( - "wire_100_GeV", std::array{0.f, 0.f, 0.f}, + "wire_100_GeV_muon", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{-1.f, 1.f}, - std::array{0.f, 2.0f * detray::constant::pi}, 100, - 100))); + std::array{0.f, 2.0f * detray::constant::pi}, -1.f, + 100, 100))); + +INSTANTIATE_TEST_SUITE_P( + KalmanFitWireChamberValidation3, KalmanFittingWireChamberTests, + ::testing::Values(std::make_tuple( + "wire_2_GeV_anti_muon", std::array{0.f, 0.f, 0.f}, + std::array{0.f, 0.f, 0.f}, std::array{2.f, 2.f}, + std::array{-1.f, 1.f}, + std::array{0.f, 2.0f * detray::constant::pi}, 1.f, + 100, 100))); \ No newline at end of file diff --git a/tests/cpu/test_simulation.cpp b/tests/cpu/test_simulation.cpp index 368d05284c..4dddf8753b 100644 --- a/tests/cpu/test_simulation.cpp +++ b/tests/cpu/test_simulation.cpp @@ -187,10 +187,10 @@ GTEST_TEST(detray_simulation, toy_geometry_simulation) { } } -// Test parameters: +// Test parameters: class TelescopeDetectorSimulation - : public ::testing::TestWithParam> { -}; + : public ::testing::TestWithParam< + std::tuple> {}; TEST_P(TelescopeDetectorSimulation, telescope_detector_simulation) { @@ -229,6 +229,8 @@ TEST_P(TelescopeDetectorSimulation, telescope_detector_simulation) { const vector3 ori{0.f, 0.f, 0.f}; const scalar theta = std::get<2>(GetParam()); + const scalar charge = std::get<3>(GetParam()); + // Track generator using generator_type = detray::uniform_track_generator; @@ -238,6 +240,7 @@ TEST_P(TelescopeDetectorSimulation, telescope_detector_simulation) { gen_cfg.origin(ori); gen_cfg.theta_range(theta, theta); gen_cfg.p_mag(mom); + gen_cfg.charge(charge); generator_type generator(gen_cfg); // Create smearer @@ -282,31 +285,31 @@ TEST_P(TelescopeDetectorSimulation, telescope_detector_simulation) { INSTANTIATE_TEST_SUITE_P( Simulation, TelescopeDetectorSimulation, ::testing::Values( - std::make_tuple("0", 0.1f * detray::unit::GeV, 0.01f), - std::make_tuple("1", 1.f * detray::unit::GeV, 0.01f), - std::make_tuple("2", 10.f * detray::unit::GeV, 0.01f), - std::make_tuple("3", 100.f * detray::unit::GeV, 0.01f), + std::make_tuple("0", 0.1f * detray::unit::GeV, 0.01f, -1.f), + std::make_tuple("1", 1.f * detray::unit::GeV, 0.01f, -1.f), + std::make_tuple("2", 10.f * detray::unit::GeV, 0.01f, -1.f), + std::make_tuple("3", 100.f * detray::unit::GeV, 0.01f, -1.f), std::make_tuple("4", 0.1f * detray::unit::GeV, - detray::constant::pi / 12.f), + detray::constant::pi / 12.f, 1.f), std::make_tuple("5", 1.f * detray::unit::GeV, - detray::constant::pi / 12.f), + detray::constant::pi / 12.f, 1.f), std::make_tuple("6", 10.f * detray::unit::GeV, - detray::constant::pi / 12.f), + detray::constant::pi / 12.f, 1.f), std::make_tuple("7", 100.f * detray::unit::GeV, - detray::constant::pi / 12.f), + detray::constant::pi / 12.f, 1.f), std::make_tuple("8", 0.1f * detray::unit::GeV, - detray::constant::pi / 8.f), + detray::constant::pi / 8.f, -1.f), std::make_tuple("9", 1.f * detray::unit::GeV, - detray::constant::pi / 8.f), + detray::constant::pi / 8.f, -1.f), std::make_tuple("10", 10.f * detray::unit::GeV, - detray::constant::pi / 8.f), + detray::constant::pi / 8.f, -1.f), std::make_tuple("11", 100.f * detray::unit::GeV, - detray::constant::pi / 8.f), + detray::constant::pi / 8.f, -1.f), std::make_tuple("12", 0.1f * detray::unit::GeV, - detray::constant::pi / 6.f), + detray::constant::pi / 6.f, 1.f), std::make_tuple("13", 1.f * detray::unit::GeV, - detray::constant::pi / 6.f), + detray::constant::pi / 6.f, 1.f), std::make_tuple("14", 10.f * detray::unit::GeV, - detray::constant::pi / 6.f), + detray::constant::pi / 6.f, 1.f), std::make_tuple("15", 100.f * detray::unit::GeV, - detray::constant::pi / 6.f))); \ No newline at end of file + detray::constant::pi / 6.f, 1.f))); \ No newline at end of file diff --git a/tests/cuda/test_ckf_sparse_tracks_telescope.cpp b/tests/cuda/test_ckf_sparse_tracks_telescope.cpp index cb5c271cb8..369142f2de 100644 --- a/tests/cuda/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cuda/test_ckf_sparse_tracks_telescope.cpp @@ -54,8 +54,9 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { const std::array eta_range = std::get<4>(GetParam()); const std::array theta_range = eta_to_theta_range(eta_range); const std::array phi_range = std::get<5>(GetParam()); - const unsigned int n_truth_tracks = std::get<6>(GetParam()); - const unsigned int n_events = std::get<7>(GetParam()); + const scalar charge = std::get<6>(GetParam()); + const unsigned int n_truth_tracks = std::get<7>(GetParam()); + const unsigned int n_events = std::get<8>(GetParam()); // Performance writer traccc::fitting_performance_writer::config fit_writer_cfg; @@ -116,6 +117,7 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { gen_cfg.phi_range(phi_range[0], phi_range[1]); gen_cfg.theta_range(theta_range[0], theta_range[1]); gen_cfg.mom_range(mom_range[0], mom_range[1]); + gen_cfg.charge(charge); generator_type generator(gen_cfg); // Smearing value for measurements @@ -289,7 +291,7 @@ INSTANTIATE_TEST_SUITE_P( "cuda_telescope_single_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 1, 5000))); + std::array{0.f, 0.f}, -1.f, 1, 5000))); INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation1, CkfSparseTrackTelescopeTests, @@ -297,7 +299,7 @@ INSTANTIATE_TEST_SUITE_P( "cuda_telescope_double_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 2, 2500))); + std::array{0.f, 0.f}, -1.f, 2, 2500))); INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation2, CkfSparseTrackTelescopeTests, @@ -305,7 +307,7 @@ INSTANTIATE_TEST_SUITE_P( "cuda_telescope_quadra_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 4, 1250))); + std::array{0.f, 0.f}, -1.f, 4, 1250))); INSTANTIATE_TEST_SUITE_P( CkfSparseTrackTelescopeValidation3, CkfSparseTrackTelescopeTests, @@ -313,4 +315,4 @@ INSTANTIATE_TEST_SUITE_P( "cuda_telescope_decade_tracks", std::array{0.f, 0.f, 0.f}, std::array{0.f, 200.f, 200.f}, std::array{1.f, 1.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 10, 500))); \ No newline at end of file + std::array{0.f, 0.f}, -1.f, 10, 500))); \ No newline at end of file diff --git a/tests/cuda/test_kalman_fitter_telescope.cpp b/tests/cuda/test_kalman_fitter_telescope.cpp index 1e60eccc66..9ee10cbc0c 100644 --- a/tests/cuda/test_kalman_fitter_telescope.cpp +++ b/tests/cuda/test_kalman_fitter_telescope.cpp @@ -55,8 +55,9 @@ TEST_P(KalmanFittingTelescopeTests, Run) { const std::array eta_range = std::get<4>(GetParam()); const std::array theta_range = eta_to_theta_range(eta_range); const std::array phi_range = std::get<5>(GetParam()); - const unsigned int n_truth_tracks = std::get<6>(GetParam()); - const unsigned int n_events = std::get<7>(GetParam()); + const scalar charge = std::get<6>(GetParam()); + const unsigned int n_truth_tracks = std::get<7>(GetParam()); + const unsigned int n_events = std::get<8>(GetParam()); // Performance writer traccc::fitting_performance_writer::config fit_writer_cfg; @@ -102,6 +103,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { gen_cfg.phi_range(phi_range[0], phi_range[1]); gen_cfg.theta_range(theta_range[0], theta_range[1]); gen_cfg.mom_range(mom_range[0], mom_range[1]); + gen_cfg.charge(charge); generator_type generator(gen_cfg); // Smearing value for measurements @@ -224,8 +226,8 @@ INSTANTIATE_TEST_SUITE_P( ::testing::Values(std::make_tuple( "cuda_telescope_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, - std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, - 100))); + std::array{0.f, 0.f}, std::array{0.f, 0.f}, + -1.f, 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, @@ -233,7 +235,7 @@ INSTANTIATE_TEST_SUITE_P( "cuda_telescope_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 100, 100))); + std::array{0.f, 0.f}, -1.f, 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, @@ -241,4 +243,4 @@ INSTANTIATE_TEST_SUITE_P( "cuda_telescope_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 100, 100))); + std::array{0.f, 0.f}, -1.f, 100, 100))); diff --git a/tests/sycl/test_kalman_fitter_telescope.sycl b/tests/sycl/test_kalman_fitter_telescope.sycl index 2ebde173aa..d2f5aeb77c 100644 --- a/tests/sycl/test_kalman_fitter_telescope.sycl +++ b/tests/sycl/test_kalman_fitter_telescope.sycl @@ -71,8 +71,9 @@ TEST_P(KalmanFittingTelescopeTests, Run) { const std::array eta_range = std::get<4>(GetParam()); const std::array theta_range = eta_to_theta_range(eta_range); const std::array phi_range = std::get<5>(GetParam()); - const unsigned int n_truth_tracks = std::get<6>(GetParam()); - const unsigned int n_events = std::get<7>(GetParam()); + const scalar charge = std::get<6>(GetParam()); + const unsigned int n_truth_tracks = std::get<7>(GetParam()); + const unsigned int n_events = std::get<8>(GetParam()); // Performance writer traccc::fitting_performance_writer::config fit_writer_cfg; @@ -122,6 +123,7 @@ TEST_P(KalmanFittingTelescopeTests, Run) { gen_cfg.phi_range(phi_range[0], phi_range[1]); gen_cfg.theta_range(theta_range[0], theta_range[1]); gen_cfg.mom_range(mom_range[0], mom_range[1]); + gen_cfg.charge(charge); generator_type generator(gen_cfg); // Smearing value for measurements @@ -231,8 +233,8 @@ INSTANTIATE_TEST_SUITE_P( ::testing::Values(std::make_tuple( "sycl_telescope_1_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{1.f, 1.f}, - std::array{0.f, 0.f}, std::array{0.f, 0.f}, 100, - 100))); + std::array{0.f, 0.f}, std::array{0.f, 0.f}, + -1.f, 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation1, KalmanFittingTelescopeTests, @@ -240,7 +242,7 @@ INSTANTIATE_TEST_SUITE_P( "sycl_telescope_10_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{10.f, 10.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 100, 100))); + std::array{0.f, 0.f}, -1.f, 100, 100))); INSTANTIATE_TEST_SUITE_P( KalmanFitTelescopeValidation2, KalmanFittingTelescopeTests, @@ -248,4 +250,4 @@ INSTANTIATE_TEST_SUITE_P( "sycl_telescope_100_GeV_0_phi", std::array{0.f, 0.f, 0.f}, std::array{0.f, 0.f, 0.f}, std::array{100.f, 100.f}, std::array{0.f, 0.f}, - std::array{0.f, 0.f}, 100, 100))); + std::array{0.f, 0.f}, -1.f, 100, 100))); From c5a54b50823b61912aba4b733af8c4a4e2d7b808 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Fri, 24 Nov 2023 21:50:34 +0100 Subject: [PATCH 10/33] Remove unused geometry writing part and keep the simulation files after test --- tests/cpu/test_ckf_sparse_tracks_telescope.cpp | 18 ------------------ tests/cpu/test_kalman_fitter_telescope.cpp | 3 --- tests/cpu/test_kalman_fitter_wire_chamber.cpp | 3 --- .../cuda/test_ckf_sparse_tracks_telescope.cpp | 18 ------------------ tests/cuda/test_kalman_fitter_telescope.cpp | 3 --- tests/sycl/test_kalman_fitter_telescope.sycl | 3 --- 6 files changed, 48 deletions(-) diff --git a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp index cde4899e3e..49fb9508fa 100644 --- a/tests/cpu/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cpu/test_ckf_sparse_tracks_telescope.cpp @@ -63,21 +63,6 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { // Memory resources used by the application. vecmem::host_memory_resource host_mr; - detray::tel_det_config<> tel_cfg{rectangle}; - tel_cfg.positions(plane_positions); - tel_cfg.module_material(mat); - tel_cfg.mat_thickness(thickness); - tel_cfg.pilot_track(traj); - - // Create telescope detector - auto [det, name_map] = create_telescope_detector(host_mr, tel_cfg); - - // Write detector file - auto writer_cfg = detray::io::detector_writer_config{} - .format(detray::io::format::json) - .replace_files(true); - detray::io::write_detector(det, name_map, writer_cfg); - // Read back detector file const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; @@ -214,9 +199,6 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { static_cast(n_success) / (n_truth_tracks * n_events); ASSERT_FLOAT_EQ(success_rate, 1.00f); - - // Remove the data - std::filesystem::remove_all(full_path); } INSTANTIATE_TEST_SUITE_P( diff --git a/tests/cpu/test_kalman_fitter_telescope.cpp b/tests/cpu/test_kalman_fitter_telescope.cpp index e6e75ea006..45fa52a495 100644 --- a/tests/cpu/test_kalman_fitter_telescope.cpp +++ b/tests/cpu/test_kalman_fitter_telescope.cpp @@ -171,9 +171,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { static_cast(n_success) / (n_truth_tracks * n_events); ASSERT_FLOAT_EQ(success_rate, 1.00f); - - // Remove the data - std::filesystem::remove_all(full_path); } INSTANTIATE_TEST_SUITE_P( diff --git a/tests/cpu/test_kalman_fitter_wire_chamber.cpp b/tests/cpu/test_kalman_fitter_wire_chamber.cpp index 0b5f772cbd..2ad1cf8414 100644 --- a/tests/cpu/test_kalman_fitter_wire_chamber.cpp +++ b/tests/cpu/test_kalman_fitter_wire_chamber.cpp @@ -179,9 +179,6 @@ TEST_P(KalmanFittingWireChamberTests, Run) { ASSERT_GE(success_rate, 0.99f); ASSERT_LE(success_rate, 1.00f); - - // Remove the data - std::filesystem::remove_all(full_path); } INSTANTIATE_TEST_SUITE_P( diff --git a/tests/cuda/test_ckf_sparse_tracks_telescope.cpp b/tests/cuda/test_ckf_sparse_tracks_telescope.cpp index 369142f2de..ce68a390f7 100644 --- a/tests/cuda/test_ckf_sparse_tracks_telescope.cpp +++ b/tests/cuda/test_ckf_sparse_tracks_telescope.cpp @@ -73,21 +73,6 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { traccc::memory_resource mr{device_mr, &host_mr}; vecmem::cuda::managed_memory_resource mng_mr; - detray::tel_det_config<> tel_cfg{rectangle}; - tel_cfg.positions(plane_positions); - tel_cfg.module_material(mat); - tel_cfg.mat_thickness(thickness); - tel_cfg.pilot_track(traj); - - // Create telescope detector - auto [det, name_map] = create_telescope_detector(host_mr, tel_cfg); - - // Write detector file - auto writer_cfg = detray::io::detector_writer_config{} - .format(detray::io::format::json) - .replace_files(true); - detray::io::write_detector(det, name_map, writer_cfg); - // Read back detector file const std::string path = name + "/"; detray::io::detector_reader_config reader_cfg{}; @@ -280,9 +265,6 @@ TEST_P(CkfSparseTrackTelescopeTests, Run) { static_cast(n_success) / (n_truth_tracks * n_events); ASSERT_FLOAT_EQ(success_rate, 1.00f); - - // Remove the data - std::filesystem::remove_all(full_path); } INSTANTIATE_TEST_SUITE_P( diff --git a/tests/cuda/test_kalman_fitter_telescope.cpp b/tests/cuda/test_kalman_fitter_telescope.cpp index 9ee10cbc0c..e90df1535c 100644 --- a/tests/cuda/test_kalman_fitter_telescope.cpp +++ b/tests/cuda/test_kalman_fitter_telescope.cpp @@ -216,9 +216,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { static_cast(n_success) / (n_truth_tracks * n_events); ASSERT_FLOAT_EQ(success_rate, 1.00f); - - // Remove the data - std::filesystem::remove_all(full_path); } INSTANTIATE_TEST_SUITE_P( diff --git a/tests/sycl/test_kalman_fitter_telescope.sycl b/tests/sycl/test_kalman_fitter_telescope.sycl index d2f5aeb77c..364eb34e8a 100644 --- a/tests/sycl/test_kalman_fitter_telescope.sycl +++ b/tests/sycl/test_kalman_fitter_telescope.sycl @@ -223,9 +223,6 @@ TEST_P(KalmanFittingTelescopeTests, Run) { static const std::vector pull_names{ "pull_d0", "pull_z0", "pull_phi", "pull_theta", "pull_qop"}; pull_value_tests(fit_writer_cfg.file_path, pull_names); - - // Remove the data - std::filesystem::remove_all(full_path); } INSTANTIATE_TEST_SUITE_P( From ca5618d07192b2e0fd538a2d0458b049516a2060 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Fri, 24 Nov 2023 01:25:16 +0100 Subject: [PATCH 11/33] Remove the unused pair integer comparator --- core/CMakeLists.txt | 1 - .../traccc/finding/finding_algorithm.ipp | 4 +- core/include/traccc/utils/compare.hpp | 24 -------- tests/core/CMakeLists.txt | 1 - tests/core/test_compare.cpp | 55 ------------------- 5 files changed, 1 insertion(+), 84 deletions(-) delete mode 100644 core/include/traccc/utils/compare.hpp delete mode 100644 tests/core/test_compare.cpp diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index 51fc5a805a..2dd605dee2 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -35,7 +35,6 @@ traccc_add_library( traccc_core core TYPE SHARED "include/traccc/geometry/pixel_data.hpp" # Utilities. "include/traccc/utils/algorithm.hpp" - "include/traccc/utils/compare.hpp" "include/traccc/utils/type_traits.hpp" "include/traccc/utils/memory_resource.hpp" "include/traccc/utils/seed_generator.hpp" diff --git a/core/include/traccc/finding/finding_algorithm.ipp b/core/include/traccc/finding/finding_algorithm.ipp index e3a45ffcec..d1608dfb7d 100644 --- a/core/include/traccc/finding/finding_algorithm.ipp +++ b/core/include/traccc/finding/finding_algorithm.ipp @@ -7,11 +7,9 @@ // Project include(s). #include "traccc/finding/candidate_link.hpp" -#include "traccc/utils/compare.hpp" // System include #include -#include #include namespace traccc { @@ -35,7 +33,7 @@ finding_algorithm::operator()( auto end = std::unique_copy(measurements.begin(), measurements.end(), uniques.begin(), measurement_equal_comp()); - unsigned int n_modules = end - uniques.begin(); + const unsigned int n_modules = end - uniques.begin(); // Get upper bounds of unique elements std::vector upper_bounds; diff --git a/core/include/traccc/utils/compare.hpp b/core/include/traccc/utils/compare.hpp deleted file mode 100644 index 865d02174a..0000000000 --- a/core/include/traccc/utils/compare.hpp +++ /dev/null @@ -1,24 +0,0 @@ -/** - * TRACCC library, part of the ACTS project (R&D line) - * - * (c) 2023 CERN for the benefit of the ACTS project - * - * Mozilla Public License Version 2.0 - */ - -#pragma once - -// Project include(s). -#include "traccc/definitions/qualifiers.hpp" - -namespace traccc { - -template