Merge 91022e7 into sapling-pr-archive-ktf

Author: ktf

Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackInfoExt.h

@@ -35,6 +35,7 @@ struct TrackInfoExt {
   float ttime = 0;
   float ttimeE = 0;
   float xmin = 0;
+  float chi2TPC = 0.f;
   float chi2ITSTPC = 0.f;
   float q2ptITS = 0.f;
   float q2ptTPC = 0.f;
@@ -56,7 +57,7 @@ struct TrackInfoExt {
   float getTPCInY0() const { return innerTPCPos0[1]; }
   float getTPCInZ0() const { return innerTPCPos0[2]; }
 
-  ClassDefNV(TrackInfoExt, 5);
+  ClassDefNV(TrackInfoExt, 6);
 };
 
 } // namespace dataformats

Detectors/GlobalTrackingWorkflow/study/src/TrackingStudy.cxx

@@ -503,6 +503,7 @@ void TrackingStudySpec::process(o2::globaltracking::RecoContainer& recoData)
               }
             }
             fillTPCClInfo(*tpcTr, trcExt, tsuse);
+            trcExt.chi2TPC = tpcTr->getChi2();
           }
           auto gidRefs = recoData.getSingleDetectorRefs(vid);
           if (gidRefs[GTrackID::ITS].isIndexSet()) {

Framework/Core/include/Framework/TMessageSerializer.h

@@ -40,6 +40,8 @@ class FairOutputTBuffer : public TBufferFile
     : TBufferFile(TBuffer::kWrite, msg.GetSize() - sizeof(char*), embedInItself(msg), false, fairMQrealloc)
   {
   }
+
+  ~FairOutputTBuffer() override;
   // Helper function to keep track of the FairMQ message that holds the data
   // in the data itself. We can use this to make sure the message can be reallocated
   // even if we simply have a pointer to the data. Hopefully ROOT will not play dirty
@@ -60,6 +62,7 @@ class FairInputTBuffer : public TBufferFile
     : TBufferFile(TBuffer::kRead, size - sizeof(char*), data + sizeof(char*), false, nullptr)
   {
   }
+  ~FairInputTBuffer() override;
 };
 
 struct TMessageSerializer {

Framework/Core/src/TMessageSerializer.cxx

@@ -15,6 +15,9 @@
 
 using namespace o2::framework;
 
+FairOutputTBuffer::~FairOutputTBuffer() = default;
+FairInputTBuffer::~FairInputTBuffer() = default;
+
 void* FairOutputTBuffer::embedInItself(fair::mq::Message& msg)
 {
   // The first bytes of the message are used to store the pointer to the message itself

GPU/Common/GPUCommonAlgorithm.h

@@ -282,30 +282,27 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp)
 {
 #ifndef GPUCA_GPUCODE
   GPUCommonAlgorithm::sort(begin, end, comp);
+#elif defined(GPUCA_DETERMINISTIC_MODE) // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression
+  if (get_local_id(0) == 0) {
+    GPUCommonAlgorithm::sort(begin, end, comp);
+  }
+  GPUbarrier();
 #else
-  GPUCA_DETERMINISTIC_CODE( // clang-format off
-    GPUbarrier();
-    if (get_local_id(0) == 0) {
-      GPUCommonAlgorithm::sort(begin, end, comp);
-    }
-    GPUbarrier();
-  , // !GPUCA_DETERMINISTIC_CODE
-    int32_t n = end - begin;
-    for (int32_t i = 0; i < n; i++) {
-      for (int32_t tIdx = get_local_id(0); tIdx < n; tIdx += get_local_size(0)) {
-        int32_t offset = i % 2;
-        int32_t curPos = 2 * tIdx + offset;
-        int32_t nextPos = curPos + 1;
-
-        if (nextPos < n) {
-          if (!comp(begin[curPos], begin[nextPos])) {
-            IterSwap(&begin[curPos], &begin[nextPos]);
-          }
+  int32_t n = end - begin;
+  for (int32_t i = 0; i < n; i++) {
+    for (int32_t tIdx = get_local_id(0); tIdx < n; tIdx += get_local_size(0)) {
+      int32_t offset = i % 2;
+      int32_t curPos = 2 * tIdx + offset;
+      int32_t nextPos = curPos + 1;
+
+      if (nextPos < n) {
+        if (!comp(begin[curPos], begin[nextPos])) {
+          IterSwap(&begin[curPos], &begin[nextPos]);
         }
       }
-      GPUbarrier();
     }
-  ) // clang-format on
+    GPUbarrier();
+  }
 #endif
 }
 

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDArtc.cu

@@ -14,9 +14,12 @@
 
 #define GPUCA_GPUCODE_GENRTC
 #define GPUCA_GPUCODE_COMPILEKERNELS
+
+// Keep some preprocessor calls unprocessed
 #define GPUCA_RTC_SPECIAL_CODE(...) GPUCA_RTC_SPECIAL_CODE(__VA_ARGS__)
 #define GPUCA_DETERMINISTIC_CODE(...) GPUCA_DETERMINISTIC_CODE(__VA_ARGS__)
-// GPUReconstructionCUDAIncludesSystem.h prependended without preprocessor running
+
+// GPUReconstructionCUDAIncludesSystem.h prependended by CMakewithout preprocessor running
 #include "GPUReconstructionCUDADef.h"
 #include "GPUReconstructionIncludesDeviceAll.h"
 

GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx

@@ -273,19 +273,19 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
 #ifdef GPUCA_GPUCODE
         static_assert(GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCompressionKernels_step1unattached) * 2 <= GPUCA_TPC_COMP_CHUNK_SIZE);
 #endif
-        GPUCA_DETERMINISTIC_CODE( // clang-format off
+#ifdef GPUCA_DETERMINISTIC_MODE // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression
+        CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
+#else  // GPUCA_DETERMINISTIC_MODE
+        if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZPadTime) {
           CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
-        , // !GPUCA_DETERMINISTIC_CODE
-          if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZPadTime) {
-            CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZPadTime>(clusters->clusters[iSector][iRow]));
-          } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZTimePad) {
-            CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(clusters->clusters[iSector][iRow]));
-          } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortPad) {
-            CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(clusters->clusters[iSector][iRow]));
-          } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortTime) {
-            CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(clusters->clusters[iSector][iRow]));
-          }
-        ) // clang-format on
+        } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortZTimePad) {
+          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortZTimePad>(clusters->clusters[iSector][iRow]));
+        } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortPad) {
+          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortPad>(clusters->clusters[iSector][iRow]));
+        } else if (param.rec.tpc.compressionSortOrder == GPUSettings::SortTime) {
+          CAAlgo::sortInBlock(sortBuffer, sortBuffer + count, GPUTPCCompressionKernels_Compare<GPUSettings::SortTime>(clusters->clusters[iSector][iRow]));
+        }
+#endif // GPUCA_DETERMINISTIC_MODE
         GPUbarrier();
       }