ITS: fix rare valid lookup at the edge of acceptance

Technically, bin (0,0,0,0) is a valid result for the
phi-z cluster query.
Note though, that this if at all a super rare case and
in local tests this did only showed up in the number of
tracklets.
So more a consistency fix than anything.

Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>

Author: f3sch

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu

@@ -408,7 +408,7 @@ GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel(
         const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + constants::Tolerance)}; /// protecting from overflows adding the detector resolution
         const float sigmaZ{o2::gpu::CAMath::Sqrt(math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInverseDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * MSAngle))};
         const int4 selectedBinsRect{o2::its::getBinsRect(currentCluster, layerIndex + 1, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut, *utils)};
-        if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) {
+        if (selectedBinsRect.x < 0) {
           continue;
         }
         int phiBinsNum{selectedBinsRect.w - selectedBinsRect.y + 1};

Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h

@@ -124,7 +124,7 @@ GPUhdi() int4 getBinsRect(const Cluster& currentCluster, const int layerIndex,
 
   if (zRangeMax < -utils.getLayerZ(layerIndex) ||
       zRangeMin > utils.getLayerZ(layerIndex) || zRangeMin > zRangeMax) {
-    return int4{0, 0, 0, 0};
+    return int4{-1, -1, -1, -1};
   }
 
   return int4{o2::gpu::GPUCommonMath::Max(0, utils.getZBinIndex(layerIndex, zRangeMin)),

Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h

@@ -53,12 +53,6 @@ class VertexerTraits
   VertexerTraits() = default;
   virtual ~VertexerTraits() = default;
 
-  GPUhdi() static consteval int4 getEmptyBinsRect()
-  {
-    return int4{0, 0, 0, 0};
-  }
-  GPUhd() const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi);
-  GPUhd() static const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi, const IndexTableUtilsN&);
   GPUhd() static const int2 getPhiBins(float phi, float deltaPhi, const IndexTableUtilsN&);
   GPUhd() const int2 getPhiBins(float phi, float deltaPhi) { return getPhiBins(phi, deltaPhi, mIndexTableUtils); }
 
@@ -134,34 +128,6 @@ GPUhdi() const int2 VertexerTraits<NLayers>::getPhiBins(float phi, float dPhi, c
               utils.getPhiBinIndex(math_utils::getNormalizedPhi(phi + dPhi))};
 }
 
-template <int NLayers>
-GPUhdi() const int4 VertexerTraits<NLayers>::getBinsRect(const Cluster& currentCluster, const int layerIndex,
-                                                         const float directionZIntersection, float maxdeltaz, float maxdeltaphi,
-                                                         const IndexTableUtilsN& utils)
-{
-  const float zRangeMin = directionZIntersection - 2 * maxdeltaz;
-  const float phiRangeMin = currentCluster.phi - maxdeltaphi;
-  const float zRangeMax = directionZIntersection + 2 * maxdeltaz;
-  const float phiRangeMax = currentCluster.phi + maxdeltaphi;
-
-  if (zRangeMax < -utils.getLayerZ(layerIndex + 1) ||
-      zRangeMin > utils.getLayerZ(layerIndex + 1) || zRangeMin > zRangeMax) {
-    return getEmptyBinsRect();
-  }
-
-  return int4{o2::gpu::GPUCommonMath::Max(0, utils.getZBinIndex(layerIndex + 1, zRangeMin)),
-              utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMin)),
-              o2::gpu::GPUCommonMath::Min(utils.getNzBins() - 1, utils.getZBinIndex(layerIndex + 1, zRangeMax)),
-              utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))};
-}
-
-template <int NLayers>
-GPUhdi() const int4 VertexerTraits<NLayers>::getBinsRect(const Cluster& currentCluster, const int layerIndex,
-                                                         const float directionZIntersection, float maxdeltaz, float maxdeltaphi)
-{
-  return VertexerTraits::getBinsRect(currentCluster, layerIndex, directionZIntersection, maxdeltaz, maxdeltaphi, mIndexTableUtils);
-}
-
 } // namespace its
 } // namespace o2
 #endif

Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx

@@ -114,7 +114,7 @@ void TrackerTraits<NLayers>::computeLayerTracklets(const int iteration, int iVer
           const auto bins = o2::its::getBinsRect(currentCluster, iLayer + 1, zAtRmin, zAtRmax,
                                                  sigmaZ * mTrkParams[iteration].NSigmaCut, mTimeFrame->getPhiCut(iLayer),
                                                  mTimeFrame->getIndexTableUtils());
-          if (bins.x == 0 && bins.y == 0 && bins.z == 0 && bins.w == 0) {
+          if (bins.x < 0) {
             continue;
           }
           int phiBinsNum = bins.w - bins.y + 1;

Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx

@@ -58,8 +58,8 @@ void trackleterKernelHost(
   for (int iCurrentLayerClusterIndex = 0; iCurrentLayerClusterIndex < clustersCurrentLayer.size(); ++iCurrentLayerClusterIndex) {
     int storedTracklets{0};
     const Cluster& currentCluster{clustersCurrentLayer[iCurrentLayerClusterIndex]};
-    const int4 selectedBinsRect{VertexerTraits<NLayers>::getBinsRect(currentCluster, (int)Mode, 0.f, 50.f, phiCut / 2, utils)};
-    if (selectedBinsRect.x != 0 || selectedBinsRect.y != 0 || selectedBinsRect.z != 0 || selectedBinsRect.w != 0) {
+    const int4 selectedBinsRect{o2::its::getBinsRect(currentCluster, (int)Mode + 1, 0.f, 0.f, 100.f, phiCut / 2, utils)};
+    if (selectedBinsRect.x >= 0) {
       int phiBinsNum{selectedBinsRect.w - selectedBinsRect.y + 1};
       if (phiBinsNum < 0) {
         phiBinsNum += PhiBins;