GPU: Extract unattached cluster compression to include file

davidrohr · davidrohr · commit 6385a0ef8d52 · 2022-11-06T23:25:37.000+01:00
diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt
@@ -159,6 +159,7 @@ set(HDRS_INSTALL
     DataTypes/GPUTRDInterfaceO2Track.h
     Base/GPUParam.inc
     DataCompression/TPCClusterDecompressor.inc
+    DataCompression/GPUTPCCompressionKernels.inc
     Merger/GPUTPCGMMergerTypes.h
     Global/GPUErrorCodes.h
     Global/GPUChainTrackingDefs.h
diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx
@@ -20,6 +20,7 @@
 #include "GPUTPCCompressionTrackModel.h"
 #include "GPUTPCGeometry.h"
 #include "GPUTPCClusterRejection.h"
+#include "GPUTPCCompressionKernels.inc"
 
 using namespace GPUCA_NAMESPACE::gpu;
 using namespace o2::tpc;
@@ -272,25 +273,9 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
       for (unsigned int j = get_local_id(0); j < count; j += get_local_size(0)) {
         int outidx = idOffsetOut + totalCount + j;
         const ClusterNative& GPUrestrict() orgCl = clusters->clusters[iSlice][iRow][sortBuffer[j]];
-        unsigned int lastTime = 0;
-        unsigned int lastPad = 0;
-        if (param.rec.tpc.compressionTypeMask & GPUSettings::CompressionDifferences) {
-          if (j != 0) {
-            const ClusterNative& GPUrestrict() orgClPre = clusters->clusters[iSlice][iRow][sortBuffer[j - 1]];
-            lastPad = orgClPre.padPacked;
-            lastTime = orgClPre.getTimePacked();
-          } else if (totalCount != 0) {
-            const ClusterNative& GPUrestrict() orgClPre = clusters->clusters[iSlice][iRow][smem.lastIndex];
-            lastPad = orgClPre.padPacked;
-            lastTime = orgClPre.getTimePacked();
-          }
 
-          c.padDiffU[outidx] = orgCl.padPacked - lastPad;
-          c.timeDiffU[outidx] = (orgCl.getTimePacked() - lastTime) & 0xFFFFFF;
-        } else {
-          c.padDiffU[outidx] = orgCl.padPacked;
-          c.timeDiffU[outidx] = orgCl.getTimePacked();
-        }
+        int preId = j != 0 ? (int)sortBuffer[j - 1] : (totalCount != 0 ? (int)smem.lastIndex : -1);
+        GPUTPCCompression_EncodeUnattached(param, orgCl, c, outidx, preId == -1 ? nullptr : &clusters->clusters[iSlice][iRow][preId]);
 
         unsigned short qtot = orgCl.qTot, qmax = orgCl.qMax;
         unsigned char sigmapad = orgCl.sigmaPadPacked, sigmatime = orgCl.sigmaTimePacked;
diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.inc b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.inc
@@ -0,0 +1,41 @@
+// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
+// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
+// All rights not expressly granted are reserved.
+//
+// This software is distributed under the terms of the GNU General Public
+// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
+//
+// In applying this license CERN does not waive the privileges and immunities
+// granted to it by virtue of its status as an Intergovernmental Organization
+// or submit itself to any jurisdiction.
+
+/// \file GPUTPCCompressionKernels.cxx
+/// \author David Rohr
+
+#include "GPUO2DataTypes.h"
+#include "GPUParam.h"
+#include "GPUTPCGeometry.h"
+#include "GPUTPCClusterRejection.h"
+
+namespace o2::gpu
+{
+
+template <class T>
+GPUdii() void GPUTPCCompression_EncodeUnattached(const GPUParam& param, const o2::tpc::ClusterNative& orgCl, T& c, int outidx, const o2::tpc::ClusterNative* orgClPre = nullptr)
+{
+  if (param.rec.tpc.compressionTypeMask & GPUSettings::CompressionDifferences) {
+    unsigned int lastTime = 0, lastPad = 0;
+    if (orgClPre) {
+      lastPad = orgClPre->padPacked;
+      lastTime = orgClPre->getTimePacked();
+    }
+
+    c.padDiffU[outidx] = orgCl.padPacked - lastPad;
+    c.timeDiffU[outidx] = (orgCl.getTimePacked() - lastTime) & 0xFFFFFF;
+  } else {
+    c.padDiffU[outidx] = orgCl.padPacked;
+    c.timeDiffU[outidx] = orgCl.getTimePacked();
+  }
+}
+
+} // namespace o2::gpu
