diff --git a/core/include/traccc/clusterization/detail/measurement_creation_helper.hpp b/core/include/traccc/clusterization/detail/measurement_creation_helper.hpp index 8a309ace69..d8bb2b1f37 100644 --- a/core/include/traccc/clusterization/detail/measurement_creation_helper.hpp +++ b/core/include/traccc/clusterization/detail/measurement_creation_helper.hpp @@ -79,7 +79,7 @@ TRACCC_HOST inline void calc_cluster_properties( /// @param[in] module is the cell module where the cluster belongs to /// @param[in] module_link is the module index /// -TRACCC_HOST void fill_measurement( +TRACCC_HOST inline void fill_measurement( measurement_collection_types::host& measurements, const cell_collection_types::host& cluster, const cell_module& module, const unsigned int module_link) { @@ -102,7 +102,7 @@ TRACCC_HOST void fill_measurement( if (totalWeight > 0.) { measurement m; m.module_link = module_link; - m.surface_link = detray::geometry::barcode{module.module}; + m.surface_link = module.surface_link; // normalize the cell position m.local = mean; // normalize the variance diff --git a/core/include/traccc/edm/cell.hpp b/core/include/traccc/edm/cell.hpp index dbb010726f..fbfa230a90 100644 --- a/core/include/traccc/edm/cell.hpp +++ b/core/include/traccc/edm/cell.hpp @@ -13,6 +13,9 @@ #include "traccc/edm/container.hpp" #include "traccc/geometry/pixel_data.hpp" +// Detray include(s). +#include "detray/geometry/barcode.hpp" + namespace traccc { /// Definition of a detector module @@ -23,7 +26,7 @@ namespace traccc { /// struct cell_module { - geometry_id module = 0; + detray::geometry::barcode surface_link{0u}; transform3 placement = transform3{}; scalar threshold = 0; @@ -37,7 +40,7 @@ using cell_module_collection_types = collection_types; /// Equality operator for cell module TRACCC_HOST_DEVICE inline bool operator==(const cell_module& lhs, const cell_module& rhs) { - return lhs.module == rhs.module; + return lhs.surface_link == rhs.surface_link; } /// Definition for one detector cell diff --git a/device/common/include/traccc/clusterization/device/impl/aggregate_cluster.ipp b/device/common/include/traccc/clusterization/device/impl/aggregate_cluster.ipp index 80b4c9a941..0663e43768 100644 --- a/device/common/include/traccc/clusterization/device/impl/aggregate_cluster.ipp +++ b/device/common/include/traccc/clusterization/device/impl/aggregate_cluster.ipp @@ -106,8 +106,8 @@ inline void aggregate_cluster( */ out.local = mean; out.variance = var; + out.surface_link = this_module.surface_link; out.module_link = module_link; - out.surface_link = detray::geometry::barcode{this_module.module}; // The following will need to be filled properly "soon". out.meas_dim = 2u; } diff --git a/device/common/include/traccc/clusterization/device/impl/reduce_problem_cell.ipp b/device/common/include/traccc/clusterization/device/impl/reduce_problem_cell.ipp index ba2dc90d98..10a609ace5 100644 --- a/device/common/include/traccc/clusterization/device/impl/reduce_problem_cell.ipp +++ b/device/common/include/traccc/clusterization/device/impl/reduce_problem_cell.ipp @@ -14,8 +14,8 @@ namespace traccc::device { * cluster. */ TRACCC_HOST_DEVICE -bool is_adjacent(channel_id ac0, channel_id ac1, channel_id bc0, - channel_id bc1) { +inline bool is_adjacent(channel_id ac0, channel_id ac1, channel_id bc0, + channel_id bc1) { unsigned int p0 = (ac0 - bc0); unsigned int p1 = (ac1 - bc1); diff --git a/device/cuda/CMakeLists.txt b/device/cuda/CMakeLists.txt index 863ec069e8..a8f29762c6 100644 --- a/device/cuda/CMakeLists.txt +++ b/device/cuda/CMakeLists.txt @@ -42,6 +42,8 @@ traccc_add_library( traccc_cuda cuda TYPE SHARED "src/seeding/seeding_algorithm.cpp" "src/cca/component_connection.cu" # Clusterization + "include/traccc/cuda/clusterization/experimental/clusterization_algorithm.hpp" + "src/clusterization/experimental/clusterization_algorithm.cu" "include/traccc/cuda/clusterization/clusterization_algorithm.hpp" "src/clusterization/clusterization_algorithm.cu" # Finding diff --git a/device/cuda/include/traccc/cuda/clusterization/experimental/clusterization_algorithm.hpp b/device/cuda/include/traccc/cuda/clusterization/experimental/clusterization_algorithm.hpp new file mode 100644 index 0000000000..4208ab7a46 --- /dev/null +++ b/device/cuda/include/traccc/cuda/clusterization/experimental/clusterization_algorithm.hpp @@ -0,0 +1,73 @@ +/** 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 + +// Local include(s). +#include "traccc/cuda/utils/stream.hpp" + +// Project include(s). +#include "traccc/edm/cell.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/utils/algorithm.hpp" +#include "traccc/utils/memory_resource.hpp" + +// VecMem include(s). +#include + +namespace traccc::cuda::experimental { + +/// Algorithm performing hit clusterization +/// +/// This algorithm implements hit clusterization in a massively-parallel +/// approach. Each thread handles a pre-determined number of detector cells. +/// +/// This algorithm returns a buffer which is not necessarily filled yet. A +/// synchronisation statement is required before destroying this buffer. +/// +class clusterization_algorithm + : public algorithm { + + public: + /// Constructor for clusterization algorithm + /// + /// @param mr The memory resource(s) to use in the algorithm + /// @param copy The copy object to use for copying data between device + /// and host memory blocks + /// @param str The CUDA stream to perform the operations in + /// @param target_cells_per_partition the average number of cells in each + /// partition + /// + clusterization_algorithm(const traccc::memory_resource& mr, + vecmem::copy& copy, stream& str, + const unsigned short target_cells_per_partition); + // const unsigned short target_cells_per_partition); + + /// Callable operator for clusterization algorithm + /// + /// @param cells a collection of cells + /// @param modules a collection of modules + /// @return a spacepoint collection (buffer) and a collection (buffer) + /// of links from cells to the spacepoints they belong to. + output_type operator()( + const cell_collection_types::const_view& cells, + const cell_module_collection_types::const_view& modules) const override; + + private: + /// The average number of cells in each partition + unsigned short m_target_cells_per_partition; + /// The memory resource(s) to use + traccc::memory_resource m_mr; + /// The copy object to use + vecmem::copy& m_copy; + /// The CUDA stream to use + stream& m_stream; +}; + +} // namespace traccc::cuda::experimental \ No newline at end of file diff --git a/device/cuda/src/clusterization/experimental/clusterization_algorithm.cu b/device/cuda/src/clusterization/experimental/clusterization_algorithm.cu new file mode 100644 index 0000000000..a46d15f858 --- /dev/null +++ b/device/cuda/src/clusterization/experimental/clusterization_algorithm.cu @@ -0,0 +1,147 @@ +/** TRACCC library, part of the ACTS project (R&D line) + * + * (c) 2022 CERN for the benefit of the ACTS project + * + * Mozilla Public License Version 2.0 + */ + +// CUDA Library include(s). +#include "../../utils/utils.hpp" +#include "traccc/cuda/clusterization/experimental/clusterization_algorithm.hpp" +#include "traccc/cuda/utils/barrier.hpp" +#include "traccc/cuda/utils/definitions.hpp" + +// Project include(s) +#include "traccc/clusterization/device/aggregate_cluster.hpp" +#include "traccc/clusterization/device/ccl_kernel.hpp" +#include "traccc/clusterization/device/reduce_problem_cell.hpp" + +// Vecmem include(s). +#include + +namespace traccc::cuda::experimental { + +namespace { +/// These indices in clusterization will only range from 0 to +/// max_cells_per_partition, so we only need a short. +using index_t = unsigned short; + +static constexpr int TARGET_CELLS_PER_THREAD = 8; +static constexpr int MAX_CELLS_PER_THREAD = 12; +} // namespace + +namespace kernels { + +/// CUDA kernel for running @c traccc::device::ccl_kernel +__global__ void ccl_kernel( + const cell_collection_types::const_view cells_view, + const cell_module_collection_types::const_view modules_view, + const index_t max_cells_per_partition, + const index_t target_cells_per_partition, + measurement_collection_types::view measurements_view, + unsigned int& measurement_count, + vecmem::data::vector_view cell_links) { + __shared__ unsigned int partition_start, partition_end; + __shared__ unsigned int outi; + extern __shared__ index_t shared_v[]; + index_t* f = &shared_v[0]; + index_t* f_next = &shared_v[max_cells_per_partition]; + traccc::cuda::barrier barry_r; + + device::ccl_kernel(threadIdx.x, blockDim.x, blockIdx.x, cells_view, + modules_view, max_cells_per_partition, + target_cells_per_partition, partition_start, + partition_end, outi, f, f_next, barry_r, + measurements_view, measurement_count, cell_links); +} + +} // namespace kernels + +clusterization_algorithm::clusterization_algorithm( + const traccc::memory_resource& mr, vecmem::copy& copy, stream& str, + const unsigned short target_cells_per_partition) + : m_mr(mr), + m_copy(copy), + m_stream(str), + m_target_cells_per_partition(target_cells_per_partition) {} + +clusterization_algorithm::output_type clusterization_algorithm::operator()( + const cell_collection_types::const_view& cells, + const cell_module_collection_types::const_view& modules) const { + + // Get a convenience variable for the stream that we'll be using. + cudaStream_t stream = details::get_stream(m_stream); + + // Number of cells + const cell_collection_types::view::size_type num_cells = + m_copy.get_size(cells); + + if (num_cells == 0) { + return {0, m_mr.main}; + } + + // Create result object for the CCL kernel with size overestimation + measurement_collection_types::buffer measurements_buffer(num_cells, + m_mr.main); + m_copy.setup(measurements_buffer); + + // Counter for number of measurements + vecmem::unique_alloc_ptr num_measurements_device = + vecmem::make_unique_alloc(m_mr.main); + CUDA_ERROR_CHECK(cudaMemsetAsync(num_measurements_device.get(), 0, + sizeof(unsigned int), stream)); + + const unsigned short max_cells_per_partition = + (m_target_cells_per_partition * MAX_CELLS_PER_THREAD + + TARGET_CELLS_PER_THREAD - 1) / + TARGET_CELLS_PER_THREAD; + const unsigned int threads_per_partition = + (m_target_cells_per_partition + TARGET_CELLS_PER_THREAD - 1) / + TARGET_CELLS_PER_THREAD; + const unsigned int num_partitions = + (num_cells + m_target_cells_per_partition - 1) / + m_target_cells_per_partition; + + // Create buffer for linking cells to their spacepoints. + vecmem::data::vector_buffer cell_links(num_cells, m_mr.main); + m_copy.setup(cell_links); + + // Launch ccl kernel. Each thread will handle a single cell. + kernels:: + ccl_kernel<<>>( + cells, modules, max_cells_per_partition, + m_target_cells_per_partition, measurements_buffer, + *num_measurements_device, cell_links); + + CUDA_ERROR_CHECK(cudaGetLastError()); + + // Copy number of measurements to host + vecmem::unique_alloc_ptr num_measurements_host = + vecmem::make_unique_alloc( + (m_mr.host != nullptr) ? *(m_mr.host) : m_mr.main); + CUDA_ERROR_CHECK(cudaMemcpyAsync( + num_measurements_host.get(), num_measurements_device.get(), + sizeof(unsigned int), cudaMemcpyDeviceToHost, stream)); + m_stream.synchronize(); + + // Create a new measurement buffer with a right size + measurement_collection_types::buffer new_measurements_buffer( + *num_measurements_host, m_mr.main); + m_copy.setup(new_measurements_buffer); + + vecmem::device_vector measurements_device(measurements_buffer); + vecmem::device_vector new_measurements_device( + new_measurements_buffer); + + CUDA_ERROR_CHECK(cudaMemcpyAsync( + new_measurements_device.begin(), measurements_device.begin(), + sizeof(measurement) * (*num_measurements_host), + cudaMemcpyDeviceToDevice, stream)); + + m_stream.synchronize(); + + return new_measurements_buffer; +} + +} // namespace traccc::cuda::experimental diff --git a/device/sycl/CMakeLists.txt b/device/sycl/CMakeLists.txt index 14c6054615..c83e02f341 100644 --- a/device/sycl/CMakeLists.txt +++ b/device/sycl/CMakeLists.txt @@ -15,6 +15,7 @@ enable_language( SYCL ) traccc_add_library( traccc_sycl sycl TYPE SHARED # header files "include/traccc/sycl/clusterization/clusterization_algorithm.hpp" + "include/traccc/sycl/clusterization/experimental/clusterization_algorithm.hpp" "include/traccc/sycl/fitting/fitting_algorithm.hpp" "include/traccc/sycl/seeding/experimental/spacepoint_formation.hpp" "include/traccc/sycl/seeding/seeding_algorithm.hpp" @@ -27,6 +28,7 @@ traccc_add_library( traccc_sycl sycl TYPE SHARED "include/traccc/sycl/utils/barrier.hpp" # implementation files "src/clusterization/clusterization_algorithm.sycl" + "src/clusterization/experimental/clusterization_algorithm.sycl" "src/fitting/fitting_algorithm.sycl" "src/seeding/experimental/spacepoint_formation.sycl" "src/seeding/seed_finding.sycl" diff --git a/device/sycl/include/traccc/sycl/clusterization/experimental/clusterization_algorithm.hpp b/device/sycl/include/traccc/sycl/clusterization/experimental/clusterization_algorithm.hpp new file mode 100644 index 0000000000..2c60b558d1 --- /dev/null +++ b/device/sycl/include/traccc/sycl/clusterization/experimental/clusterization_algorithm.hpp @@ -0,0 +1,76 @@ +/** 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 + +// SYCL library include(s). +#include "traccc/sycl/utils/queue_wrapper.hpp" + +// Project include(s). +#include "traccc/edm/cell.hpp" +#include "traccc/edm/measurement.hpp" +#include "traccc/utils/algorithm.hpp" +#include "traccc/utils/memory_resource.hpp" + +// VecMem include(s). +#include + +namespace traccc::sycl::experimental { + +/// Algorithm performing hit clusterization +/// +/// This algorithm implements hit clusterization in a massively-parallel +/// approach. Each thread handles a pre-determined number of detector cells. +/// +/// This algorithm returns a buffer which is not necessarily filled yet. A +/// synchronisation statement is required before destroying this buffer. +/// +class clusterization_algorithm + : public algorithm { + + public: + /// Constructor for clusterization algorithm + /// + /// @param mr The memory resource(s) to use in the algorithm + /// @param copy The copy object to use for copying data between device + /// and host memory blocks + /// @param str The CUDA stream to perform the operations in + /// @param target_cells_per_partition the average number of cells in each + /// partition + /// + clusterization_algorithm(const traccc::memory_resource& mr, + vecmem::copy& copy, queue_wrapper queue, + const unsigned short target_cells_per_partition); + // const unsigned short target_cells_per_partition); + + /// Callable operator for clusterization algorithm + /// + /// @param cells a collection of cells + /// @param modules a collection of modules + /// @return a spacepoint collection (buffer) and a collection (buffer) + /// of links from cells to the spacepoints they belong to. + output_type operator()( + const cell_collection_types::const_view& cells, + const cell_module_collection_types::const_view& modules) const override; + + private: + /// The average number of cells in each partition + unsigned short m_target_cells_per_partition; + /// The maximum number of threads in a work group + unsigned int m_max_work_group_size; + + /// The memory resource(s) to use + traccc::memory_resource m_mr; + /// The copy object to use + vecmem::copy& m_copy; + /// The SYCL queue object + mutable queue_wrapper m_queue; +}; + +} // namespace traccc::sycl::experimental \ No newline at end of file diff --git a/device/sycl/src/clusterization/experimental/clusterization_algorithm.sycl b/device/sycl/src/clusterization/experimental/clusterization_algorithm.sycl new file mode 100644 index 0000000000..4f81352885 --- /dev/null +++ b/device/sycl/src/clusterization/experimental/clusterization_algorithm.sycl @@ -0,0 +1,165 @@ +/** 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 + */ + +// Local include(s). +#include "../../utils/get_queue.hpp" +#include "traccc/sycl/clusterization/experimental/clusterization_algorithm.hpp" +#include "traccc/sycl/utils/barrier.hpp" +#include "traccc/sycl/utils/calculate1DimNdRange.hpp" + +// Project include(s) +#include "traccc/clusterization/device/aggregate_cluster.hpp" +#include "traccc/clusterization/device/ccl_kernel.hpp" +#include "traccc/clusterization/device/reduce_problem_cell.hpp" + +// Vecmem include(s). +#include +#include +#include + +// System include(s). +#include + +namespace traccc::sycl::experimental { + +namespace { +/// These indices in clusterization will only range from 0 to +/// max_cells_per_partition, so we only need a short +using index_t = unsigned short; + +static constexpr int TARGET_CELLS_PER_THREAD = 8; +static constexpr int MAX_CELLS_PER_THREAD = 12; +} // namespace + +namespace kernels { + +/// Class identifying the kernel running @c traccc::device::ccl_kernel +class ccl_kernel; + +} // namespace kernels + +clusterization_algorithm::clusterization_algorithm( + const traccc::memory_resource& mr, vecmem::copy& copy, queue_wrapper queue, + const unsigned short target_cells_per_partition) + : m_target_cells_per_partition(target_cells_per_partition), + m_max_work_group_size( + details::get_queue(queue) + .get_device() + .get_info<::sycl::info::device::max_work_group_size>()), + m_mr(mr), + m_copy(copy), + m_queue(queue) {} + +clusterization_algorithm::output_type clusterization_algorithm::operator()( + const cell_collection_types::const_view& cells, + const cell_module_collection_types::const_view& modules) const { + + // Number of cells + const cell_collection_types::view::size_type num_cells = + m_copy.get_size(cells); + + if (num_cells == 0) { + return {0, m_mr.main}; + } + + // Create result object for the CCL kernel with size overestimation + measurement_collection_types::buffer measurements_buffer(num_cells, + m_mr.main); + m_copy.setup(measurements_buffer)->wait(); + measurement_collection_types::view measurements_view(measurements_buffer); + + // Counter for number of measurements + vecmem::unique_alloc_ptr num_measurements_device = + vecmem::make_unique_alloc(m_mr.main); + details::get_queue(m_queue) + .memset(num_measurements_device.get(), 0, sizeof(unsigned int)) + .wait_and_throw(); + + const unsigned short max_cells_per_partition = + (m_target_cells_per_partition * MAX_CELLS_PER_THREAD + + TARGET_CELLS_PER_THREAD - 1) / + TARGET_CELLS_PER_THREAD; + const unsigned int threads_per_partition = + (m_target_cells_per_partition + TARGET_CELLS_PER_THREAD - 1) / + TARGET_CELLS_PER_THREAD; + const unsigned int num_partitions = + (num_cells + m_target_cells_per_partition - 1) / + m_target_cells_per_partition; + const unsigned int target_cells_per_partition = + m_target_cells_per_partition; + + ::sycl::nd_range cclKernelRange( + ::sycl::range<1>(num_partitions * threads_per_partition), + ::sycl::range<1>(threads_per_partition)); + + // Check if device is capable of allocating sufficient local memory + assert(sizeof(index_t) * 2 * max_cells_per_partition + + 3 * sizeof(unsigned int) < + details::get_queue(m_queue) + .get_device() + .get_info<::sycl::info::device::local_mem_size>()); + + // Create buffer for linking cells to their spacepoints. + vecmem::data::vector_buffer cell_links(num_cells, m_mr.main); + m_copy.setup(cell_links)->wait(); + vecmem::data::vector_view cell_links_view(cell_links); + + auto aux_num_measurements_device = num_measurements_device.get(); + // Run ccl kernel + details::get_queue(m_queue) + .submit([&](::sycl::handler& h) { + vecmem::sycl::local_accessor shared_uint(3, h); + vecmem::sycl::local_accessor shared_idx( + 2 * max_cells_per_partition, h); + + h.parallel_for( + cclKernelRange, [=](::sycl::nd_item<1> item) { + index_t* f = &shared_idx[0]; + index_t* f_next = &shared_idx[max_cells_per_partition]; + unsigned int& partition_start = shared_uint[0]; + unsigned int& partition_end = shared_uint[1]; + unsigned int& outi = shared_uint[2]; + traccc::sycl::barrier barry_r(item); + + device::ccl_kernel( + item.get_local_linear_id(), item.get_local_range(0), + item.get_group_linear_id(), cells, modules, + max_cells_per_partition, target_cells_per_partition, + partition_start, partition_end, outi, f, f_next, + barry_r, measurements_view, + *aux_num_measurements_device, cell_links_view); + }); + }) + .wait_and_throw(); + + // Copy number of measurements to host + vecmem::unique_alloc_ptr num_measurements_host = + vecmem::make_unique_alloc( + (m_mr.host != nullptr) ? *(m_mr.host) : m_mr.main); + details::get_queue(m_queue) + .memcpy(num_measurements_host.get(), num_measurements_device.get(), + sizeof(unsigned int)) + .wait_and_throw(); + + // Create a new measurement buffer with a right size + measurement_collection_types::buffer new_measurements_buffer( + *num_measurements_host, m_mr.main); + m_copy.setup(new_measurements_buffer); + + vecmem::device_vector measurements_device(measurements_buffer); + vecmem::device_vector new_measurements_device( + new_measurements_buffer); + + details::get_queue(m_queue) + .memcpy(new_measurements_device.begin(), measurements_device.begin(), + sizeof(measurement) * (*num_measurements_host)) + .wait_and_throw(); + + return new_measurements_buffer; +} + +} // namespace traccc::sycl::experimental \ No newline at end of file diff --git a/io/src/csv/read_cells.cpp b/io/src/csv/read_cells.cpp index 7f24b50007..f3a8be2838 100644 --- a/io/src/csv/read_cells.cpp +++ b/io/src/csv/read_cells.cpp @@ -32,21 +32,20 @@ traccc::cell_module get_module(traccc::io::csv::cell c, const traccc::digitization_config* dconfig) { traccc::cell_module result; - - result.module = c.geometry_id; + result.surface_link = detray::geometry::barcode{c.geometry_id}; // Find/set the 3D position of the detector module. if (geom != nullptr) { // Check if the module ID is known. - if (!geom->contains(result.module)) { + if (!geom->contains(result.surface_link.value())) { throw std::runtime_error( "Could not find placement for geometry ID " + - std::to_string(result.module)); + std::to_string(result.surface_link.value())); } // Set the value on the module description. - result.placement = (*geom)[result.module]; + result.placement = (*geom)[result.surface_link.value()]; } // Find/set the digitization configuration of the detector module. @@ -54,11 +53,11 @@ traccc::cell_module get_module(traccc::io::csv::cell c, // Check if the module ID is known. const traccc::digitization_config::Iterator geo_it = - dconfig->find(result.module); + dconfig->find(result.surface_link.value()); if (geo_it == dconfig->end()) { throw std::runtime_error( "Could not find digitization config for geometry ID " + - std::to_string(result.module)); + std::to_string(result.surface_link.value())); } // Set the value on the module description. @@ -100,7 +99,8 @@ void read_cells(cell_reader_output& out, std::string_view filename, // Look for current module in cell counter vector. auto rit = std::find_if(result_modules.rbegin(), result_modules.rend(), [&iocell](const cell_module& mod) { - return mod.module == iocell.geometry_id; + return mod.surface_link.value() == + iocell.geometry_id; }); if (rit == result_modules.rend()) { // Add new cell and new cell counter if a new module is found diff --git a/io/src/csv/read_measurements.cpp b/io/src/csv/read_measurements.cpp index 0677311946..127df51f06 100644 --- a/io/src/csv/read_measurements.cpp +++ b/io/src/csv/read_measurements.cpp @@ -10,6 +10,9 @@ #include "make_measurement_reader.hpp" +// Detray include(s). +#include "detray/geometry/barcode.hpp" + // System include(s). #include @@ -40,7 +43,7 @@ void read_measurements(measurement_reader_output& out, link = result_modules.size(); m[iomeas.geometry_id] = link; cell_module mod; - mod.module = iomeas.geometry_id; + mod.surface_link = detray::geometry::barcode{iomeas.geometry_id}; result_modules.push_back(mod); } @@ -72,7 +75,7 @@ measurement_container_types::host read_measurements_container( // Construct the module ID for the measurement. cell_module module; - module.module = iomeas.geometry_id; + module.surface_link = detray::geometry::barcode{iomeas.geometry_id}; // Construct the measurement object. const traccc::measurement meas{ diff --git a/io/src/csv/read_spacepoints.cpp b/io/src/csv/read_spacepoints.cpp index d2d29a6c30..fe269a4cbb 100644 --- a/io/src/csv/read_spacepoints.cpp +++ b/io/src/csv/read_spacepoints.cpp @@ -12,6 +12,9 @@ #include "make_measurement_hit_id_reader.hpp" #include "read_measurements.hpp" +// Detray include(s). +#include "detray/geometry/barcode.hpp" + // System include(s). #include #include @@ -55,7 +58,7 @@ void read_spacepoints(spacepoint_reader_output& out, std::string_view filename, link = result_modules.size(); m[iohit.geometry_id] = link; cell_module mod; - mod.module = iohit.geometry_id; + mod.surface_link = detray::geometry::barcode{iohit.geometry_id}; mod.placement = geom[iohit.geometry_id]; result_modules.push_back(mod); } diff --git a/io/src/mapper.cpp b/io/src/mapper.cpp index 9625bb6a43..291f43cf7d 100644 --- a/io/src/mapper.cpp +++ b/io/src/mapper.cpp @@ -242,7 +242,7 @@ measurement_particle_map generate_measurement_particle_map( geoId_link_map link_map; for (unsigned int i = 0; i < modules.size(); ++i) { - link_map[modules[i].module] = i; + link_map[modules[i].surface_link.value()] = i; } // generate cell particle map @@ -279,7 +279,7 @@ measurement_particle_map generate_measurement_particle_map( geoId_link_map link_map; for (unsigned int i = 0; i < modules.size(); ++i) { - link_map[modules[i].module] = i; + link_map[modules[i].surface_link.value()] = i; } auto h_p_map = diff --git a/tests/cpu/test_cca.cpp b/tests/cpu/test_cca.cpp index ed680eb8b7..430d52f8af 100644 --- a/tests/cpu/test_cca.cpp +++ b/tests/cpu/test_cca.cpp @@ -34,8 +34,8 @@ cca_function_t f = [](const traccc::cell_collection_types::host& cells, auto measurements = ca(cells, modules); for (std::size_t i = 0; i < measurements.size(); i++) { - result[modules.at(measurements.at(i).module_link).module].push_back( - measurements.at(i)); + result[modules.at(measurements.at(i).module_link).surface_link.value()] + .push_back(measurements.at(i)); } return result; diff --git a/tests/cpu/test_clusterization_resolution.cpp b/tests/cpu/test_clusterization_resolution.cpp index 25a5b380b8..e4b9b617c1 100644 --- a/tests/cpu/test_clusterization_resolution.cpp +++ b/tests/cpu/test_clusterization_resolution.cpp @@ -76,8 +76,8 @@ TEST_P(SurfaceBinningTests, Run) { const auto& sp_truth = spacepoints_truth[i]; // Check that the spacepoints belong to the same module - EXPECT_EQ(modules.at(sp_recon.meas.module_link).module, - modules_2.at(sp_truth.meas.module_link).module); + EXPECT_EQ(modules.at(sp_recon.meas.module_link).surface_link, + modules_2.at(sp_truth.meas.module_link).surface_link); // Make sure that the difference in spacepoint position is less than // 1% diff --git a/tests/cuda/CMakeLists.txt b/tests/cuda/CMakeLists.txt index ee9bcf6319..54447ba707 100644 --- a/tests/cuda/CMakeLists.txt +++ b/tests/cuda/CMakeLists.txt @@ -30,6 +30,7 @@ traccc_add_test( test_basic.cu test_cca.cpp test_ckf_sparse_tracks.cpp + test_clusterization.cpp test_copy.cu test_kalman_filter.cpp test_spacepoint_formation.cpp diff --git a/tests/cuda/test_cca.cpp b/tests/cuda/test_cca.cpp index 76b4d7cb00..42b57231be 100644 --- a/tests/cuda/test_cca.cpp +++ b/tests/cuda/test_cca.cpp @@ -23,8 +23,8 @@ cca_function_t f = [](const traccc::cell_collection_types::host& cells, auto measurements = cc(cells); for (std::size_t i = 0; i < measurements.size(); i++) { - result[modules.at(measurements.at(i).module_link).module].push_back( - measurements.at(i)); + result[modules.at(measurements.at(i).module_link).surface_link.value()] + .push_back(measurements.at(i)); } return result; diff --git a/tests/cuda/test_clusterization.cpp b/tests/cuda/test_clusterization.cpp new file mode 100644 index 0000000000..cf4425acb8 --- /dev/null +++ b/tests/cuda/test_clusterization.cpp @@ -0,0 +1,71 @@ +/** 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 + */ + +// Project include(s). +#include "traccc/cuda/clusterization/experimental/clusterization_algorithm.hpp" +#include "traccc/definitions/common.hpp" + +// VecMem include(s). +#include +#include + +// GTest include(s). +#include + +using namespace traccc; + +TEST(clusterization, cuda) { + + // Memory resource used by the EDM. + vecmem::cuda::managed_memory_resource mng_mr; + traccc::memory_resource mr{mng_mr}; + + // Cuda stream + traccc::cuda::stream stream; + + // Cuda copy objects + vecmem::cuda::async_copy copy{stream.cudaStream()}; + + // Create cell collection + traccc::cell_collection_types::host cells{&mng_mr}; + + cells.push_back({1u, 2u, 1.f, 0, 0}); + cells.push_back({2u, 2u, 1.f, 0, 0}); + cells.push_back({3u, 2u, 1.f, 0, 0}); + + cells.push_back({5u, 5u, 1.f, 0, 0}); + cells.push_back({6u, 4u, 1.f, 0, 0}); + cells.push_back({6u, 5u, 1.f, 0, 0}); + cells.push_back({6u, 6u, 1.f, 0, 0}); + cells.push_back({7u, 5u, 1.f, 0, 0}); + + // Create module collection + traccc::cell_module_collection_types::host modules{&mng_mr}; + modules.push_back({}); + + // Run Clusterization + cuda::experimental::clusterization_algorithm ca_cuda(mr, copy, stream, + 1024); + + auto measurements_buffer = + ca_cuda(vecmem::get_data(cells), vecmem::get_data(modules)); + + measurement_collection_types::device measurements(measurements_buffer); + + // Check the results + EXPECT_EQ(copy.get_size(measurements_buffer), 2u); + std::set test; + test.insert(measurements[0]); + test.insert(measurements[1]); + + std::set ref; + ref.insert({{2.f, 2.f}, {0.75, 0.0833333}, detray::geometry::barcode{0u}}); + ref.insert( + {{6.f, 5.f}, {0.483333, 0.483333}, detray::geometry::barcode{0u}}); + + EXPECT_EQ(test, ref); +} \ No newline at end of file diff --git a/tests/io/test_binary.cpp b/tests/io/test_binary.cpp index a07637e21f..2689c53c8f 100644 --- a/tests/io/test_binary.cpp +++ b/tests/io/test_binary.cpp @@ -93,7 +93,7 @@ TEST(io_binary, cell) { ASSERT_EQ(cells_csv[i], cells_binary[i]); } for (std::size_t i = 0; i < modules_csv.size(); i++) { - ASSERT_EQ(modules_csv[i].module, modules_binary[i].module); + ASSERT_EQ(modules_csv[i].surface_link, modules_binary[i].surface_link); ASSERT_EQ(modules_csv[i].placement, modules_binary[i].placement); } } @@ -226,6 +226,6 @@ TEST(io_binary, measurement) { ASSERT_EQ(measurements_csv[i], measurements_binary[i]); } for (std::size_t i = 0; i < modules_csv.size(); i++) { - ASSERT_EQ(modules_csv[i].module, modules_binary[i].module); + ASSERT_EQ(modules_csv[i].surface_link, modules_binary[i].surface_link); } } \ No newline at end of file diff --git a/tests/io/test_csv.cpp b/tests/io/test_csv.cpp index 80fa53953e..42395901c0 100644 --- a/tests/io/test_csv.cpp +++ b/tests/io/test_csv.cpp @@ -38,7 +38,7 @@ TEST_F(io, csv_read_single_module) { ASSERT_EQ(modules.size(), 1u); auto module = single_module_cells.modules.at(0); - ASSERT_EQ(module.module, 0u); + ASSERT_EQ(module.surface_link.value(), 0u); ASSERT_EQ(cells.at(0).channel0, 123u); ASSERT_EQ(cells.at(0).channel1, 32u); ASSERT_EQ(cells.at(5).channel0, 174u); @@ -65,7 +65,7 @@ TEST_F(io, csv_read_two_modules) { ASSERT_EQ(cells.at(5).channel1, 880u); ASSERT_EQ(cells.at(5).module_link, 0u); - ASSERT_EQ(modules.at(0u).module, 0u); + ASSERT_EQ(modules.at(0u).surface_link.value(), 0u); // Check cells in second module ASSERT_EQ(cells.at(6).channel0, 0u); @@ -75,7 +75,7 @@ TEST_F(io, csv_read_two_modules) { ASSERT_EQ(cells.at(13).channel1, 98u); ASSERT_EQ(cells.at(13).module_link, 1u); - ASSERT_EQ(modules.at(1u).module, 1u); + ASSERT_EQ(modules.at(1u).surface_link.value(), 1u); } // This reads in the tml pixel barrel first event diff --git a/tests/sycl/CMakeLists.txt b/tests/sycl/CMakeLists.txt index c34fb325be..a627496acc 100644 --- a/tests/sycl/CMakeLists.txt +++ b/tests/sycl/CMakeLists.txt @@ -12,6 +12,7 @@ traccc_add_test( sycl # Define the sources for the test. + test_clusterization.sycl test_kalman_filter.sycl test_spacepoint_formation.sycl diff --git a/tests/sycl/test_clusterization.sycl b/tests/sycl/test_clusterization.sycl new file mode 100644 index 0000000000..a95559226e --- /dev/null +++ b/tests/sycl/test_clusterization.sycl @@ -0,0 +1,85 @@ +/** 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 + */ + +// Project include(s). +#include "traccc/definitions/common.hpp" +#include "traccc/sycl/clusterization/experimental/clusterization_algorithm.hpp" + +// VecMem include(s). +#include +#include + +// GTest include(s). +#include + +using namespace traccc; + +// Simple asynchronous handler function +auto handle_async_error = [](::sycl::exception_list elist) { + for (auto& e : elist) { + try { + std::rethrow_exception(e); + } catch (::sycl::exception& e) { + std::cout << "ASYNC EXCEPTION!!\n"; + std::cout << e.what() << "\n"; + } + } +}; + +TEST(clusterization, sycl) { + + // Memory resource used by the EDM. + vecmem::sycl::shared_memory_resource shared_mr; + traccc::memory_resource mr{shared_mr}; + + // Creating SYCL queue object + ::sycl::queue q(handle_async_error); + std::cout << "Running Seeding on device: " + << q.get_device().get_info<::sycl::info::device::name>() << "\n"; + + // Copy object + vecmem::sycl::copy copy{&q}; + + // Create cell collection + traccc::cell_collection_types::host cells{&shared_mr}; + + cells.push_back({1u, 2u, 1.f, 0, 0}); + cells.push_back({2u, 2u, 1.f, 0, 0}); + cells.push_back({3u, 2u, 1.f, 0, 0}); + + cells.push_back({5u, 5u, 1.f, 0, 0}); + cells.push_back({6u, 4u, 1.f, 0, 0}); + cells.push_back({6u, 5u, 1.f, 0, 0}); + cells.push_back({6u, 6u, 1.f, 0, 0}); + cells.push_back({7u, 5u, 1.f, 0, 0}); + + // Create module collection + traccc::cell_module_collection_types::host modules{&shared_mr}; + modules.push_back({}); + + // Run Clusterization + traccc::sycl::experimental::clusterization_algorithm ca_sycl(mr, copy, &q, + 1024); + + auto measurements_buffer = + ca_sycl(vecmem::get_data(cells), vecmem::get_data(modules)); + + measurement_collection_types::device measurements(measurements_buffer); + + // Check the results + EXPECT_EQ(copy.get_size(measurements_buffer), 2u); + std::set test; + test.insert(measurements[0]); + test.insert(measurements[1]); + + std::set ref; + ref.insert({{2.f, 2.f}, {0.75, 0.0833333}, detray::geometry::barcode{0u}}); + ref.insert( + {{6.f, 5.f}, {0.483333, 0.483333}, detray::geometry::barcode{0u}}); + + EXPECT_EQ(test, ref); +} \ No newline at end of file