diff --git a/core/include/traccc/finding/finding_algorithm.ipp b/core/include/traccc/finding/finding_algorithm.ipp index 2401615e00..e3a45ffcec 100644 --- a/core/include/traccc/finding/finding_algorithm.ipp +++ b/core/include/traccc/finding/finding_algorithm.ipp @@ -29,30 +29,21 @@ finding_algorithm::operator()( * Measurement Operations *****************************************************************/ - // Copy the measurements - measurement_collection_types::host sorted_measurements = measurements; - - // Sort the measurements w.r.t geometry barcode - std::sort(sorted_measurements.begin(), sorted_measurements.end(), - measurement_sort_comp()); - // Get copy of barcode uniques std::vector uniques; - uniques.resize(sorted_measurements.size()); + uniques.resize(measurements.size()); - auto end = - std::unique_copy(sorted_measurements.begin(), sorted_measurements.end(), - uniques.begin(), measurement_equal_comp()); + auto end = std::unique_copy(measurements.begin(), measurements.end(), + uniques.begin(), measurement_equal_comp()); unsigned int n_modules = end - uniques.begin(); // Get upper bounds of unique elements std::vector upper_bounds; upper_bounds.reserve(n_modules); for (unsigned int i = 0; i < n_modules; i++) { - auto up = std::upper_bound(sorted_measurements.begin(), - sorted_measurements.end(), uniques[i], - measurement_sort_comp()); - upper_bounds.push_back(std::distance(sorted_measurements.begin(), up)); + auto up = std::upper_bound(measurements.begin(), measurements.end(), + uniques[i], measurement_sort_comp()); + upper_bounds.push_back(std::distance(measurements.begin(), up)); } // Get the number of measurements of each module @@ -178,7 +169,7 @@ finding_algorithm::operator()( bound_track_parameters bound_param(in_param.surface_link(), in_param.vector(), in_param.covariance()); - const auto& meas = sorted_measurements[item_id]; + const auto& meas = measurements[item_id]; track_state trk_state(meas); @@ -281,7 +272,7 @@ finding_algorithm::operator()( auto& cand = *it; - cand = sorted_measurements.at(L.meas_idx); + cand = measurements.at(L.meas_idx); // Break the loop if the iterator is at the first candidate and // fill the seed diff --git a/device/cuda/src/clusterization/experimental/clusterization_algorithm.cu b/device/cuda/src/clusterization/experimental/clusterization_algorithm.cu index a46d15f858..ce1e635f56 100644 --- a/device/cuda/src/clusterization/experimental/clusterization_algorithm.cu +++ b/device/cuda/src/clusterization/experimental/clusterization_algorithm.cu @@ -19,6 +19,10 @@ // Vecmem include(s). #include +// Thrust include(s). +#include +#include + namespace traccc::cuda::experimental { namespace { @@ -141,6 +145,10 @@ clusterization_algorithm::output_type clusterization_algorithm::operator()( m_stream.synchronize(); + // Sort the measurements w.r.t geometry barcode + thrust::sort(thrust::cuda::par.on(stream), new_measurements_device.begin(), + new_measurements_device.end(), measurement_sort_comp()); + return new_measurements_buffer; } diff --git a/device/cuda/src/finding/finding_algorithm.cu b/device/cuda/src/finding/finding_algorithm.cu index cc4cdfa03c..48a34421e1 100644 --- a/device/cuda/src/finding/finding_algorithm.cu +++ b/device/cuda/src/finding/finding_algorithm.cu @@ -225,27 +225,16 @@ finding_algorithm::operator()( * Measurement Operations *****************************************************************/ - // Copy the measurements - measurement_collection_types::buffer sorted_measurements_buffer( - m_copy->get_size(measurements), m_mr.main); - measurement_collection_types::device sorted_measurements( - sorted_measurements_buffer); measurement_collection_types::const_device measurements_device( measurements); - thrust::copy(thrust::device, measurements_device.begin(), - measurements_device.end(), sorted_measurements.begin()); - - // Sort the measurements w.r.t geometry barcode - thrust::sort(thrust::device, sorted_measurements.begin(), - sorted_measurements.end(), measurement_sort_comp()); // Get copy of barcode uniques measurement_collection_types::buffer uniques_buffer{ - sorted_measurements.size(), m_mr.main}; + measurements_device.size(), m_mr.main}; measurement_collection_types::device uniques(uniques_buffer); measurement* end = thrust::unique_copy( - thrust::device, sorted_measurements.begin(), sorted_measurements.end(), + thrust::device, measurements_device.begin(), measurements_device.end(), uniques.begin(), measurement_equal_comp()); unsigned int n_modules = end - uniques.begin(); @@ -254,8 +243,8 @@ finding_algorithm::operator()( m_mr.main}; vecmem::device_vector upper_bounds(upper_bounds_buffer); - thrust::upper_bound(thrust::device, sorted_measurements.begin(), - sorted_measurements.end(), uniques.begin(), + thrust::upper_bound(thrust::device, measurements_device.begin(), + measurements_device.end(), uniques.begin(), uniques.begin() + n_modules, upper_bounds.begin(), measurement_sort_comp()); @@ -267,7 +256,7 @@ finding_algorithm::operator()( upper_bounds.end(), sizes.begin()); // Number of total measurements - const unsigned int n_total_measurements = sorted_measurements.size(); + const unsigned int n_total_measurements = measurements_device.size(); /***************************************************************** * Kernel1: Create barcode sequence @@ -365,10 +354,9 @@ finding_algorithm::operator()( if (nBlocks > 0) { kernels::find_tracks <<>>( - m_cfg, det_view, sorted_measurements_buffer, - barcodes_buffer, upper_bounds_buffer, in_params_buffer, - n_threads_buffer, step, - (*global_counter_device).n_measurements_per_thread, + 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, updated_params_buffer, link_map[step], (*global_counter_device).n_candidates); @@ -512,8 +500,8 @@ finding_algorithm::operator()( nThreads = WARP_SIZE * 2; nBlocks = (n_tips_total + nThreads - 1) / nThreads; kernels::build_tracks<<>>( - sorted_measurements_buffer, seeds_buffer, links_buffer, - param_to_link_buffer, tips_buffer, track_candidates_buffer); + measurements, seeds_buffer, links_buffer, param_to_link_buffer, + tips_buffer, track_candidates_buffer); CUDA_ERROR_CHECK(cudaGetLastError()); CUDA_ERROR_CHECK(cudaDeviceSynchronize()); diff --git a/device/sycl/src/clusterization/experimental/clusterization_algorithm.sycl b/device/sycl/src/clusterization/experimental/clusterization_algorithm.sycl index 4f81352885..3a276bb24a 100644 --- a/device/sycl/src/clusterization/experimental/clusterization_algorithm.sycl +++ b/device/sycl/src/clusterization/experimental/clusterization_algorithm.sycl @@ -159,6 +159,12 @@ clusterization_algorithm::output_type clusterization_algorithm::operator()( sizeof(measurement) * (*num_measurements_host)) .wait_and_throw(); + // @NOTE Uncomment once the onedpl is available + // oneapi::dpl::experimental::sort_async( + // oneapi::dpl::execution::dpcpp_default, + // new_measurements_device.begin(), new_measurements_device.end(), + // measurement_sort_comp()); + return new_measurements_buffer; } diff --git a/io/src/csv/read_measurements.cpp b/io/src/csv/read_measurements.cpp index b9e39b6c20..d5e262d78a 100644 --- a/io/src/csv/read_measurements.cpp +++ b/io/src/csv/read_measurements.cpp @@ -81,6 +81,9 @@ void read_measurements(measurement_reader_output& out, result_measurements.push_back(meas); } + + std::sort(result_measurements.begin(), result_measurements.end(), + measurement_sort_comp()); } measurement_container_types::host read_measurements_container(