Fix for TPC edge clusters in CTF decoding

Author: shahor02

GPU/GPUTracking/DataCompression/GPUTPCDecompressionKernels.cxx

@@ -44,6 +44,7 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste
   CompressedClusters& GPUrestrict() cmprClusters = decompressor.mInputGPU;
   ClusterNative* GPUrestrict() clusterBuffer = decompressor.mNativeClustersBuffer;
   const ClusterNativeAccess* outputAccess = decompressor.mClusterNativeAccess;
+  constexpr GPUTPCGeometry geo;
   uint32_t* offsets = decompressor.mUnattachedClustersOffsets;
   for (int32_t i = get_global_id(0); i < GPUCA_ROW_COUNT * nSectors; i += get_global_size(0)) {
     uint32_t iRow = i % GPUCA_ROW_COUNT;
@@ -57,6 +58,18 @@ GPUdii() void GPUTPCDecompressionKernels::Thread<GPUTPCDecompressionKernels::ste
     ClusterNative* clout = buffer + decompressor.mNativeClustersIndex[linearIndex];
     uint32_t end = offsets[linearIndex] + ((linearIndex >= decompressor.mInputGPU.nSliceRows) ? 0 : decompressor.mInputGPU.nSliceRowClusters[linearIndex]);
     TPCClusterDecompressionCore::decompressHits(cmprClusters, offsets[linearIndex], end, clout);
+    if (processors.param.rec.tpc.clustersEdgeFixDistance > 0.f) {
+      for (uint32_t k = 0; k < outputAccess->nClusters[iSector][iRow]; k++) {
+        auto& cluster = buffer[k];
+        if (cluster.getFlags() & ClusterNative::flagEdge) {
+          auto padF = cluster.getPad();
+          float distEdge = padF < geo.NPads(iRow) / 2 ? padF : geo.NPads(iRow) - 1 - padF;
+          if (distEdge > processors.param.rec.tpc.clustersEdgeFixDistance) {
+            cluster.setFlags(cluster.getFlags() ^ ClusterNative::flagEdge);
+          }
+        }
+      }
+    }
     if (processors.param.rec.tpc.clustersShiftTimebins != 0.f) {
       for (uint32_t k = 0; k < outputAccess->nClusters[iSector][iRow]; k++) {
         auto& cl = buffer[k];

GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc

@@ -137,7 +137,14 @@ class TPCClusterDecompressionCore
         time = cmprClusters.timeA[trackIndex];
         pad = cmprClusters.padA[trackIndex];
       }
-      const auto cluster = decompressTrackStore(cmprClusters, clusterOffset, slice, row, pad, time, args...);
+      auto cluster = decompressTrackStore(cmprClusters, clusterOffset, slice, row, pad, time, args...);
+      if ((cluster.getFlags() & ClusterNative::flagEdge) && param.rec.tpc.clustersShiftTimebins != 0.f) {
+        auto padF = cluster.getPad();
+        float distEdge = padF < geo.NPads(row) / 2 ? padF : geo.NPads(row) - 1 - padF;
+        if (distEdge > param.rec.tpc.clustersShiftTimebins) {
+          cluster.setFlags(cluster.getFlags() ^ ClusterNative::flagEdge);
+        }
+      }
       float y = track.LinearPad2Y(slice, cluster.getPad(), geo.PadWidth(row), geo.NPads(row));
       float z = geo.LinearTime2Z(slice, cluster.getTime());
       if (clusterIndex == 0) {
@@ -152,7 +159,7 @@ class TPCClusterDecompressionCore
     clusterOffset += cmprClusters.nTrackClusters[trackIndex] - clusterIndex;
   }
 
-  GPUdi() static const auto& decompressHitsStore(const CompressedClusters& cmprClusters, uint32_t k, uint32_t time, uint16_t pad, ClusterNative*& clusterBuffer)
+  GPUdi() static auto& decompressHitsStore(const CompressedClusters& cmprClusters, uint32_t k, uint32_t time, uint16_t pad, ClusterNative*& clusterBuffer)
   {
     return ((*(clusterBuffer++) = ClusterNative(time, cmprClusters.flagsU[k], pad, cmprClusters.sigmaTimeU[k], cmprClusters.sigmaPadU[k], cmprClusters.qMaxU[k], cmprClusters.qTotU[k])));
   }
@@ -183,7 +190,6 @@ class TPCClusterDecompressionCore
         time = cmprClusters.timeDiffU[k];
         pad = cmprClusters.padDiffU[k];
       }
-      decompressHitsStore(cmprClusters, k, time, pad, args...);
     }
   }
 };

