From f2e7956bbfe9bdd661acbb25cdef722773513ddc Mon Sep 17 00:00:00 2001
From: czhang <chi.zhang@cern.ch>
Date: Mon, 27 Jan 2025 15:27:12 +0000
Subject: [PATCH 1/8] Add termination for record writer to save record file
 correctly

---
 Detectors/MUON/MCH/Align/src/Aligner.cxx | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/Detectors/MUON/MCH/Align/src/Aligner.cxx b/Detectors/MUON/MCH/Align/src/Aligner.cxx
index 71bafad5b9ff3..a19b7f602cdb8 100644
--- a/Detectors/MUON/MCH/Align/src/Aligner.cxx
+++ b/Detectors/MUON/MCH/Align/src/Aligner.cxx
@@ -359,14 +359,17 @@ void Aligner::init(TString DataRecFName, TString ConsRecFName)
 void Aligner::terminate()
 {
   fInitialized = kFALSE;
-  LOG(info) << "Closing Evaluation TFile";
   if (fDoEvaluation) {
+    LOG(info) << "Closing Evaluation TFile";
     if (fTFile && fTTree) {
       fTFile->cd();
       fTTree->Write();
       fTFile->Close();
     }
   }
+  if (!fDisableRecordWriter) {
+    mRecordWriter->terminate();
+  }
 }
 
 //_____________________________________________________

From 8baefd989c16b9d926a2a69efe985909cbaa06d4 Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Tue, 28 Jan 2025 14:26:08 +0100
Subject: [PATCH 2/8] DPL Analysis: out of line HistogramRegistry population
 methods (#13906)

---
 .../include/Framework/HistogramRegistry.h     | 72 ++++++++-----------
 Framework/Core/src/HistogramRegistry.cxx      | 70 ++++++++++++++++++
 2 files changed, 101 insertions(+), 41 deletions(-)

diff --git a/Framework/Core/include/Framework/HistogramRegistry.h b/Framework/Core/include/Framework/HistogramRegistry.h
index 0801064b6f4cc..9f272be38da0c 100644
--- a/Framework/Core/include/Framework/HistogramRegistry.h
+++ b/Framework/Core/include/Framework/HistogramRegistry.h
@@ -22,10 +22,12 @@
 #include "Framework/SerializationMethods.h"
 #include "Framework/TableBuilder.h"
 #include "Framework/RuntimeError.h"
+#include "StepTHn.h"
 
 #include <TDataMember.h>
 #include <TDataType.h>
 #include <TArrayL.h>
+#include <THnSparse.h>
 #include <TProfile2D.h>
 #include <fmt/core.h>
 
@@ -403,28 +405,6 @@ constexpr HistogramRegistry::HistName::HistName(const ConstStr<chars...>& hashed
 {
 }
 
-template <typename T>
-std::shared_ptr<T> HistogramRegistry::add(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2)
-{
-  auto histVariant = add(name, title, histConfigSpec, callSumw2);
-  if (auto histPtr = std::get_if<std::shared_ptr<T>>(&histVariant)) {
-    return *histPtr;
-  } else {
-    throw runtime_error_f(R"(Histogram type specified in add<>("%s") does not match the actual type of the histogram!)", name);
-  }
-}
-
-template <typename T>
-std::shared_ptr<T> HistogramRegistry::add(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2)
-{
-  auto histVariant = add(name, title, histType, axes, callSumw2);
-  if (auto histPtr = std::get_if<std::shared_ptr<T>>(&histVariant)) {
-    return *histPtr;
-  } else {
-    throw runtime_error_f(R"(Histogram type specified in add<>("%s") does not match the actual type of the histogram!)", name);
-  }
-}
-
 template <typename T>
 std::shared_ptr<T> HistogramRegistry::add(const std::string& name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2)
 {
@@ -447,25 +427,6 @@ std::shared_ptr<T> HistogramRegistry::operator()(const HistName& histName)
   return get<T>(histName);
 }
 
-template <typename T>
-HistPtr HistogramRegistry::insertClone(const HistName& histName, const std::shared_ptr<T> originalHist)
-{
-  validateHistName(histName.str, histName.hash);
-  for (auto i = 0u; i < MAX_REGISTRY_SIZE; ++i) {
-    TObject* rawPtr = nullptr;
-    std::visit([&](const auto& sharedPtr) { rawPtr = sharedPtr.get(); }, mRegistryValue[imask(histName.idx + i)]);
-    if (!rawPtr) {
-      registerName(histName.str);
-      mRegistryKey[imask(histName.idx + i)] = histName.hash;
-      mRegistryValue[imask(histName.idx + i)] = std::shared_ptr<T>(static_cast<T*>(originalHist->Clone(histName.str)));
-      lookup += i;
-      return mRegistryValue[imask(histName.idx + i)];
-    }
-  }
-  LOGF(fatal, R"(Internal array of HistogramRegistry "%s" is full.)", mName);
-  return HistPtr();
-}
-
 template <typename T>
 uint32_t HistogramRegistry::getHistIndex(const T& histName)
 {
@@ -491,6 +452,35 @@ extern template void HistogramRegistry::fill(const HistName& histName, double);
 extern template void HistogramRegistry::fill(const HistName& histName, float);
 extern template void HistogramRegistry::fill(const HistName& histName, int);
 
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TH1>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TH2>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TH3>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TProfile>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TProfile2D>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TProfile3D>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<THnSparse>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<THn>);
+extern template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<StepTHn>);
+
+extern template std::shared_ptr<TH1> HistogramRegistry::add<TH1>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<TH1> HistogramRegistry::add<TH1>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<TH2> HistogramRegistry::add<TH2>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<TH2> HistogramRegistry::add<TH2>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<TH3> HistogramRegistry::add<TH3>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<TH3> HistogramRegistry::add<TH3>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<TProfile> HistogramRegistry::add<TProfile>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<TProfile> HistogramRegistry::add<TProfile>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<TProfile2D> HistogramRegistry::add<TProfile2D>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<TProfile2D> HistogramRegistry::add<TProfile2D>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<TProfile3D> HistogramRegistry::add<TProfile3D>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<TProfile3D> HistogramRegistry::add<TProfile3D>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<THn> HistogramRegistry::add<THn>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<THn> HistogramRegistry::add<THn>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<THnSparse> HistogramRegistry::add<THnSparse>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<THnSparse> HistogramRegistry::add<THnSparse>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+extern template std::shared_ptr<StepTHn> HistogramRegistry::add<StepTHn>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+extern template std::shared_ptr<StepTHn> HistogramRegistry::add<StepTHn>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+
 template <typename... Cs, typename T>
 void HistogramRegistry::fill(const HistName& histName, const T& table, const o2::framework::expressions::Filter& filter)
 {
diff --git a/Framework/Core/src/HistogramRegistry.cxx b/Framework/Core/src/HistogramRegistry.cxx
index 0836e72ffa935..0a0cc1fc3a690 100644
--- a/Framework/Core/src/HistogramRegistry.cxx
+++ b/Framework/Core/src/HistogramRegistry.cxx
@@ -414,4 +414,74 @@ void HistFiller::badHistogramFill(char const* name)
   LOGF(fatal, "The number of arguments in fill function called for histogram %s is incompatible with histogram dimensions.", name);
 }
 
+template <typename T>
+HistPtr HistogramRegistry::insertClone(const HistName& histName, const std::shared_ptr<T> originalHist)
+{
+  validateHistName(histName.str, histName.hash);
+  for (auto i = 0u; i < MAX_REGISTRY_SIZE; ++i) {
+    TObject* rawPtr = nullptr;
+    std::visit([&](const auto& sharedPtr) { rawPtr = sharedPtr.get(); }, mRegistryValue[imask(histName.idx + i)]);
+    if (!rawPtr) {
+      registerName(histName.str);
+      mRegistryKey[imask(histName.idx + i)] = histName.hash;
+      mRegistryValue[imask(histName.idx + i)] = std::shared_ptr<T>(static_cast<T*>(originalHist->Clone(histName.str)));
+      lookup += i;
+      return mRegistryValue[imask(histName.idx + i)];
+    }
+  }
+  LOGF(fatal, R"(Internal array of HistogramRegistry "%s" is full.)", mName);
+  return HistPtr();
+}
+
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TH1>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TH2>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TH3>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TProfile>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TProfile2D>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<TProfile3D>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<THnSparse>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<THn>);
+template HistPtr HistogramRegistry::insertClone(const HistName&, const std::shared_ptr<StepTHn>);
+
+template <typename T>
+std::shared_ptr<T> HistogramRegistry::add(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2)
+{
+  auto histVariant = add(name, title, histConfigSpec, callSumw2);
+  if (auto histPtr = std::get_if<std::shared_ptr<T>>(&histVariant)) {
+    return *histPtr;
+  } else {
+    throw runtime_error_f(R"(Histogram type specified in add<>("%s") does not match the actual type of the histogram!)", name);
+  }
+}
+
+template <typename T>
+std::shared_ptr<T> HistogramRegistry::add(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2)
+{
+  auto histVariant = add(name, title, histType, axes, callSumw2);
+  if (auto histPtr = std::get_if<std::shared_ptr<T>>(&histVariant)) {
+    return *histPtr;
+  } else {
+    throw runtime_error_f(R"(Histogram type specified in add<>("%s") does not match the actual type of the histogram!)", name);
+  }
+}
+
+template std::shared_ptr<TH1> HistogramRegistry::add<TH1>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<TH1> HistogramRegistry::add<TH1>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<TH2> HistogramRegistry::add<TH2>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<TH2> HistogramRegistry::add<TH2>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<TH3> HistogramRegistry::add<TH3>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<TH3> HistogramRegistry::add<TH3>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<TProfile> HistogramRegistry::add<TProfile>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<TProfile> HistogramRegistry::add<TProfile>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<TProfile2D> HistogramRegistry::add<TProfile2D>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<TProfile2D> HistogramRegistry::add<TProfile2D>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<TProfile3D> HistogramRegistry::add<TProfile3D>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<TProfile3D> HistogramRegistry::add<TProfile3D>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<THn> HistogramRegistry::add<THn>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<THn> HistogramRegistry::add<THn>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<THnSparse> HistogramRegistry::add<THnSparse>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<THnSparse> HistogramRegistry::add<THnSparse>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+template std::shared_ptr<StepTHn> HistogramRegistry::add<StepTHn>(char const* const name, char const* const title, const HistogramConfigSpec& histConfigSpec, bool callSumw2);
+template std::shared_ptr<StepTHn> HistogramRegistry::add<StepTHn>(char const* const name, char const* const title, HistType histType, const std::vector<AxisSpec>& axes, bool callSumw2);
+
 } // namespace o2::framework

