tests/cuda/test_ambiguity_resolution.cpp

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2025-2026 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 // Project include(s).
0009 #include "traccc/ambiguity_resolution/ambiguity_resolution_config.hpp"
0010 #include "traccc/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp"
0011 #include "traccc/cuda/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.hpp"
0012 #include "traccc/device/container_d2h_copy_alg.hpp"
0013 #include "traccc/device/container_h2d_copy_alg.hpp"
0014 #include "traccc/utils/memory_resource.hpp"
0015
0016 // VecMem include(s).
0017 #include <vecmem/memory/cuda/device_memory_resource.hpp>
0018 #include <vecmem/memory/cuda/managed_memory_resource.hpp>
0019 #include <vecmem/memory/host_memory_resource.hpp>
0020 #include <vecmem/utils/cuda/async_copy.hpp>
0021 #include <vecmem/utils/cuda/stream_wrapper.hpp>
0022
0023 // GTest include(s).
0024 #include <gtest/gtest.h>
0025
0026 // System include(s).
0027 #include <chrono>
0028 #include <random>
0029 #include <thread>
0030
0031 using namespace traccc;
0032
0033 void fill_measurements(edm::measurement_collection::host& measurements,
0034                        const measurement_id_type max_meas_id) {
0035   measurements.reserve(max_meas_id + 1);
0036   for (measurement_id_type i = 0; i <= max_meas_id; i++) {
0037     measurements.push_back({});
0038     measurements.at(measurements.size() - 1).identifier() = i;
0039   }
0040 }
0041
0042 void fill_pattern(edm::track_container<default_algebra>::host& track_candidates,
0043                   const traccc::scalar pval,
0044                   const std::vector<measurement_id_type>& pattern) {
0045   track_candidates.tracks.resize(track_candidates.tracks.size() + 1u);
0046   track_candidates.tracks.pval().back() = pval;
0047
0048   edm::measurement_collection::const_device measurements{
0049       track_candidates.measurements};
0050
0051   for (const auto& meas_id : pattern) {
0052     const auto meas_iter =
0053         std::lower_bound(measurements.identifier().begin(),
0054                          measurements.identifier().end(), meas_id);
0055
0056     const auto meas_idx =
0057         std::distance(measurements.identifier().begin(), meas_iter);
0058     track_candidates.tracks.constituent_links().back().push_back(
0059         {edm::track_constituent_link::measurement,
0060          static_cast<measurement_id_type>(meas_idx)});
0061   }
0062 }
0063
0064 bool find_pattern(
0065     const edm::track_container<default_algebra>::const_device& tracks,
0066     const std::vector<measurement_id_type>& pattern) {
0067   const auto n_tracks = tracks.tracks.size();
0068   for (unsigned int i = 0; i < n_tracks; i++) {
0069     std::vector<measurement_id_type> ids;
0070     for (const auto& [type, meas_idx] :
0071          tracks.tracks.constituent_links().at(i)) {
0072       assert(type == edm::track_constituent_link::measurement);
0073       ids.push_back(tracks.measurements.at(meas_idx).identifier());
0074     }
0075     if (pattern == ids) {
0076       return true;
0077     }
0078   }
0079   return false;
0080 }
0081
0082 std::vector<measurement_id_type> get_pattern(
0083     const edm::track_container<default_algebra>::host& track_candidates,
0084     const std::size_t idx) {
0085   edm::measurement_collection::const_device measurements{
0086       track_candidates.measurements};
0087   std::vector<measurement_id_type> ret;
0088   // A const reference would be fine here. But GCC fears that that would lead
0089   // to a dangling reference...
0090   const auto meas_links = track_candidates.tracks.at(idx).constituent_links();
0091   for (const auto& [type, meas_idx] : meas_links) {
0092     assert(type == edm::track_constituent_link::measurement);
0093     ret.push_back(measurements.at(meas_idx).identifier());
0094   }
0095
0096   return ret;
0097 }
0098
0099 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest0) {
0100   // Memory resource used by the EDM.
0101   vecmem::cuda::managed_memory_resource mng_mr;
0102
0103   // Cuda stream
0104   vecmem::cuda::stream_wrapper vecmem_stream;
0105   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0106
0107   // Cuda copy objects
0108   vecmem::cuda::async_copy copy{stream.cudaStream()};
0109
0110   edm::measurement_collection::host measurements{mng_mr};
0111   fill_measurements(measurements, 100);
0112
0113   edm::track_container<default_algebra>::host trk_cands{
0114       mng_mr, vecmem::get_data(measurements)};
0115   fill_pattern(trk_cands, 0.23f, {5, 1, 11, 3});
0116   fill_pattern(trk_cands, 0.85f, {12, 10, 9, 8, 7, 6});
0117   fill_pattern(trk_cands, 0.42f, {4, 2, 13});
0118
0119   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0120       resolution_config;
0121
0122   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0123       resolution_config, {mng_mr}, copy, stream);
0124   {
0125     resolution_alg_cuda.get_config().min_meas_per_track = 3;
0126     auto res_trk_cands_buffer = resolution_alg_cuda(
0127         edm::track_container<default_algebra>::const_data(trk_cands));
0128     stream.synchronize();
0129     edm::track_container<default_algebra>::const_device res_trk_cands(
0130         res_trk_cands_buffer);
0131     // All tracks are accepted as they have more than three measurements
0132     EXPECT_EQ(res_trk_cands.tracks.size(), 3u);
0133     ASSERT_TRUE(find_pattern(res_trk_cands, {5, 1, 11, 3}));
0134     ASSERT_TRUE(find_pattern(res_trk_cands, {12, 10, 9, 8, 7, 6}));
0135     ASSERT_TRUE(find_pattern(res_trk_cands, {4, 2, 13}));
0136   }
0137
0138   {
0139     resolution_alg_cuda.get_config().min_meas_per_track = 5;
0140     auto res_trk_cands_buffer = resolution_alg_cuda(
0141         edm::track_container<default_algebra>::const_data(trk_cands));
0142     stream.synchronize();
0143     edm::track_container<default_algebra>::const_device res_trk_cands(
0144         res_trk_cands_buffer);
0145     // Only the second track with six measurements is accepted
0146     ASSERT_EQ(res_trk_cands.tracks.size(), 1u);
0147     ASSERT_TRUE(find_pattern(res_trk_cands, {12, 10, 9, 8, 7, 6}));
0148   }
0149 }
0150
0151 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest1) {
0152   // Memory resource used by the EDM.
0153   vecmem::cuda::managed_memory_resource mng_mr;
0154
0155   // Cuda stream
0156   vecmem::cuda::stream_wrapper vecmem_stream;
0157   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0158
0159   // Cuda copy objects
0160   vecmem::cuda::async_copy copy{stream.cudaStream()};
0161
0162   edm::measurement_collection::host measurements{mng_mr};
0163   fill_measurements(measurements, 100);
0164
0165   edm::track_container<default_algebra>::host trk_cands{
0166       mng_mr, vecmem::get_data(measurements)};
0167   fill_pattern(trk_cands, 0.12f, {5, 14, 1, 11, 18, 16, 3});
0168   fill_pattern(trk_cands, 0.53f, {3, 6, 5, 13});
0169
0170   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0171       resolution_config;
0172
0173   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0174       resolution_config, {mng_mr}, copy, stream);
0175
0176   resolution_alg_cuda.get_config().min_meas_per_track = 3;
0177   auto res_trk_cands_buffer = resolution_alg_cuda(
0178       edm::track_container<default_algebra>::const_data(trk_cands));
0179   stream.synchronize();
0180   edm::track_container<default_algebra>::const_device res_trk_cands(
0181       res_trk_cands_buffer);
0182   // All tracks are accepted as they have more than three measurements
0183   ASSERT_EQ(res_trk_cands.tracks.size(), 1u);
0184
0185   // The first track is selected over the second one as its relative
0186   // shared measurement (2/7) is lower than the one of the second track
0187   // (2/4)
0188   ASSERT_TRUE(find_pattern(res_trk_cands, {5, 14, 1, 11, 18, 16, 3}));
0189 }
0190
0191 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest2) {
0192   // Memory resource used by the EDM.
0193   vecmem::cuda::managed_memory_resource mng_mr;
0194
0195   // Cuda stream
0196   vecmem::cuda::stream_wrapper vecmem_stream;
0197   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0198
0199   // Cuda copy objects
0200   vecmem::cuda::async_copy copy{stream.cudaStream()};
0201
0202   edm::measurement_collection::host measurements{mng_mr};
0203   fill_measurements(measurements, 100);
0204
0205   edm::track_container<default_algebra>::host trk_cands{
0206       mng_mr, vecmem::get_data(measurements)};
0207   fill_pattern(trk_cands, 0.8f, {1, 3, 5, 11});
0208   fill_pattern(trk_cands, 0.9f, {3, 5, 6, 13});
0209
0210   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0211       resolution_config;
0212   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0213       resolution_config, {mng_mr}, copy, stream);
0214
0215   auto res_trk_cands_buffer = resolution_alg_cuda(
0216       edm::track_container<default_algebra>::const_data(trk_cands));
0217   stream.synchronize();
0218   edm::track_container<default_algebra>::const_device res_trk_cands(
0219       res_trk_cands_buffer);
0220   ASSERT_EQ(res_trk_cands.tracks.size(), 1u);
0221
0222   // The second track is selected over the first one as their relative
0223   // shared measurement (2/4) is the same but its p-value is higher
0224   ASSERT_TRUE(find_pattern(res_trk_cands, {3, 5, 6, 13}));
0225 }
0226
0227 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest3) {
0228   // Memory resource used by the EDM.
0229   vecmem::cuda::managed_memory_resource mng_mr;
0230
0231   // Cuda stream
0232   vecmem::cuda::stream_wrapper vecmem_stream;
0233   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0234
0235   // Cuda copy objects
0236   vecmem::cuda::async_copy copy{stream.cudaStream()};
0237
0238   edm::measurement_collection::host measurements{mng_mr};
0239   fill_measurements(measurements, 100);
0240
0241   edm::track_container<default_algebra>::host trk_cands{
0242       mng_mr, vecmem::get_data(measurements)};
0243   fill_pattern(trk_cands, 0.2f, {5, 1, 11, 3});
0244   fill_pattern(trk_cands, 0.5f, {6, 2});
0245   fill_pattern(trk_cands, 0.4f, {3, 21, 12, 6, 19, 14});
0246   fill_pattern(trk_cands, 0.1f, {13, 16, 2, 7, 11});
0247   fill_pattern(trk_cands, 0.3f, {1, 7, 8});
0248   fill_pattern(trk_cands, 0.6f, {1, 3, 11, 22});
0249
0250   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0251       resolution_config;
0252   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0253       resolution_config, {mng_mr}, copy, stream);
0254
0255   auto res_trk_cands_buffer = resolution_alg_cuda(
0256       edm::track_container<default_algebra>::const_data(trk_cands));
0257   stream.synchronize();
0258   edm::track_container<default_algebra>::const_device res_trk_cands(
0259       res_trk_cands_buffer);
0260   ASSERT_EQ(res_trk_cands.tracks.size(), 2u);
0261
0262   ASSERT_TRUE(find_pattern(res_trk_cands, {3, 21, 12, 6, 19, 14}));
0263   ASSERT_TRUE(find_pattern(res_trk_cands, {13, 16, 2, 7, 11}));
0264 }
0265
0266 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest5) {
0267   // Memory resource used by the EDM.
0268   vecmem::cuda::managed_memory_resource mng_mr;
0269
0270   // Cuda stream
0271   vecmem::cuda::stream_wrapper vecmem_stream;
0272   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0273
0274   // Cuda copy objects
0275   vecmem::cuda::async_copy copy{stream.cudaStream()};
0276
0277   edm::measurement_collection::host measurements{mng_mr};
0278   fill_measurements(measurements, 100);
0279
0280   edm::track_container<default_algebra>::host trk_cands{
0281       mng_mr, vecmem::get_data(measurements)};
0282   fill_pattern(trk_cands, 0.2f, {1, 2, 1, 1});
0283   fill_pattern(trk_cands, 0.5f, {3, 2, 1});
0284   fill_pattern(trk_cands, 0.4f, {2, 4, 5, 7, 2});
0285   fill_pattern(trk_cands, 0.1f, {6, 6, 6, 6});
0286
0287   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0288       resolution_config;
0289   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0290       resolution_config, {mng_mr}, copy, stream);
0291
0292   auto res_trk_cands_buffer = resolution_alg_cuda(
0293       edm::track_container<default_algebra>::const_data(trk_cands));
0294   stream.synchronize();
0295   edm::track_container<default_algebra>::const_device res_trk_cands(
0296       res_trk_cands_buffer);
0297   ASSERT_EQ(res_trk_cands.tracks.size(), 2u);
0298
0299   ASSERT_TRUE(find_pattern(res_trk_cands, {3, 2, 1}));
0300   ASSERT_TRUE(find_pattern(res_trk_cands, {6, 6, 6, 6}));
0301 }
0302
0303 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest6) {
0304   // Memory resource used by the EDM.
0305   vecmem::cuda::managed_memory_resource mng_mr;
0306
0307   // Cuda stream
0308   vecmem::cuda::stream_wrapper vecmem_stream;
0309   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0310
0311   // Cuda copy objects
0312   vecmem::cuda::async_copy copy{stream.cudaStream()};
0313
0314   edm::measurement_collection::host measurements{mng_mr};
0315   fill_measurements(measurements, 100);
0316
0317   edm::track_container<default_algebra>::host trk_cands{
0318       mng_mr, vecmem::get_data(measurements)};
0319   fill_pattern(trk_cands, 0.2f, {7, 3, 5, 7, 7, 7, 2});
0320   fill_pattern(trk_cands, 0.5f, {2});
0321   fill_pattern(trk_cands, 0.4f, {8, 9, 7, 2, 3, 4, 3, 7});
0322   fill_pattern(trk_cands, 0.1f, {8, 9, 0, 8, 1, 4, 6});
0323   fill_pattern(trk_cands, 0.9f, {10, 3, 2});
0324
0325   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0326       resolution_config;
0327   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0328       resolution_config, {mng_mr}, copy, stream);
0329
0330   auto res_trk_cands_buffer = resolution_alg_cuda(
0331       edm::track_container<default_algebra>::const_data(trk_cands));
0332   stream.synchronize();
0333   edm::track_container<default_algebra>::const_device res_trk_cands(
0334       res_trk_cands_buffer);
0335   ASSERT_EQ(res_trk_cands.tracks.size(), 2u);
0336
0337   ASSERT_TRUE(find_pattern(res_trk_cands, {7, 3, 5, 7, 7, 7, 2}));
0338   ASSERT_TRUE(find_pattern(res_trk_cands, {8, 9, 0, 8, 1, 4, 6}));
0339 }
0340
0341 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest7) {
0342   // Memory resource used by the EDM.
0343   vecmem::cuda::managed_memory_resource mng_mr;
0344
0345   // Cuda stream
0346   vecmem::cuda::stream_wrapper vecmem_stream;
0347   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0348
0349   // Cuda copy objects
0350   vecmem::cuda::async_copy copy{stream.cudaStream()};
0351
0352   edm::measurement_collection::host measurements{mng_mr};
0353   fill_measurements(measurements, 100);
0354
0355   edm::track_container<default_algebra>::host trk_cands{
0356       mng_mr, vecmem::get_data(measurements)};
0357   fill_pattern(trk_cands, 0.173853f, {10, 3, 6, 8});
0358   fill_pattern(trk_cands, 0.548019f, {3, 3, 1});
0359   fill_pattern(trk_cands, 0.276757f, {2, 8, 5, 4});
0360
0361   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0362       resolution_config;
0363   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0364       resolution_config, {mng_mr}, copy, stream);
0365
0366   auto res_trk_cands_buffer = resolution_alg_cuda(
0367       edm::track_container<default_algebra>::const_data(trk_cands));
0368   stream.synchronize();
0369   edm::track_container<default_algebra>::const_device res_trk_cands(
0370       res_trk_cands_buffer);
0371   ASSERT_EQ(res_trk_cands.tracks.size(), 1u);
0372
0373   ASSERT_TRUE(find_pattern(res_trk_cands, {2, 8, 5, 4}));
0374 }
0375
0376 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest8) {
0377   // Memory resource used by the EDM.
0378   vecmem::cuda::managed_memory_resource mng_mr;
0379
0380   // Cuda stream
0381   vecmem::cuda::stream_wrapper vecmem_stream;
0382   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0383
0384   // Cuda copy objects
0385   vecmem::cuda::async_copy copy{stream.cudaStream()};
0386
0387   edm::measurement_collection::host measurements{mng_mr};
0388   fill_measurements(measurements, 100);
0389
0390   edm::track_container<default_algebra>::host trk_cands{
0391       mng_mr, vecmem::get_data(measurements)};
0392   fill_pattern(trk_cands, 0.0623132f, {10, 4});
0393   fill_pattern(trk_cands, 0.207417f, {6, 7, 5});
0394   fill_pattern(trk_cands, 0.325736f, {8, 2, 2});
0395   fill_pattern(trk_cands, 0.581643f, {5, 7, 9, 7});
0396   fill_pattern(trk_cands, 0.389551f, {1, 9, 3, 0});
0397
0398   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0399       resolution_config;
0400   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0401       resolution_config, {mng_mr}, copy, stream);
0402
0403   auto res_trk_cands_buffer = resolution_alg_cuda(
0404       edm::track_container<default_algebra>::const_data(trk_cands));
0405   stream.synchronize();
0406   edm::track_container<default_algebra>::const_device res_trk_cands(
0407       res_trk_cands_buffer);
0408   ASSERT_EQ(res_trk_cands.tracks.size(), 3u);
0409
0410   ASSERT_TRUE(find_pattern(res_trk_cands, {6, 7, 5}));
0411   ASSERT_TRUE(find_pattern(res_trk_cands, {8, 2, 2}));
0412   ASSERT_TRUE(find_pattern(res_trk_cands, {1, 9, 3, 0}));
0413 }
0414
0415 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest9) {
0416   // Memory resource used by the EDM.
0417   vecmem::cuda::managed_memory_resource mng_mr;
0418
0419   // Cuda stream
0420   vecmem::cuda::stream_wrapper vecmem_stream;
0421   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0422
0423   // Cuda copy objects
0424   vecmem::cuda::async_copy copy{stream.cudaStream()};
0425
0426   edm::measurement_collection::host measurements{mng_mr};
0427   fill_measurements(measurements, 100);
0428
0429   edm::track_container<default_algebra>::host trk_cands{
0430       mng_mr, vecmem::get_data(measurements)};
0431   fill_pattern(trk_cands, 0.542984f, {0, 4, 8, 1, 1});
0432   fill_pattern(trk_cands, 0.583695f, {10, 6, 8, 7});
0433   fill_pattern(trk_cands, 0.280232f, {4, 1, 8, 10});
0434
0435   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0436       resolution_config;
0437   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0438       resolution_config, {mng_mr}, copy, stream);
0439
0440   auto res_trk_cands_buffer = resolution_alg_cuda(
0441       edm::track_container<default_algebra>::const_data(trk_cands));
0442   stream.synchronize();
0443   edm::track_container<default_algebra>::const_device res_trk_cands(
0444       res_trk_cands_buffer);
0445   ASSERT_EQ(res_trk_cands.tracks.size(), 1u);
0446
0447   ASSERT_TRUE(find_pattern(res_trk_cands, {0, 4, 8, 1, 1}));
0448 }
0449
0450 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest10) {
0451   // Memory resource used by the EDM.
0452   vecmem::cuda::managed_memory_resource mng_mr;
0453
0454   // Cuda stream
0455   vecmem::cuda::stream_wrapper vecmem_stream;
0456   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0457
0458   // Cuda copy objects
0459   vecmem::cuda::async_copy copy{stream.cudaStream()};
0460
0461   edm::measurement_collection::host measurements{mng_mr};
0462   fill_measurements(measurements, 100);
0463
0464   edm::track_container<default_algebra>::host trk_cands{
0465       mng_mr, vecmem::get_data(measurements)};
0466   fill_pattern(trk_cands, 0.399106f, {14, 51});
0467   fill_pattern(trk_cands, 0.43899f, {80, 35, 41, 55});
0468   fill_pattern(trk_cands, 0.0954247f, {73, 63, 49, 89});
0469   fill_pattern(trk_cands, 0.158046f, {81, 22, 58, 54, 91});
0470   fill_pattern(trk_cands, 0.349878f, {97, 89, 80});
0471
0472   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0473       resolution_config;
0474   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0475       resolution_config, {mng_mr}, copy, stream);
0476
0477   auto res_trk_cands_buffer = resolution_alg_cuda(
0478       edm::track_container<default_algebra>::const_data(trk_cands));
0479   stream.synchronize();
0480   edm::track_container<default_algebra>::const_device res_trk_cands(
0481       res_trk_cands_buffer);
0482   ASSERT_EQ(res_trk_cands.tracks.size(), 3u);
0483
0484   ASSERT_TRUE(find_pattern(res_trk_cands, {80, 35, 41, 55}));
0485   ASSERT_TRUE(find_pattern(res_trk_cands, {73, 63, 49, 89}));
0486   ASSERT_TRUE(find_pattern(res_trk_cands, {81, 22, 58, 54, 91}));
0487 }
0488
0489 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest11) {
0490   // Memory resource used by the EDM.
0491   vecmem::cuda::managed_memory_resource mng_mr;
0492
0493   // Cuda stream
0494   vecmem::cuda::stream_wrapper vecmem_stream;
0495   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0496
0497   // Cuda copy objects
0498   vecmem::cuda::async_copy copy{stream.cudaStream()};
0499
0500   edm::measurement_collection::host measurements{mng_mr};
0501   fill_measurements(measurements, 100);
0502
0503   edm::track_container<default_algebra>::host trk_cands{
0504       mng_mr, vecmem::get_data(measurements)};
0505   fill_pattern(trk_cands, 0.95f, {56, 87});
0506   fill_pattern(trk_cands, 0.894f, {64, 63});
0507   fill_pattern(trk_cands, 0.824f, {70, 17});
0508   fill_pattern(trk_cands, 0.862f, {27, 0});
0509   fill_pattern(trk_cands, 0.871f, {27, 19});
0510
0511   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0512       resolution_config;
0513   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0514       resolution_config, {mng_mr}, copy, stream);
0515
0516   auto res_trk_cands_buffer = resolution_alg_cuda(
0517       edm::track_container<default_algebra>::const_data(trk_cands));
0518   stream.synchronize();
0519   edm::track_container<default_algebra>::const_device res_trk_cands(
0520       res_trk_cands_buffer);
0521   ASSERT_EQ(res_trk_cands.tracks.size(), 0u);
0522 }
0523
0524 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest12) {
0525   // Memory resource used by the EDM.
0526   vecmem::cuda::managed_memory_resource mng_mr;
0527
0528   // Cuda stream
0529   vecmem::cuda::stream_wrapper vecmem_stream;
0530   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0531
0532   // Cuda copy objects
0533   vecmem::cuda::async_copy copy{stream.cudaStream()};
0534
0535   edm::measurement_collection::host measurements{mng_mr};
0536   fill_measurements(measurements, 100);
0537
0538   edm::track_container<default_algebra>::host trk_cands{
0539       mng_mr, vecmem::get_data(measurements)};
0540   fill_pattern(trk_cands, 0.948f, {17, 6, 1, 69, 78});  // 69
0541   fill_pattern(trk_cands, 0.609f, {17});
0542   fill_pattern(trk_cands, 0.453f, {84, 45, 81, 69});      // 84, 69
0543   fill_pattern(trk_cands, 0.910f, {54, 64, 49, 96, 40});  // 64
0544   fill_pattern(trk_cands, 0.153f, {59, 57, 84, 27, 64});  // 84, 64
0545
0546   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0547       resolution_config;
0548   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0549       resolution_config, {mng_mr}, copy, stream);
0550
0551   auto res_trk_cands_buffer = resolution_alg_cuda(
0552       edm::track_container<default_algebra>::const_data(trk_cands));
0553   stream.synchronize();
0554   edm::track_container<default_algebra>::const_device res_trk_cands(
0555       res_trk_cands_buffer);
0556   ASSERT_EQ(res_trk_cands.tracks.size(), 2u);
0557
0558   ASSERT_TRUE(find_pattern(res_trk_cands, {17, 6, 1, 69, 78}));
0559   ASSERT_TRUE(find_pattern(res_trk_cands, {54, 64, 49, 96, 40}));
0560 }
0561
0562 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest13) {
0563   // Memory resource used by the EDM.
0564   vecmem::cuda::managed_memory_resource mng_mr;
0565
0566   // Cuda stream
0567   vecmem::cuda::stream_wrapper vecmem_stream;
0568   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0569
0570   // Cuda copy objects
0571   vecmem::cuda::async_copy copy{stream.cudaStream()};
0572
0573   edm::measurement_collection::host measurements{mng_mr};
0574   fill_measurements(measurements, 100);
0575
0576   edm::track_container<default_algebra>::host trk_cands{
0577       mng_mr, vecmem::get_data(measurements)};
0578   fill_pattern(trk_cands, 0.211f, {46, 92, 74, 58});
0579   fill_pattern(trk_cands, 0.694f, {15, 78, 9});
0580   fill_pattern(trk_cands, 0.432f, {15, 4, 58, 68});
0581   fill_pattern(trk_cands, 0.958f, {38, 93, 68});
0582   fill_pattern(trk_cands, 0.203f, {57, 64, 57, 36});
0583   fill_pattern(trk_cands, 0.118f, {4, 85, 65, 14});
0584
0585   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0586       resolution_config;
0587   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0588       resolution_config, {mng_mr}, copy, stream);
0589
0590   auto res_trk_cands_buffer = resolution_alg_cuda(
0591       edm::track_container<default_algebra>::const_data(trk_cands));
0592   stream.synchronize();
0593   edm::track_container<default_algebra>::const_device res_trk_cands(
0594       res_trk_cands_buffer);
0595   ASSERT_EQ(res_trk_cands.tracks.size(), 5u);
0596
0597   ASSERT_TRUE(find_pattern(res_trk_cands, {46, 92, 74, 58}));
0598   ASSERT_TRUE(find_pattern(res_trk_cands, {15, 78, 9}));
0599   ASSERT_TRUE(find_pattern(res_trk_cands, {38, 93, 68}));
0600   ASSERT_TRUE(find_pattern(res_trk_cands, {57, 64, 57, 36}));
0601   ASSERT_TRUE(find_pattern(res_trk_cands, {4, 85, 65, 14}));
0602 }
0603
0604 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest14) {
0605   // Memory resource used by the EDM.
0606   vecmem::cuda::managed_memory_resource mng_mr;
0607
0608   // Cuda stream
0609   vecmem::cuda::stream_wrapper vecmem_stream;
0610   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0611
0612   // Cuda copy objects
0613   vecmem::cuda::async_copy copy{stream.cudaStream()};
0614
0615   edm::measurement_collection::host measurements{mng_mr};
0616   fill_measurements(measurements, 100);
0617
0618   edm::track_container<default_algebra>::host trk_cands{
0619       mng_mr, vecmem::get_data(measurements)};
0620   fill_pattern(trk_cands, 0.932f, {8, 4, 3});
0621   fill_pattern(trk_cands, 0.263f, {1, 1, 9, 3});
0622   fill_pattern(trk_cands, 0.876f, {1, 2, 5});
0623   fill_pattern(trk_cands, 0.058f, {2, 0, 4, 7});
0624
0625   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0626       resolution_config;
0627   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0628       resolution_config, {mng_mr}, copy, stream);
0629
0630   auto res_trk_cands_buffer = resolution_alg_cuda(
0631       edm::track_container<default_algebra>::const_data(trk_cands));
0632   stream.synchronize();
0633   edm::track_container<default_algebra>::const_device res_trk_cands(
0634       res_trk_cands_buffer);
0635   ASSERT_EQ(res_trk_cands.tracks.size(), 2u);
0636
0637   ASSERT_TRUE(find_pattern(res_trk_cands, {8, 4, 3}));
0638   ASSERT_TRUE(find_pattern(res_trk_cands, {1, 2, 5}));
0639 }
0640
0641 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest15) {
0642   // Memory resource used by the EDM.
0643   vecmem::cuda::managed_memory_resource mng_mr;
0644
0645   // Cuda stream
0646   vecmem::cuda::stream_wrapper vecmem_stream;
0647   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0648
0649   // Cuda copy objects
0650   vecmem::cuda::async_copy copy{stream.cudaStream()};
0651
0652   edm::measurement_collection::host measurements{mng_mr};
0653   fill_measurements(measurements, 100);
0654
0655   edm::track_container<default_algebra>::host trk_cands{
0656       mng_mr, vecmem::get_data(measurements)};
0657   fill_pattern(trk_cands, 0.293f, {2, 0, 4});
0658   fill_pattern(trk_cands, 0.362f, {8, 4, 9, 3});
0659   fill_pattern(trk_cands, 0.011f, {9, 4, 8, 4});
0660   fill_pattern(trk_cands, 0.843f, {8, 7, 1});
0661
0662   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0663       resolution_config;
0664   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0665       resolution_config, {mng_mr}, copy, stream);
0666
0667   auto res_trk_cands_buffer = resolution_alg_cuda(
0668       edm::track_container<default_algebra>::const_data(trk_cands));
0669   stream.synchronize();
0670   edm::track_container<default_algebra>::const_device res_trk_cands(
0671       res_trk_cands_buffer);
0672   ASSERT_EQ(res_trk_cands.tracks.size(), 2u);
0673
0674   ASSERT_TRUE(find_pattern(res_trk_cands, {2, 0, 4}));
0675   ASSERT_TRUE(find_pattern(res_trk_cands, {8, 7, 1}));
0676 }
0677
0678 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest16) {
0679   // Memory resource used by the EDM.
0680   vecmem::cuda::managed_memory_resource mng_mr;
0681
0682   // Cuda stream
0683   vecmem::cuda::stream_wrapper vecmem_stream;
0684   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0685
0686   // Cuda copy objects
0687   vecmem::cuda::async_copy copy{stream.cudaStream()};
0688
0689   edm::measurement_collection::host measurements{mng_mr};
0690   fill_measurements(measurements, 100);
0691
0692   edm::track_container<default_algebra>::host trk_cands{
0693       mng_mr, vecmem::get_data(measurements)};
0694   fill_pattern(trk_cands, 0.622598f, {95, 24, 62, 83, 67});
0695   fill_pattern(trk_cands, 0.541774f, {6, 52, 57, 87, 75});
0696   fill_pattern(trk_cands, 0.361033f, {14, 52, 29, 79, 89});
0697   fill_pattern(trk_cands, 0.622598f, {57, 85, 63, 90});
0698   fill_pattern(trk_cands, 0.481157f, {80, 45, 94});
0699
0700   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0701       resolution_config;
0702   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0703       resolution_config, {mng_mr}, copy, stream);
0704
0705   auto res_trk_cands_buffer = resolution_alg_cuda(
0706       edm::track_container<default_algebra>::const_data(trk_cands));
0707   stream.synchronize();
0708   edm::track_container<default_algebra>::const_device res_trk_cands(
0709       res_trk_cands_buffer);
0710   ASSERT_EQ(res_trk_cands.tracks.size(), 4u);
0711
0712   ASSERT_TRUE(find_pattern(res_trk_cands, {95, 24, 62, 83, 67}));
0713   ASSERT_TRUE(find_pattern(res_trk_cands, {14, 52, 29, 79, 89}));
0714   ASSERT_TRUE(find_pattern(res_trk_cands, {57, 85, 63, 90}));
0715   ASSERT_TRUE(find_pattern(res_trk_cands, {80, 45, 94}));
0716 }
0717
0718 TEST(CUDAAmbiguitySolverTests, GreedyResolverTest17) {
0719   // Memory resource used by the EDM.
0720   vecmem::cuda::managed_memory_resource mng_mr;
0721
0722   // Cuda stream
0723   vecmem::cuda::stream_wrapper vecmem_stream;
0724   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0725
0726   // Cuda copy objects
0727   vecmem::cuda::async_copy copy{stream.cudaStream()};
0728
0729   edm::measurement_collection::host measurements{mng_mr};
0730   fill_measurements(measurements, 100);
0731
0732   edm::track_container<default_algebra>::host trk_cands{
0733       mng_mr, vecmem::get_data(measurements)};
0734   fill_pattern(trk_cands, 0.17975f, {7, 4, 10, 3, 0});
0735   fill_pattern(trk_cands, 0.924326f, {0, 0, 9});
0736   fill_pattern(trk_cands, 0.0832954f, {0, 2, 0});
0737   fill_pattern(trk_cands, 0.303148f, {0, 6, 4, 5, 5});
0738
0739   traccc::cuda::greedy_ambiguity_resolution_algorithm::config_type
0740       resolution_config;
0741   traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0742       resolution_config, {mng_mr}, copy, stream);
0743
0744   auto res_trk_cands_buffer = resolution_alg_cuda(
0745       edm::track_container<default_algebra>::const_data(trk_cands));
0746   stream.synchronize();
0747   edm::track_container<default_algebra>::const_device res_trk_cands(
0748       res_trk_cands_buffer);
0749   ASSERT_EQ(res_trk_cands.tracks.size(), 1u);
0750
0751   ASSERT_TRUE(find_pattern(res_trk_cands, {0, 6, 4, 5, 5}));
0752 }
0753
0754 // Test class for the ambiguity resolution comparison with CPU implementation
0755 // Input tuple: < n_event, n_tracks, track_length_range , max_meas_id,
0756 // allow_duplicate >
0757 class CUDAGreedyResolutionCompareToCPU
0758     : public ::testing::TestWithParam<
0759           std::tuple<std::size_t, std::size_t, std::array<std::size_t, 2u>,
0760                      measurement_id_type, bool>> {};
0761
0762 TEST_P(CUDAGreedyResolutionCompareToCPU, Comparison) {
0763   const std::size_t n_events = std::get<0>(GetParam());
0764   const std::size_t n_tracks = std::get<1>(GetParam());
0765   const std::array<std::size_t, 2u> trk_length_range = std::get<2>(GetParam());
0766   const measurement_id_type max_meas_id = std::get<3>(GetParam());
0767   const bool allow_duplicate = std::get<4>(GetParam());
0768
0769   // Memory resource used by the EDM.
0770   vecmem::cuda::device_memory_resource device_mr;
0771   vecmem::host_memory_resource host_mr;
0772   traccc::memory_resource mr{device_mr, &host_mr};
0773
0774   // Cuda stream
0775   vecmem::cuda::stream_wrapper vecmem_stream;
0776   traccc::cuda::stream_wrapper stream{vecmem_stream.stream()};
0777
0778   // Cuda copy objects
0779   vecmem::cuda::async_copy copy{stream.cudaStream()};
0780
0781   for (std::size_t i_evt = 0u; i_evt < n_events; i_evt++) {
0782     std::size_t sd = 42u + i_evt;
0783     std::mt19937 gen(sd);
0784     std::cout << "Event: " << i_evt << " Seed: " << sd << std::endl;
0785
0786     edm::measurement_collection::host measurements{host_mr};
0787     fill_measurements(measurements, max_meas_id);
0788     edm::track_container<default_algebra>::host trk_cands{
0789         host_mr, vecmem::get_data(measurements)};
0790
0791     for (std::size_t i = 0; i < n_tracks; i++) {
0792       std::uniform_int_distribution<std::size_t> track_length_dist(
0793           trk_length_range[0], trk_length_range[1]);
0794       std::uniform_int_distribution<measurement_id_type> meas_id_dist(
0795           0, max_meas_id);
0796       std::uniform_real_distribution<traccc::scalar> pval_dist(0.0f, 1.0f);
0797
0798       const std::size_t track_length = track_length_dist(gen);
0799       const traccc::scalar pval = pval_dist(gen);
0800       std::vector<measurement_id_type> pattern;
0801       // std::cout << pval << std::endl;
0802       while (pattern.size() < track_length) {
0803         auto mid = meas_id_dist(gen);
0804         if (!allow_duplicate) {
0805           while (std::find(pattern.begin(), pattern.end(), mid) !=
0806                  pattern.end()) {
0807             mid = meas_id_dist(gen);
0808           }
0809         }
0810         // std::cout << mid << ", ";
0811         pattern.push_back(mid);
0812       }
0813       // std::cout << std::endl;
0814
0815       // Make sure that partern size is eqaul to the track length
0816       ASSERT_EQ(pattern.size(), track_length);
0817
0818       // Fill the pattern
0819       fill_pattern(trk_cands, pval, pattern);
0820     }
0821
0822     // CPU algorithm
0823     traccc::host::greedy_ambiguity_resolution_algorithm::config_type
0824         resolution_config;
0825     traccc::host::greedy_ambiguity_resolution_algorithm resolution_alg_cpu(
0826         resolution_config, host_mr);
0827
0828     auto start_cpu = std::chrono::high_resolution_clock::now();
0829
0830     auto res_trk_cands_cpu = resolution_alg_cpu(
0831         edm::track_container<default_algebra>::const_data(trk_cands));
0832
0833     auto end_cpu = std::chrono::high_resolution_clock::now();
0834     auto duration_cpu = std::chrono::duration_cast<std::chrono::milliseconds>(
0835         end_cpu - start_cpu);
0836     std::cout << " Time for the cpu method " << duration_cpu.count() << " ms"
0837               << std::endl;
0838
0839     // CUDA algorithm
0840     traccc::cuda::greedy_ambiguity_resolution_algorithm resolution_alg_cuda(
0841         resolution_config, mr, copy, stream);
0842
0843     // H2D transfer
0844     edm::measurement_collection::buffer measurements_buffer =
0845         copy.to(vecmem::get_data(measurements), device_mr, &host_mr,
0846                 vecmem::copy::type::host_to_device);
0847     traccc::edm::track_container<default_algebra>::buffer trk_cands_buffer{
0848         copy.to(vecmem::get_data(trk_cands.tracks), device_mr, &host_mr,
0849                 vecmem::copy::type::host_to_device),
0850         {},
0851         measurements_buffer};
0852
0853     auto start_cuda = std::chrono::high_resolution_clock::now();
0854
0855     // Instantiate output cuda containers/collections
0856     auto res_trk_cands_buffer = resolution_alg_cuda(trk_cands_buffer);
0857     stream.synchronize();
0858
0859     auto end_cuda = std::chrono::high_resolution_clock::now();
0860     auto duration_cuda = std::chrono::duration_cast<std::chrono::milliseconds>(
0861         end_cuda - start_cuda);
0862     std::cout << " Time for the cuda method " << duration_cuda.count() << " ms"
0863               << std::endl;
0864
0865     traccc::edm::track_container<default_algebra>::buffer res_trk_cands_cuda{
0866         copy.to(res_trk_cands_buffer.tracks, host_mr, nullptr,
0867                 vecmem::copy::type::device_to_host),
0868         {},
0869         vecmem::get_data(measurements)};
0870
0871     const auto n_tracks_cpu = res_trk_cands_cpu.tracks.size();
0872     ASSERT_EQ(n_tracks_cpu, res_trk_cands_cuda.tracks.capacity());
0873
0874     // Make sure that CPU and CUDA track candidates have same
0875     // patterns
0876     edm::track_container<default_algebra>::const_device
0877         res_trk_cands_cuda_device{res_trk_cands_cuda};
0878     for (unsigned int i = 0; i < n_tracks_cpu; i++) {
0879       ASSERT_TRUE(find_pattern(res_trk_cands_cuda_device,
0880                                get_pattern(res_trk_cands_cpu, i)));
0881     }
0882   }
0883 };
0884
0885 INSTANTIATE_TEST_SUITE_P(
0886     CUDAStandard, CUDAGreedyResolutionCompareToCPU,
0887     ::testing::Values(std::make_tuple(5u, 50000u,
0888                                       std::array<std::size_t, 2u>{1u, 10u},
0889                                       20000u, true),
0890                       std::make_tuple(5u, 50000u,
0891                                       std::array<std::size_t, 2u>{1u, 10u},
0892                                       20000u, false)));
0893
0894 INSTANTIATE_TEST_SUITE_P(
0895     CUDASparse, CUDAGreedyResolutionCompareToCPU,
0896     ::testing::Values(std::make_tuple(3u, 5000u,
0897                                       std::array<std::size_t, 2u>{3u, 10u},
0898                                       1000000u, true),
0899                       std::make_tuple(3u, 5000u,
0900                                       std::array<std::size_t, 2u>{3u, 10u},
0901                                       1000000u, false)));
0902
0903 INSTANTIATE_TEST_SUITE_P(
0904     CUDADense, CUDAGreedyResolutionCompareToCPU,
0905     ::testing::Values(std::make_tuple(3u, 5000u,
0906                                       std::array<std::size_t, 2u>{3u, 10u},
0907                                       100u, true),
0908                       std::make_tuple(3u, 5000u,
0909                                       std::array<std::size_t, 2u>{3u, 10u},
0910                                       100u, false)));
0911
0912 INSTANTIATE_TEST_SUITE_P(
0913     CUDALong, CUDAGreedyResolutionCompareToCPU,
0914     ::testing::Values(std::make_tuple(3u, 10000u,
0915                                       std::array<std::size_t, 2u>{3u, 500u},
0916                                       10000u, true),
0917                       std::make_tuple(3u, 10000u,
0918                                       std::array<std::size_t, 2u>{3u, 500u},
0919                                       10000u, false)));
0920
0921 INSTANTIATE_TEST_SUITE_P(
0922     CUDASimple, CUDAGreedyResolutionCompareToCPU,
0923     ::testing::Values(
0924         std::make_tuple(3u, 5u, std::array<std::size_t, 2u>{3u, 5u}, 10u, true),
0925         std::make_tuple(3u, 5u, std::array<std::size_t, 2u>{3u, 5u}, 10u,
0926                         false)));