diff --git a/core/include/traccc/seeding/seed_selecting_helper.hpp b/core/include/traccc/seeding/seed_selecting_helper.hpp index d860b30501..51580cb1a3 100644 --- a/core/include/traccc/seeding/seed_selecting_helper.hpp +++ b/core/include/traccc/seeding/seed_selecting_helper.hpp @@ -70,16 +70,12 @@ struct seed_selecting_helper { /// @param seed current seed to possibly cut /// /// @return boolean value + template static TRACCC_HOST_DEVICE bool cut_per_middle_sp( - const seedfilter_config& filter_config, - const edm::spacepoint_collection::const_device& spacepoints, - const details::spacepoint_grid_types::const_device& grid, - const triplet& seed) { + const seedfilter_config& filter_config, const spacepoint_type& spB, + const scalar weight) { - const edm::spacepoint_collection::const_device::const_proxy_type spB = - spacepoints.at(grid.bin(seed.sp1.bin_idx)[seed.sp1.sp_idx]); - - return (seed.weight > filter_config.seed_min_weight || + return (weight > filter_config.seed_min_weight || spB.radius() > filter_config.spB_min_radius); } }; diff --git a/core/src/seeding/seed_filtering.cpp b/core/src/seeding/seed_filtering.cpp index c0fb75d418..7f56e36ea6 100644 --- a/core/src/seeding/seed_filtering.cpp +++ b/core/src/seeding/seed_filtering.cpp @@ -95,9 +95,14 @@ void seed_filtering::operator()( std::min(triplets_passing_single_seed_cuts.size(), static_cast(m_finder_config.maxSeedsPerSpM)); for (std::size_t i = 1; i < itLength; ++i) { + const traccc::details::spacepoint_grid_types::const_device + sp_grid_accessor(sp_grid_data); + const auto& this_seed = triplets_passing_single_seed_cuts[i].get(); if (seed_selecting_helper::cut_per_middle_sp( - m_filter_config, spacepoints, sp_grid_data, - triplets_passing_single_seed_cuts[i])) { + m_filter_config, + spacepoints.at(sp_grid_accessor.bin( + this_seed.sp1.bin_idx)[this_seed.sp1.sp_idx]), + this_seed.weight)) { triplets_passing_final_cuts.push_back( triplets_passing_single_seed_cuts[i]); } diff --git a/device/alpaka/src/seeding/seed_finding.cpp b/device/alpaka/src/seeding/seed_finding.cpp index 6a8f27dc24..adf478b115 100644 --- a/device/alpaka/src/seeding/seed_finding.cpp +++ b/device/alpaka/src/seeding/seed_finding.cpp @@ -134,7 +134,6 @@ struct UpdateTripletWeights { ALPAKA_FN_ACC void operator()( TAcc const& acc, seedfilter_config filter_config, edm::spacepoint_collection::const_view spacepoints, - traccc::details::spacepoint_grid_types::const_view sp_grid, device::triplet_counter_spM_collection_types::const_view spM_tc, device::triplet_counter_collection_types::const_view midBot_tc, device::device_triplet_collection_types::view triplet_view) const { @@ -151,8 +150,8 @@ struct UpdateTripletWeights { scalar* dataPos = &data[localThreadIdx * filter_config.compatSeedLimit]; device::update_triplet_weights(globalThreadIdx, filter_config, - spacepoints, sp_grid, spM_tc, midBot_tc, - dataPos, triplet_view); + spacepoints, spM_tc, midBot_tc, dataPos, + triplet_view); } }; @@ -175,10 +174,12 @@ struct SelectSeeds { // Array for temporary storage of quality parameters for comparing // triplets within weight updating kernel - triplet* const data = ::alpaka::getDynSharedMem(acc); + device::device_triplet* const data = + ::alpaka::getDynSharedMem(acc); - // Each thread uses maxSeedsPerSpM elements of the array - triplet* dataPos = &data[localThreadIdx * finder_config.maxSeedsPerSpM]; + // Each thread uses max_triplets_per_spM elements of the array + device::device_triplet* dataPos = + &data[localThreadIdx * finder_config.maxSeedsPerSpM]; device::select_seeds(globalThreadIdx, finder_config, filter_config, spacepoints, sp_view, spM_tc, midBot_tc, @@ -360,7 +361,7 @@ edm::seed_collection::buffer seed_finding::operator()( // Update the weights of all spacepoint triplets. ::alpaka::exec(queue, workDiv, kernels::UpdateTripletWeights{}, - m_seedfilter_config, spacepoints_view, g2_view, + m_seedfilter_config, spacepoints_view, vecmem::get_data(triplet_counter_spM_buffer), vecmem::get_data(triplet_counter_midBot_buffer), vecmem::get_data(triplet_buffer)); @@ -419,7 +420,7 @@ struct BlockSharedMemDynSizeBytes { ) -> std::size_t { return static_cast(finder_config.maxSeedsPerSpM * blockThreadExtent.prod()) * - sizeof(traccc::triplet); + sizeof(traccc::device::device_triplet); } }; diff --git a/device/common/include/traccc/edm/device/device_triplet.hpp b/device/common/include/traccc/edm/device/device_triplet.hpp index b3f351bbf5..a73376f983 100644 --- a/device/common/include/traccc/edm/device/device_triplet.hpp +++ b/device/common/include/traccc/edm/device/device_triplet.hpp @@ -17,7 +17,7 @@ namespace traccc::device { /// Triplets of bottom, middle and top spacepoints struct device_triplet { // top spacepoint location in internal spacepoint container - sp_location spT; + unsigned int spB, spM, spT; using link_type = device::triplet_counter_collection_types::host::size_type; /// Link to triplet counter where the middle and bottom spacepoints are diff --git a/device/common/include/traccc/seeding/device/impl/find_triplets.ipp b/device/common/include/traccc/seeding/device/impl/find_triplets.ipp index f85ab337fe..3475876154 100644 --- a/device/common/include/traccc/seeding/device/impl/find_triplets.ipp +++ b/device/common/include/traccc/seeding/device/impl/find_triplets.ipp @@ -56,12 +56,14 @@ inline void find_triplets( const sp_location spB_loc = mid_bot_counter.spB; // middle spacepoint + const unsigned int spM_idx = sp_grid.bin(spM_loc.bin_idx)[spM_loc.sp_idx]; const edm::spacepoint_collection::const_device::const_proxy_type spM = - spacepoints.at(sp_grid.bin(spM_loc.bin_idx)[spM_loc.sp_idx]); + spacepoints.at(spM_idx); // bottom spacepoint + const unsigned int spB_idx = sp_grid.bin(spB_loc.bin_idx)[spB_loc.sp_idx]; const edm::spacepoint_collection::const_device::const_proxy_type spB = - spacepoints.at(sp_grid.bin(spB_loc.bin_idx)[spB_loc.sp_idx]); + spacepoints.at(spB_idx); // Set up the device result collection device_triplet_collection_types::device triplets(triplet_view); @@ -95,8 +97,10 @@ inline void find_triplets( for (unsigned int i = mt_start_idx; i < mt_end_idx; ++i) { const sp_location spT_loc = mid_top_doublet_device[i].sp2; + const unsigned int spT_idx = + sp_grid.bin(spT_loc.bin_idx)[spT_loc.sp_idx]; const edm::spacepoint_collection::const_device::const_proxy_type spT = - spacepoints.at(sp_grid.bin(spT_loc.bin_idx)[spT_loc.sp_idx]); + spacepoints.at(spT_idx); // Apply the conformal transformation to middle-top doublet const traccc::lin_circle lt = @@ -110,7 +114,7 @@ inline void find_triplets( // Add triplet to jagged vector triplets.at(posTriplets++) = device_triplet( - {spT_loc, globalIndex, curvature, + {spB_idx, spM_idx, spT_idx, globalIndex, curvature, -impact_parameter * filter_config.impactWeightFactor, lb.Zo()}); } diff --git a/device/common/include/traccc/seeding/device/impl/select_seeds.ipp b/device/common/include/traccc/seeding/device/impl/select_seeds.ipp index 0d67eb01aa..786331cfa6 100644 --- a/device/common/include/traccc/seeding/device/impl/select_seeds.ipp +++ b/device/common/include/traccc/seeding/device/impl/select_seeds.ipp @@ -20,7 +20,7 @@ namespace traccc::device { namespace details { // Finding minimum element algorithm template -TRACCC_HOST_DEVICE std::size_t min_elem(const triplet* arr, +TRACCC_HOST_DEVICE std::size_t min_elem(const device_triplet* arr, const std::size_t begin_idx, const std::size_t end_idx, Comparator comp) { @@ -38,11 +38,11 @@ TRACCC_HOST_DEVICE std::size_t min_elem(const triplet* arr, // Sorting algorithm for sorting seeds in the local memory template -TRACCC_HOST_DEVICE void insertionSort(triplet* arr, +TRACCC_HOST_DEVICE void insertionSort(device_triplet* arr, const unsigned int begin_idx, const unsigned int n, Comparator comp) { int j = 0; - triplet key = arr[begin_idx]; + device_triplet key = arr[begin_idx]; for (unsigned int i = 0; i < n; ++i) { key = arr[begin_idx + i]; j = static_cast(i) - 1; @@ -67,7 +67,7 @@ inline void select_seeds( const triplet_counter_spM_collection_types::const_view& spM_tc_view, const triplet_counter_collection_types::const_view& tc_view, const device_triplet_collection_types::const_view& triplet_view, - triplet* data, edm::seed_collection::view seed_view) { + device_triplet* data, edm::seed_collection::view seed_view) { // Check if anything needs to be done. const triplet_counter_spM_collection_types::const_device triplet_counts_spM( @@ -90,8 +90,9 @@ inline void select_seeds( // Current work item = middle spacepoint const triplet_counter_spM spM_counter = triplet_counts_spM.at(globalIndex); const sp_location spM_loc = spM_counter.spM; + const unsigned int spM_idx = sp_device.bin(spM_loc.bin_idx)[spM_loc.sp_idx]; const edm::spacepoint_collection::const_device::const_proxy_type spM = - spacepoints.at(sp_device.bin(spM_loc.bin_idx)[spM_loc.sp_idx]); + spacepoints.at(spM_idx); // Number of triplets added for this spM unsigned int n_triplets_per_spM = 0; @@ -103,14 +104,12 @@ inline void select_seeds( device_triplet aTriplet = triplets[i]; // spacepoints bottom and top for this triplet - const sp_location spB_loc = - triplet_counts.at(static_cast(aTriplet.counter_link)) - .spB; - const sp_location spT_loc = aTriplet.spT; + const unsigned int spB_idx = aTriplet.spB; const edm::spacepoint_collection::const_device::const_proxy_type spB = - spacepoints.at(sp_device.bin(spB_loc.bin_idx)[spB_loc.sp_idx]); + spacepoints.at(spB_idx); + const unsigned int spT_idx = aTriplet.spT; const edm::spacepoint_collection::const_device::const_proxy_type spT = - spacepoints.at(sp_device.bin(spT_loc.bin_idx)[spT_loc.sp_idx]); + spacepoints.at(spT_idx); // update weight of triplet seed_selecting_helper::seed_weight(filter_config, spM, spB, spT, @@ -126,27 +125,23 @@ inline void select_seeds( // the triplet with the lowest weight is removed if (n_triplets_per_spM >= finder_config.maxSeedsPerSpM) { - const std::size_t min_index = - details::min_elem(data, 0, finder_config.maxSeedsPerSpM, - [](const triplet lhs, const triplet rhs) { - return lhs.weight > rhs.weight; - }); + const std::size_t min_index = details::min_elem( + data, 0, finder_config.maxSeedsPerSpM, + [](const device_triplet& lhs, const device_triplet& rhs) { + return lhs.weight > rhs.weight; + }); const scalar& min_weight = data[min_index].weight; if (aTriplet.weight > min_weight) { - data[min_index] = {spB_loc, spM_loc, - spT_loc, aTriplet.curvature, - aTriplet.weight, aTriplet.z_vertex}; + data[min_index] = aTriplet; } } // if the number of good triplets is below the threshold, add // the current triplet to the array else if (n_triplets_per_spM < finder_config.maxSeedsPerSpM) { - data[n_triplets_per_spM] = {spB_loc, spM_loc, - spT_loc, aTriplet.curvature, - aTriplet.weight, aTriplet.z_vertex}; + data[n_triplets_per_spM] = aTriplet; n_triplets_per_spM++; } } @@ -154,16 +149,16 @@ inline void select_seeds( // sort the triplets per spM details::insertionSort( data, 0, n_triplets_per_spM, - traccc::details::triplet_sorter{spacepoints, sp_device}); + [](const device_triplet& lhs, const device_triplet& rhs) { + return lhs.weight > rhs.weight; + }); // the number of good seed per compatible middle spacepoint unsigned int n_seeds_per_spM = 0; // iterate over the good triplets for final selection of seeds for (unsigned int i = 0; i < n_triplets_per_spM; ++i) { - const triplet& aTriplet = data[i]; - const sp_location& spB_loc = aTriplet.sp1; - const sp_location& spT_loc = aTriplet.sp3; + const device_triplet& aTriplet = data[i]; // if the number of seeds reaches the threshold, break if (n_seeds_per_spM >= finder_config.maxSeedsPerSpM + 1) { @@ -171,16 +166,13 @@ inline void select_seeds( } // check if it is a good triplet - if (seed_selecting_helper::cut_per_middle_sp(filter_config, spacepoints, - sp_device, aTriplet) || + if (seed_selecting_helper::cut_per_middle_sp( + filter_config, spacepoints.at(aTriplet.spB), aTriplet.weight) || n_seeds_per_spM == 0) { n_seeds_per_spM++; - seeds_device.push_back( - {sp_device.bin(spB_loc.bin_idx)[spB_loc.sp_idx], - sp_device.bin(spM_loc.bin_idx)[spM_loc.sp_idx], - sp_device.bin(spT_loc.bin_idx)[spT_loc.sp_idx]}); + seeds_device.push_back({aTriplet.spB, aTriplet.spM, aTriplet.spT}); } } } diff --git a/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp b/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp index 42fcf1d52c..62ef23cda8 100644 --- a/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp +++ b/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp @@ -19,7 +19,6 @@ TRACCC_HOST_DEVICE inline void update_triplet_weights( const global_index_t globalIndex, const seedfilter_config& filter_config, const edm::spacepoint_collection::const_view& spacepoints_view, - const traccc::details::spacepoint_grid_types::const_view& sp_view, const triplet_counter_spM_collection_types::const_view& spM_tc_view, const triplet_counter_collection_types::const_view& tc_view, scalar* data, device_triplet_collection_types::view triplet_view) { @@ -33,7 +32,6 @@ inline void update_triplet_weights( // Set up the device containers const edm::spacepoint_collection::const_device spacepoints{ spacepoints_view}; - const traccc::details::spacepoint_grid_types::const_device sp_grid(sp_view); const triplet_counter_spM_collection_types::const_device triplet_counts_spM( spM_tc_view); const triplet_counter_collection_types::const_device triplet_counts( @@ -42,11 +40,8 @@ inline void update_triplet_weights( // Current work item device_triplet this_triplet = triplets.at(globalIndex); - const sp_location& spT_idx = this_triplet.spT; - const edm::spacepoint_collection::const_device::const_proxy_type - current_spT = - spacepoints.at(sp_grid.bin(spT_idx.bin_idx)[spT_idx.sp_idx]); + current_spT = spacepoints.at(this_triplet.spT); const scalar currentTop_r = current_spT.radius(); @@ -82,10 +77,8 @@ inline void update_triplet_weights( } const device_triplet other_triplet = triplets[i]; - const sp_location other_spT_idx = other_triplet.spT; const edm::spacepoint_collection::const_device::const_proxy_type - other_spT = spacepoints.at( - sp_grid.bin(other_spT_idx.bin_idx)[other_spT_idx.sp_idx]); + other_spT = spacepoints.at(other_triplet.spT); // compared top SP should have at least deltaRMin distance const scalar otherTop_r = other_spT.radius(); diff --git a/device/common/include/traccc/seeding/device/update_triplet_weights.hpp b/device/common/include/traccc/seeding/device/update_triplet_weights.hpp index a0bb68dbbd..c33a591ab3 100644 --- a/device/common/include/traccc/seeding/device/update_triplet_weights.hpp +++ b/device/common/include/traccc/seeding/device/update_triplet_weights.hpp @@ -37,7 +37,6 @@ TRACCC_HOST_DEVICE inline void update_triplet_weights( global_index_t globalIndex, const seedfilter_config& filter_config, const edm::spacepoint_collection::const_view& spacepoints, - const traccc::details::spacepoint_grid_types::const_view& sp_view, const triplet_counter_spM_collection_types::const_view& spM_tc_view, const triplet_counter_collection_types::const_view& tc_view, scalar* data, device_triplet_collection_types::view triplet_view); diff --git a/device/cuda/src/seeding/seed_finding.cu b/device/cuda/src/seeding/seed_finding.cu index 6c74d3d341..ac0decf2df 100644 --- a/device/cuda/src/seeding/seed_finding.cu +++ b/device/cuda/src/seeding/seed_finding.cu @@ -111,7 +111,6 @@ __global__ void find_triplets( __global__ void update_triplet_weights( seedfilter_config filter_config, edm::spacepoint_collection::const_view spacepoints, - traccc::details::spacepoint_grid_types::const_view sp_grid, device::triplet_counter_spM_collection_types::const_view spM_tc, device::triplet_counter_collection_types::const_view midBot_tc, device::device_triplet_collection_types::view triplet_view) { @@ -123,8 +122,8 @@ __global__ void update_triplet_weights( scalar* dataPos = &data[threadIdx.x * filter_config.compatSeedLimit]; device::update_triplet_weights(details::global_index1(), filter_config, - spacepoints, sp_grid, spM_tc, midBot_tc, - dataPos, triplet_view); + spacepoints, spM_tc, midBot_tc, dataPos, + triplet_view); } /// CUDA kernel for running @c traccc::device::select_seeds @@ -139,9 +138,10 @@ __global__ void select_seeds( // Array for temporary storage of triplets for comparing within seed // selecting kernel - extern __shared__ triplet data2[]; - // Each thread uses maxSeedsPerSpM elements of the array - triplet* dataPos = &data2[threadIdx.x * finder_config.maxSeedsPerSpM]; + extern __shared__ device::device_triplet data2[]; + // Each thread uses max_triplets_per_spM elements of the array + device::device_triplet* dataPos = + &data2[threadIdx.x * finder_config.maxSeedsPerSpM]; device::select_seeds(details::global_index1(), finder_config, filter_config, spacepoints, sp_view, spM_tc, midBot_tc, triplet_view, @@ -336,7 +336,7 @@ edm::seed_collection::buffer seed_finding::operator()( nWeightUpdatingBlocks, nWeightUpdatingThreads, sizeof(scalar) * m_seedfilter_config.compatSeedLimit * nWeightUpdatingThreads, - stream>>>(m_seedfilter_config, spacepoints_view, g2_view, + stream>>>(m_seedfilter_config, spacepoints_view, triplet_counter_spM_buffer, triplet_counter_midBot_buffer, triplet_buffer); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); @@ -355,14 +355,14 @@ edm::seed_collection::buffer seed_finding::operator()( nSeedSelectingThreads; // Create seeds out of selected triplets - kernels:: - select_seeds<<>>( - m_seedfinder_config, m_seedfilter_config, spacepoints_view, g2_view, - triplet_counter_spM_buffer, triplet_counter_midBot_buffer, - triplet_buffer, seed_buffer); + kernels::select_seeds<<>>( + m_seedfinder_config, m_seedfilter_config, spacepoints_view, g2_view, + triplet_counter_spM_buffer, triplet_counter_midBot_buffer, + triplet_buffer, seed_buffer); TRACCC_CUDA_ERROR_CHECK(cudaGetLastError()); return seed_buffer; diff --git a/device/sycl/src/seeding/seed_finding.sycl b/device/sycl/src/seeding/seed_finding.sycl index 9224d5f147..cf90d62e1d 100644 --- a/device/sycl/src/seeding/seed_finding.sycl +++ b/device/sycl/src/seeding/seed_finding.sycl @@ -308,7 +308,7 @@ edm::seed_collection::buffer seed_finding::operator()( h.parallel_for( weightUpdatingRange, - [filter_config = m_seedfilter_config, spacepoints_view, g2_view, + [filter_config = m_seedfilter_config, spacepoints_view, triplet_counter_spM_view, triplet_counter_midBot_view, local_mem, triplet_view](::sycl::nd_item<1> item) { // Each thread uses compatSeedLimit elements of the array @@ -317,7 +317,7 @@ edm::seed_collection::buffer seed_finding::operator()( device::update_triplet_weights( details::global_index(item), filter_config, - spacepoints_view, g2_view, triplet_counter_spM_view, + spacepoints_view, triplet_counter_spM_view, triplet_counter_midBot_view, dataPos, triplet_view); }); }); @@ -349,7 +349,7 @@ edm::seed_collection::buffer seed_finding::operator()( .submit([&](::sycl::handler& h) { // Array for temporary storage of triplets for comparing within // kernel - vecmem::sycl::local_accessor local_mem( + vecmem::sycl::local_accessor local_mem( m_seedfinder_config.maxSeedsPerSpM * seedSelectingLocalSize, h); h.parallel_for( @@ -359,8 +359,9 @@ edm::seed_collection::buffer seed_finding::operator()( triplet_counter_spM_view, triplet_counter_midBot_view, triplet_view, local_mem, seed_view](::sycl::nd_item<1> item) { // Each thread uses compatSeedLimit elements of the array - triplet* dataPos = &local_mem[item.get_local_id() * - finder_config.maxSeedsPerSpM]; + device::device_triplet* dataPos = + &local_mem[item.get_local_id() * + finder_config.maxSeedsPerSpM]; device::select_seeds( details::global_index(item), finder_config,