From 53ed758f32eaf659b7a9dc0a7c701fcd300a7544 Mon Sep 17 00:00:00 2001
From: Anton Alkin <anton.alkin@cern.ch>
Date: Tue, 28 Jan 2025 14:26:35 +0100
Subject: [PATCH 3/8] DPL Analysis: remove unnecessary instances of
 selected_pack (#13892)

---
 Framework/Core/include/Framework/ASoA.h       | 115 ++++++++++++------
 .../Core/include/Framework/Configurable.h     |   5 +-
 .../include/Framework/GroupedCombinations.h   |  22 +++-
 3 files changed, 101 insertions(+), 41 deletions(-)

diff --git a/Framework/Core/include/Framework/ASoA.h b/Framework/Core/include/Framework/ASoA.h
index 8af872a64176d..8ef7ed9539ec7 100644
--- a/Framework/Core/include/Framework/ASoA.h
+++ b/Framework/Core/include/Framework/ASoA.h
@@ -208,8 +208,6 @@ template <typename D, typename... Cs>
 struct TableMetadata {
   using columns = framework::pack<Cs...>;
   using persistent_columns_t = framework::selected_pack<soa::is_persistent_column_t, Cs...>;
-  using external_index_columns_t = framework::selected_pack<soa::is_external_index_t, Cs...>;
-  using internal_index_columns_t = framework::selected_pack<soa::is_self_index_t, Cs...>;
 
   template <typename Key, typename... PCs>
   static consteval std::array<bool, sizeof...(PCs)> getMap(framework::pack<PCs...>)
@@ -806,9 +804,6 @@ concept is_dynamic_column = requires(C& c) {
 template <typename C>
 concept is_marker_column = requires { &C::mark; };
 
-template <typename T>
-using is_dynamic_t = std::conditional_t<is_dynamic_column<T>, std::true_type, std::false_type>;
-
 template <typename T>
 concept is_column = is_persistent_column<T> || is_dynamic_column<T> || is_indexing_column<T> || is_marker_column<T>;
 
@@ -1031,6 +1026,17 @@ concept can_bind = requires(T&& t) {
 template <typename... C>
 concept has_index = (is_indexing_column<C> || ...);
 
+template <is_index_column C>
+  requires(!is_self_index_column<C>)
+consteval auto getBinding() -> typename C::binding_t
+{
+}
+
+template <typename C>
+consteval auto getBinding() -> void
+{
+}
+
 template <typename D, typename O, typename IP, typename... C>
 struct TableIterator : IP, C... {
  public:
@@ -1038,9 +1044,9 @@ struct TableIterator : IP, C... {
   using policy_t = IP;
   using all_columns = framework::pack<C...>;
   using persistent_columns_t = framework::selected_pack<soa::is_persistent_column_t, C...>;
-  using external_index_columns_t = framework::selected_pack<soa::is_external_index_t, C...>;
-  using internal_index_columns_t = framework::selected_pack<soa::is_self_index_t, C...>;
-  using bindings_pack_t = decltype([]<typename... Cs>(framework::pack<Cs...>) -> framework::pack<typename Cs::binding_t...> {}(external_index_columns_t{})); // decltype(extractBindings(external_index_columns_t{}));
+  using bindings_pack_t = decltype([]<typename... Cs>(framework::pack<Cs...>) {
+    return framework::pack<decltype(getBinding<Cs>())...>{};
+  }(all_columns{}));
 
   TableIterator(arrow::ChunkedArray* columnData[sizeof...(C)], IP&& policy)
     : IP{policy},
@@ -1133,7 +1139,13 @@ struct TableIterator : IP, C... {
   template <typename... CL, typename TA>
   void doSetCurrentIndex(framework::pack<CL...>, TA* current)
   {
-    (CL::setCurrent(current), ...);
+    (framework::overloaded{
+      [&current, this]<is_index_column CI>
+        requires(!is_self_index_column<CI>)
+      () { CI::setCurrent(current); },
+      []<typename CI>() {}}
+        .template operator()<CL>(),
+      ...);
   }
 
   template <typename CL>
@@ -1145,24 +1157,40 @@ struct TableIterator : IP, C... {
   template <typename... Cs>
   auto getIndexBindingsImpl(framework::pack<Cs...>) const
   {
-    return std::vector<o2::soa::Binding>{static_cast<Cs const&>(*this).getCurrentRaw()...};
+    std::vector<o2::soa::Binding> result;
+    (framework::overloaded{
+      [this, &result]<is_index_column CI>
+        requires(!is_self_index_column<CI>)
+      () mutable {
+        result.emplace_back(CI::getCurrentRaw());
+      },
+      []<typename CI>() {}}
+        .template operator()<Cs>(),
+      ...);
+    return result;
   }
 
   auto getIndexBindings() const
   {
-    return getIndexBindingsImpl(external_index_columns_t{});
+    return getIndexBindingsImpl(all_columns{});
   }
 
   template <typename... TA>
   void bindExternalIndices(TA*... current)
   {
-    (doSetCurrentIndex(external_index_columns_t{}, current), ...);
+    (doSetCurrentIndex(all_columns{}, current), ...);
   }
 
   template <typename... Cs>
   void doSetCurrentIndexRaw(framework::pack<Cs...> p, std::vector<o2::soa::Binding>&& ptrs)
   {
-    (Cs::setCurrentRaw(ptrs[framework::has_type_at_v<Cs>(p)]), ...);
+    (framework::overloaded{
+      [&ptrs, p, this]<is_self_index_column CI>
+        requires(!is_self_index_column<CI>)
+      () { CI::setCurrentRaw(ptrs[framework::has_type_at_v<CI>(p)]); },
+      []<typename CI>() {}}
+        .template operator()<Cs>(),
+      ...);
   }
 
   template <typename... Cs, typename I>
@@ -1170,18 +1198,22 @@ struct TableIterator : IP, C... {
   {
     o2::soa::Binding b;
     b.bind(ptr);
-    (Cs::setCurrentRaw(b), ...);
+    (framework::overloaded{
+       [&ptr, &b, this]<is_self_index_column CI>() { CI::setCurrentRaw(b); },
+       []<typename CI>() {}}
+       .template operator()<Cs>(),
+     ...);
   }
 
   void bindExternalIndicesRaw(std::vector<o2::soa::Binding>&& ptrs)
   {
-    doSetCurrentIndexRaw(external_index_columns_t{}, std::forward<std::vector<o2::soa::Binding>>(ptrs));
+    doSetCurrentIndexRaw(all_columns{}, std::forward<std::vector<o2::soa::Binding>>(ptrs));
   }
 
   template <typename I>
   void bindInternalIndices(I const* table)
   {
-    doSetCurrentInternal(internal_index_columns_t{}, table);
+    doSetCurrentInternal(all_columns{}, table);
   }
 
  private:
@@ -1365,25 +1397,37 @@ static constexpr std::string getLabelFromTypeForKey(std::string const& key)
 template <typename B, typename... C>
 consteval static bool hasIndexTo(framework::pack<C...>&&)
 {
-  return (o2::soa::is_binding_compatible_v<B, typename C::binding_t>() || ...);
+  return (framework::overloaded{
+    []<is_index_column CI>
+      requires(!is_self_index_column<CI>)
+    () { return o2::soa::is_binding_compatible_v<B, typename CI::binding_t>(); },
+    []<typename CI>() { return false; }}
+      .template operator()<C>() ||
+    ...);
 }
 
 template <typename B, typename... C>
 consteval static bool hasSortedIndexTo(framework::pack<C...>&&)
 {
-  return ((C::sorted && o2::soa::is_binding_compatible_v<B, typename C::binding_t>()) || ...);
+  return (framework::overloaded{
+    []<is_index_column CI>
+      requires(!is_self_index_column<CI>)
+    () { return (CI::sorted && o2::soa::is_binding_compatible_v<B, typename CI::binding_t>()); },
+    []<typename CI>() {}}
+      .template operator()<C>() ||
+    ...);
 }
 
 template <typename B, typename Z>
 consteval static bool relatedByIndex()
 {
-  return hasIndexTo<B>(typename Z::table_t::external_index_columns_t{});
+  return hasIndexTo<B>(typename Z::table_t::columns_t{});
 }
 
 template <typename B, typename Z>
 consteval static bool relatedBySortedIndex()
 {
-  return hasSortedIndexTo<B>(typename Z::table_t::external_index_columns_t{});
+  return hasSortedIndexTo<B>(typename Z::table_t::columns_t{});
 }
 } // namespace o2::soa
 
@@ -1726,16 +1770,13 @@ class Table
   using persistent_columns_t = decltype([]<typename... C>(framework::pack<C...>&&) -> framework::selected_pack<soa::is_persistent_column_t, C...> {}(columns_t{}));
   using column_types = decltype([]<typename... C>(framework::pack<C...>) -> framework::pack<typename C::type...> {}(persistent_columns_t{}));
 
-  using external_index_columns_t = decltype([]<typename... C>(framework::pack<C...>&&) -> framework::selected_pack<soa::is_external_index_t, C...> {}(columns_t{}));
-  using internal_index_columns_t = decltype([]<typename... C>(framework::pack<C...>&&) -> framework::selected_pack<soa::is_self_index_t, C...> {}(columns_t{}));
   template <typename IP>
   using base_iterator = decltype(base_iter<D, O, IP>(columns_t{}));
 
   template <typename IP, typename Parent, typename... T>
   struct TableIteratorBase : base_iterator<IP> {
     using columns_t = typename Parent::columns_t;
-    using external_index_columns_t = typename Parent::external_index_columns_t;
-    using bindings_pack_t = decltype([]<typename... C>(framework::pack<C...>) -> framework::pack<typename C::binding_t...> {}(external_index_columns_t{}));
+    using bindings_pack_t = typename base_iterator<IP>::bindings_pack_t;
     // static constexpr const std::array<TableRef, sizeof...(T)> originals{T::ref...};
     static constexpr auto originals = Parent::originals;
     using policy_t = IP;
@@ -1828,7 +1869,7 @@ class Table
       using decayed = std::decay_t<TI>;
       if constexpr (framework::has_type<decayed>(bindings_pack_t{})) { // index to another table
         constexpr auto idx = framework::has_type_at_v<decayed>(bindings_pack_t{});
-        return framework::pack_element_t<idx, external_index_columns_t>::getId();
+        return framework::pack_element_t<idx, columns_t>::getId();
       } else if constexpr (std::same_as<decayed, Parent>) { // self index
         return this->globalIndex();
       } else if constexpr (is_indexing_column<decayed>) { // soa::Index<>
@@ -1838,20 +1879,17 @@ class Table
       }
     }
 
-    template <typename CD, typename... CDArgs>
+    template <soa::is_dynamic_column CD, typename... CDArgs>
     auto getDynamicColumn() const
     {
-      using decayed = std::decay_t<CD>;
-      static_assert(is_dynamic_t<decayed>(), "Requested column is not a dynamic column");
-      return static_cast<decayed>(*this).template getDynamicValue<CDArgs...>();
+      return static_cast<std::decay_t<CD>>(*this).template getDynamicValue<CDArgs...>();
     }
 
     template <typename B, typename CC>
+      requires(is_dynamic_column<CC> || is_persistent_column<CC>)
     auto getValue() const
     {
-      using COL = std::decay_t<CC>;
-      static_assert(is_dynamic_t<COL>() || soa::is_persistent_column<COL>, "Should be persistent or dynamic column with no argument that has a return type convertable to float");
-      return static_cast<B>(static_cast<COL>(*this).get());
+      return static_cast<B>(static_cast<std::decay_t<CC>>(*this).get());
     }
 
     template <typename B, typename... CCs>
@@ -2056,13 +2094,17 @@ class Table
 
   void bindInternalIndicesExplicit(o2::soa::Binding binding)
   {
-    doBindInternalIndicesExplicit(internal_index_columns_t{}, binding);
+    doBindInternalIndicesExplicit(columns_t{}, binding);
   }
 
   template <typename... Cs>
   void doBindInternalIndicesExplicit(framework::pack<Cs...>, o2::soa::Binding binding)
   {
-    (static_cast<Cs>(mBegin).setCurrentRaw(binding), ...);
+    (framework::overloaded{
+       [this, &binding]<is_self_index_column CI>() { static_cast<CI>(mBegin).setCurrentRaw(binding); },
+       []<typename CI>() {}}
+       .template operator()<Cs>(),
+     ...);
   }
 
   void bindExternalIndicesRaw(std::vector<o2::soa::Binding>&& ptrs)
@@ -2079,7 +2121,7 @@ class Table
   template <typename T>
   void copyIndexBindings(T& dest) const
   {
-    doCopyIndexBindings(external_index_columns_t{}, dest);
+    doCopyIndexBindings(columns_t{}, dest);
   }
 
   auto select(framework::expressions::Filter const& f) const
@@ -3298,7 +3340,6 @@ class FilteredBase : public T
   using T::originals;
   using columns_t = typename T::columns_t;
   using persistent_columns_t = typename T::persistent_columns_t;
-  using external_index_columns_t = typename T::external_index_columns_t;
 
   using iterator = T::template iterator_template_o<FilteredIndexPolicy, self_t>;
   using unfiltered_iterator = T::template iterator_template_o<DefaultIndexPolicy, self_t>;
@@ -3444,7 +3485,7 @@ class FilteredBase : public T
   template <typename T1>
   void copyIndexBindings(T1& dest) const
   {
-    doCopyIndexBindings(external_index_columns_t{}, dest);
+    doCopyIndexBindings(columns_t{}, dest);
   }
 
   template <typename T1>
diff --git a/Framework/Core/include/Framework/Configurable.h b/Framework/Core/include/Framework/Configurable.h
index 88e50cf3c7c26..930c37e700105 100644
--- a/Framework/Core/include/Framework/Configurable.h
+++ b/Framework/Core/include/Framework/Configurable.h
@@ -83,6 +83,9 @@ struct Configurable : IP {
 template <typename T, ConfigParamKind K = ConfigParamKind::kGeneric>
 using MutableConfigurable = Configurable<T, K, ConfigurablePolicyMutable<T, K>>;
 
+template <typename T>
+concept is_configurable = requires(T& t) { &T::operator typename T::type; };
+
 using ConfigurableAxis = Configurable<std::vector<double>, ConfigParamKind::kAxisSpec, ConfigurablePolicyConst<std::vector<double>, ConfigParamKind::kAxisSpec>>;
 
 template <typename R, typename T, typename... As>
@@ -97,7 +100,7 @@ struct ProcessConfigurable : Configurable<bool, ConfigParamKind::kProcessFlag> {
 };
 
 template <typename T>
-concept is_process_configurable = base_of_template<ProcessConfigurable, T>;
+concept is_process_configurable = is_configurable<T> && requires(T& t) { t.process; };
 
 #define PROCESS_SWITCH(_Class_, _Name_, _Help_, _Default_) \
   decltype(ProcessConfigurable{&_Class_ ::_Name_, #_Name_, _Default_, _Help_}) do##_Name_ = ProcessConfigurable{&_Class_ ::_Name_, #_Name_, _Default_, _Help_};
diff --git a/Framework/Core/include/Framework/GroupedCombinations.h b/Framework/Core/include/Framework/GroupedCombinations.h
index 21d8384e3aa6e..9f450489ac50f 100644
--- a/Framework/Core/include/Framework/GroupedCombinations.h
+++ b/Framework/Core/include/Framework/GroupedCombinations.h
@@ -34,14 +34,30 @@ auto interleaveTuples(std::tuple<T1s...>& t1, std::tuple<T2s...>& t2)
   return interleaveTuplesImpl(t1, t2, std::index_sequence_for<T1s...>());
 }
 
+template <soa::is_index_column T, typename G>
+  requires(!soa::is_self_index_column<T>)
+consteval auto isIndexTo()
+{
+  if constexpr (o2::soa::is_binding_compatible_v<G, typename T::binding_t>()) {
+    return std::true_type{};
+  } else {
+    return std::false_type{};
+  }
+}
+
+template <typename T, typename G>
+consteval auto isIndexTo()
+{
+  return std::false_type{};
+}
+
 template <typename T, typename G>
-using is_index_to_g_t = typename std::conditional<o2::soa::is_binding_compatible_v<G, typename T::binding_t>(), std::true_type, std::false_type>::type;
+using is_index_to_g_t = decltype(isIndexTo<T, G>());
 
 template <typename G, typename A>
 expressions::BindingNode getMatchingIndexNode()
 {
-  using external_index_columns_pack = typename A::external_index_columns_t;
-  using selected_indices_t = selected_pack_multicondition<is_index_to_g_t, pack<G>, external_index_columns_pack>;
+  using selected_indices_t = selected_pack_multicondition<is_index_to_g_t, pack<G>, typename A::columns_t>;
   static_assert(pack_size(selected_indices_t{}) == 1, "No matching index column from associated to grouping");
   using index_column_t = pack_head_t<selected_indices_t>;
   return expressions::BindingNode{index_column_t::mLabel, o2::framework::TypeIdHelpers::uniqueId<typename index_column_t::column_t>(), expressions::selectArrowType<typename index_column_t::type>()};

From 1640a6274e0b4eb1ecf25d0a418cc3aae0a849ea Mon Sep 17 00:00:00 2001
From: swenzel <sandro.wenzel@cern.ch>
Date: Tue, 28 Jan 2025 14:37:14 +0100
Subject: [PATCH 4/8] DigitizationContext: Prevent a segfault in QED filling

---
 DataFormats/simulation/src/DigitizationContext.cxx | 5 +++++
 1 file changed, 5 insertions(+)

diff --git a/DataFormats/simulation/src/DigitizationContext.cxx b/DataFormats/simulation/src/DigitizationContext.cxx
index 975458c41fcb3..e875cb61951a9 100644
--- a/DataFormats/simulation/src/DigitizationContext.cxx
+++ b/DataFormats/simulation/src/DigitizationContext.cxx
@@ -289,6 +289,11 @@ DigitizationContext* DigitizationContext::loadFromFile(std::string_view filename
 
 void DigitizationContext::fillQED(std::string_view QEDprefix, int max_events, double qedrate)
 {
+  if (mEventRecords.size() <= 1) {
+    // nothing to do
+    return;
+  }
+
   o2::steer::InteractionSampler qedInteractionSampler;
   qedInteractionSampler.setBunchFilling(mBCFilling);
 

From 0a74715ccddebbf7bc2cf23152f9884e3c5adccd Mon Sep 17 00:00:00 2001
From: shahoian <ruben.shahoyan@cern.ch>
Date: Tue, 28 Jan 2025 11:34:42 +0100
Subject: [PATCH 5/8] Add to matching debug info TPC inermost cl. row/pad

---
 .../include/GlobalTracking/MatchTPCITS.h      |  6 +++-
 Detectors/GlobalTracking/src/MatchTPCITS.cxx  | 32 ++++++++++++++-----
 2 files changed, 29 insertions(+), 9 deletions(-)

diff --git a/Detectors/GlobalTracking/include/GlobalTracking/MatchTPCITS.h b/Detectors/GlobalTracking/include/GlobalTracking/MatchTPCITS.h
index 8a8dbcb8fa9ae..00f2fc157a5ec 100644
--- a/Detectors/GlobalTracking/include/GlobalTracking/MatchTPCITS.h
+++ b/Detectors/GlobalTracking/include/GlobalTracking/MatchTPCITS.h
@@ -50,6 +50,7 @@
 #include "ITSReconstruction/RecoGeomHelper.h"
 #include "TPCFastTransform.h"
 #include "GPUO2InterfaceRefit.h"
+#include "GPUTPCGeometry.h"
 #include "GlobalTracking/MatchTPCITSParams.h"
 #include "DataFormatsITSMFT/TopologyDictionary.h"
 #include "DataFormatsITSMFT/TrkClusRef.h"
@@ -132,6 +133,8 @@ struct TrackLocTPC : public o2::track::TrackParCov {
   int sourceID = 0;                     ///< TPC track origin in
   o2::dataformats::GlobalTrackID gid{}; // global track source ID (TPC track may be part of it)
   int matchID = MinusOne;               ///< entry (non if MinusOne) of its matchTPC struct in the mMatchesTPC
+  uint8_t lowestRow = -1;
+  uint8_t padFromEdge = -1;
   Constraint_t constraint{Constrained};
 
   float getCorrectedTime(float dt) const // return time0 corrected for extra drift (to match certain Z)
@@ -143,7 +146,7 @@ struct TrackLocTPC : public o2::track::TrackParCov {
     return constraint == Constrained ? 0.f : (constraint == ASide ? dt : -dt);
   }
 
-  ClassDefNV(TrackLocTPC, 2);
+  ClassDefNV(TrackLocTPC, 3);
 };
 
 ///< ITS track outward parameters propagated to reference X, with time bracket and index of
@@ -738,6 +741,7 @@ class MatchTPCITS
   static constexpr float MaxSnp = 0.9;                 // max snp of ITS or TPC track at xRef to be matched
   static constexpr float MaxTgp = 2.064;               // max tg corresponting to MaxSnp = MaxSnp/std::sqrt(1.-MaxSnp^2)
   static constexpr float MinTBToCleanCache = 600.;     // keep in AB ITS cluster refs cache at most this number of TPC bins
+  static const o2::gpu::GPUTPCGeometry TPCGeometry;
 
   enum TimerIDs { SWTot,
                   SWPrepITS,
diff --git a/Detectors/GlobalTracking/src/MatchTPCITS.cxx b/Detectors/GlobalTracking/src/MatchTPCITS.cxx
index eb2a2212edb30..436a12df51352 100644
--- a/Detectors/GlobalTracking/src/MatchTPCITS.cxx
+++ b/Detectors/GlobalTracking/src/MatchTPCITS.cxx
@@ -9,6 +9,13 @@
 // granted to it by virtue of its status as an Intergovernmental Organization
 // or submit itself to any jurisdiction.
 
+#include "GPUO2Interface.h" // Needed for propper settings in GPUParam.h
+#include "GPUParam.h"
+#include "GPUParam.inc"
+#ifdef WITH_OPENMP
+#include <omp.h>
+#endif
+
 #include <TTree.h>
 #include <cassert>
 #include <algorithm>
@@ -50,13 +57,6 @@
 #include "ITS3Reconstruction/IOUtils.h"
 #endif
 
-#include "GPUO2Interface.h" // Needed for propper settings in GPUParam.h
-#include "GPUParam.h"
-#include "GPUParam.inc"
-#ifdef WITH_OPENMP
-#include <omp.h>
-#endif
-
 using namespace o2::globaltracking;
 
 using MatrixDSym4 = ROOT::Math::SMatrix<double, 4, 4, ROOT::Math::MatRepSym<double, 4>>;
@@ -68,6 +68,8 @@ constexpr float MatchTPCITS::Tan70, MatchTPCITS::Cos70I2, MatchTPCITS::MaxSnp, M
 
 LinksPoolMT* TPCABSeed::gLinksPool = nullptr;
 
+const o2::gpu::GPUTPCGeometry MatchTPCITS::TPCGeometry{};
+
 //______________________________________________
 MatchTPCITS::MatchTPCITS() = default;
 
@@ -428,6 +430,12 @@ int MatchTPCITS::addTPCSeed(const o2::track::TrackParCov& _tr, float t0, float t
   if (clRow > mParams->askMinTPCRow[clSect]) {
     return -9;
   }
+  const auto& clus = mTPCClusterIdxStruct->clusters[clSect][clRow][clIdx];
+  uint8_t padFromEdge = uint8_t(clus.getPad());
+  if (padFromEdge > TPCGeometry.NPads(clRow) / 2) {
+    padFromEdge = TPCGeometry.NPads(clRow) - 1 - padFromEdge;
+  }
+
   // create working copy of track param
   bool extConstrained = srcGID.getSource() != GTrackID::TPC;
   if (extConstrained) {
@@ -442,6 +450,8 @@ int MatchTPCITS::addTPCSeed(const o2::track::TrackParCov& _tr, float t0, float t
                 tpcID,
                 srcGID,
                 MinusOne,
+                clRow,
+                padFromEdge,
                 (extConstrained || tpcOrig.hasBothSidesClusters()) ? TrackLocTPC::Constrained : (tpcOrig.hasASideClustersOnly() ? TrackLocTPC::ASide : TrackLocTPC::CSide)});
   // propagate to matching Xref
   const auto& trackTune = TrackTuneParams::Instance();
@@ -2872,7 +2882,7 @@ void MatchTPCITS::dumpTPCOrig(bool acc, int tpcIndex)
   ///< fill debug tree for TPC original tracks (passing pT cut)
   mTimer[SWDBG].Start(false);
   const auto& tpcOrig = mTPCTracksArray[tpcIndex];
-  uint8_t clSect = 0, clRow = 0, prevRow = 0xff;
+  uint8_t clSect = 0, clRow = 0, prevRow = 0xff, padFromEdge = -1;
   uint32_t clIdx = 0;
   int nshared = 0;
   std::array<bool, 152> shMap{};
@@ -2888,6 +2898,11 @@ void MatchTPCITS::dumpTPCOrig(bool acc, int tpcIndex)
       prevRawShared = true;
     }
   }
+  const auto& clus = mTPCClusterIdxStruct->clusters[clSect][clRow][clIdx];
+  padFromEdge = uint8_t(clus.getPad());
+  if (padFromEdge > TPCGeometry.NPads(clRow) / 2) {
+    padFromEdge = TPCGeometry.NPads(clRow) - 1 - padFromEdge;
+  }
   int tb = tpcOrig.getTime0() * mNTPCOccBinLengthInv;
   float mltTPC = tb < 0 ? mTBinClOcc[0] : (tb >= mTBinClOcc.size() ? mTBinClOcc.back() : mTBinClOcc[tb]);
   (*mDBGOut) << "tpcOrig"
@@ -2900,6 +2915,7 @@ void MatchTPCITS::dumpTPCOrig(bool acc, int tpcIndex)
              << "time0=" << tpcOrig.getTime0()
              << "trc=" << ((o2::track::TrackParCov&)tpcOrig)
              << "minRow=" << clRow
+             << "padFromEdge=" << padFromEdge
              << "multTPC=" << mltTPC;
   if (mMCTruthON) {
     (*mDBGOut) << "tpcOrig"

From e7179fcb15cb0b9852f3abc22b4ebaa883d55cb2 Mon Sep 17 00:00:00 2001
From: Matteo Concas <matteo.concas@cern.ch>
Date: Wed, 29 Jan 2025 08:41:50 +0100
Subject: [PATCH 6/8] ITS-GPU: Cleanup for some host code (#13907)

* Cleanup

* Fix nCells printout
---
 .../GPU/ITStrackingGPU/TimeFrameGPU.h         | 12 ++++-
 .../GPU/ITStrackingGPU/TrackingKernels.h      | 26 +++++-----
 .../ITS/tracking/GPU/cuda/TimeFrameGPU.cu     | 14 +++++-
 .../tracking/GPU/cuda/TrackerTraitsGPU.cxx    | 48 ++++++++-----------
 .../ITS/tracking/GPU/cuda/TrackingKernels.cu  | 46 +++++++++---------
 5 files changed, 81 insertions(+), 65 deletions(-)

diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h
index 066bef7631415..100e49def0d50 100644
--- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h
+++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h
@@ -77,7 +77,8 @@ class TimeFrameGPU : public TimeFrame
   void createCellsDevice();
   void createCellsLUTDevice();
   void createNeighboursIndexTablesDevice();
-  void createNeighboursDevice(const unsigned int& layer, std::vector<std::pair<int, int>>& neighbours);
+  void createNeighboursDevice(const unsigned int layer, const unsigned int nNeighbours);
+  void createNeighboursDevice(const unsigned int layer, std::vector<std::pair<int, int>>& neighbours);
   void createNeighboursLUTDevice(const int, const unsigned int);
   void createNeighboursDeviceArray();
   void createTrackITSExtDevice(std::vector<CellSeed>&);
@@ -151,6 +152,9 @@ class TimeFrameGPU : public TimeFrame
   gsl::span<Tracklet*> getDeviceTracklet() { return mTrackletsDevice; }
   gsl::span<CellSeed*> getDeviceCells() { return mCellsDevice; }
 
+  // Overridden getters
+  int getNumberOfCells() const;
+
  private:
   void allocMemAsync(void**, size_t, Stream*, bool); // Abstract owned and unowned memory allocations
   bool mHostRegistered = false;
@@ -252,6 +256,12 @@ inline std::vector<unsigned int> TimeFrameGPU<nLayers>::getClusterSizes()
   return sizes;
 }
 
+template <int nLayers>
+inline int TimeFrameGPU<nLayers>::getNumberOfCells() const
+{
+  return std::accumulate(mNCells.begin(), mNCells.end(), 0);
+}
+
 } // namespace gpu
 } // namespace its
 } // namespace o2
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h
index 78636d00788bf..720867ddaba29 100644
--- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h
+++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h
@@ -148,19 +148,19 @@ void computeCellsHandler(const Cluster** sortedClusters,
                          const int nBlocks,
                          const int nThreads);
 
-void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
-                                int* neighboursLUTs,
-                                int** cellsLUTs,
-                                gpuPair<int, int>* cellNeighbours,
-                                int* neighboursIndexTable,
-                                const float maxChi2ClusterAttachment,
-                                const float bz,
-                                const int layerIndex,
-                                const unsigned int nCells,
-                                const unsigned int nCellsNext,
-                                const int maxCellNeighbours,
-                                const int nBlocks,
-                                const int nThreads);
+unsigned int countCellNeighboursHandler(CellSeed** cellsLayersDevice,
+                                        int* neighboursLUTs,
+                                        int** cellsLUTs,
+                                        gpuPair<int, int>* cellNeighbours,
+                                        int* neighboursIndexTable,
+                                        const float maxChi2ClusterAttachment,
+                                        const float bz,
+                                        const int layerIndex,
+                                        const unsigned int nCells,
+                                        const unsigned int nCellsNext,
+                                        const int maxCellNeighbours,
+                                        const int nBlocks,
+                                        const int nThreads);
 
 void computeCellNeighboursHandler(CellSeed** cellsLayersDevice,
                                   int* neighboursLUTs,
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu
index fd067b9930fd0..b1aa55f533c34 100644
--- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu
+++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu
@@ -406,7 +406,19 @@ void TimeFrameGPU<nLayers>::loadTrackSeedsDevice(std::vector<CellSeed>& seeds)
 }
 
 template <int nLayers>
-void TimeFrameGPU<nLayers>::createNeighboursDevice(const unsigned int& layer, std::vector<std::pair<int, int>>& neighbours)
+void TimeFrameGPU<nLayers>::createNeighboursDevice(const unsigned int layer, const unsigned int nNeighbours)
+{
+  START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "reserving neighbours");
+  LOGP(debug, "gpu-allocation: reserving {} neighbours (pairs), for {} MB.", nNeighbours, nNeighbours * sizeof(gpuPair<int, int>) / MB);
+  allocMemAsync(reinterpret_cast<void**>(&mNeighbourPairsDevice[layer]), nNeighbours * sizeof(gpuPair<int, int>), &(mGpuStreams[0]), getExtAllocator());
+  checkGPUError(cudaMemsetAsync(mNeighbourPairsDevice[layer], -1, nNeighbours * sizeof(gpuPair<int, int>), mGpuStreams[0].get()));
+  LOGP(debug, "gpu-allocation: reserving {} neighbours, for {} MB.", nNeighbours, nNeighbours * sizeof(gpuPair<int, int>) / MB);
+  allocMemAsync(reinterpret_cast<void**>(&mNeighboursDevice[layer]), nNeighbours * sizeof(int), &(mGpuStreams[0]), getExtAllocator());
+  STOP_GPU_STREAM_TIMER(mGpuStreams[0].get());
+}
+
+template <int nLayers>
+void TimeFrameGPU<nLayers>::createNeighboursDevice(const unsigned int layer, std::vector<std::pair<int, int>>& neighbours)
 {
   START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "reserving neighbours");
   mCellsNeighbours[layer].clear();
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
index 4821ebb636f54..3c65faddcff71 100644
--- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
+++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
@@ -212,36 +212,30 @@ void TrackerTraitsGPU<nLayers>::findCellsNeighboursHybrid(const int iteration)
 {
   mTimeFrameGPU->createNeighboursIndexTablesDevice();
   auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
-  std::vector<std::vector<std::pair<int, int>>> cellsNeighboursLayer(mTrkParams[iteration].CellsPerRoad() - 1);
   for (int iLayer{0}; iLayer < mTrkParams[iteration].CellsPerRoad() - 1; ++iLayer) {
     const int nextLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer + 1])};
-    mTimeFrameGPU->getCellsNeighboursLUT()[iLayer].clear();
-    mTimeFrameGPU->getCellsNeighboursLUT()[iLayer].resize(nextLayerCellsNum, 0);
 
-    // if (mTimeFrameGPU->getCells()[iLayer + 1].empty() ||
-    //     mTimeFrameGPU->getCellsLookupTable()[iLayer].empty()) {
-    //   mTimeFrameGPU->getCellsNeighbours()[iLayer].clear();
-    //   continue;
-    // }
+    if (!nextLayerCellsNum) {
+      continue;
+    }
 
     mTimeFrameGPU->createNeighboursLUTDevice(iLayer, nextLayerCellsNum);
-    countCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
-                               mTimeFrameGPU->getDeviceNeighboursLUT(iLayer), // LUT is initialised here.
-                               mTimeFrameGPU->getDeviceArrayCellsLUT(),
-                               mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
-                               mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer),
-                               mTrkParams[0].MaxChi2ClusterAttachment,
-                               mBz,
-                               iLayer,
-                               mTimeFrameGPU->getNCells()[iLayer],
-                               nextLayerCellsNum,
-                               1e2,
-                               conf.nBlocks,
-                               conf.nThreads);
-    mTimeFrameGPU->downloadNeighboursLUTDevice(mTimeFrameGPU->getCellsNeighboursLUT()[iLayer], iLayer);
-    // Get the number of found cells from LUT
-    cellsNeighboursLayer[iLayer].resize(mTimeFrameGPU->getCellsNeighboursLUT()[iLayer].back());
-    mTimeFrameGPU->createNeighboursDevice(iLayer, cellsNeighboursLayer[iLayer]);
+    unsigned int nNeigh = countCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
+                                                     mTimeFrameGPU->getDeviceNeighboursLUT(iLayer), // LUT is initialised here.
+                                                     mTimeFrameGPU->getDeviceArrayCellsLUT(),
+                                                     mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
+                                                     mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer),
+                                                     mTrkParams[0].MaxChi2ClusterAttachment,
+                                                     mBz,
+                                                     iLayer,
+                                                     mTimeFrameGPU->getNCells()[iLayer],
+                                                     nextLayerCellsNum,
+                                                     1e2,
+                                                     conf.nBlocks,
+                                                     conf.nThreads);
+
+    mTimeFrameGPU->createNeighboursDevice(iLayer, nNeigh);
+
     computeCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
                                  mTimeFrameGPU->getDeviceNeighboursLUT(iLayer),
                                  mTimeFrameGPU->getDeviceArrayCellsLUT(),
@@ -255,13 +249,11 @@ void TrackerTraitsGPU<nLayers>::findCellsNeighboursHybrid(const int iteration)
                                  1e2,
                                  conf.nBlocks,
                                  conf.nThreads);
-    mTimeFrameGPU->getCellsNeighbours()[iLayer].clear();
-    mTimeFrameGPU->getCellsNeighbours()[iLayer].reserve(cellsNeighboursLayer[iLayer].size());
 
     filterCellNeighboursHandler(mTimeFrameGPU->getCellsNeighbours()[iLayer],
                                 mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
                                 mTimeFrameGPU->getDeviceNeighbours(iLayer),
-                                cellsNeighboursLayer[iLayer].size());
+                                nNeigh);
   }
   mTimeFrameGPU->createNeighboursDeviceArray();
   mTimeFrameGPU->unregisterRest();
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
index 10459cf800b6c..4fa7913c10e82 100644
--- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
+++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
@@ -1066,19 +1066,19 @@ void computeCellsHandler(
     nSigmaCut);               // const float
 }
 
