src/Seeding/DupletSearch.hpp

0001 // This file is part of the Acts project.
0002 //
0003 // Copyright (C) 2020-2021 CERN for the benefit of the Acts project
0004 //
0005 // This Source Code Form is subject to the terms of the Mozilla Public
0006 // License, v. 2.0. If a copy of the MPL was not distributed with this
0007 // file, You can obtain one at http://mozilla.org/MPL/2.0/.
0008
0009 #pragma once
0010
0011 // Local include(s).
0012 #include "Acts/Plugins/Sycl/Seeding/detail/Types.hpp"
0013
0014 #include "../Utilities/Arrays.hpp"
0015 #include "SpacePointType.hpp"
0016
0017 // VecMem include(s).
0018 #include "vecmem/containers/data/jagged_vector_view.hpp"
0019 #include "vecmem/containers/data/vector_view.hpp"
0020 #include "vecmem/containers/device_vector.hpp"
0021 #include "vecmem/containers/jagged_device_vector.hpp"
0022
0023 // SYCL include(s).
0024 #include <CL/sycl.hpp>
0025
0026 // System include(s).
0027 #include <cstdint>
0028
0029 namespace Acts::Sycl::detail {
0030
0031 /// Functor taking care of finding viable spacepoint duplets
0032 template <SpacePointType OtherSPType>
0033 class DupletSearch {
0034   // Sanity check(s).
0035   static_assert((OtherSPType == SpacePointType::Bottom) ||
0036                     (OtherSPType == SpacePointType::Top),
0037                 "Class must be instantiated with either "
0038                 "Acts::Sycl::detail::SpacePointType::Bottom or "
0039                 "Acts::Sycl::detail::SpacePointType::Top");
0040
0041  public:
0042   /// Constructor with all the necessary arguments
0043   DupletSearch(vecmem::data::vector_view<const DeviceSpacePoint> middleSPs,
0044                vecmem::data::vector_view<const DeviceSpacePoint> otherSPs,
0045                vecmem::data::jagged_vector_view<uint32_t> middleOtherSPIndices,
0046                const DeviceSeedFinderConfig& config)
0047       : m_middleSPs(middleSPs),
0048         m_otherSPs(otherSPs),
0049         m_middleOtherSPIndices(middleOtherSPIndices),
0050         m_config(config) {}
0051
0052   /// Operator performing the duplet search
0053   void operator()(cl::sycl::nd_item<2> item) const {
0054     // Get the indices of the spacepoints to evaluate.
0055     const auto middleIndex = item.get_global_id(0);
0056     const auto otherIndex = item.get_global_id(1);
0057
0058     // We check whether this thread actually makes sense (within bounds).
0059     // The number of threads is usually a factor of 2, or 3*2^k (k \in N), etc.
0060     // Without this check we may index out of arrays.
0061     if ((middleIndex >= m_middleSPs.size()) ||
0062         (otherIndex >= m_otherSPs.size())) {
0063       return;
0064     }
0065
0066     // Create a copy of the spacepoint objects for the current thread. On
0067     // dedicated GPUs this provides a better performance than accessing
0068     // variables one-by-one from global device memory.
0069     const vecmem::device_vector<const DeviceSpacePoint> middleSPs(m_middleSPs);
0070     const DeviceSpacePoint middleSP = middleSPs[middleIndex];
0071     const vecmem::device_vector<const DeviceSpacePoint> otherSPs(m_otherSPs);
0072     const DeviceSpacePoint otherSP = otherSPs[otherIndex];
0073
0074     // Calculate the variables that the duplet quality selection are based on.
0075     // Note that the asserts of the functor make sure that 'OtherSPType' must be
0076     // either SpacePointType::Bottom or SpacePointType::Top.
0077     float deltaR = 0.0f, cotTheta = 0.0f;
0078     if constexpr (OtherSPType == SpacePointType::Bottom) {
0079       deltaR = middleSP.r - otherSP.r;
0080       cotTheta = (middleSP.z - otherSP.z) / deltaR;
0081     } else {
0082       deltaR = otherSP.r - middleSP.r;
0083       cotTheta = (otherSP.z - middleSP.z) / deltaR;
0084     }
0085     const float zOrigin = middleSP.z - middleSP.r * cotTheta;
0086
0087     // Check if the duplet passes our quality requirements.
0088     if ((deltaR >= m_config.deltaRMin) && (deltaR <= m_config.deltaRMax) &&
0089         (cl::sycl::abs(cotTheta) <= m_config.cotThetaMax) &&
0090         (zOrigin >= m_config.collisionRegionMin) &&
0091         (zOrigin <= m_config.collisionRegionMax)) {
0092       // Create device vector based on the view and push to it
0093       vecmem::jagged_device_vector<uint32_t> middleOtherSPIndices(
0094           m_middleOtherSPIndices);
0095       middleOtherSPIndices.at(middleIndex).push_back(otherIndex);
0096     }
0097   }
0098
0099  private:
0100   /// Middle spacepoints
0101   vecmem::data::vector_view<const DeviceSpacePoint> m_middleSPs;
0102   /// "Other" (bottom or top) spacepoints
0103   vecmem::data::vector_view<const DeviceSpacePoint> m_otherSPs;
0104
0105   /// The 2D array storing the compatible middle-other spacepoint indices
0106   vecmem::data::jagged_vector_view<uint32_t> m_middleOtherSPIndices;
0107
0108   /// Configuration for the seed finding
0109   DeviceSeedFinderConfig m_config;
0110
0111 };  // struct DupletSearch
0112
0113 }  // namespace Acts::Sycl::detail