GPU/GPUTracking/DataCompression/TPCClusterDecompressor.cxx

@@ -85,6 +85,7 @@ int32_t TPCClusterDecompressor::decompress(const CompressedClusters* clustersCom
   }
   clustersNative.clustersLinear = clusterBuffer;
   clustersNative.setOffsetPtrs();
+  constexpr GPUTPCGeometry geo;
   tbb::parallel_for<uint32_t>(0, NSECTORS, [&](auto i) {
     for (uint32_t j = 0; j < GPUCA_ROW_COUNT; j++) {
       ClusterNative* buffer = &clusterBuffer[clustersNative.clusterOffset[i][j]];
@@ -94,6 +95,18 @@ int32_t TPCClusterDecompressor::decompress(const CompressedClusters* clustersCom
       ClusterNative* clout = buffer + clusters[i][j].size();
       uint32_t end = offsets[i][j] + ((i * GPUCA_ROW_COUNT + j >= clustersCompressed->nSliceRows) ? 0 : clustersCompressed->nSliceRowClusters[i * GPUCA_ROW_COUNT + j]);
       TPCClusterDecompressionCore::decompressHits(*clustersCompressed, offsets[i][j], end, clout);
+      if (param.rec.tpc.clustersEdgeFixDistance > 0.f) {
+        for (uint32_t k = 0; k < clustersNative.nClusters[i][j]; k++) {
+          auto& cluster = buffer[k];
+          if (cluster.getFlags() & ClusterNative::flagEdge) {
+            auto padF = cluster.getPad();
+            float distEdge = padF < geo.NPads(j) / 2 ? padF : geo.NPads(j) - 1 - padF;
+            if (distEdge > param.rec.tpc.clustersEdgeFixDistance) {
+              cluster.setFlags(cluster.getFlags() ^ ClusterNative::flagEdge);
+            }
+          }
+        }
+      }
       if (param.rec.tpc.clustersShiftTimebins != 0.f) {
         for (uint32_t k = 0; k < clustersNative.nClusters[i][j]; k++) {
           auto& cl = buffer[k];

GPU/GPUTracking/Definitions/GPUSettingsList.h

@@ -72,6 +72,7 @@ AddOptionRTC(tubeChi2, float, 5.f * 5.f, "", 0, "Max chi2 to mark cluster adjace
 AddOptionRTC(tubeMaxSize2, float, 2.5f * 2.5f, "", 0, "Square of max tube size (normally derrived from tpcTubeChi2)")
 AddOptionRTC(clustersShiftTimebins, float, 0, "", 0, "Shift of TPC clusters (applied during CTF cluster decoding)")
 AddOptionRTC(clustersShiftTimebinsClusterizer, float, 0, "", 0, "Shift of TPC clusters (applied during CTF clusterization)")
+AddOptionRTC(clustersEdgeFixDistance, float, 0.f, "", 0, "If >0, revert cluster.flag edge bit distance to edge exceeds this parameter (fixed during CTF decoding)")
 AddOptionRTC(defaultZOffsetOverR, float, 0.5210953f, "", 0, "Shift of TPC clusters (applied during CTF cluster decoding)")
 AddOptionRTC(PID_EKrangeMin, float, 0.47f, "", 0, "min P of electron/K BB bands crossing")
 AddOptionRTC(PID_EKrangeMax, float, 0.57f, "", 0, "max P of electron/K BB bands crossing")