-void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
-                                int* neighboursLUT,
-                                int** cellsLUTs,
-                                gpuPair<int, int>* cellNeighbours,
-                                int* neighboursIndexTable,
-                                const float maxChi2ClusterAttachment,
-                                const float bz,
-                                const int layerIndex,
-                                const unsigned int nCells,
-                                const unsigned int nCellsNext,
-                                const int maxCellNeighbours,
-                                const int nBlocks,
-                                const int nThreads)
+unsigned int countCellNeighboursHandler(CellSeed** cellsLayersDevice,
+                                        int* neighboursLUT,
+                                        int** cellsLUTs,
+                                        gpuPair<int, int>* cellNeighbours,
+                                        int* neighboursIndexTable,
+                                        const float maxChi2ClusterAttachment,
+                                        const float bz,
+                                        const int layerIndex,
+                                        const unsigned int nCells,
+                                        const unsigned int nCellsNext,
+                                        const int maxCellNeighbours,
+                                        const int nBlocks,
+                                        const int nThreads)
 {
   gpu::computeLayerCellNeighboursKernel<true><<<nBlocks, nThreads>>>(
     cellsLayersDevice,
@@ -1091,8 +1091,7 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
     layerIndex,
     nCells,
     maxCellNeighbours);
-  // gpuCheckError(cudaPeekAtLastError());
-  // gpuCheckError(cudaDeviceSynchronize());
+
   void *d_temp_storage = nullptr, *d_temp_storage_2 = nullptr;
   size_t temp_storage_bytes = 0, temp_storage_bytes_2 = 0;
   gpuCheckError(cub::DeviceScan::InclusiveSum(d_temp_storage,     // d_temp_storage
@@ -1102,17 +1101,19 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
                                               nCellsNext));       // num_items
 
   discardResult(cudaMalloc(&d_temp_storage, temp_storage_bytes));
-  gpuCheckError(cub::DeviceScan::InclusiveSum(d_temp_storage,       // d_temp_storage
-                                              temp_storage_bytes,   // temp_storage_bytes
-                                              neighboursLUT,        // d_in
-                                              neighboursLUT,        // d_out
-                                              nCellsNext));         // num_items
+  gpuCheckError(cub::DeviceScan::InclusiveSum(d_temp_storage,     // d_temp_storage
+                                              temp_storage_bytes, // temp_storage_bytes
+                                              neighboursLUT,      // d_in
+                                              neighboursLUT,      // d_out
+                                              nCellsNext));       // num_items
+
   gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage_2,     // d_temp_storage
                                               temp_storage_bytes_2, // temp_storage_bytes
                                               neighboursIndexTable, // d_in
                                               neighboursIndexTable, // d_out
                                               nCells + 1,           // num_items
                                               0));                  // NOLINT: this is the offset of the sum, not a pointer
