Analysis Software
Documentation for sPHENIX simulation software
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
TripletFilter.hpp
Go to the documentation of this file. Or view the newest version in sPHENIX GitHub for file TripletFilter.hpp
1 // This file is part of the Acts project.
2 //
3 // Copyright (C) 2020-2021 CERN for the benefit of the Acts project
4 //
5 // This Source Code Form is subject to the terms of the Mozilla Public
6 // License, v. 2.0. If a copy of the MPL was not distributed with this
7 // file, You can obtain one at http://mozilla.org/MPL/2.0/.
8 
9 #pragma once
10 
11 // Local include(s).
12 #include "Acts/Plugins/Sycl/Seeding/detail/Types.hpp"
13 
14 #include "../Utilities/Arrays.hpp"
15 #include "SpacePointType.hpp"
16 
17 // VecMem include(s).
18 #include "vecmem/containers/data/jagged_vector_view.hpp"
19 #include "vecmem/containers/data/vector_view.hpp"
20 #include "vecmem/containers/device_vector.hpp"
21 #include "vecmem/containers/jagged_device_vector.hpp"
22 
23 // SYCL include(s).
24 #include <CL/sycl.hpp>
25 
26 // System include(s).
27 #include <cstdint>
28 
29 namespace Acts::Sycl::detail {
30 
32 class TripletFilter {
33  public:
35  TripletFilter(
36  const uint32_t numTripletFilterThreads,
37  vecmem::data::vector_view<uint32_t> sumBotMidView,
38  const uint32_t firstMiddle,
39  vecmem::data::vector_view<uint32_t> indMidBotCompView,
40  vecmem::data::vector_view<uint32_t> indBotDupletView,
41  vecmem::data::vector_view<uint32_t> sumBotTopCombView,
42  vecmem::data::jagged_vector_view<uint32_t> midTopDupletView,
43  vecmem::data::vector_view<detail::DeviceTriplet> curvImpactView,
44  vecmem::data::vector_view<const detail::DeviceSpacePoint> topSPsView,
45  vecmem::data::vector_view<const detail::DeviceSpacePoint> middleSPsView,
46  vecmem::data::vector_view<const detail::DeviceSpacePoint> bottomSPsView,
47  vecmem::data::vector_view<uint32_t> countTripletsView,
48  vecmem::data::vector_view<detail::SeedData> seedArrayView,
49  const DeviceSeedFinderConfig& config, const DeviceExperimentCuts& cuts)
50  : m_numTripletFilterThreads(numTripletFilterThreads),
51  m_sumBotMidView(sumBotMidView),
52  m_firstMiddle(firstMiddle),
53  m_indMidBotCompView(indMidBotCompView),
54  m_indBotDupletView(indBotDupletView),
55  m_sumBotTopCombView(sumBotTopCombView),
56  m_midTopDupletView(midTopDupletView),
57  m_curvImpactView(curvImpactView),
58  m_topSPsView(topSPsView),
59  m_middleSPsView(middleSPsView),
60  m_bottomSPsView(bottomSPsView),
61  m_countTripletsView(countTripletsView),
62  m_seedArrayView(seedArrayView),
63  m_config(config),
64  m_cuts(cuts) {}
65 
67  void operator()(cl::sycl::nd_item<1> item) const {
68  if (item.get_global_linear_id() < m_numTripletFilterThreads) {
69  vecmem::device_vector<uint32_t> sumBotMidPrefix(m_sumBotMidView),
70  deviceIndMidBot(m_indMidBotCompView),
71  deviceIndBotDuplets(m_indBotDupletView),
72  sumBotTopCombPrefix(m_sumBotTopCombView),
73  deviceCountTriplets(m_countTripletsView);
74  vecmem::jagged_device_vector<uint32_t> midTopDuplets(m_midTopDupletView);
75  const auto idx =
76  sumBotMidPrefix[m_firstMiddle] + item.get_global_linear_id();
77  const auto mid = deviceIndMidBot[idx];
78  const auto bot = deviceIndBotDuplets[idx];
79  const auto sumCombUptoFirstMiddle = sumBotTopCombPrefix[m_firstMiddle];
80  const auto tripletBegin =
81  sumBotTopCombPrefix[mid] - sumCombUptoFirstMiddle +
82  (idx - sumBotMidPrefix[mid]) * midTopDuplets.at(mid).size();
83  const auto tripletEnd = tripletBegin + deviceCountTriplets[idx];
84  const vecmem::device_vector<detail::DeviceTriplet> deviceCurvImpactConst(
85  m_curvImpactView);
86  for (auto i1 = tripletBegin; i1 < tripletEnd; ++i1) {
87  const auto current = deviceCurvImpactConst[i1];
88  const auto top = current.topSPIndex;
89 
90  const auto invHelixDiameter = current.curvature;
91  const auto lowerLimitCurv =
92  invHelixDiameter - m_config.deltaInvHelixDiameter;
93  const auto upperLimitCurv =
94  invHelixDiameter + m_config.deltaInvHelixDiameter;
95  const vecmem::device_vector<const detail::DeviceSpacePoint> topSPs(
96  m_topSPsView);
97  const auto currentTop_r = topSPs[top].r;
98  auto weight = -(current.impact * m_config.impactWeightFactor);
99 
100  uint32_t compatCounter = 0;
101  // By default compatSeedLimit is 2 -> 2 is
102  // currently hard coded, because variable length
103  // arrays are not supported in SYCL kernels.
104  float compatibleSeedR[2];
105  for (auto i2 = tripletBegin;
106  i2 < tripletEnd && compatCounter < m_config.compatSeedLimit;
107  ++i2) {
108  const auto other = deviceCurvImpactConst[i2];
109 
110  const auto otherCurv = other.curvature;
111  const auto otherTop_r = topSPs[other.topSPIndex].r;
112  const float deltaR = cl::sycl::abs(currentTop_r - otherTop_r);
113  if (deltaR >= m_config.filterDeltaRMin &&
114  otherCurv >= lowerLimitCurv && otherCurv <= upperLimitCurv) {
115  uint32_t c = 0;
116  for (; c < compatCounter &&
117  cl::sycl::abs(compatibleSeedR[c] - otherTop_r) >=
118  m_config.filterDeltaRMin;
119  ++c) {
120  }
121  if (c == compatCounter) {
122  compatibleSeedR[c] = otherTop_r;
123  ++compatCounter;
124  }
125  }
126  }
127  weight += compatCounter * m_config.compatSeedWeight;
128  const vecmem::device_vector<const detail::DeviceSpacePoint> middleSPs(
129  m_middleSPsView),
130  bottomSPs(m_bottomSPsView);
131  const auto bottomSP = bottomSPs[bot];
132  const auto middleSP = middleSPs[mid];
133  const auto topSP = topSPs[top];
134 
135  weight += m_cuts.seedWeight(bottomSP, middleSP, topSP);
136 
137  if (m_cuts.singleSeedCut(weight, bottomSP, middleSP, topSP)) {
138  detail::SeedData D;
139  D.bottom = bot;
140  D.top = top;
141  D.middle = mid;
142  D.weight = weight;
143  vecmem::device_vector<detail::SeedData> seedArray(m_seedArrayView);
144  seedArray.push_back(D);
145  }
146  }
147  }
148  }
149 
150  private:
151  const uint32_t m_numTripletFilterThreads;
152  vecmem::data::vector_view<uint32_t> m_sumBotMidView;
153  const uint32_t m_firstMiddle;
154  vecmem::data::vector_view<uint32_t> m_indMidBotCompView;
155  vecmem::data::vector_view<uint32_t> m_indBotDupletView;
156  vecmem::data::vector_view<uint32_t> m_sumBotTopCombView;
157  vecmem::data::jagged_vector_view<uint32_t> m_midTopDupletView;
158  vecmem::data::vector_view<detail::DeviceTriplet> m_curvImpactView;
159  vecmem::data::vector_view<const detail::DeviceSpacePoint> m_topSPsView;
160  vecmem::data::vector_view<const detail::DeviceSpacePoint> m_middleSPsView;
161  vecmem::data::vector_view<const detail::DeviceSpacePoint> m_bottomSPsView;
162  vecmem::data::vector_view<uint32_t> m_countTripletsView;
163  vecmem::data::vector_view<detail::SeedData> m_seedArrayView;
164  DeviceSeedFinderConfig m_config;
165  DeviceExperimentCuts m_cuts;
166 }; // struct TripletFilter
167 } // namespace Acts::Sycl::detail