+
   discardResult(cudaMalloc(&d_temp_storage_2, temp_storage_bytes_2));
   gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage_2,     // d_temp_storage
                                               temp_storage_bytes_2, // temp_storage_bytes
@@ -1120,10 +1121,11 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
                                               neighboursIndexTable, // d_out
                                               nCells + 1,           // num_items
                                               0));                  // NOLINT: this is the offset of the sum, not a pointer
+  unsigned int nNeighbours;
+  gpuCheckError(cudaMemcpy(&nNeighbours, &neighboursLUT[nCellsNext - 1], sizeof(unsigned int), cudaMemcpyDeviceToHost));
   gpuCheckError(cudaFree(d_temp_storage));
   gpuCheckError(cudaFree(d_temp_storage_2));
-  gpuCheckError(cudaPeekAtLastError());
-  gpuCheckError(cudaDeviceSynchronize());
+  return nNeighbours;
 }
 
 void computeCellNeighboursHandler(CellSeed** cellsLayersDevice,

From dff469f9af9f2a8665b7680015402ba4b5e1637b Mon Sep 17 00:00:00 2001
From: shahoian <ruben.shahoyan@cern.ch>
Date: Tue, 28 Jan 2025 14:35:16 +0100
Subject: [PATCH 7/8] Fix fake/correct status for barrel tracks

The fMcMask bit 15 (fake global track label or TOF_label != TPC_lable) was
wrong since original TOF cluster label (set in the reconstruction) was compared
with TPC remapped label prepared for AOD storage.
In fact, we don't need to consider separately the global_label.isFake and TOF-TPC mismach:
TOF is the last detector in the matching process and the global track label is determined
by the TPC track label. Hence, if the TOF match is present but its cluster is not contributed
by the TPC track, the global label isFake will be necessarilly true, and vice versa.

Also, the status of bit 13 (flagging ITS-TPC mismatch) was covering only track-to-track
matches but not those from the afterburner.

Now settings of fakeness relies on the isFake status from the reconstruction.
---
 .../GlobalTracking/src/RecoContainer.cxx      |  2 +-
 .../AODProducerWorkflowSpec.h                 |  2 -
 Detectors/AOD/src/AODProducerWorkflowSpec.cxx | 56 +++++--------------
 3 files changed, 15 insertions(+), 45 deletions(-)

diff --git a/DataFormats/Detectors/GlobalTracking/src/RecoContainer.cxx b/DataFormats/Detectors/GlobalTracking/src/RecoContainer.cxx
index 60c18b966abed..c26de2bfda896 100644
--- a/DataFormats/Detectors/GlobalTracking/src/RecoContainer.cxx
+++ b/DataFormats/Detectors/GlobalTracking/src/RecoContainer.cxx
@@ -1447,7 +1447,7 @@ RecoContainer::GlobalIDSet RecoContainer::getSingleDetectorRefs(GTrackID gidx) c
     table[GTrackID::MCH] = parent0.getMCHRef();
     table[GTrackID::MID] = parent0.getMIDRef();
   }
-  return std::move(table);
+  return table;
 }
 
 //________________________________________________________
diff --git a/Detectors/AOD/include/AODProducerWorkflow/AODProducerWorkflowSpec.h b/Detectors/AOD/include/AODProducerWorkflow/AODProducerWorkflowSpec.h
index 241846f1a9270..ae866b2006e7f 100644
--- a/Detectors/AOD/include/AODProducerWorkflow/AODProducerWorkflowSpec.h
+++ b/Detectors/AOD/include/AODProducerWorkflow/AODProducerWorkflowSpec.h
@@ -482,8 +482,6 @@ class AODProducerWorkflowDPL : public Task
   // using -1 as dummies for AOD
   struct MCLabels {
     uint32_t labelID = -1;
-    uint32_t labelITS = -1;
-    uint32_t labelTPC = -1;
     uint16_t labelMask = 0;
     uint8_t fwdLabelMask = 0;
   };
diff --git a/Detectors/AOD/src/AODProducerWorkflowSpec.cxx b/Detectors/AOD/src/AODProducerWorkflowSpec.cxx
index 6083052eb1168..c8e3209cc0e29 100644
--- a/Detectors/AOD/src/AODProducerWorkflowSpec.cxx
+++ b/Detectors/AOD/src/AODProducerWorkflowSpec.cxx
@@ -109,6 +109,7 @@ using PVertex = o2::dataformats::PrimaryVertex;
 using GIndex = o2::dataformats::VtxTrackIndex;
 using DataRequest = o2::globaltracking::DataRequest;
 using GID = o2::dataformats::GlobalTrackID;
+using DetID = o2::detectors::DetID;
 using SMatrix55Sym = ROOT::Math::SMatrix<double, 5, 5, ROOT::Math::MatRepSym<double, 5>>;
 
 namespace o2::aodproducer
@@ -1058,9 +1059,9 @@ void AODProducerWorkflowDPL::fillMCTrackLabelsTable(MCTrackLabelCursorType& mcTr
                                                     int vertexId)
 {
   // labelMask (temporary) usage:
-  //   bit 13 -- ITS/TPC or TPC/TOF labels are not equal
+  //   bit 13 -- ITS/TPC with ITS label (track of AB tracklet) different from TPC
   //   bit 14 -- isNoise() == true
-  //   bit 15 -- isFake() == true
+  //   bit 15 -- isFake() == true (defined by the fakeness of the top level global track, i.e. if TOF is present, fake means that the track of the TPC label does not contribute to TOF cluster)
   // labelID = -1 -- label is not set
 
   for (int src = GIndex::NSources; src--;) {
@@ -1084,7 +1085,7 @@ void AODProducerWorkflowDPL::fillMCTrackLabelsTable(MCTrackLabelCursorType& mcTr
 
       if (GIndex::includesSource(src, mInputSources)) {
         auto mcTruth = data.getTrackMCLabel(trackIndex);
-        MCLabels labelHolder;
+        MCLabels labelHolder{};
         if ((src == GIndex::Source::MFT) || (src == GIndex::Source::MFTMCH) || (src == GIndex::Source::MCH) || (src == GIndex::Source::MCHMID)) { // treating mft and fwd labels separately
           if (!needToStore(src == GIndex::Source::MFT ? mGIDToTableMFTID : mGIDToTableFwdID)) {
             continue;
@@ -1110,51 +1111,22 @@ void AODProducerWorkflowDPL::fillMCTrackLabelsTable(MCTrackLabelCursorType& mcTr
             continue;
           }
           if (mcTruth.isValid()) { // if not set, -1 will be stored
-            labelHolder.labelID = (mToStore[mcTruth.getSourceID()][mcTruth.getEventID()])[mcTruth.getTrackID()];
-          }
-          // treating possible mismatches and fakes for global tracks
-          auto contributorsGID = data.getSingleDetectorRefs(trackIndex);
-          bool isSetTPC = contributorsGID[GIndex::Source::TPC].isIndexSet();
-          bool isSetITS = contributorsGID[GIndex::Source::ITS].isIndexSet();
-          bool isSetTOF = contributorsGID[GIndex::Source::TOF].isIndexSet();
-          bool isTOFFake = true;
-          if (isSetTPC && (isSetITS || isSetTOF)) {
-            auto mcTruthTPC = data.getTrackMCLabel(contributorsGID[GIndex::Source::TPC]);
-            if (mcTruthTPC.isValid()) {
-              labelHolder.labelTPC = (mToStore[mcTruthTPC.getSourceID()][mcTruthTPC.getEventID()])[mcTruthTPC.getTrackID()];
-              labelHolder.labelID = labelHolder.labelTPC;
-            }
-            if (isSetITS) {
-              auto mcTruthITS = data.getTrackMCLabel(contributorsGID[GIndex::Source::ITS]);
-              if (mcTruthITS.isValid()) {
-                labelHolder.labelITS = (mToStore[mcTruthITS.getSourceID()][mcTruthITS.getEventID()])[mcTruthITS.getTrackID()];
-              }
-              if (labelHolder.labelITS != labelHolder.labelTPC) {
-                LOG(debug) << "ITS-TPC MCTruth: labelIDs do not match at " << trackIndex.getIndex() << ", src = " << src;
-                labelHolder.labelMask |= (0x1 << 13);
-              }
+            labelHolder.labelID = (mToStore[mcTruth.getSourceID()][mcTruth.getEventID()])[mcTruth.getTrackID()]; // defined by TPC if it contributes, otherwise: by ITS
+            if (mcTruth.isFake()) {
+              labelHolder.labelMask |= (0x1 << 15);
             }
-            if (isSetTOF) {
-              const auto& labelsTOF = data.getTOFClustersMCLabels()->getLabels(contributorsGID[GIndex::Source::TOF]);
-              for (auto& mcLabel : labelsTOF) {
-                if (!mcLabel.isValid()) {
-                  continue;
-                }
-                if (mcLabel == labelHolder.labelTPC) {
-                  isTOFFake = false;
-                  break;
+            if (trackIndex.includesDet(DetID::TPC) && trackIndex.getSource() != GIndex::Source::TPC) { // this is global track
+              auto contributorsGID = data.getSingleDetectorRefs(trackIndex);
+              if (contributorsGID[GIndex::Source::ITSTPC].isIndexSet()) { // there is a match to ITS tracks or ITSAB tracklet!
+                if (data.getTrackMCLabel(contributorsGID[GIndex::Source::ITSTPC]).isFake()) {
+                  labelHolder.labelMask |= (0x1 << 13);
                 }
               }
             }
-          }
-          if (mcTruth.isFake() || (isSetTOF && isTOFFake)) {
-            labelHolder.labelMask |= (0x1 << 15);
-          }
-          if (mcTruth.isNoise()) {
+          } else if (mcTruth.isNoise()) {
             labelHolder.labelMask |= (0x1 << 14);
           }
-          mcTrackLabelCursor(labelHolder.labelID,
-                             labelHolder.labelMask);
+          mcTrackLabelCursor(labelHolder.labelID, labelHolder.labelMask);
         }
       }
     }

From 895dc910cdf77b919cfa06c82fa27105a79e4ce5 Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Wed, 29 Jan 2025 10:42:46 +0100
Subject: [PATCH 8/8] DPL: drop duplicated code

---
 Framework/Core/src/DPLMonitoringBackend.cxx | 9 +--------
 1 file changed, 1 insertion(+), 8 deletions(-)

diff --git a/Framework/Core/src/DPLMonitoringBackend.cxx b/Framework/Core/src/DPLMonitoringBackend.cxx
index 354c9fceef963..5fc0ff99703ba 100644
--- a/Framework/Core/src/DPLMonitoringBackend.cxx
+++ b/Framework/Core/src/DPLMonitoringBackend.cxx
@@ -13,20 +13,13 @@
 #include "Framework/DriverClient.h"
 #include "Framework/ServiceRegistry.h"
 #include "Framework/RuntimeError.h"
+#include "Framework/VariantHelpers.h"
 #include <fmt/format.h>
 #include <sstream>
 
 namespace o2::framework
 {
 
-template <class... Ts>
-struct overloaded : Ts... {
-  using Ts::operator()...;
-};
-template <class... Ts>
-overloaded(Ts...) -> overloaded<Ts...>;
-
-
 DPLMonitoringBackend::DPLMonitoringBackend(ServiceRegistryRef registry)
   : mRegistry{registry}
 {