From 2f52be1afa9d96fc02148b3fee7eb73fb1ef7661 Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Tue, 19 May 2026 11:33:06 +0200
Subject: [PATCH 01/17] DPL MCP: support multiple workflows

Rather than having to specify a single workflow in the
MCP configuration, allow it to connect to any running workflow.
---
 .../scripts/dpl-mcp-server/dpl_mcp_server.py  | 351 +++++++++++-------
 1 file changed, 223 insertions(+), 128 deletions(-)

diff --git a/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py b/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py
index 3900a646632a1..dca5058b01dcd 100644
--- a/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py
+++ b/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py
@@ -14,19 +14,21 @@
 Bridges the DPL driver /status WebSocket endpoint to MCP tools so that an
 AI assistant (e.g. Claude) can inspect and monitor a running DPL workflow.
 
+Supports multiple concurrent workflows.  Use the ``connect`` tool to attach
+to a running topology by port or PID, then pass the returned workflow name
+to every other tool.
+
 Usage
 -----
-    python3 dpl_mcp_server.py --port 8080
-    python3 dpl_mcp_server.py --pid 12345       # port derived as 8080 + pid % 30000
-    DPL_STATUS_PORT=8080 python3 dpl_mcp_server.py
+    python3 dpl_mcp_server.py
 
-Wire protocol (client → driver)
+Wire protocol (client -> driver)
 --------------------------------
     {"cmd":"list_metrics","device":"<name>"}
     {"cmd":"subscribe","device":"<name>","metrics":["m1","m2"]}
     {"cmd":"unsubscribe","device":"<name>","metrics":["m1"]}
 
-Wire protocol (driver → client)
+Wire protocol (driver -> client)
 --------------------------------
     {"type":"snapshot","devices":[{"name","pid","active","streamingState","deviceState"},...]}
     {"type":"update","device":<idx>,"name":"<name>","metrics":{<name:value,...>}}
@@ -35,80 +37,115 @@
 
 from __future__ import annotations
 
-import argparse
 import asyncio
 import json
-import os
-import sys
 from typing import Any
 
 import websockets
 from mcp.server.fastmcp import FastMCP
 
+
 # ---------------------------------------------------------------------------
-# Global connection state (all access from the single asyncio event loop)
+# Per-workflow connection state
 # ---------------------------------------------------------------------------
-_port: int = 8080
-_ws: Any = None
-_reader_task: asyncio.Task | None = None
-_snapshot: dict = {}
-_updates: list[dict] = []
-_logs: list[dict] = []
-_metrics_lists: dict[str, list[str]] = {}
-
-
-async def _ensure_connected() -> None:
-    """Connect (or reconnect) to the driver's /status WebSocket."""
-    global _ws, _reader_task
-
-    # Check liveness of existing connection.
-    if _ws is not None:
+class WorkflowConnection:
+    """Holds WebSocket connection and buffered state for one DPL workflow."""
+
+    def __init__(self, port: int, name: str):
+        self.port = port
+        self.name = name
+        self.ws: Any = None
+        self.reader_task: asyncio.Task | None = None
+        self.snapshot: dict = {}
+        self.updates: list[dict] = []
+        self.logs: list[dict] = []
+        self.metrics_lists: dict[str, list[str]] = {}
+
+    async def ensure_connected(self) -> None:
+        """Connect (or reconnect) to the driver's /status WebSocket."""
+        if self.ws is not None:
+            try:
+                pong = await asyncio.wait_for(self.ws.ping(), timeout=2.0)
+                await pong
+                return
+            except Exception:
+                old_ws = self.ws
+                self.ws = None
+                if self.reader_task is not None and not self.reader_task.done():
+                    self.reader_task.cancel()
+                    try:
+                        await self.reader_task
+                    except (asyncio.CancelledError, Exception):
+                        pass
+                self.reader_task = None
+                try:
+                    await old_ws.close()
+                except Exception:
+                    pass
+
+        url = f"ws://localhost:{self.port}/status"
+        self.ws = await websockets.connect(url, subprotocols=["dpl"])
+        if self.reader_task is None or self.reader_task.done():
+            self.reader_task = asyncio.create_task(self._reader())
+
+    async def _reader(self) -> None:
+        """Background task: read frames from the driver and buffer them."""
         try:
-            pong = await asyncio.wait_for(_ws.ping(), timeout=2.0)
-            await pong
-            return
+            async for raw in self.ws:
+                try:
+                    msg = json.loads(raw)
+                except json.JSONDecodeError:
+                    continue
+                t = msg.get("type")
+                if t == "snapshot":
+                    self.snapshot = msg
+                    self.metrics_lists.clear()
+                elif t == "update":
+                    self.updates.append(msg)
+                elif t == "log":
+                    self.logs.append(msg)
+                elif t == "metrics_list":
+                    device = msg.get("device", "")
+                    self.metrics_lists[device] = msg.get("metrics", [])
         except Exception:
-            _ws = None
-            if _reader_task is not None and not _reader_task.done():
-                _reader_task.cancel()
-            _reader_task = None
-
-    url = f"ws://localhost:{_port}/status"
-    _ws = await websockets.connect(url, subprotocols=["dpl"])
-    if _reader_task is None or _reader_task.done():
-        _reader_task = asyncio.create_task(_reader())
-
-
-async def _reader() -> None:
-    """Background task: read frames from the driver and buffer them."""
-    global _ws, _snapshot, _updates, _logs, _metrics_lists
-    try:
-        async for raw in _ws:
+            pass
+        finally:
+            self.ws = None
+
+    async def send(self, obj: dict) -> None:
+        await self.ensure_connected()
+        await self.ws.send(json.dumps(obj, separators=(",", ":")))
+
+    async def close(self) -> None:
+        ws = self.ws
+        self.ws = None
+        if self.reader_task is not None and not self.reader_task.done():
+            self.reader_task.cancel()
             try:
-                msg = json.loads(raw)
-            except json.JSONDecodeError:
-                continue
-            t = msg.get("type")
-            if t == "snapshot":
-                _snapshot = msg
-                # Clear stale metric lists from a previous driver instance.
-                _metrics_lists.clear()
-            elif t == "update":
-                _updates.append(msg)
-            elif t == "log":
-                _logs.append(msg)
-            elif t == "metrics_list":
-                device = msg.get("device", "")
-                _metrics_lists[device] = msg.get("metrics", [])
-    except Exception:
-        pass
-    finally:
-        _ws = None
-
-
-async def _send(obj: dict) -> None:
-    await _ensure_connected()
-    await _ws.send(json.dumps(obj, separators=(",", ":")))
+                await self.reader_task
+            except (asyncio.CancelledError, Exception):
+                pass
+        self.reader_task = None
+        if ws is not None:
+            await ws.close()
+
+
+# ---------------------------------------------------------------------------
+# Workflow registry
+# ---------------------------------------------------------------------------
+_workflows: dict[str, WorkflowConnection] = {}
+
+
+def _get(workflow: str) -> WorkflowConnection:
+    """Look up a workflow by name, raising a clear error if not found."""
+    conn = _workflows.get(workflow)
+    if conn is None:
+        available = ", ".join(_workflows.keys()) if _workflows else "(none)"
+        raise ValueError(
+            f"No workflow named '{workflow}'. Connected workflows: {available}. "
+            f"Use the connect tool first."
+        )
+    return conn
 
 
 # ---------------------------------------------------------------------------
@@ -118,16 +155,81 @@ async def _send(obj: dict) -> None:
 
 
 @mcp.tool()
-async def list_devices() -> str:
+async def connect(port: int = 0, pid: int = 0, name: str = "") -> str:
+    """Connect to a running DPL workflow.
+
+    Provide either ``port`` (the driver's WebSocket port) or ``pid`` (the
+    driver PID, port derived as 8080 + pid % 30000).  An optional ``name``
+    gives the workflow a human-friendly label; if omitted the port number is
+    used.
+
+    Args:
+        port: TCP port of the DPL driver status WebSocket.
+        pid:  PID of the DPL driver process (alternative to port).
+        name: Optional human-friendly name for this workflow.
+    """
+    if pid:
+        port = 8080 + pid % 30000
+    if not port:
+        return "Provide either port or pid."
+
+    wf_name = name or str(port)
+    if wf_name in _workflows:
+        old = _workflows[wf_name]
+        await old.close()
+
+    conn = WorkflowConnection(port, wf_name)
+    await conn.ensure_connected()
+    _workflows[wf_name] = conn
+
+    devices = conn.snapshot.get("devices", [])
+    return (
+        f"Connected to workflow '{wf_name}' on port {port} "
+        f"({len(devices)} device(s))."
+    )
+
+
+@mcp.tool()
+async def disconnect(workflow: str) -> str:
+    """Disconnect from a DPL workflow and release its resources.
+
+    Args:
+        workflow: Workflow name as returned by connect.
+    """
+    conn = _get(workflow)
+    await conn.close()
+    del _workflows[workflow]
+    return f"Disconnected from workflow '{workflow}'."
+
+
+@mcp.tool()
+async def list_workflows() -> str:
+    """List all currently connected DPL workflows."""
+    if not _workflows:
+        return "No workflows connected. Use the connect tool first."
+    lines = []
+    for wf_name, conn in _workflows.items():
+        n = len(conn.snapshot.get("devices", []))
+        status = "connected" if conn.ws is not None else "disconnected"
+        lines.append(f"{wf_name}: port={conn.port} devices={n} status={status}")
+    return "\n".join(lines)
+
+
+@mcp.tool()
+async def list_devices(workflow: str) -> str:
     """List all DPL devices with their current status.
 
     Returns each device's name, PID, active flag, streaming state, and device
     state as reported by the driver snapshot.
+
+    Args:
+        workflow: Workflow name as returned by connect.
     """
-    await _ensure_connected()
-    if not _snapshot:
-        return "No snapshot received yet — the driver may still be starting."
-    devices = _snapshot.get("devices", [])
+    conn = _get(workflow)
+    await conn.ensure_connected()
+    if not conn.snapshot:
+        return "No snapshot received yet -- the driver may still be starting."
+    devices = conn.snapshot.get("devices", [])
     if not devices:
         return "No devices in snapshot."
     lines = []
@@ -140,7 +242,7 @@ async def list_devices() -> str:
 
 
 @mcp.tool()
-async def list_metrics(device: str) -> str:
+async def list_metrics(workflow: str, device: str) -> str:
     """List the available numeric metrics for a DPL device.
 
     Sends a list_metrics command to the driver and waits up to 3 seconds for
@@ -148,15 +250,16 @@ async def list_metrics(device: str) -> str:
     and enum metrics are excluded.
 
     Args:
+        workflow: Workflow name as returned by connect.
         device: Device name exactly as shown by list_devices.
     """
-    # Remove any stale cached result so we can detect the fresh reply.
-    _metrics_lists.pop(device, None)
-    await _send({"cmd": "list_metrics", "device": device})
+    conn = _get(workflow)
+    conn.metrics_lists.pop(device, None)
+    await conn.send({"cmd": "list_metrics", "device": device})
     for _ in range(60):          # up to 3 s
         await asyncio.sleep(0.05)
-        if device in _metrics_lists:
-            names = _metrics_lists[device]
+        if device in conn.metrics_lists:
+            names = conn.metrics_lists[device]
             if not names:
                 return f"Device '{device}' has no numeric metrics yet."
             return f"{len(names)} metric(s): " + ", ".join(names)
@@ -164,7 +267,7 @@ async def list_metrics(device: str) -> str:
 
 
 @mcp.tool()
-async def subscribe(device: str, metrics: list[str]) -> str:
+async def subscribe(workflow: str, device: str, metrics: list[str]) -> str:
     """Subscribe to one or more metrics for a DPL device.
 
     After subscribing, the driver will push update frames for the device
@@ -172,60 +275,70 @@ async def subscribe(device: str, metrics: list[str]) -> str:
     the buffer.
 
     Args:
+        workflow: Workflow name as returned by connect.
         device: Device name exactly as shown by list_devices.
         metrics: List of metric names to subscribe to (from list_metrics).
     """
-    await _send({"cmd": "subscribe", "device": device, "metrics": metrics})
+    conn = _get(workflow)
+    await conn.send({"cmd": "subscribe", "device": device, "metrics": metrics})
     return f"Subscribed to {len(metrics)} metric(s) for '{device}': {', '.join(metrics)}"
 
 
 @mcp.tool()
-async def unsubscribe(device: str, metrics: list[str]) -> str:
+async def unsubscribe(workflow: str, device: str, metrics: list[str]) -> str:
     """Stop receiving updates for specific metrics of a DPL device.
 
     Args:
+        workflow: Workflow name as returned by connect.
         device: Device name exactly as shown by list_devices.
         metrics: List of metric names to unsubscribe from.
     """
-    await _send({"cmd": "unsubscribe", "device": device, "metrics": metrics})
+    conn = _get(workflow)
+    await conn.send({"cmd": "unsubscribe", "device": device, "metrics": metrics})
     return f"Unsubscribed from {len(metrics)} metric(s) for '{device}'."
 
 
 @mcp.tool()
-async def subscribe_logs(device: str) -> str:
+async def subscribe_logs(workflow: str, device: str) -> str:
     """Subscribe to log output for a DPL device.
 
     After subscribing, new log lines from the device will be buffered and
     can be retrieved with get_logs().
 
     Args:
+        workflow: Workflow name as returned by connect.
         device: Device name exactly as shown by list_devices.
     """
-    await _send({"cmd": "subscribe_logs", "device": device})
+    conn = _get(workflow)
+    await conn.send({"cmd": "subscribe_logs", "device": device})
     return f"Subscribed to logs for '{device}'."
 
 
 @mcp.tool()
-async def unsubscribe_logs(device: str) -> str:
+async def unsubscribe_logs(workflow: str, device: str) -> str:
     """Stop receiving log output for a DPL device.
 
     Args:
+        workflow: Workflow name as returned by connect.
         device: Device name exactly as shown by list_devices.
     """
-    await _send({"cmd": "unsubscribe_logs", "device": device})
+    conn = _get(workflow)
+    await conn.send({"cmd": "unsubscribe_logs", "device": device})
     return f"Unsubscribed from logs for '{device}'."
 
 
 @mcp.tool()
-async def get_logs(max_lines: int = 100) -> str:
+async def get_logs(workflow: str, max_lines: int = 100) -> str:
     """Drain and return buffered log lines received since the last call.
 
     Args:
+        workflow: Workflow name as returned by connect.
         max_lines: Maximum number of log lines to return (default 100).
     """
-    await _ensure_connected()
-    batch = _logs[:max_lines]
-    del _logs[:max_lines]
+    conn = _get(workflow)
+    await conn.ensure_connected()
+    batch = conn.logs[:max_lines]
+    del conn.logs[:max_lines]
     if not batch:
         return "No buffered log lines."
     lines = []
@@ -238,17 +351,21 @@ async def get_logs(max_lines: int = 100) -> str:
 
 
 @mcp.tool()
-async def start_devices() -> str:
+async def start_devices(workflow: str) -> str:
     """Resume all stopped DPL devices (send SIGCONT).
 
     Use this when the workflow was started with -s (all devices paused).
+
+    Args:
+        workflow: Workflow name as returned by connect.
     """
-    await _send({"cmd": "start_devices"})
+    conn = _get(workflow)
+    await conn.send({"cmd": "start_devices"})
     return "Sent SIGCONT to all active devices."
 
 
 @mcp.tool()
-async def enable_signpost(device: str, streams: list[str]) -> str:
+async def enable_signpost(workflow: str, device: str, streams: list[str]) -> str:
     """Enable one or more signpost log streams for a DPL device.
 
     Signpost streams produce detailed trace output visible in the device logs.
@@ -259,27 +376,31 @@ async def enable_signpost(device: str, streams: list[str]) -> str:
     ch.cern.aliceo2.data_processor_context, ch.cern.aliceo2.stream_context.
 
     Args:
+        workflow: Workflow name as returned by connect.
         device: Device name as shown by list_devices, or "" for the driver.
         streams: List of full signpost log names to enable.
     """
-    await _send({"cmd": "enable_signpost", "device": device, "streams": streams})
+    conn = _get(workflow)
+    await conn.send({"cmd": "enable_signpost", "device": device, "streams": streams})
     return f"Enabled {len(streams)} signpost stream(s) for '{device or 'driver'}': {', '.join(streams)}"
 
 
 @mcp.tool()
-async def disable_signpost(device: str, streams: list[str]) -> str:
+async def disable_signpost(workflow: str, device: str, streams: list[str]) -> str:
     """Disable one or more signpost log streams for a DPL device.
 
     Args:
+        workflow: Workflow name as returned by connect.
         device: Device name as shown by list_devices, or "" for the driver.
         streams: List of full signpost log names to disable.
     """
-    await _send({"cmd": "disable_signpost", "device": device, "streams": streams})
+    conn = _get(workflow)
+    await conn.send({"cmd": "disable_signpost", "device": device, "streams": streams})
     return f"Disabled {len(streams)} signpost stream(s) for '{device or 'driver'}': {', '.join(streams)}"
 
 
 @mcp.tool()
-async def get_updates(max_updates: int = 50) -> str:
+async def get_updates(workflow: str, max_updates: int = 50) -> str:
     """Drain and return buffered metric update frames received since the last call.
 
     Each frame contains the latest values of all subscribed metrics that
@@ -287,11 +408,13 @@ async def get_updates(max_updates: int = 50) -> str:
     time-ordered view of metric evolution.
 
     Args:
+        workflow: Workflow name as returned by connect.
         max_updates: Maximum number of update frames to return (default 50).
     """
-    await _ensure_connected()
-    batch = _updates[:max_updates]
-    del _updates[:max_updates]
+    conn = _get(workflow)
+    await conn.ensure_connected()
+    batch = conn.updates[:max_updates]
+    del conn.updates[:max_updates]
     if not batch:
         return "No buffered updates."
     lines = []
@@ -310,34 +433,6 @@ async def get_updates(max_updates: int = 50) -> str:
 # Entry point
 # ---------------------------------------------------------------------------
 def main() -> None:
-    global _port
-
-    parser = argparse.ArgumentParser(
-        description="DPL status MCP server — expose DPL driver metrics via MCP tools"
-    )
-    group = parser.add_mutually_exclusive_group()
-    group.add_argument(
-        "--port",
-        type=int,
-        default=None,
-        help="TCP port of the DPL driver status WebSocket (default: 8080 or DPL_STATUS_PORT env var)",
-    )
-    group.add_argument(
-        "--pid",
-        type=int,
-        default=None,
-        help="PID of the DPL driver process; port is derived as 8080 + pid %% 30000",
-    )
-    args = parser.parse_args()
-
-    if args.pid is not None:
-        _port = 8080 + args.pid % 30000
-    elif args.port is not None:
-        _port = args.port
-    elif "DPL_STATUS_PORT" in os.environ:
-        _port = int(os.environ["DPL_STATUS_PORT"])
-    # else leave _port at the default 8080
-
     mcp.run()
 
 

From 3ee5c9fe4f3e8b26cab7c37198fab29582d6019b Mon Sep 17 00:00:00 2001
From: mcoquet642 <74600025+mcoquet642@users.noreply.github.com>
Date: Tue, 19 May 2026 22:05:49 +0200
Subject: [PATCH 02/17] [MFT] Fixing number of links per RU/zone in mapping
 (#15410)

Co-authored-by: Maurice Coquet <mcoquet@alicecerno2.cern.ch>
---
 .../ITSMFTReconstruction/ChipMappingMFT.h       | 17 ++++++++---------
 .../reconstruction/src/ChipMappingMFT.cxx       |  6 +++---
 2 files changed, 11 insertions(+), 12 deletions(-)

diff --git a/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/ChipMappingMFT.h b/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/ChipMappingMFT.h
index eee9bdbb6a4dc..63d37a25ffbc9 100644
--- a/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/ChipMappingMFT.h
+++ b/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/ChipMappingMFT.h
@@ -73,16 +73,15 @@ class ChipMappingMFT
   ///< total number of RUs
   static constexpr Int_t getNRUs() { return NRUs; }
 
-  ///< get FEEId of the RU (software id of the RU), read via given link
+  ///< get software id of the RU, from first 8 bits of FEEID (HW id of RU)
   uint8_t FEEId2RUSW(uint16_t hw) const { return mFEEId2RUSW[hw & 0xff]; }
 
-  ///< get HW id of the RU (software id of the RU)
+  ///< get FEEID, from software id of the RU and link number
   uint16_t RUSW2FEEId(uint16_t sw, uint16_t linkID = 0) const { return ((linkID << 8) + mRUInfo[sw].idHW); }
 
   ///< compose FEEid for given stave (ru) relative to layer and link, see documentation in the constructor
   uint16_t composeFEEId(uint16_t layer, uint16_t ruOnLayer, uint16_t link) const
   {
-    // only one link is used
     // ruOnLayer is 0, 1, 2, 3 for half = 0
     //              4, 5, 6, 7            1
     auto dhalf = std::div(ruOnLayer, 4);
@@ -114,7 +113,7 @@ class ChipMappingMFT
     face = (feeID >> 2) & 0x1;
   }
 
-  ///< get info on sw RU
+  ///< get info on sw RU corresponding to given FEEID
   const RUInfo* getRUInfoFEEId(Int_t feeID) const { return &mRUInfo[FEEId2RUSW(feeID)]; }
 
   ///< get number of chips served by single cable on given RU type
@@ -123,13 +122,13 @@ class ChipMappingMFT
     return ((0x1 << 7) + (cableHW & 0x1f));
   }
 
-  ///< convert HW cable ID to its position on the ActiveLanes word in the GBT.header for given RU type
+  ///< convert HW cable ID to its position on the ActiveLanes word in the GBT.header for given RU type (note: this position is equal to the HW cable ID)
   uint8_t cableHW2Pos(uint8_t ruType, uint8_t hwid) const { return mCableHW2Pos[ruType][hwid]; }
 
   ///< convert HW cable ID to SW ID for give RU type
   uint8_t cableHW2SW(uint8_t ruType, uint8_t hwid) const { return hwid < mCableHW2SW[ruType].size() ? mCableHW2SW[ruType][hwid] : 0xff; }
 
-  ///< convert cable iterator ID to its position on the ActiveLanes word in the GBT.header for given RU type
+  ///< convert cable iterator ID (i.e. chipOnModule) to its position on the ActiveLanes word in the GBT.header for given RU type (note: this position is equal to the HW cable ID)
   uint8_t cablePos(uint8_t ruType, uint8_t id) const { return mCablePos[ruType][id]; }
 
   ///< get chipID on module from chip global SW ID, cable SW ID and stave (RU) info
@@ -139,7 +138,7 @@ class ChipMappingMFT
     return 0xffff;
   }
 
-  ///< get chip global SW ID from chipID on module, cable SW ID and stave (RU) info
+  ///< get chip global SW ID from cable HW ID and stave (RU) info (note: chOnModuleHW is unused)
   uint16_t getGlobalChipID(uint16_t chOnModuleHW, int cableHW, const RUInfo& ruInfo) const
   {
     auto chipOnRU = cableHW2SW(ruInfo.ruType, cableHW);
@@ -393,11 +392,11 @@ class ChipMappingMFT
 
  private:
   Int_t invalid() const;
-  static constexpr Int_t NRUs = NLayers * NZonesPerLayer;
+  static constexpr Int_t NRUs = NLayers * NZonesPerLayer; // 10 layers * 8 zones per layer
   static constexpr Int_t NModules = 280;
   static constexpr Int_t NChipsInfo = 7 + 8 + 9 + 10 + 11 + 12 + 13 + 14 + 16 + 17 + 18 + 19 + 14;
   static constexpr Int_t NChipsPerCable = 1;
-  static constexpr Int_t NLinks = 1;
+  static constexpr Int_t NLinks = 3;
   static constexpr Int_t NConnectors = 5;
   static constexpr Int_t NMaxChipsPerLadder = 5;
   static constexpr Int_t NRUCables = 25;
diff --git a/Detectors/ITSMFT/common/reconstruction/src/ChipMappingMFT.cxx b/Detectors/ITSMFT/common/reconstruction/src/ChipMappingMFT.cxx
index de2358469e894..b79c529bef803 100644
--- a/Detectors/ITSMFT/common/reconstruction/src/ChipMappingMFT.cxx
+++ b/Detectors/ITSMFT/common/reconstruction/src/ChipMappingMFT.cxx
@@ -1624,7 +1624,7 @@ ChipMappingMFT::ChipMappingMFT()
 {
   // init chips info
 
-  uint32_t maxRUHW = composeFEEId(NLayers - 1, NZonesPerLayer - 1, NLinks - 1); // Max possible FEE ID
+  uint32_t maxRUHW = composeFEEId(NLayers - 1, NZonesPerLayer - 1, 0); // Max possible RU HW ID (first 8 bits of max FEEID, while link stored in 9th and 10th bit of FEEID)
   mFEEId2RUSW.resize(maxRUHW + 1, 0xff);
 
   int curLayer = -1, curZone = -1, curHalf = -1;
@@ -1698,8 +1698,8 @@ ChipMappingMFT::ChipMappingMFT()
       auto& ruInfo = mRUInfo[ctrRU];
       ruInfo.idSW = ctrRU++;
 
-      // map FEEIds (RU read out by at most 3 GBT links) to SW ID
-      ruInfo.idHW = composeFEEId(iLayer, iZone, 0); // FEEId for link 0
+      // map RU HW ID (RU read out by at most 3 GBT links) to SW ID
+      ruInfo.idHW = composeFEEId(iLayer, iZone, 0); // RU HW ID (first 8 bits of FEEID)
       mFEEId2RUSW[ruInfo.idHW] = ruInfo.idSW;
       ruInfo.layer = iLayer;
       ruInfo.ruType = ZoneRUType[iZone % 4][iLayer / 2];

From 2e26434c440f559711b7a4e0aa576f0ab5ded134 Mon Sep 17 00:00:00 2001
From: shahoian <ruben.shahoyan@cern.ch>
Date: Tue, 19 May 2026 17:00:48 +0200
Subject: [PATCH 03/17] Fixes for ROFs downscaling for ITS tracking

---
 .../ITSMFT/ITS/tracking/src/FastMultEst.cxx   | 86 +++++--------------
 1 file changed, 21 insertions(+), 65 deletions(-)

diff --git a/Detectors/ITSMFT/ITS/tracking/src/FastMultEst.cxx b/Detectors/ITSMFT/ITS/tracking/src/FastMultEst.cxx
index cb831d7db71d0..cfbfdd8a9150e 100644
--- a/Detectors/ITSMFT/ITS/tracking/src/FastMultEst.cxx
+++ b/Detectors/ITSMFT/ITS/tracking/src/FastMultEst.cxx
@@ -168,85 +168,41 @@ int FastMultEst::selectROFs(const std::array<gsl::span<const o2::itsmft::ROFReco
                             ROFMaskTableN& sel)
 {
   const auto& multEstConf = FastMultEstConfig::Instance(); // parameters for mult estimation and cuts
-  const int selectionLayer = overlapView.getClock();
-  int multLayer = std::clamp(multEstConf.cutMultClusLayer, 0, NLayers - 1);
-  if (doStaggering && rofs[multLayer].empty()) {
-    LOGP(info, "FastMultEst multiplicity layer {} has no ROFs, falling back to selection layer {}", multLayer, selectionLayer);
-    multLayer = selectionLayer;
-  }
-
-  const auto multCounts = buildMultiplicityCounts<NLayers>(rofs, clus, doStaggering, multLayer);
+  const int selectionLayer = multEstConf.isMultCutRequested() ? std::clamp(multEstConf.cutMultClusLayer, 0, NLayers - 1) : overlapView.getClock();
+  const auto multCounts = buildMultiplicityCounts<NLayers>(rofs, clus, doStaggering, selectionLayer);
   const int selectionRofCount = doStaggering ? static_cast<int>(rofs[selectionLayer].size()) : static_cast<int>(rofs[0].size());
 
   sel.resetMask();
   lastRandomSeed = gRandom->GetSeed();
   const o2::InteractionRecord tfStartIR{0, firstTForbit};
-
-  if (!trig.empty()) {
+  // mask ROFs which are not good from the multiplicity selection (if any) POV
+  struct ROFStatus {
+    int entry = 0, priority = 0;
+  };
+  std::vector<ROFStatus> selROFs;
+  selROFs.reserve(selectionRofCount);
+  bool selmult = multEstConf.isMultCutRequested();
+  for (int selectionRof = 0; selectionRof < selectionRofCount; ++selectionRof) {
+    selROFs.emplace_back(selectionRof, (selmult && !multEstConf.isPassingMultCut(process(multCounts[selectionRof]))) ? -1 : 0);
+  }
+  if (!trig.empty() && multEstConf.preferTriggered) {
     const auto& selectionLayerTiming = overlapView.getLayer(selectionLayer);
-    const auto& multLayerTiming = overlapView.getLayer(multLayer);
-
     for (const auto& trigger : trig) {
       const int selectionRof = findROFForIR(trigger.ir, tfStartIR, selectionLayerTiming);
-      if (selectionRof < 0) {
-        continue;
-      }
-      if (multEstConf.cutRandomFraction > 0.f && gRandom->Rndm() < multEstConf.cutRandomFraction) {
-        continue;
-      }
-      if (multEstConf.isMultCutRequested()) {
-        const int triggerMultRof = doStaggering ? findROFForIR(trigger.ir, tfStartIR, multLayerTiming) : selectionRof;
-        if (triggerMultRof < 0 || triggerMultRof >= static_cast<int>(multCounts.size())) {
-          continue;
-        }
-        if (!multEstConf.isPassingMultCut(process(multCounts[triggerMultRof]))) {
-          continue;
-        }
-      }
-      enableCompatibleROFs<NLayers>(selectionLayer, selectionRof, overlapView, sel);
-    }
-  } else {
-    LOGP(info, "FastMultEst received no physics/TRD triggers, falling back to ROF-driven filtering on layer {}", selectionLayer);
-    for (int selectionRof = 0; selectionRof < selectionRofCount; ++selectionRof) {
-      if (multEstConf.isMultCutRequested()) {
-        bool passes = false;
-        if (!doStaggering || selectionLayer == multLayer) {
-          if (selectionRof < static_cast<int>(multCounts.size())) {
-            passes = multEstConf.isPassingMultCut(process(multCounts[selectionRof]));
-          }
-        } else {
-          const auto& overlap = overlapView.getOverlap(selectionLayer, multLayer, selectionRof);
-          for (int rof = overlap.getFirstEntry(); rof < overlap.getEntriesBound(); ++rof) {
-            if (rof < static_cast<int>(multCounts.size())) {
-              if (multEstConf.isPassingMultCut(process(multCounts[rof]))) {
-                passes = true;
-                break;
-              }
-            }
-          }
-        }
-        if (!passes) {
-          continue;
-        }
-      }
-      if (multEstConf.cutRandomFraction > 0.f && gRandom->Rndm() < multEstConf.cutRandomFraction) {
+      if (selectionRof < 0 || selROFs[selectionRof].priority < 0) {
         continue;
       }
-      enableCompatibleROFs<NLayers>(selectionLayer, selectionRof, overlapView, sel);
+      selROFs[selectionRof].priority++; // increment trigger counter
     }
+    sort(selROFs.begin(), selROFs.end(), [](const ROFStatus& a, const ROFStatus& b) { return a.priority > b.priority; }); // order in number of triggers, masked will go to the end
   }
-
-  const auto selView = sel.getView();
   int nsel = 0;
-  for (int irof = 0; irof < selectionRofCount; ++irof) {
-    nsel += selView.isROFEnabled(selectionLayer, irof);
-  }
-
-  if (!trig.empty() && multEstConf.preferTriggered) {
-    LOGP(debug, "FastMultEst preferTriggered is ignored in trigger-driven mask mode");
+  for (auto& rof : selROFs) {
+    if (rof.priority >= 0 && (multEstConf.cutRandomFraction <= 0.f || (gRandom->Rndm() > multEstConf.cutRandomFraction))) {
+      enableCompatibleROFs<NLayers>(selectionLayer, rof.entry, overlapView, sel);
+      nsel++;
+    }
   }
-
   LOGP(debug, "NSel = {} of {} rofs on layer {} Seeds: before {} after {}", nsel, selectionRofCount, selectionLayer, lastRandomSeed, gRandom->GetSeed());
-
   return nsel;
 }

From f8f5d1eb6c74aa9d5840407a082bef9af3f56555 Mon Sep 17 00:00:00 2001
From: Martin Eide <43970264+mrtineide@users.noreply.github.com>
Date: Wed, 20 May 2026 14:06:56 +0200
Subject: [PATCH 04/17] Add token for the test CCDB instance

Before the test instance did not need a token. This has changed, which
broke tests that relied on hardcoded HTTP, and not HTTPS.
Now the test instance redirects to HTTPS and we need a token.
---
 CCDB/src/CcdbApi.cxx | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/CCDB/src/CcdbApi.cxx b/CCDB/src/CcdbApi.cxx
index 42bc13904bf61..93a79ad56c477 100644
--- a/CCDB/src/CcdbApi.cxx
+++ b/CCDB/src/CcdbApi.cxx
@@ -213,8 +213,7 @@ void CcdbApi::init(std::string const& host)
     snapshotReport += ')';
   }
 
-  mNeedAlienToken = (host.find("https://") != std::string::npos) || (host.find("alice-ccdb.cern.ch") != std::string::npos);
-
+  mNeedAlienToken = (host.find("https://") != std::string::npos) || (host.find("alice-ccdb.cern.ch") != std::string::npos) || (host.find("ccdb-test.cern.ch") != std::string::npos);
   // Set the curl timeout. It can be forced with an env var or it has different defaults based on the deployment mode.
   if (getenv("ALICEO2_CCDB_CURL_TIMEOUT_DOWNLOAD")) {
     auto timeout = atoi(getenv("ALICEO2_CCDB_CURL_TIMEOUT_DOWNLOAD"));

From 084c3660c7bce5e8fa011fa0fee40a32d0d86f84 Mon Sep 17 00:00:00 2001
From: Matthias Kleiner <matthias.kleiner@cern.ch>
Date: Wed, 20 May 2026 15:23:30 +0200
Subject: [PATCH 05/17] TPC: add option for disabling corrections

- add option for using static corrections instead of main corrections
---
 .../calibration/src/CorrectionMapsLoader.cxx  | 13 +++-
 .../calibration/src/CorrectionMapsOptions.cxx |  4 +-
 Detectors/TPC/workflow/src/TPCScalerSpec.cxx  | 63 ++++++++++---------
 .../CorrectionMapsHelper.cxx                  | 43 ++++++++++++-
 .../CorrectionMapsTypes.h                     | 14 +++--
 5 files changed, 96 insertions(+), 41 deletions(-)

diff --git a/Detectors/TPC/calibration/src/CorrectionMapsLoader.cxx b/Detectors/TPC/calibration/src/CorrectionMapsLoader.cxx
index 9569e0eb8abd2..c8bdfa0f99350 100644
--- a/Detectors/TPC/calibration/src/CorrectionMapsLoader.cxx
+++ b/Detectors/TPC/calibration/src/CorrectionMapsLoader.cxx
@@ -28,8 +28,13 @@ using namespace o2::framework;
 void CorrectionMapsLoader::extractCCDBInputs(ProcessingContext& pc, float tpcScaler)
 {
   pc.inputs().get<o2::tpc::CorrMapParam*>("tpcCorrPar");
-  pc.inputs().get<o2::gpu::TPCFastTransform*>("tpcCorrMap");
-  pc.inputs().get<o2::gpu::TPCFastTransform*>("tpcCorrMapRef");
+  const auto lumiMode = getLumiScaleMode();
+  if (lumiMode != LumiScaleMode::NoCorrection && lumiMode != LumiScaleMode::StaticMapOnly) {
+    pc.inputs().get<o2::gpu::TPCFastTransform*>("tpcCorrMap");
+  }
+  if (lumiMode != LumiScaleMode::NoCorrection) {
+    pc.inputs().get<o2::gpu::TPCFastTransform*>("tpcCorrMapRef");
+  }
   const int maxDumRep = 5;
   int dumRep = 0;
   o2::ctp::LumiInfo lumiObj;
@@ -97,6 +102,10 @@ void CorrectionMapsLoader::requestCCDBInputs(std::vector<InputSpec>& inputs, con
     // for MC corrections
     addInput(inputs, {"tpcCorrMap", "TPC", "CorrMap", 0, Lifetime::Condition, ccdbParamSpec(CDBTypeMap.at(CDBType::CalCorrMapMC), {}, 1)});            // time-dependent
     addInput(inputs, {"tpcCorrMapRef", "TPC", "CorrMapRef", 0, Lifetime::Condition, ccdbParamSpec(CDBTypeMap.at(CDBType::CalCorrDerivMapMC), {}, 1)}); // time-dependent
+  } else if (gloOpts.lumiMode == LumiScaleMode::NoCorrection) {
+    // no correction maps needed — a dummy map is created at runtime
+  } else if (gloOpts.lumiMode == LumiScaleMode::StaticMapOnly) {
+    addInput(inputs, {"tpcCorrMapRef", "TPC", "CorrMapRef", 0, Lifetime::Condition, ccdbParamSpec(CDBTypeMap.at(CDBType::CalCorrMapRef), {}, 0)}); // load once
   } else {
     LOG(fatal) << "Correction mode unknown! Choose either 0 (default) or 1 (derivative map) for flag corrmap-lumi-mode.";
   }
diff --git a/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx b/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx
index 604b7c680385b..45c3771db57bf 100644
--- a/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx
+++ b/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx
@@ -21,13 +21,13 @@ CorrectionMapsGloOpts CorrectionMapsOptions::parseGlobalOptions(const o2::framew
 {
   CorrectionMapsGloOpts tpcopt;
   auto lumiTypeVal = opts.get<int>("lumi-type");
-  if (lumiTypeVal < -1 || lumiTypeVal > 2) {
+  if (lumiTypeVal < static_cast<int>(LumiScaleType::Unset) || lumiTypeVal >= static_cast<int>(LumiScaleType::Count)) {
     LOGP(fatal, "Invalid lumi-type value: {}", lumiTypeVal);
   }
   tpcopt.lumiType = static_cast<LumiScaleType>(lumiTypeVal);
 
   auto lumiModeVal = opts.get<int>("corrmap-lumi-mode");
-  if (lumiModeVal < -1 || lumiModeVal > 2) {
+  if (lumiModeVal < static_cast<int>(LumiScaleMode::Unset) || lumiModeVal >= static_cast<int>(LumiScaleMode::Count)) {
     LOGP(fatal, "Invalid corrmap-lumi-mode value: {}", lumiModeVal);
   }
   tpcopt.lumiMode = static_cast<LumiScaleMode>(lumiModeVal);
diff --git a/Detectors/TPC/workflow/src/TPCScalerSpec.cxx b/Detectors/TPC/workflow/src/TPCScalerSpec.cxx
index 8e2a78d69757b..1df192dd5ec00 100644
--- a/Detectors/TPC/workflow/src/TPCScalerSpec.cxx
+++ b/Detectors/TPC/workflow/src/TPCScalerSpec.cxx
@@ -183,40 +183,45 @@ class TPCScalerSpec : public Task
 
   void buildMap(ProcessingContext& pc)
   {
-    // reference map
-    auto* corrMap = mTPCCorrMapsLoader.getCorrMap();
-
-    // // new correction map
+    const auto lumiMode = mTPCCorrMapsLoader.getLumiScaleMode();
     o2::gpu::TPCFastTransform finalMap;
-    finalMap.cloneFromObject(*corrMap, nullptr);
-    finalMap.setApplyCorrectionOn();
-
-    const auto* corrMapRef = mTPCCorrMapsLoader.getCorrMapRef();
-    const float lumiScale = mTPCCorrMapsLoader.getLumiScale();
     std::vector<std::pair<const o2::gpu::TPCFastSpaceChargeCorrection*, float>> additionalCorrections;
 
-    // if standard scaling is used: map(lumi) = (mean_map - ref_map) * lumiScale + ref_map
-    if (mTPCCorrMapsLoader.getLumiScaleMode() == LumiScaleMode::Linear) {
-      const std::vector<std::pair<const o2::gpu::TPCFastSpaceChargeCorrection*, float>> step0{{&(corrMapRef->getCorrection()), -1.f}};
-      // finalMap = (mean_map - finalMap)
-      TPCFastSpaceChargeCorrectionHelper::instance()->mergeCorrections(finalMap.getCorrection(), 1, step0, true);
-
-      // finalMap = finalMap * lumiScale + ref_map
-      const std::vector<std::pair<const o2::gpu::TPCFastSpaceChargeCorrection*, float>> step1{{&(corrMapRef->getCorrection()), 1.f}};
-      TPCFastSpaceChargeCorrectionHelper::instance()->mergeCorrections(finalMap.getCorrection(), lumiScale, step1, true);
-
-    } else if (mTPCCorrMapsLoader.getLumiScaleMode() == LumiScaleMode::DerivativeMap || mTPCCorrMapsLoader.getLumiScaleMode() == LumiScaleMode::DerivativeMapMC) {
-      additionalCorrections.emplace_back(&(corrMapRef->getCorrection()), lumiScale);
-    }
+    if (lumiMode == LumiScaleMode::NoCorrection) {
+      std::unique_ptr<o2::gpu::TPCFastTransform> dummy(TPCFastTransformHelperO2::instance()->create(0));
+      finalMap.cloneFromObject(*dummy, nullptr);
+      finalMap.setApplyCorrectionOff();
+    } else {
+      auto* corrMap = mTPCCorrMapsLoader.getCorrMap();
+      const auto* corrMapRef = mTPCCorrMapsLoader.getCorrMapRef();
+      finalMap.cloneFromObject(lumiMode == LumiScaleMode::StaticMapOnly && corrMapRef ? *corrMapRef : *corrMap, nullptr);
+      finalMap.setApplyCorrectionOn();
+
+      const float lumiScale = mTPCCorrMapsLoader.getLumiScale();
+
+      // if standard scaling is used: map(lumi) = (mean_map - ref_map) * lumiScale + ref_map
+      if (lumiMode == LumiScaleMode::Linear) {
+        const std::vector<std::pair<const o2::gpu::TPCFastSpaceChargeCorrection*, float>> step0{{&(corrMapRef->getCorrection()), -1.f}};
+        // finalMap = (mean_map - finalMap)
+        TPCFastSpaceChargeCorrectionHelper::instance()->mergeCorrections(finalMap.getCorrection(), 1, step0, true);
+
+        // finalMap = finalMap * lumiScale + ref_map
+        const std::vector<std::pair<const o2::gpu::TPCFastSpaceChargeCorrection*, float>> step1{{&(corrMapRef->getCorrection()), 1.f}};
+        TPCFastSpaceChargeCorrectionHelper::instance()->mergeCorrections(finalMap.getCorrection(), lumiScale, step1, true);
+
+      } else if (lumiMode == LumiScaleMode::DerivativeMap || lumiMode == LumiScaleMode::DerivativeMapMC) {
+        additionalCorrections.emplace_back(&(corrMapRef->getCorrection()), lumiScale);
+      }
 
-    // if mshape map valid
-    if (!mTPCCorrMapsLoader.isCorrMapMShapeDummy()) {
-      LOGP(info, "Adding M-shape correction to the final map with scaling factor {}", mMShapeScalingFac);
-      additionalCorrections.emplace_back(&(mTPCCorrMapsLoader.getCorrMapMShape()->getCorrection()), 1.f);
-    }
+      // if mshape map valid
+      if (!mTPCCorrMapsLoader.isCorrMapMShapeDummy()) {
+        LOGP(info, "Adding M-shape correction to the final map with scaling factor {}", mMShapeScalingFac);
+        additionalCorrections.emplace_back(&(mTPCCorrMapsLoader.getCorrMapMShape()->getCorrection()), 1.f);
+      }
 
-    if (!additionalCorrections.empty()) {
-      TPCFastSpaceChargeCorrectionHelper::instance()->mergeCorrections(finalMap.getCorrection(), 1, additionalCorrections, true);
+      if (!additionalCorrections.empty()) {
+        TPCFastSpaceChargeCorrectionHelper::instance()->mergeCorrections(finalMap.getCorrection(), 1, additionalCorrections, true);
+      }
     }
 
     Output corrMapOutput{header::gDataOriginTPC, "TPCCORRMAP", 0};
diff --git a/GPU/TPCFastTransformation/CorrectionMapsHelper.cxx b/GPU/TPCFastTransformation/CorrectionMapsHelper.cxx
index 7f7deddafe1c8..4bfedc117dec7 100644
--- a/GPU/TPCFastTransformation/CorrectionMapsHelper.cxx
+++ b/GPU/TPCFastTransformation/CorrectionMapsHelper.cxx
@@ -38,6 +38,10 @@ void CorrectionMapsHelper::setCorrMapMShape(std::unique_ptr<TPCFastTransform>&&
 void CorrectionMapsHelper::updateLumiScale(bool report)
 {
   if (!canUseCorrections()) {
+    if (mLumiScaleMode != LumiScaleMode::NoCorrection) {
+      LOGP(warning, "Negative meanLumi={} detected, switching to NoCorrection mode for backward compatibility", mMeanLumi);
+      mLumiScaleMode = LumiScaleMode::NoCorrection;
+    }
     mLumiScale = -1.f;
   } else if ((mLumiScaleMode == LumiScaleMode::DerivativeMap) || (mLumiScaleMode == LumiScaleMode::DerivativeMapMC)) {
     mLumiScale = mMeanLumiRef ? (mInstLumi - mMeanLumi) / mMeanLumiRef : 0.f;
@@ -54,7 +58,40 @@ void CorrectionMapsHelper::updateLumiScale(bool report)
 //________________________________________________________
 void CorrectionMapsHelper::reportScaling()
 {
-  LOGP(info, "Map scaling update: LumiScaleType={} instLumi(CTP)={} instLumi(scaling)={} meanLumiRef={}, meanLumi={} -> LumiScale={} lumiScaleMode={}, M-Shape map valid: {}, M-Shape default: {}",
-       mLumiScaleType == LumiScaleType::NoScaling ? "NoScaling" : (mLumiScaleType == LumiScaleType::CTPLumi ? "LumiCTP" : "TPCScaler"), getInstLumiCTP(), getInstLumi(), getMeanLumiRef(), getMeanLumi(), getLumiScale(),
-       mLumiScaleMode == LumiScaleMode::Linear ? "Linear" : "Derivative", (mCorrMapMShape != nullptr), isCorrMapMShapeDummy());
+  auto lumiTypeName = [](LumiScaleType t) {
+    switch (t) {
+      case LumiScaleType::NoScaling:
+        return "NoScaling";
+      case LumiScaleType::CTPLumi:
+        return "CTPLumi";
+      case LumiScaleType::TPCScaler:
+        return "TPCScaler";
+      default:
+        return "Unknown";
+    }
+  };
+
+  const bool mshapeValid = (mCorrMapMShape != nullptr) && !isCorrMapMShapeDummy();
+
+  if (mLumiScaleMode == LumiScaleMode::NoCorrection) {
+    LOGP(info, "Map scaling update: mode=NoCorrection (corrections disabled, dummy map in use)");
+  } else if (mLumiScaleMode == LumiScaleMode::StaticMapOnly) {
+    LOGP(info, "Map scaling update: mode=StaticMapOnly (static reference map, no lumi scaling), M-Shape correction: {}", mshapeValid ? "applied" : "not applied");
+  } else {
+    auto lumiModeName = [](LumiScaleMode m) {
+      switch (m) {
+        case LumiScaleMode::Linear:
+          return "Linear";
+        case LumiScaleMode::DerivativeMap:
+          return "DerivativeMap";
+        case LumiScaleMode::DerivativeMapMC:
+          return "DerivativeMapMC";
+        default:
+          return "Unknown";
+      }
+    };
+    LOGP(info, "Map scaling update: LumiScaleType={} instLumi(CTP)={} instLumi(scaling)={} meanLumiRef={} meanLumi={} -> LumiScale={} lumiScaleMode={}, M-Shape correction: {}",
+         lumiTypeName(mLumiScaleType), getInstLumiCTP(), getInstLumi(), getMeanLumiRef(), getMeanLumi(), getLumiScale(),
+         lumiModeName(mLumiScaleMode), mshapeValid ? "applied" : "not applied");
+  }
 }
diff --git a/GPU/TPCFastTransformation/CorrectionMapsTypes.h b/GPU/TPCFastTransformation/CorrectionMapsTypes.h
index e239b668ab751..092a2927ebe3e 100644
--- a/GPU/TPCFastTransformation/CorrectionMapsTypes.h
+++ b/GPU/TPCFastTransformation/CorrectionMapsTypes.h
@@ -22,14 +22,18 @@ enum class LumiScaleType : int {
   Unset = -1,    ///< init value
   NoScaling = 0, ///< no scaling, use map as is
   CTPLumi = 1,   ///< use CTP luminosity for scaling
-  TPCScaler = 2  ///< use TPC scaler for scaling
+  TPCScaler = 2, ///< use TPC scaler for scaling
+  Count          ///< sentinel - keep last
 };
 
 enum class LumiScaleMode : int {
-  Unset = -1,         ///< init value
-  Linear = 0,         ///< map(lumi) = (mean_map - referenceMap) * lumiScale + referenceMap
-  DerivativeMap = 1,  ///< map(lumi) = mean_map + lumiScale * (derivativeMap) where derivativeMap = (mean_map_A - mean_map_B)
-  DerivativeMapMC = 2 ///< same DerivativeMap, but for MC
+  Unset = -1,          ///< init value
+  Linear = 0,          ///< map(lumi) = (mean_map - referenceMap) * lumiScale + referenceMap
+  DerivativeMap = 1,   ///< map(lumi) = mean_map + lumiScale * (derivativeMap) where derivativeMap = (mean_map_A - mean_map_B)
+  DerivativeMapMC = 2, ///< same DerivativeMap, but for MC
+  NoCorrection = 3,    ///< no corrections at all
+  StaticMapOnly = 4,   ///< use only static map instead of main map
+  Count                ///< sentinel - keep last
 };
 
 struct CorrectionMapsGloOpts {

From aa96c1a9fd301f775c2ab3a04d49fe890c03782d Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Nicol=C3=B2=20Jacazio?= <njacazio@users.noreply.github.com>
Date: Thu, 21 May 2026 12:00:34 +0200
Subject: [PATCH 06/17] IOTOF: align geometry to specs (#15414)

- add macro to draw geometry
- streamline setup of IOTOF active layers
---
 .../ALICE3/IOTOF/macros/CMakeLists.txt        |  3 +
 .../ALICE3/IOTOF/macros/drawTOFGeometry.C     | 90 +++++++++++++++++++
 .../ALICE3/IOTOF/simulation/src/Detector.cxx  | 29 ++++--
 .../ALICE3/IOTOF/simulation/src/Layer.cxx     | 53 ++++++++---
 4 files changed, 156 insertions(+), 19 deletions(-)
 create mode 100644 Detectors/Upgrades/ALICE3/IOTOF/macros/drawTOFGeometry.C

diff --git a/Detectors/Upgrades/ALICE3/IOTOF/macros/CMakeLists.txt b/Detectors/Upgrades/ALICE3/IOTOF/macros/CMakeLists.txt
index b2f1857186c0b..41b800ed114b4 100644
--- a/Detectors/Upgrades/ALICE3/IOTOF/macros/CMakeLists.txt
+++ b/Detectors/Upgrades/ALICE3/IOTOF/macros/CMakeLists.txt
@@ -11,3 +11,6 @@
 
 o2_add_test_root_macro(defineIOTOFGeo.C
                        LABELS alice3)
+
+o2_add_test_root_macro(drawTOFGeometry.C
+                       LABELS alice3)
diff --git a/Detectors/Upgrades/ALICE3/IOTOF/macros/drawTOFGeometry.C b/Detectors/Upgrades/ALICE3/IOTOF/macros/drawTOFGeometry.C
new file mode 100644
index 0000000000000..4e58fb54fbf6e
--- /dev/null
+++ b/Detectors/Upgrades/ALICE3/IOTOF/macros/drawTOFGeometry.C
@@ -0,0 +1,90 @@
+// 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.
+
+#include "IOTOFBase/GeometryTGeo.h"
+#include "IOTOFSimulation/Layer.h"
+
+#include <TCanvas.h>
+#include <TGeoManager.h>
+#include <TGeoMaterial.h>
+#include <TGeoMatrix.h>
+#include <TGeoVolume.h>
+#include <TStyle.h>
+
+#include <iostream>
+
+namespace
+{
+void ensureMedium(const char* name, int id, double a, double z, double density)
+{
+  if (!gGeoManager->GetMedium(name)) {
+    auto* mat = new TGeoMaterial(name, a, z, density);
+    new TGeoMedium(name, id, mat);
+  }
+}
+
+void prepareMinimalMedia()
+{
+  ensureMedium("VACUUM$", 0, 1., 1., 1.e-16);
+  ensureMedium("TF3_AIR$", 1, 14.61, 7.3, 1.20479e-3);
+  ensureMedium("TF3_SILICON$", 3, 28.086, 14., 2.33);
+}
+} // namespace
+
+void drawTOFGeometry(double x2x0 = 0.02,
+                     double sensorThickness = 0.005,
+                     bool checkOverlaps = true,
+                     double overlapToleranceCm = 0.01)
+{
+  gStyle->SetOptStat(0);
+
+  if (gGeoManager) {
+    delete gGeoManager;
+  }
+
+  auto* geo = new TGeoManager("IOTOFGeomFromLayer", "Geometry built from Layer.h classes");
+  prepareMinimalMedia();
+
+  auto* top = geo->MakeBox("TOP", geo->GetMedium("VACUUM$"), 1200., 1200., 1200.);
+  geo->SetTopVolume(top);
+
+  auto* mother = new TGeoVolumeAssembly("IOTOFMacroVol");
+  top->AddNode(mother, 1, new TGeoTranslation(0., 0., 0.));
+
+  // Build using the same classes and createLayer() used by detector geometry code.
+  o2::iotof::ITOFLayer itof(o2::iotof::GeometryTGeo::getITOFLayerPattern(),
+                            21.f, 0.f, 129.f, 0.f, x2x0,
+                            o2::iotof::Layer::kBarrelSegmented,
+                            24, 5.42, 3.0, 10, sensorThickness);
+
+  o2::iotof::OTOFLayer otof(o2::iotof::GeometryTGeo::getOTOFLayerPattern(),
+                            92.f, 0.f, 680.f, 0.f, x2x0,
+                            o2::iotof::Layer::kBarrelSegmented,
+                            62, 9.74, 5.0, 54, sensorThickness);
+
+  itof.createLayer(mother);
+  otof.createLayer(mother);
+
+  geo->CloseGeometry();
+
+  std::cout << "Built geometry from Layer.h classes with x2x0=" << x2x0
+            << " and sensorThickness=" << sensorThickness << " cm\n";
+  std::cout << "ITOF sensitive volumes: " << o2::iotof::ITOFLayer::mRegister.size() << "\n";
+  std::cout << "OTOF sensitive volumes: " << o2::iotof::OTOFLayer::mRegister.size() << "\n";
+
+  if (checkOverlaps) {
+    std::cout << "Checking overlaps with tolerance=" << overlapToleranceCm << " cm\n";
+    geo->CheckOverlaps(overlapToleranceCm);
+    geo->PrintOverlaps();
+  }
+
+  top->Draw("ogl");
+}
diff --git a/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Detector.cxx b/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Detector.cxx
index bed8cbfd6dfac..ab9a68bd401ec 100644
--- a/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Detector.cxx
+++ b/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Detector.cxx
@@ -200,28 +200,47 @@ void Detector::defineSensitiveVolumes()
   TGeoManager* geoManager = gGeoManager;
   TGeoVolume* v;
 
-  // The names of the IOTOF sensitive volumes have the format: IOTOFLayer(0...mLayers.size()-1)
   auto& iotofPars = IOTOFBaseParam::Instance();
-  if (iotofPars.enableInnerTOF) {
+  const bool itof = iotofPars.enableInnerTOF;
+  const bool otof = iotofPars.enableOuterTOF;
+  bool ftof = iotofPars.enableForwardTOF;
+  bool btof = iotofPars.enableBackwardTOF;
+  const std::string pattern = iotofPars.detectorPattern;
+  if (pattern == "") {
+    LOG(info) << "Default pattern";
+  } else if (pattern == "v3b") {
+    ftof = false;
+    btof = false;
+  } else if (pattern == "v3b1a") {
+  } else if (pattern == "v3b1b") {
+  } else if (pattern == "v3b2a") {
+  } else if (pattern == "v3b2b") {
+  } else if (pattern == "v3b3") {
+  } else {
+    LOG(fatal) << "IOTOF layer pattern " << pattern << " not recognized, exiting";
+  }
+
+  // The names of the IOTOF sensitive volumes have the format: IOTOFLayer(0...mLayers.size()-1)
+  if (itof) {
     for (const std::string& itofSensor : ITOFLayer::mRegister) {
       v = geoManager->GetVolume(itofSensor.c_str());
       LOGP(info, "Adding IOTOF Sensitive Volume {}", v->GetName());
       AddSensitiveVolume(v);
     }
   }
-  if (iotofPars.enableOuterTOF) {
+  if (otof) {
     for (const std::string& otofSensor : OTOFLayer::mRegister) {
       v = geoManager->GetVolume(otofSensor.c_str());
       LOGP(info, "Adding IOTOF Sensitive Volume {}", v->GetName());
       AddSensitiveVolume(v);
     }
   }
-  if (iotofPars.enableForwardTOF) {
+  if (ftof) {
     v = geoManager->GetVolume(GeometryTGeo::getFTOFSensorPattern());
     LOGP(info, "Adding IOTOF Sensitive Volume {}", v->GetName());
     AddSensitiveVolume(v);
   }
-  if (iotofPars.enableBackwardTOF) {
+  if (btof) {
     v = geoManager->GetVolume(GeometryTGeo::getBTOFSensorPattern());
     LOGP(info, "Adding IOTOF Sensitive Volume {}", v->GetName());
     AddSensitiveVolume(v);
diff --git a/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Layer.cxx b/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Layer.cxx
index 627fb599ff8ae..f2e42e1bce172 100644
--- a/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Layer.cxx
+++ b/Detectors/Upgrades/ALICE3/IOTOF/simulation/src/Layer.cxx
@@ -296,12 +296,24 @@ void OTOFLayer::createLayer(TGeoVolume* motherVolume)
     case kBarrelSegmented: {
       // First we create the volume for the whole layer, which will be used as mother volume for the segments
       const double avgRadius = 0.5 * (mInnerRadius + mOuterRadius);
-      const double staveSizeX = mStaves.second;                                                                                                                    // cm
-      const double staveSizeY = mOuterRadius - mInnerRadius;                                                                                                       // cm
-      const double staveSizeZ = mZLength;                                                                                                                          // cm
-      const double deltaForTilt = 0.5 * (std::sin(TMath::DegToRad() * mTiltAngle) * staveSizeX + std::cos(TMath::DegToRad() * mTiltAngle) * staveSizeY);           // we increase the size of the layer to account for the tilt of the staves
-      const double radiusMax = std::sqrt(avgRadius * avgRadius + 0.25 * staveSizeX * staveSizeX + 0.25 * staveSizeY * staveSizeY + avgRadius * 2. * deltaForTilt); // we increase the outer radius to account for the tilt of the staves
-      const double radiusMin = std::sqrt(avgRadius * avgRadius + 0.25 * staveSizeX * staveSizeX + 0.25 * staveSizeY * staveSizeY - avgRadius * 2. * deltaForTilt); // we decrease the inner radius to account for the tilt of the staves
+      const double staveSizeX = mStaves.second;              // cm, tangential stave size
+      const double staveSizeY = mOuterRadius - mInnerRadius; // cm, radial stave size
+      const double staveSizeZ = mZLength;                    // cm
+
+      // Build the mother layer tube from the exact inscribed/outscribed radii of a tilted stave rectangle.
+      const double alpha = mTiltAngle * TMath::DegToRad();
+      const double u0 = -avgRadius * std::cos(alpha);
+      const double v0 = avgRadius * std::sin(alpha);
+      const double uClamped = std::max(-0.5 * staveSizeY, std::min(0.5 * staveSizeY, u0));
+      const double vClamped = std::max(-0.5 * staveSizeX, std::min(0.5 * staveSizeX, v0));
+      const double radiusMin = std::hypot(uClamped - u0, vClamped - v0);
+
+      const double uCorners[4] = {-0.5 * staveSizeY, 0.5 * staveSizeY, 0.5 * staveSizeY, -0.5 * staveSizeY};
+      const double vCorners[4] = {-0.5 * staveSizeX, -0.5 * staveSizeX, 0.5 * staveSizeX, 0.5 * staveSizeX};
+      double radiusMax = 0.0;
+      for (int i = 0; i < 4; ++i) {
+        radiusMax = std::max(radiusMax, std::hypot(uCorners[i] - u0, vCorners[i] - v0));
+      }
       TGeoTube* layer = new TGeoTube(radiusMin, radiusMax, mZLength / 2);
       TGeoVolume* layerVol = new TGeoVolume(mLayerName.c_str(), layer, medAir);
       setLayerStyle(layerVol);
@@ -312,10 +324,21 @@ void OTOFLayer::createLayer(TGeoVolume* motherVolume)
       setStaveStyle(staveVol);
 
       // Now we create the volume for a single module (sensor + chip)
-      const int modulesPerStaveX = 1;                           // we assume that each stave is divided in 2 modules along the x direction
-      const double moduleSizeX = staveSizeX / modulesPerStaveX; // cm
-      const double moduleSizeY = staveSizeY;                    // cm
-      const double moduleSizeZ = staveSizeZ / mModulesPerStave; // cm
+      // oTOF V2 is a 2xN matrix of modules per stave with overlap along z.
+      const int modulesPerStaveX = 2;
+      if (mModulesPerStave % modulesPerStaveX != 0) {
+        LOG(fatal) << "Invalid oTOF module layout: total modules per stave " << mModulesPerStave
+                   << " is not divisible by modulesPerStaveX=" << modulesPerStaveX;
+      }
+      const int modulesPerStaveZ = mModulesPerStave / modulesPerStaveX;
+      const double moduleOverlapZ = 0.7; // cm, 7 mm longitudinal overlap from oTOF V2 specs
+      const double moduleSizeX = staveSizeX / modulesPerStaveX;
+      const double moduleSizeY = staveSizeY;
+      const double moduleSizeZ = (staveSizeZ + (modulesPerStaveZ - 1) * moduleOverlapZ) / modulesPerStaveZ;
+      const double modulePitchZ = moduleSizeZ - moduleOverlapZ;
+      if (modulePitchZ <= 0.0) {
+        LOG(fatal) << "Invalid oTOF module overlap " << moduleOverlapZ << " cm for module size " << moduleSizeZ << " cm";
+      }
       TGeoBBox* module = new TGeoBBox(moduleSizeX * 0.5, moduleSizeY * 0.5, moduleSizeZ * 0.5);
       TGeoVolume* moduleVol = new TGeoVolume(moduleName, module, medAir);
       setModuleStyle(moduleVol);
@@ -363,10 +386,12 @@ void OTOFLayer::createLayer(TGeoVolume* motherVolume)
 
       // Now we build a stave from modules
       for (int i = 0; i < modulesPerStaveX; ++i) {
-        for (int j = 0; j < mModulesPerStave; ++j) {
-          LOGP(info, "oTOF: Creating module {}/{} for stave {}/{}", i + 1, modulesPerStaveX, j + 1, mModulesPerStave);
-          auto* translation = new TGeoTranslation((i + 0.5) * moduleSizeX - 0.5 * staveSizeX, 0, (j + 0.5) * moduleSizeZ - 0.5 * staveSizeZ);
-          staveVol->AddNode(moduleVol, 1 + i * mModulesPerStave + j, translation);
+        for (int j = 0; j < modulesPerStaveZ; ++j) {
+          LOGP(info, "oTOF: Creating module {}/{} for stave {}/{}", i + 1, modulesPerStaveX, j + 1, modulesPerStaveZ);
+          const double tx = (i + 0.5) * moduleSizeX - 0.5 * staveSizeX;
+          const double tz = -0.5 * staveSizeZ + 0.5 * moduleSizeZ + j * modulePitchZ;
+          auto* translation = new TGeoTranslation(tx, 0, tz);
+          staveVol->AddNode(moduleVol, 1 + i * modulesPerStaveZ + j, translation);
         }
       }
 

From f017263d5da7a08e808cbb061139195cd2bf4a31 Mon Sep 17 00:00:00 2001
From: Maximiliano Puccio <mpuccio@users.noreply.github.com>
Date: Thu, 21 May 2026 12:18:14 +0200
Subject: [PATCH 07/17] ALICE3: factor GPU tracking into dynamically loaded
 CUDA/HIP backend (#15420)

---
 .../reconstruction/CMakeLists.txt             |  33 +--
 ...lAllocator.cxx => GPUExternalAllocator.cu} |  37 +--
 .../workflow/CMakeLists.txt                   |  34 +++
 .../TrackerSpec.h                             |  10 +
 .../TrackerSpecImpl.h                         | 226 ++++++++++++++++
 .../workflow/src/TrackerSpec.cxx              | 250 +++---------------
 .../workflow/src/TrackerSpecGPU.cxx           |  28 ++
 7 files changed, 342 insertions(+), 276 deletions(-)
 rename Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/{GPUExternalAllocator.cxx => GPUExternalAllocator.cu} (81%)
 create mode 100644 Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h
 create mode 100644 Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpecGPU.cxx

diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt b/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt
index 8805c1885b079..1dfcb7a22f725 100644
--- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt
+++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/CMakeLists.txt
@@ -13,24 +13,9 @@ if(Acts_FOUND)
   set(actsTarget Acts::Core)
 endif()
 
-set(alice3GlobalRecoGpuSources "")
-set(alice3GlobalRecoGpuTargets "")
-set(alice3GlobalRecoGpuPrivateTargets "")
-if(CUDA_ENABLED)
-  find_package(CUDAToolkit REQUIRED)
-  list(APPEND alice3GlobalRecoGpuSources src/TimeFrameGPU.cxx src/GPUExternalAllocator.cxx)
-  list(APPEND alice3GlobalRecoGpuTargets O2::ITStrackingCUDA)
-  list(APPEND alice3GlobalRecoGpuPrivateTargets CUDA::cudart)
-elseif(HIP_ENABLED)
-  list(APPEND alice3GlobalRecoGpuSources src/TimeFrameGPU.cxx src/GPUExternalAllocator.cxx)
-  list(APPEND alice3GlobalRecoGpuTargets O2::ITStrackingHIP)
-  list(APPEND alice3GlobalRecoGpuPrivateTargets hip::host)
-endif()
-
 o2_add_library(ALICE3GlobalReconstruction
                TARGETVARNAME targetName
                SOURCES src/TimeFrame.cxx
-                       ${alice3GlobalRecoGpuSources}
                        $<$<BOOL:${Acts_FOUND}>:src/TrackerACTS.cxx>
                PUBLIC_LINK_LIBRARIES
                        O2::ITStracking
@@ -48,26 +33,10 @@ o2_add_library(ALICE3GlobalReconstruction
                        O2::TRKReconstruction
                        O2::TRKSimulation
                        nlohmann_json::nlohmann_json
-                       ${alice3GlobalRecoGpuTargets}
                        ${actsTarget}
                PRIVATE_LINK_LIBRARIES
                        O2::Steer
-                       TBB::tbb
-                       ${alice3GlobalRecoGpuPrivateTargets})
-
-if(alice3GlobalRecoGpuTargets)
-  target_compile_definitions(${targetName} PUBLIC TRK_HAS_GPU_TRACKING)
-endif()
-
-if(CUDA_ENABLED)
-  target_include_directories(${targetName} PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
-endif()
-
-if(CUDA_ENABLED)
-  target_compile_definitions(${targetName} PUBLIC TRK_HAS_CUDA_TRACKING)
-elseif(HIP_ENABLED)
-  target_compile_definitions(${targetName} PUBLIC TRK_HAS_HIP_TRACKING)
-endif()
+                       TBB::tbb)
 
 if(Acts_FOUND)
   target_compile_definitions(${targetName} PUBLIC O2_WITH_ACTS)
diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cxx b/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cu
similarity index 81%
rename from Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cxx
rename to Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cu
index df2a2c30b037a..c7b5f1cee50f5 100644
--- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cxx
+++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/reconstruction/src/GPUExternalAllocator.cu
@@ -9,11 +9,9 @@
 // granted to it by virtue of its status as an Intergovernmental Organization
 // or submit itself to any jurisdiction.
 
-#if defined(TRK_HAS_CUDA_TRACKING)
+#define GPUCA_GPUCODE_HOSTONLY
+
 #include <cuda_runtime.h>
-#elif defined(TRK_HAS_HIP_TRACKING)
-#include <hip/hip_runtime.h>
-#endif
 
 #include "ALICE3GlobalReconstruction/GPUExternalAllocator.h"
 
@@ -23,21 +21,12 @@
 
 namespace
 {
-#if defined(TRK_HAS_CUDA_TRACKING)
 void checkGpuError(cudaError_t error, const char* call)
 {
   if (error != cudaSuccess) {
     throw std::runtime_error(std::string(call) + ": " + cudaGetErrorString(error));
   }
 }
-#elif defined(TRK_HAS_HIP_TRACKING)
-void checkGpuError(hipError_t error, const char* call)
-{
-  if (error != hipSuccess) {
-    throw std::runtime_error(std::string(call) + ": " + hipGetErrorString(error));
-  }
-}
-#endif
 } // namespace
 
 namespace o2::trk
@@ -147,26 +136,14 @@ void GPUExternalAllocator::releaseAll()
 void* GPUExternalAllocator::allocateHost(size_t size)
 {
   void* ptr = nullptr;
-#if defined(TRK_HAS_CUDA_TRACKING)
   checkGpuError(cudaHostAlloc(&ptr, size, cudaHostAllocPortable), "cudaHostAlloc");
-#elif defined(TRK_HAS_HIP_TRACKING)
-  checkGpuError(hipHostMalloc(&ptr, size, hipHostMallocPortable), "hipHostMalloc");
-#else
-  throw std::runtime_error("GPUExternalAllocator built without a GPU backend");
-#endif
   return ptr;
 }
 
 void* GPUExternalAllocator::allocateDevice(size_t size)
 {
   void* ptr = nullptr;
-#if defined(TRK_HAS_CUDA_TRACKING)
   checkGpuError(cudaMalloc(&ptr, size), "cudaMalloc");
-#elif defined(TRK_HAS_HIP_TRACKING)
-  checkGpuError(hipMalloc(&ptr, size), "hipMalloc");
-#else
-  throw std::runtime_error("GPUExternalAllocator built without a GPU backend");
-#endif
   return ptr;
 }
 
@@ -176,21 +153,11 @@ void GPUExternalAllocator::freeAllocation(void* ptr, AllocationSpace space)
     return;
   }
 
-#if defined(TRK_HAS_CUDA_TRACKING)
   if (space == AllocationSpace::Host) {
     checkGpuError(cudaFreeHost(ptr), "cudaFreeHost");
   } else {
     checkGpuError(cudaFree(ptr), "cudaFree");
   }
-#elif defined(TRK_HAS_HIP_TRACKING)
-  if (space == AllocationSpace::Host) {
-    checkGpuError(hipHostFree(ptr), "hipHostFree");
-  } else {
-    checkGpuError(hipFree(ptr), "hipFree");
-  }
-#else
-  (void)space;
-#endif
 }
 
 void GPUExternalAllocator::removeFromTagLocked(uint64_t tag, void* ptr)
diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt
index be6add9c03483..6a4994e11467b 100644
--- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt
+++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/CMakeLists.txt
@@ -24,8 +24,42 @@ o2_add_library(ALICE3GlobalReconstructionWorkflow
                                      O2::TRKBase
                                      O2::TRKSimulation
                                      O2::ALICE3GlobalReconstruction
+                                     O2::CommonUtils
                                      nlohmann_json::nlohmann_json)
 
+if(CUDA_ENABLED OR HIP_ENABLED)
+  target_compile_definitions(${targetName} PUBLIC TRK_HAS_GPU_TRACKING)
+endif()
+
+if(CUDA_ENABLED)
+  find_package(CUDAToolkit REQUIRED)
+  target_compile_definitions(${targetName} PUBLIC TRK_HAS_CUDA_TRACKING)
+  o2_add_library(ALICE3GlobalReconstructionWorkflowCUDA
+                 TARGETVARNAME cudaTargetName
+                 SOURCES src/TrackerSpecGPU.cxx
+                         ../reconstruction/src/TimeFrameGPU.cxx
+                         ../reconstruction/src/GPUExternalAllocator.cu
+                 PUBLIC_LINK_LIBRARIES
+                         O2::ALICE3GlobalReconstructionWorkflow
+                         O2::ITStrackingCUDA
+                 PRIVATE_LINK_LIBRARIES
+                         CUDA::cudart)
+  target_include_directories(${cudaTargetName} PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
+endif()
+
+if(HIP_ENABLED)
+  target_compile_definitions(${targetName} PUBLIC TRK_HAS_HIP_TRACKING)
+  o2_add_hipified_library(ALICE3GlobalReconstructionWorkflowHIP
+                          SOURCES src/TrackerSpecGPU.cxx
+                                  ../reconstruction/src/TimeFrameGPU.cxx
+                                  ../reconstruction/src/GPUExternalAllocator.cu
+                          PUBLIC_LINK_LIBRARIES
+                                  O2::ALICE3GlobalReconstructionWorkflow
+                                  O2::ITStrackingHIP
+                          PRIVATE_LINK_LIBRARIES
+                                  hip::host)
+endif()
+
 o2_add_executable(reco-workflow
                   SOURCES src/alice3-global-reconstruction-workflow.cxx
                   COMPONENT_NAME alice3-global-reconstruction
diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h
index 006bb4cbf5260..c1e7e051fb3f1 100644
--- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h
+++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpec.h
@@ -32,6 +32,10 @@
 
 #include <nlohmann/json.hpp>
 
+#include <memory>
+#include <utility>
+#include <vector>
+
 namespace o2::trk
 {
 class TrackerDPL : public framework::Task
@@ -48,10 +52,15 @@ class TrackerDPL : public framework::Task
   void endOfStream(framework::EndOfStreamContext& ec) final;
   //   void finaliseCCDB(framework::ConcreteDataMatcher& matcher, void* obj) final;
   void stop() final;
+  template <typename TimeFrameT, typename TrackerTraitsT>
+  void runTracking(framework::ProcessingContext& pc, TimeFrameT& timeFrame, TrackerTraitsT& trackerTraits);
+  const std::shared_ptr<its::ExternalAllocator>& getGPUAllocator() const noexcept { return mGPUAllocator; }
+  void setGPUAllocator(std::shared_ptr<its::ExternalAllocator> allocator) { mGPUAllocator = std::move(allocator); }
 
  private:
   void updateTimeDependentParams(framework::ProcessingContext& pc);
   std::vector<o2::its::TrackingParameters> createTrackingParamsFromConfig();
+  void runGPUTracking(framework::ProcessingContext& pc);
   //   std::unique_ptr<o2::gpu::GPUReconstruction> mRecChain = nullptr;
   //   std::unique_ptr<o2::gpu::GPUChainITS> mChainITS = nullptr;
   //   std::shared_ptr<o2::base::GRPGeomRequest> mGGCCDBRequest;
@@ -61,6 +70,7 @@ class TrackerDPL : public framework::Task
   std::shared_ptr<its::BoundedMemoryResource> mMemoryPool;
   std::shared_ptr<its::ExternalAllocator> mGPUAllocator;
   std::shared_ptr<tbb::task_arena> mTaskArena;
+  std::vector<o2::its::TrackingParameters> mTrackingParams;
   nlohmann::json mHitRecoConfig;
   nlohmann::json mClusterRecoConfig;
   TStopwatch mTimer;
diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h
new file mode 100644
index 0000000000000..f6221e485f369
--- /dev/null
+++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/include/ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h
@@ -0,0 +1,226 @@
+// 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.
+
+#ifndef O2_TRK_TRACKERSPECIMPL_H
+#define O2_TRK_TRACKERSPECIMPL_H
+
+#include "ALICE3GlobalReconstructionWorkflow/TrackerSpec.h"
+
+#include "CommonDataFormat/IRFrame.h"
+#include "DataFormatsTRK/Cluster.h"
+#include "DataFormatsTRK/ROFRecord.h"
+#include "DetectorsBase/GeometryManager.h"
+#include "Field/MagFieldParam.h"
+#include "Field/MagneticField.h"
+#include "Framework/ControlService.h"
+#include "ITStracking/Tracker.h"
+#include "SimulationDataFormat/MCCompLabel.h"
+#include "SimulationDataFormat/MCEventHeader.h"
+#include "SimulationDataFormat/MCTruthContainer.h"
+#include "TRKBase/GeometryTGeo.h"
+#include "TRKSimulation/Hit.h"
+
+#include <TFile.h>
+#include <TGeoGlobalMagField.h>
+#include <TTree.h>
+
+#include <algorithm>
+#include <array>
+#include <chrono>
+#include <format>
+#include <numeric>
+#include <vector>
+
+namespace o2::trk
+{
+
+template <typename TimeFrameT, typename TrackerTraitsT>
+void TrackerDPL::runTracking(framework::ProcessingContext& pc, TimeFrameT& timeFrame, TrackerTraitsT& trackerTraits)
+{
+  o2::its::Tracker<11> itsTracker(&trackerTraits);
+  timeFrame.setMemoryPool(mMemoryPool);
+  trackerTraits.setMemoryPool(mMemoryPool);
+  trackerTraits.setNThreads(mTaskArena->max_concurrency(), mTaskArena);
+  trackerTraits.adoptTimeFrame(static_cast<o2::its::TimeFrame<11>*>(&timeFrame));
+  itsTracker.adoptTimeFrame(timeFrame);
+  trackerTraits.updateTrackingParameters(mTrackingParams);
+  timeFrame.initTrackerTopologies(mTrackingParams, 11);
+
+  int nRofs{0};
+  if (!mHitRecoConfig.empty()) {
+    TFile hitsFile(mHitRecoConfig["inputfiles"]["hits"].get<std::string>().c_str(), "READ");
+    TFile mcHeaderFile(mHitRecoConfig["inputfiles"]["mcHeader"].get<std::string>().c_str(), "READ");
+    TTree* hitsTree = hitsFile.Get<TTree>("o2sim");
+    std::vector<o2::trk::Hit>* trkHit = nullptr;
+    hitsTree->SetBranchAddress("TRKHit", &trkHit);
+
+    TTree* mcHeaderTree = mcHeaderFile.Get<TTree>("o2sim");
+    auto mcheader = new o2::dataformats::MCEventHeader;
+    mcHeaderTree->SetBranchAddress("MCEventHeader.", &mcheader);
+
+    o2::base::GeometryManager::loadGeometry(mHitRecoConfig["inputfiles"]["geometry"].get<std::string>().c_str(), false, true);
+    auto* gman = o2::trk::GeometryTGeo::Instance();
+
+    const Long64_t nEvents{hitsTree->GetEntries()};
+    LOGP(info, "Starting {} reconstruction from hits for {} events", trackerTraits.getName(), nEvents);
+
+    trackerTraits.setBz(mHitRecoConfig["geometry"]["bz"].get<float>());
+    auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mHitRecoConfig["geometry"]["bz"].get<float>() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform);
+    TGeoGlobalMagField::Instance()->SetField(field);
+    TGeoGlobalMagField::Instance()->Lock();
+
+    nRofs = timeFrame.loadROFsFromHitTree(hitsTree, gman, mHitRecoConfig);
+    const int inROFpileup{mHitRecoConfig.contains("inROFpileup") ? mHitRecoConfig["inROFpileup"].get<int>() : 1};
+    timeFrame.getPrimaryVerticesFromMC(mcHeaderTree, nRofs, nEvents, inROFpileup);
+  } else if (!mClusterRecoConfig.empty()) {
+    LOGP(info, "Starting {} reconstruction from clusters", trackerTraits.getName());
+
+    o2::base::GeometryManager::loadGeometry(mClusterRecoConfig["inputfiles"]["geometry"].get<std::string>().c_str(), false, true);
+    o2::trk::GeometryTGeo::Instance();
+
+    trackerTraits.setBz(mClusterRecoConfig["geometry"]["bz"].get<float>());
+    auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mClusterRecoConfig["geometry"]["bz"].get<float>() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform);
+    TGeoGlobalMagField::Instance()->SetField(field);
+    TGeoGlobalMagField::Instance()->Lock();
+
+    constexpr int nLayers{11};
+    std::array<gsl::span<const o2::trk::Cluster>, nLayers> layerClusters;
+    std::array<gsl::span<const unsigned char>, nLayers> layerPatterns;
+    std::array<gsl::span<const o2::trk::ROFRecord>, nLayers> layerROFs;
+    std::array<const dataformats::MCTruthContainer<MCCompLabel>*, nLayers> layerLabels{};
+
+    size_t nInputRofs{0};
+    for (int iLayer = 0; iLayer < nLayers; ++iLayer) {
+      layerClusters[iLayer] = pc.inputs().get<gsl::span<o2::trk::Cluster>>(std::format("compClusters_{}", iLayer));
+      layerPatterns[iLayer] = pc.inputs().get<gsl::span<unsigned char>>(std::format("patterns_{}", iLayer));
+      layerROFs[iLayer] = pc.inputs().get<gsl::span<o2::trk::ROFRecord>>(std::format("ROframes_{}", iLayer));
+      nInputRofs = std::max(nInputRofs, layerROFs[iLayer].size());
+      if (mIsMC) {
+        layerLabels[iLayer] = pc.inputs().get<const dataformats::MCTruthContainer<MCCompLabel>*>(std::format("trkmclabels_{}", iLayer)).release();
+      }
+    }
+
+    timeFrame.deriveAndInitTiming(layerROFs);
+
+    const float yPlaneMLOT = 0.0010f;
+    nRofs = timeFrame.loadROFrameData(layerROFs, layerClusters, layerPatterns, mIsMC ? &layerLabels : nullptr, yPlaneMLOT);
+    timeFrame.addTruthSeedingVertices();
+  }
+
+  const auto trackingLoopStart = std::chrono::steady_clock::now();
+  for (size_t iter{0}; iter < mTrackingParams.size(); ++iter) {
+    LOGP(info, "{}", mTrackingParams[iter].asString());
+    trackerTraits.initialiseTimeFrame(iter);
+    trackerTraits.computeLayerTracklets(iter, -1);
+    LOGP(info, "Number of tracklets in iteration {}: {}", iter, timeFrame.getNumberOfTracklets());
+    trackerTraits.computeLayerCells(iter);
+    LOGP(info, "Number of cells in iteration {}: {}", iter, timeFrame.getNumberOfCells());
+    trackerTraits.findCellsNeighbours(iter);
+    LOGP(info, "Number of cell neighbours in iteration {}: {}", iter, timeFrame.getNumberOfNeighbours());
+    trackerTraits.findRoads(iter);
+    LOGP(info, "Number of roads in iteration {}: {}", iter, timeFrame.getNumberOfTracks());
+  }
+  const auto trackingLoopElapsedMs = std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::steady_clock::now() - trackingLoopStart).count();
+  LOGP(info, "Tracking iterations block took {} ms", trackingLoopElapsedMs);
+
+  if (mIsMC) {
+    itsTracker.computeTracksMClabels();
+  }
+
+  const auto& tracks = timeFrame.getTracks();
+  const auto& labels = timeFrame.getTracksLabel();
+  std::vector<o2::its::TrackITS> allTracks(tracks.begin(), tracks.end());
+  std::vector<o2::MCCompLabel> allLabels;
+
+  int totalTracks = allTracks.size();
+  int goodTracks = 0;
+  int fakeTracks = 0;
+
+  if (mIsMC) {
+    allLabels.assign(labels.begin(), labels.end());
+    for (const auto& label : allLabels) {
+      if (label.isFake()) {
+        ++fakeTracks;
+      } else {
+        ++goodTracks;
+      }
+    }
+  }
+
+  LOGP(info, "=== Tracking Summary ===");
+  LOGP(info, "Total tracks reconstructed: {}", totalTracks);
+  LOGP(info, "Good tracks: {} ({:.1f}%)", goodTracks, totalTracks > 0 ? 100.0 * goodTracks / totalTracks : 0);
+  LOGP(info, "Fake tracks: {} ({:.1f}%)", fakeTracks, totalTracks > 0 ? 100.0 * fakeTracks / totalTracks : 0);
+
+  const auto& rofView = timeFrame.getROFOverlapTableView();
+  const auto& clockLayer = rofView.getClockLayer();
+  const int clockLayerId = rofView.getClock();
+  const int64_t anchorBC = timeFrame.getTFAnchorIR().toLong();
+
+  int highestROF = static_cast<int>(clockLayer.mNROFsTF);
+  for (const auto& trc : allTracks) {
+    highestROF = std::max(highestROF, static_cast<int>(clockLayer.getROF(trc.getTimeStamp())));
+  }
+  for (const auto& vtx : timeFrame.getPrimaryVertices()) {
+    highestROF = std::max(highestROF, static_cast<int>(clockLayer.getROF(vtx.getTimeStamp().lower())));
+  }
+
+  std::vector<o2::trk::ROFRecord> allTrackROFs(highestROF);
+  for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) {
+    auto& rof = allTrackROFs[iROF];
+    o2::InteractionRecord ir;
+    ir.setFromLong(anchorBC + static_cast<int64_t>(clockLayer.getROFStartInBC(iROF)));
+    rof.setBCData(ir);
+    rof.setROFrame(iROF);
+    rof.setFirstEntry(0);
+    rof.setNEntries(0);
+  }
+
+  std::vector<int> rofEntries(highestROF + 1, 0);
+  for (const auto& trc : allTracks) {
+    const int rof = static_cast<int>(clockLayer.getROF(trc.getTimeStamp()));
+    if (rof >= 0 && rof < highestROF) {
+      ++rofEntries[rof];
+    }
+  }
+  std::exclusive_scan(rofEntries.begin(), rofEntries.end(), rofEntries.begin(), 0);
+
+  std::vector<o2::dataformats::IRFrame> irFrames;
+  irFrames.reserve(allTrackROFs.size());
+  const auto& maskView = timeFrame.getROFMaskView();
+  const auto rofLenMinus1 = clockLayer.mROFLength > 0 ? clockLayer.mROFLength - 1 : 0;
+  for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) {
+    allTrackROFs[iROF].setFirstEntry(rofEntries[iROF]);
+    allTrackROFs[iROF].setNEntries(rofEntries[iROF + 1] - rofEntries[iROF]);
+    if (maskView.isROFEnabled(clockLayerId, static_cast<int>(iROF))) {
+      const auto& bcStart = allTrackROFs[iROF].getBCData();
+      auto& irFrame = irFrames.emplace_back(bcStart, bcStart + rofLenMinus1);
+      irFrame.info = allTrackROFs[iROF].getNEntries();
+    }
+  }
+
+  pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKS", 0}, allTracks);
+  pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSROF", 0}, allTrackROFs);
+  pc.outputs().snapshot(o2::framework::Output{"TRK", "IRFRAMES", 0}, irFrames);
+  if (mIsMC) {
+    pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSMCTR", 0}, allLabels);
+  }
+
+  LOGP(info, "TRK pushed {} tracks in {} ROFs and {} IR frames{}",
+       allTracks.size(), allTrackROFs.size(), irFrames.size(),
+       mIsMC ? " (with MC labels)" : "");
+
+  timeFrame.wipe();
+}
+
+} // namespace o2::trk
+
+#endif // O2_TRK_TRACKERSPECIMPL_H
diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx
index 9fb2899ab3ef5..6f9f5561a5ef6 100644
--- a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx
+++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpec.cxx
@@ -17,6 +17,7 @@
 #include <fstream>
 #include <numeric>
 
+#include "CommonUtils/DLLoaderBase.h"
 #include "CommonDataFormat/IRFrame.h"
 #include "DataFormatsTRK/Cluster.h"
 #include "DataFormatsTRK/ROFRecord.h"
@@ -36,12 +37,8 @@
 #include "TRKBase/SegmentationChip.h"
 #include "TRKSimulation/Hit.h"
 #include "ALICE3GlobalReconstruction/TimeFrame.h"
-#ifdef TRK_HAS_GPU_TRACKING
-#include "ALICE3GlobalReconstruction/TimeFrameGPU.h"
-#include "ALICE3GlobalReconstruction/GPUExternalAllocator.h"
-#include "ITStrackingGPU/TrackerTraitsGPU.h"
-#endif
 #include "ALICE3GlobalReconstructionWorkflow/TrackerSpec.h"
+#include "ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h"
 #include <TGeoGlobalMagField.h>
 
 #ifdef O2_WITH_ACTS
@@ -58,6 +55,18 @@ namespace trk
 {
 using Vertex = o2::dataformats::Vertex<o2::dataformats::TimeStamp<int>>;
 
+namespace
+{
+class ALICE3TrackingBackendLoader : public o2::utils::DLLoaderBase<ALICE3TrackingBackendLoader>
+{
+  O2DLLoaderDef(ALICE3TrackingBackendLoader)
+};
+
+O2DLLoaderImpl(ALICE3TrackingBackendLoader)
+
+  constexpr const char* kGPUBackendFunction = "runALICE3GPUTracking";
+} // namespace
+
 TrackerDPL::TrackerDPL(std::shared_ptr<o2::base::GRPGeomRequest> gr,
                        bool isMC,
                        const std::string& hitRecoConfigFileName,
@@ -249,220 +258,20 @@ void TrackerDPL::run(ProcessingContext& pc)
     mTaskArena = std::make_shared<tbb::task_arena>(1); /// TODO: make it configurable
   }
 
-  auto trackingParams = createTrackingParamsFromConfig();
+  mTrackingParams = createTrackingParamsFromConfig();
 
   auto cput = mTimer.CpuTime();
   auto realt = mTimer.RealTime();
   mTimer.Start(false);
 
   const bool useGPU = mDeviceType != o2::gpu::gpudatatypes::DeviceType::CPU;
-#ifndef TRK_HAS_GPU_TRACKING
-  if (useGPU) {
-    LOGP(fatal, "TRK GPU tracking was requested but this build has no TRK GPU tracking backend");
-  }
-#else
-#ifdef TRK_HAS_CUDA_TRACKING
-  if (useGPU && mDeviceType != o2::gpu::gpudatatypes::DeviceType::CUDA) {
-    LOGP(fatal, "This build provides the CUDA TRK tracking backend only, but device type {} was requested", static_cast<int>(mDeviceType));
-  }
-#elif defined(TRK_HAS_HIP_TRACKING)
-  if (useGPU && mDeviceType != o2::gpu::gpudatatypes::DeviceType::HIP) {
-    LOGP(fatal, "This build provides the HIP TRK tracking backend only, but device type {} was requested", static_cast<int>(mDeviceType));
-  }
-#endif
-#endif
-
-  auto runTracking = [&](auto& timeFrame, auto& trackerTraits) {
-    o2::its::Tracker<11> itsTracker(&trackerTraits);
-    timeFrame.setMemoryPool(mMemoryPool);
-    trackerTraits.setMemoryPool(mMemoryPool);
-    trackerTraits.setNThreads(mTaskArena->max_concurrency(), mTaskArena);
-    trackerTraits.adoptTimeFrame(static_cast<o2::its::TimeFrame<11>*>(&timeFrame));
-    itsTracker.adoptTimeFrame(timeFrame);
-    trackerTraits.updateTrackingParameters(trackingParams);
-
-    int nRofs{0};
-    if (!mHitRecoConfig.empty()) {
-      TFile hitsFile(mHitRecoConfig["inputfiles"]["hits"].get<std::string>().c_str(), "READ");
-      TFile mcHeaderFile(mHitRecoConfig["inputfiles"]["mcHeader"].get<std::string>().c_str(), "READ");
-      TTree* hitsTree = hitsFile.Get<TTree>("o2sim");
-      std::vector<o2::trk::Hit>* trkHit = nullptr;
-      hitsTree->SetBranchAddress("TRKHit", &trkHit);
-
-      TTree* mcHeaderTree = mcHeaderFile.Get<TTree>("o2sim");
-      auto mcheader = new o2::dataformats::MCEventHeader;
-      mcHeaderTree->SetBranchAddress("MCEventHeader.", &mcheader);
-
-      o2::base::GeometryManager::loadGeometry(mHitRecoConfig["inputfiles"]["geometry"].get<std::string>().c_str(), false, true);
-      auto* gman = o2::trk::GeometryTGeo::Instance();
-
-      const Long64_t nEvents{hitsTree->GetEntries()};
-      LOGP(info, "Starting {} reconstruction from hits for {} events", trackerTraits.getName(), nEvents);
-
-      trackerTraits.setBz(mHitRecoConfig["geometry"]["bz"].get<float>());
-      auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mHitRecoConfig["geometry"]["bz"].get<float>() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform);
-      TGeoGlobalMagField::Instance()->SetField(field);
-      TGeoGlobalMagField::Instance()->Lock();
-
-      nRofs = timeFrame.loadROFsFromHitTree(hitsTree, gman, mHitRecoConfig);
-      const int inROFpileup{mHitRecoConfig.contains("inROFpileup") ? mHitRecoConfig["inROFpileup"].get<int>() : 1};
-      timeFrame.getPrimaryVerticesFromMC(mcHeaderTree, nRofs, nEvents, inROFpileup);
-    } else if (!mClusterRecoConfig.empty()) {
-      LOGP(info, "Starting {} reconstruction from clusters", trackerTraits.getName());
-
-      o2::base::GeometryManager::loadGeometry(mClusterRecoConfig["inputfiles"]["geometry"].get<std::string>().c_str(), false, true);
-      o2::trk::GeometryTGeo::Instance();
-
-      trackerTraits.setBz(mClusterRecoConfig["geometry"]["bz"].get<float>());
-      auto field = new field::MagneticField("ALICE3Mag", "ALICE 3 Magnetic Field", mClusterRecoConfig["geometry"]["bz"].get<float>() / 5.f, 0.0, o2::field::MagFieldParam::k5kGUniform);
-      TGeoGlobalMagField::Instance()->SetField(field);
-      TGeoGlobalMagField::Instance()->Lock();
-
-      constexpr int nLayers{11};
-      std::array<gsl::span<const o2::trk::Cluster>, nLayers> layerClusters;
-      std::array<gsl::span<const unsigned char>, nLayers> layerPatterns;
-      std::array<gsl::span<const o2::trk::ROFRecord>, nLayers> layerROFs;
-      std::array<const dataformats::MCTruthContainer<MCCompLabel>*, nLayers> layerLabels{};
-
-      size_t nInputRofs{0};
-      for (int iLayer = 0; iLayer < nLayers; ++iLayer) {
-        layerClusters[iLayer] = pc.inputs().get<gsl::span<o2::trk::Cluster>>(std::format("compClusters_{}", iLayer));
-        layerPatterns[iLayer] = pc.inputs().get<gsl::span<unsigned char>>(std::format("patterns_{}", iLayer));
-        layerROFs[iLayer] = pc.inputs().get<gsl::span<o2::trk::ROFRecord>>(std::format("ROframes_{}", iLayer));
-        nInputRofs = std::max(nInputRofs, layerROFs[iLayer].size());
-        if (mIsMC) {
-          layerLabels[iLayer] = pc.inputs().get<const dataformats::MCTruthContainer<MCCompLabel>*>(std::format("trkmclabels_{}", iLayer)).release();
-        }
-      }
-
-      timeFrame.deriveAndInitTiming(layerROFs);
-
-      const float yPlaneMLOT = 0.0010f;
-      nRofs = timeFrame.loadROFrameData(layerROFs, layerClusters, layerPatterns, mIsMC ? &layerLabels : nullptr, yPlaneMLOT);
-      timeFrame.addTruthSeedingVertices();
-    }
-
-    const auto trackingLoopStart = std::chrono::steady_clock::now();
-    for (size_t iter{0}; iter < trackingParams.size(); ++iter) {
-      LOGP(info, "{}", trackingParams[iter].asString());
-      trackerTraits.initialiseTimeFrame(iter);
-      trackerTraits.computeLayerTracklets(iter, -1);
-      LOGP(info, "Number of tracklets in iteration {}: {}", iter, timeFrame.getNumberOfTracklets());
-      trackerTraits.computeLayerCells(iter);
-      LOGP(info, "Number of cells in iteration {}: {}", iter, timeFrame.getNumberOfCells());
-      trackerTraits.findCellsNeighbours(iter);
-      LOGP(info, "Number of cell neighbours in iteration {}: {}", iter, timeFrame.getNumberOfNeighbours());
-      trackerTraits.findRoads(iter);
-      LOGP(info, "Number of roads in iteration {}: {}", iter, timeFrame.getNumberOfTracks());
-    }
-    const auto trackingLoopElapsedMs = std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::steady_clock::now() - trackingLoopStart).count();
-    LOGP(info, "Tracking iterations block took {} ms", trackingLoopElapsedMs);
-
-    if (mIsMC) {
-      itsTracker.computeTracksMClabels();
-    }
-
-    const auto& tracks = timeFrame.getTracks();
-    const auto& labels = timeFrame.getTracksLabel();
-    std::vector<o2::its::TrackITS> allTracks(tracks.begin(), tracks.end());
-    std::vector<o2::MCCompLabel> allLabels;
-
-    int totalTracks = allTracks.size();
-    int goodTracks = 0;
-    int fakeTracks = 0;
-
-    if (mIsMC) {
-      allLabels.assign(labels.begin(), labels.end());
-      for (const auto& label : allLabels) {
-        if (label.isFake()) {
-          ++fakeTracks;
-        } else {
-          ++goodTracks;
-        }
-      }
-    }
-
-    LOGP(info, "=== Tracking Summary ===");
-    LOGP(info, "Total tracks reconstructed: {}", totalTracks);
-    LOGP(info, "Good tracks: {} ({:.1f}%)", goodTracks, totalTracks > 0 ? 100.0 * goodTracks / totalTracks : 0);
-    LOGP(info, "Fake tracks: {} ({:.1f}%)", fakeTracks, totalTracks > 0 ? 100.0 * fakeTracks / totalTracks : 0);
-
-    const auto& rofView = timeFrame.getROFOverlapTableView();
-    const auto& clockLayer = rofView.getClockLayer();
-    const int clockLayerId = rofView.getClock();
-    const int64_t anchorBC = timeFrame.getTFAnchorIR().toLong();
 
-    int highestROF = static_cast<int>(clockLayer.mNROFsTF);
-    for (const auto& trc : allTracks) {
-      highestROF = std::max(highestROF, static_cast<int>(clockLayer.getROF(trc.getTimeStamp())));
-    }
-    for (const auto& vtx : timeFrame.getPrimaryVertices()) {
-      highestROF = std::max(highestROF, static_cast<int>(clockLayer.getROF(vtx.getTimeStamp().lower())));
-    }
-
-    std::vector<o2::trk::ROFRecord> allTrackROFs(highestROF);
-    for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) {
-      auto& rof = allTrackROFs[iROF];
-      o2::InteractionRecord ir;
-      ir.setFromLong(anchorBC + static_cast<int64_t>(clockLayer.getROFStartInBC(iROF)));
-      rof.setBCData(ir);
-      rof.setROFrame(iROF);
-      rof.setFirstEntry(0);
-      rof.setNEntries(0);
-    }
-
-    std::vector<int> rofEntries(highestROF + 1, 0);
-    for (const auto& trc : allTracks) {
-      const int rof = static_cast<int>(clockLayer.getROF(trc.getTimeStamp()));
-      if (rof >= 0 && rof < highestROF) {
-        ++rofEntries[rof];
-      }
-    }
-    std::exclusive_scan(rofEntries.begin(), rofEntries.end(), rofEntries.begin(), 0);
-
-    std::vector<o2::dataformats::IRFrame> irFrames;
-    irFrames.reserve(allTrackROFs.size());
-    const auto& maskView = timeFrame.getROFMaskView();
-    const auto rofLenMinus1 = clockLayer.mROFLength > 0 ? clockLayer.mROFLength - 1 : 0;
-    for (size_t iROF = 0; iROF < allTrackROFs.size(); ++iROF) {
-      allTrackROFs[iROF].setFirstEntry(rofEntries[iROF]);
-      allTrackROFs[iROF].setNEntries(rofEntries[iROF + 1] - rofEntries[iROF]);
-      if (maskView.isROFEnabled(clockLayerId, static_cast<int>(iROF))) {
-        const auto& bcStart = allTrackROFs[iROF].getBCData();
-        auto& irFrame = irFrames.emplace_back(bcStart, bcStart + rofLenMinus1);
-        irFrame.info = allTrackROFs[iROF].getNEntries();
-      }
-    }
-
-    pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKS", 0}, allTracks);
-    pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSROF", 0}, allTrackROFs);
-    pc.outputs().snapshot(o2::framework::Output{"TRK", "IRFRAMES", 0}, irFrames);
-    if (mIsMC) {
-      pc.outputs().snapshot(o2::framework::Output{"TRK", "TRACKSMCTR", 0}, allLabels);
-    }
-
-    LOGP(info, "TRK pushed {} tracks in {} ROFs and {} IR frames{}",
-         allTracks.size(), allTrackROFs.size(), irFrames.size(),
-         mIsMC ? " (with MC labels)" : "");
-
-    timeFrame.wipe();
-  };
-
-#ifdef TRK_HAS_GPU_TRACKING
   if (useGPU) {
-    o2::trk::TimeFrameGPU<11> timeFrame;
-    o2::its::TrackerTraitsGPU<11> itsTrackerTraits;
-    if (!mGPUAllocator) {
-      mGPUAllocator = std::make_shared<o2::trk::GPUExternalAllocator>();
-    }
-    timeFrame.setFrameworkAllocator(mGPUAllocator.get());
-    runTracking(timeFrame, itsTrackerTraits);
-  } else
-#endif
-  {
+    runGPUTracking(pc);
+  } else {
     o2::trk::TimeFrame<11> timeFrame;
     o2::its::TrackerTraits<11> itsTrackerTraits;
-    runTracking(timeFrame, itsTrackerTraits);
+    runTracking(pc, timeFrame, itsTrackerTraits);
   }
 
   pc.services().get<o2::framework::ControlService>().endOfStream();
@@ -472,6 +281,29 @@ void TrackerDPL::run(ProcessingContext& pc)
   LOGP(info, "CPU Reconstruction time for this TF {} s (cpu), {} s (wall)", mTimer.CpuTime() - cput, mTimer.RealTime() - realt);
 }
 
+void TrackerDPL::runGPUTracking(ProcessingContext& pc)
+{
+  auto& loader = ALICE3TrackingBackendLoader::Instance();
+  switch (mDeviceType) {
+    case o2::gpu::gpudatatypes::DeviceType::CUDA:
+#ifdef TRK_HAS_CUDA_TRACKING
+      loader.executeFunctionAlias<int, TrackerDPL*, ProcessingContext*>("O2ALICE3GlobalReconstructionWorkflowCUDA", kGPUBackendFunction, this, &pc);
+      return;
+#else
+      LOGP(fatal, "CUDA TRK GPU tracking was requested but this build has no CUDA TRK GPU tracking backend");
+#endif
+    case o2::gpu::gpudatatypes::DeviceType::HIP:
+#ifdef TRK_HAS_HIP_TRACKING
+      loader.executeFunctionAlias<int, TrackerDPL*, ProcessingContext*>("O2ALICE3GlobalReconstructionWorkflowHIP", kGPUBackendFunction, this, &pc);
+      return;
+#else
+      LOGP(fatal, "HIP TRK GPU tracking was requested but this build has no HIP TRK GPU tracking backend");
+#endif
+    default:
+      LOGP(fatal, "Unsupported TRK GPU device type {}", static_cast<int>(mDeviceType));
+  }
+}
+
 void TrackerDPL::endOfStream(EndOfStreamContext& ec)
 {
   LOGF(info, "TRK CA-Tracker total timing: Cpu: %.3e Real: %.3e s in %d slots", mTimer.CpuTime(), mTimer.RealTime(), mTimer.Counter() - 1);
diff --git a/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpecGPU.cxx b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpecGPU.cxx
new file mode 100644
index 0000000000000..ea98ab3f852e5
--- /dev/null
+++ b/Detectors/Upgrades/ALICE3/GlobalReconstruction/workflow/src/TrackerSpecGPU.cxx
@@ -0,0 +1,28 @@
+// 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.
+
+#include "ALICE3GlobalReconstruction/GPUExternalAllocator.h"
+#include "ALICE3GlobalReconstruction/TimeFrameGPU.h"
+#include "ALICE3GlobalReconstructionWorkflow/TrackerSpec.h"
+#include "ALICE3GlobalReconstructionWorkflow/TrackerSpecImpl.h"
+#include "ITStrackingGPU/TrackerTraitsGPU.h"
+
+extern "C" int runALICE3GPUTracking(o2::trk::TrackerDPL* tracker, o2::framework::ProcessingContext* pc)
+{
+  o2::trk::TimeFrameGPU<11> timeFrame;
+  o2::its::TrackerTraitsGPU<11> itsTrackerTraits;
+  if (!tracker->getGPUAllocator()) {
+    tracker->setGPUAllocator(std::make_shared<o2::trk::GPUExternalAllocator>());
+  }
+  timeFrame.setFrameworkAllocator(tracker->getGPUAllocator().get());
+  tracker->runTracking(*pc, timeFrame, itsTrackerTraits);
+  return 0;
+}

From c7d5958f58ac9ac4295c8e79f9578f74ddd72bf7 Mon Sep 17 00:00:00 2001
From: Fabrizio Chinu <91954233+fchinu@users.noreply.github.com>
Date: Thu, 21 May 2026 13:36:01 +0200
Subject: [PATCH 08/17] ITS: add selections on tracks sharing clusters (#15406)

* ITS: add selections on tracks sharing clusters

* Refactor selection function, return simple int instead of reference to int

* Use isPhiDifferenceBelow, avoid variable narrowing

* Improve memory management for selection of tracks with shared clusters
---
 .../ITS/include/DataFormatsITS/TrackITS.h     |  8 ++-
 .../tracking/GPU/cuda/TrackerTraitsGPU.cxx    |  4 +-
 .../include/ITStracking/Configuration.h       |  7 ++-
 .../include/ITStracking/TrackerTraits.h       |  4 +-
 .../include/ITStracking/TrackingConfigParam.h |  5 ++
 .../ITSMFT/ITS/tracking/src/Configuration.cxx |  3 +
 .../ITSMFT/ITS/tracking/src/TrackerTraits.cxx | 61 ++++++++++++-------
 7 files changed, 63 insertions(+), 29 deletions(-)

diff --git a/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h b/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h
index a06395e76afff..89f6416c6e177 100644
--- a/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h
+++ b/DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h
@@ -192,7 +192,13 @@ class TrackITSExt : public TrackITS
     getClusterRefs().setEntries(ncl);
   }
 
-  GPUhdi() const int& getClusterIndex(int lr) const { return mIndex[lr]; }
+  GPUhdi() const int getClusterIndex(int lr) const { return mIndex[lr]; }
+
+  GPUh() const int getFirstLayerClusterIndex() const
+  {
+    int firstLayer = getFirstClusterLayer();
+    return getClusterIndex(firstLayer);
+  }
 
   GPUhdi() void setExternalClusterIndex(int layer, int idx, bool newCluster = false)
   {
diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
index f1812c9f6f764..32c46e2ea55d2 100644
--- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
+++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
@@ -387,10 +387,10 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
     mTimeFrameGPU->downloadTrackITSExtDevice();
 
     auto& tracks = mTimeFrameGPU->getTrackITSExt();
-    this->acceptTracks(iteration, tracks, firstClusters, sharedFirstClusters);
+    this->acceptTracks(iteration, tracks, firstClusters);
     mTimeFrameGPU->loadUsedClustersDevice();
   }
-  this->markTracks(iteration, sharedFirstClusters);
+  this->markTracks(iteration);
   // wipe the artefact memory
   mTimeFrameGPU->popMemoryStack(iteration);
 };
diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h
index c939f39532fdb..852c5ecd24633 100644
--- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h
+++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h
@@ -70,7 +70,6 @@ struct TrackingParameters {
   float DiamondCov[6] = {25.e-6f, 0.f, 0.f, 25.e-6f, 0.f, 36.f};
 
   /// General parameters
-  bool AllowSharingFirstCluster = false;
   int ClusterSharing = 0;
   int MinTrackLength = 7;
   int MaxHoles = 0;
@@ -98,6 +97,12 @@ struct TrackingParameters {
   bool PrintMemory = false; // print allocator usage in epilog report
   size_t MaxMemory = std::numeric_limits<size_t>::max();
   bool DropTFUponFailure = false;
+
+  // Selections on tracks sharing clusters
+  bool AllowSharingFirstCluster = false;
+  float SharedClusterMaxDeltaPhi = 0.05f; // For tracks sharing clusters, maximum allowed delta phi at the cluster position
+  float SharedClusterMaxDeltaEta = 0.03f; // For tracks sharing clusters, maximum allowed delta eta at the cluster position
+  bool SharedClusterOppositeSign = false; // For tracks sharing clusters, require opposite sign of the tracklets
 };
 
 struct VertexingParameters {
diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h
index 647403bb6b548..f536e86fe95d5 100644
--- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h
+++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h
@@ -55,8 +55,8 @@ class TrackerTraits
   template <typename InputSeed>
   void processNeighbours(int iteration, int defaultCellTopologyId, int iLevel, const bounded_vector<InputSeed>& currentCellSeed, const bounded_vector<int>& currentCellId, const bounded_vector<int>& currentCellTopologyId, bounded_vector<TrackSeedN>& updatedCellSeed, bounded_vector<int>& updatedCellId, bounded_vector<int>& updatedCellTopologyId);
 
-  void acceptTracks(int iteration, bounded_vector<TrackITSExt>& tracks, bounded_vector<bounded_vector<int>>& firstClusters, bounded_vector<bounded_vector<int>>& sharedFirstClusters);
-  void markTracks(int iteration, bounded_vector<bounded_vector<int>>& sharedFirstClusters);
+  void acceptTracks(int iteration, bounded_vector<TrackITSExt>& tracks, bounded_vector<bounded_vector<int>>& firstClusters);
+  void markTracks(int iteration);
 
   void updateTrackingParameters(const std::vector<TrackingParameters>& trkPars)
   {
diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h
index 21b4f928d5b73..69aa3c5fdaf06 100644
--- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h
+++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingConfigParam.h
@@ -102,7 +102,12 @@ struct TrackerParamConfig : public o2::conf::ConfigurableParamHelper<TrackerPara
   size_t maxMemory = std::numeric_limits<size_t>::max();
   bool dropTFUponFailure = false;
   bool fataliseUponFailure = true;       // granular management of the fatalisation in async mode
+
+  // Selections on tracks sharing clusters
   bool allowSharingFirstCluster = false; // allow first cluster sharing among tracks
+  float sharedClusterMaxDeltaPhi = 0.05f; // Maximum allowed delta phi at the cluster position
+  float sharedClusterMaxDeltaEta = 0.03f; // Maximum allowed delta eta at the cluster position
+  bool sharedClusterOppositeSign = false; // Require opposite sign of the tracklets
 
   O2ParamDef(TrackerParamConfig, "ITSCATrackerParam");
 };
diff --git a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx
index 0087da0a85ac2..0bf383c996a68 100644
--- a/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx
+++ b/Detectors/ITSMFT/ITS/tracking/src/Configuration.cxx
@@ -204,6 +204,9 @@ std::vector<TrackingParameters> TrackingMode::getTrackingParameters(TrackingMode
     p.SaveTimeBenchmarks = tc.saveTimeBenchmarks;
     p.FataliseUponFailure = tc.fataliseUponFailure;
     p.AllowSharingFirstCluster = tc.allowSharingFirstCluster;
+    p.SharedClusterMaxDeltaPhi = tc.sharedClusterMaxDeltaPhi;
+    p.SharedClusterMaxDeltaEta = tc.sharedClusterMaxDeltaEta;
+    p.SharedClusterOppositeSign = tc.sharedClusterOppositeSign;
     const auto iter = &p - trackParams.data();
     if (iter < constants::MaxIter) {
       p.MaxHoles = tc.maxHolesIter[iter];
diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx
index 9fef067559e8a..3432b60162002 100644
--- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx
+++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx
@@ -661,9 +661,7 @@ template <int NLayers>
 void TrackerTraits<NLayers>::findRoads(const int iteration)
 {
   bounded_vector<bounded_vector<int>> firstClusters(mTrkParams[iteration].NLayers, bounded_vector<int>(mMemoryPool.get()), mMemoryPool.get());
-  bounded_vector<bounded_vector<int>> sharedFirstClusters(mTrkParams[iteration].NLayers, bounded_vector<int>(mMemoryPool.get()), mMemoryPool.get());
   firstClusters.resize(mTrkParams[iteration].NLayers);
-  sharedFirstClusters.resize(mTrkParams[iteration].NLayers);
   const auto propagator = o2::base::Propagator::Instance();
   const TrackingFrameInfo* tfInfos[NLayers]{};
   const Cluster* unsortedClusters[NLayers]{};
@@ -787,13 +785,13 @@ void TrackerTraits<NLayers>::findRoads(const int iteration)
       return track::isBetter(a, b);
     });
 
-    acceptTracks(iteration, tracks, firstClusters, sharedFirstClusters);
+    acceptTracks(iteration, tracks, firstClusters);
   }
-  markTracks(iteration, sharedFirstClusters);
+  markTracks(iteration);
 }
 
 template <int NLayers>
-void TrackerTraits<NLayers>::acceptTracks(int iteration, bounded_vector<TrackITSExt>& tracks, bounded_vector<bounded_vector<int>>& firstClusters, bounded_vector<bounded_vector<int>>& sharedFirstClusters)
+void TrackerTraits<NLayers>::acceptTracks(int iteration, bounded_vector<TrackITSExt>& tracks, bounded_vector<bounded_vector<int>>& firstClusters)
 {
   auto& trks = mTimeFrame->getTracks();
   trks.reserve(trks.size() + tracks.size());
@@ -860,34 +858,51 @@ void TrackerTraits<NLayers>::acceptTracks(int iteration, bounded_vector<TrackITS
 
     if (mTrkParams[iteration].AllowSharingFirstCluster) {
       firstClusters[firstLayer].push_back(firstCluster);
-      if (isFirstShared) {
-        sharedFirstClusters[firstLayer].push_back(firstCluster);
-      }
     }
   }
 }
 
 template <int NLayers>
-void TrackerTraits<NLayers>::markTracks(int iteration, bounded_vector<bounded_vector<int>>& sharedFirstClusters)
+void TrackerTraits<NLayers>::markTracks(int iteration)
 {
   if (mTrkParams[iteration].AllowSharingFirstCluster) {
     /// Now we have to set the shared cluster flag
-    for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) {
-      std::sort(sharedFirstClusters[iLayer].begin(), sharedFirstClusters[iLayer].end());
-    }
+    auto& tracks = mTimeFrame->getTracks();
 
-    for (auto& track : mTimeFrame->getTracks()) {
-      int firstLayer{mTrkParams[iteration].NLayers}, firstCluster{constants::UnusedIndex};
-      for (int iLayer{0}; iLayer < mTrkParams[iteration].NLayers; ++iLayer) {
-        if (track.getClusterIndex(iLayer) == constants::UnusedIndex) {
-          continue;
-        }
-        firstLayer = iLayer;
-        firstCluster = track.getClusterIndex(iLayer);
-        break;
+    bounded_vector<int> fclusSort(tracks.size(), mMemoryPool.get());
+    std::iota(fclusSort.begin(), fclusSort.end(), 0);
+    std::sort(fclusSort.begin(), fclusSort.end(), [&tracks](int a, int b) {
+      return tracks[a].getFirstLayerClusterIndex() < tracks[b].getFirstLayerClusterIndex();
+    });
+
+    auto areTracksSelected = [this, iteration](const TrackITSExt& t1, const TrackITSExt& t2) {
+      const auto t1FirstLayer{t1.getFirstClusterLayer()}, t2FirstLayer{t2.getFirstClusterLayer()};
+      if (t1FirstLayer != t2FirstLayer) {
+        return false;
+      }
+      if (mTimeFrame->getClusterROF(t1FirstLayer, t1.getClusterIndex(t1FirstLayer)) != mTimeFrame->getClusterROF(t2FirstLayer, t2.getClusterIndex(t2FirstLayer))) {
+        return false;
+      }
+      if (!math_utils::isPhiDifferenceBelow(t1.getPhi(), t2.getPhi(), mTrkParams[iteration].SharedClusterMaxDeltaPhi)) {
+        return false;
+      }
+      if (std::abs(t1.getEta() - t2.getEta()) > mTrkParams[iteration].SharedClusterMaxDeltaEta) {
+        return false;
       }
-      if (std::binary_search(sharedFirstClusters[firstLayer].begin(), sharedFirstClusters[firstLayer].end(), firstCluster)) {
-        track.setSharedClusters();
+      if (mTrkParams[iteration].SharedClusterOppositeSign && t1.getSign() == t2.getSign()) {
+        return false;
+      }
+      return true;
+    };
+
+    for (int i{0}; i < static_cast<int>(fclusSort.size()); ++i) {
+      auto& track = tracks[fclusSort[i]];
+      for (int j{i + 1}; j < static_cast<int>(fclusSort.size()) && tracks[fclusSort[j]].getFirstLayerClusterIndex() == track.getFirstLayerClusterIndex(); ++j) {
+        auto& track2 = tracks[fclusSort[j]];
+        if (areTracksSelected(track, track2)) {
+          track.setSharedClusters();
+          track2.setSharedClusters();
+        }
       }
     }
   }

From 8c2216f6a5eb93d433f0a834d5c410efd2e98904 Mon Sep 17 00:00:00 2001
From: shahoian <ruben.shahoyan@cern.ch>
Date: Thu, 21 May 2026 18:04:09 +0200
Subject: [PATCH 09/17] Update CorrectionMapsOptions help hints

---
 Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx b/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx
index 45c3771db57bf..5518d680420ca 100644
--- a/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx
+++ b/Detectors/TPC/calibration/src/CorrectionMapsOptions.cxx
@@ -45,7 +45,7 @@ void CorrectionMapsOptions::addGlobalOptions(std::vector<ConfigParamSpec>& optio
 {
   // these are options which should be added at the workflow level, since they modify the inputs of the devices
   addOption(options, ConfigParamSpec{"lumi-type", o2::framework::VariantType::Int, 0, {"1 = use CTP lumi for TPC correction scaling, 2 = use TPC scalers for TPC correction scaling"}});
-  addOption(options, ConfigParamSpec{"corrmap-lumi-mode", o2::framework::VariantType::Int, 0, {"scaling mode: (default) 0 = static + scale * full; 1 = full + scale * derivative; 2 = full + scale * derivative (for MC)"}});
+  addOption(options, ConfigParamSpec{"corrmap-lumi-mode", o2::framework::VariantType::Int, 0, {"scaling mode: (default) 0 = static + scale * full; 1 = full + scale * derivative; 2 = full + scale * derivative (for MC); 3 = no correction; 4 = static only"}});
   addOption(options, ConfigParamSpec{"enable-M-shape-correction", o2::framework::VariantType::Bool, false, {"Enable M-shape distortion correction"}});
   addOption(options, ConfigParamSpec{"disable-ctp-lumi-request", o2::framework::VariantType::Bool, false, {"do not request CTP lumi (regardless what is used for corrections)"}});
   addOption(options, ConfigParamSpec{"disable-lumi-type-consistency-check", o2::framework::VariantType::Bool, false, {"disable check of selected CTP or IDC scaling source being consistent with the map"}});

From 86a999e410d7959a8a74c9ac16b42e828795c20d Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Fri, 22 May 2026 19:39:53 +0200
Subject: [PATCH 10/17] DPL MCP: allow connecting to a running Hyperloop test

---
 .../scripts/dpl-mcp-server/dpl_mcp_server.py  | 59 +++++++++++++++++--
 1 file changed, 54 insertions(+), 5 deletions(-)

diff --git a/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py b/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py
index dca5058b01dcd..ed457b8a57d9d 100644
--- a/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py
+++ b/Framework/Core/scripts/dpl-mcp-server/dpl_mcp_server.py
@@ -39,7 +39,9 @@
 
 import asyncio
 import json
+import os
 from typing import Any
+from urllib.parse import urlparse
 
 import websockets
 from mcp.server.fastmcp import FastMCP
@@ -51,9 +53,10 @@
 class WorkflowConnection:
     """Holds WebSocket connection and buffered state for one DPL workflow."""
 
-    def __init__(self, port: int, name: str):
-        self.port = port
+    def __init__(self, *, url: str, name: str, extra_headers: dict[str, str] | None = None):
+        self.url = url
         self.name = name
+        self.extra_headers = extra_headers or {}
         self.ws: Any = None
         self.reader_task: asyncio.Task | None = None
         self.snapshot: dict = {}
@@ -83,8 +86,11 @@ async def ensure_connected(self) -> None:
                 except Exception:
                     pass
 
-        url = f"ws://localhost:{self.port}/status"
-        self.ws = await websockets.connect(url, subprotocols=["dpl"])
+        self.ws = await websockets.connect(
+            self.url,
+            subprotocols=["dpl"],
+            additional_headers=self.extra_headers if self.extra_headers else None,
+        )
         if self.reader_task is None or self.reader_task.done():
             self.reader_task = asyncio.create_task(self._reader())
 
@@ -178,7 +184,8 @@ async def connect(port: int = 0, pid: int = 0, name: str = "") -> str:
         old = _workflows[wf_name]
         await old.close()
 
-    conn = WorkflowConnection(port, wf_name)
+    url = f"ws://localhost:{port}/status"
+    conn = WorkflowConnection(url=url, name=wf_name)
     await conn.ensure_connected()
     _workflows[wf_name] = conn
 
@@ -189,6 +196,48 @@ async def connect(port: int = 0, pid: int = 0, name: str = "") -> str:
     )
 
 
+@mcp.tool()
+async def connect_hyperloop(url: str, name: str = "", token: str = "") -> str:
+    """Connect to a DPL workflow running on Hyperloop via the remote proxy.
+
+    Accepts a URL like:
+        https://alimonitor.cern.ch/train-workdir/remote-gui/remote_proxy.html?<token>/<port>
+
+    and remaps it to the local WebSocket proxy endpoint.
+
+    Args:
+        url: The remote_proxy.html URL from alimonitor.
+        name: Optional human-friendly name for this workflow.
+        token: Hyperloop auth token. Falls back to HYPERLOOP_TOKEN env var.
+    """
+    token = token or os.environ.get("HYPERLOOP_TOKEN", "")
+    if not token:
+        return "No token provided and HYPERLOOP_TOKEN environment variable is not set."
+
+    parsed = urlparse(url)
+    path_suffix = parsed.query  # everything after '?'
+    if not path_suffix:
+        return f"Cannot parse token/port from URL: {url}"
+
+    ws_url = f"ws://localhost:8888/remote-mcp/o2/{path_suffix}/status"
+    wf_name = name or path_suffix.split("/")[-1]
+
+    if wf_name in _workflows:
+        old = _workflows[wf_name]
+        await old.close()
+
+    headers = {"Authorization": f"Bearer {token}"}
+    conn = WorkflowConnection(url=ws_url, name=wf_name, extra_headers=headers)
+    await conn.ensure_connected()
+    _workflows[wf_name] = conn
+
+    devices = conn.snapshot.get("devices", [])
+    return (
+        f"Connected to Hyperloop workflow '{wf_name}' via {ws_url} "
+        f"({len(devices)} device(s))."
+    )
+
+
 @mcp.tool()
 async def disconnect(workflow: str) -> str:
     """Disconnect from a DPL workflow and release its resources.

From d13ebbbc141dc9229dcbe3ec69905b167b3d5075 Mon Sep 17 00:00:00 2001
From: shahoian <ruben.shahoyan@cern.ch>
Date: Sat, 23 May 2026 00:39:54 +0200
Subject: [PATCH 11/17] Fix output file name for checkResid streamers

---
 .../GlobalTrackingWorkflow/study/src/CheckResidSpec.cxx      | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/Detectors/GlobalTrackingWorkflow/study/src/CheckResidSpec.cxx b/Detectors/GlobalTrackingWorkflow/study/src/CheckResidSpec.cxx
index 01ec999fce1eb..6a1915791a911 100644
--- a/Detectors/GlobalTrackingWorkflow/study/src/CheckResidSpec.cxx
+++ b/Detectors/GlobalTrackingWorkflow/study/src/CheckResidSpec.cxx
@@ -132,7 +132,7 @@ void CheckResidSpec::init(InitContext& ic)
   int maxLanes = ic.services().get<const o2::framework::DeviceSpec>().maxInputTimeslices;
   std::string nm = params.outname;
   if (maxLanes > 1) {
-    o2::conf::ConfigurableParam::updateFromString(fmt::format("checkresid.outname={}_{}", nm, lane));
+    o2::conf::ConfigurableParam::updateFromString(fmt::format("checkresid.outname={}_t{}", nm, lane));
   }
   if (mDraw) {
     mFillHistos = true;
@@ -173,8 +173,7 @@ void CheckResidSpec::init(InitContext& ic)
   mNThreads = 1;
 #endif
   if (mFillTree) {
-    nm += ".root";
-    mDBGOut = std::make_unique<o2::utils::TreeStreamRedirector>(nm.c_str(), "recreate");
+    mDBGOut = std::make_unique<o2::utils::TreeStreamRedirector>(fmt::format("{}.root", params.outname).c_str(), "recreate");
   }
 }
 

From c36617d3be5c39397aa9373f8b49b07d9e8fd453 Mon Sep 17 00:00:00 2001
From: Matthias Kleiner <matthias.kleiner@cern.ch>
Date: Tue, 19 May 2026 12:09:06 +0200
Subject: [PATCH 12/17] TPC timeseries: make data requests conditional on input
 sources

- do not request PV and FT0 in TPC-only mode
- do not request TPC clusters if TPC is not in input sources
- make time series work without any track input
---
 Detectors/TPC/workflow/src/TPCTimeSeriesSpec.cxx | 9 +++++----
 1 file changed, 5 insertions(+), 4 deletions(-)

diff --git a/Detectors/TPC/workflow/src/TPCTimeSeriesSpec.cxx b/Detectors/TPC/workflow/src/TPCTimeSeriesSpec.cxx
index ac3ff15fd3a29..0c0ae72056318 100644
--- a/Detectors/TPC/workflow/src/TPCTimeSeriesSpec.cxx
+++ b/Detectors/TPC/workflow/src/TPCTimeSeriesSpec.cxx
@@ -1825,15 +1825,16 @@ o2::framework::DataProcessorSpec getTPCTimeSeriesSpec(const bool disableWriter,
   auto dataRequest = std::make_shared<DataRequest>();
   bool useMC = false;
   GTrackID::mask_t srcTracks = GTrackID::getSourcesMask("TPC,ITS,ITS-TPC,ITS-TPC-TRD,ITS-TPC-TOF,ITS-TPC-TRD-TOF") & src;
-  srcTracks.set(GTrackID::TPC); // TPC must be always there
   dataRequest->requestTracks(srcTracks, useMC);
-  dataRequest->requestClusters(GTrackID::getSourcesMask("TPC"), useMC);
+  if (src[GTrackID::TPC]) {
+    dataRequest->requestClusters(GTrackID::getSourcesMask("TPC"), useMC);
+  }
 
   bool tpcOnly = srcTracks == GTrackID::getSourcesMask("TPC");
-  if (!tpcOnly) {
+  if (srcTracks.any() && !tpcOnly) {
     dataRequest->requestFT0RecPoints(useMC);
+    dataRequest->requestPrimaryVertices(useMC);
   }
-  dataRequest->requestPrimaryVertices(useMC);
 
   const bool enableAskMatLUT = matType == o2::base::Propagator::MatCorrType::USEMatCorrLUT;
   auto ccdbRequest = std::make_shared<o2::base::GRPGeomRequest>(!disableWriter,                 // orbitResetTime

From a4bd6bc0c69b2dd18b1073390a8b062002d8e360 Mon Sep 17 00:00:00 2001
From: Matthias Kleiner <Matthias.Kleiner@cern.ch>
Date: Thu, 23 Apr 2026 16:04:23 +0200
Subject: [PATCH 13/17] TPC: move nthreads to local option

---
 .../include/TPCWorkflow/TPCFourierTransformAggregatorSpec.h  | 5 ++++-
 .../TPC/workflow/src/tpc-fouriertransform-aggregator.cxx     | 3 ---
 2 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/Detectors/TPC/workflow/include/TPCWorkflow/TPCFourierTransformAggregatorSpec.h b/Detectors/TPC/workflow/include/TPCWorkflow/TPCFourierTransformAggregatorSpec.h
index 35f51dd489115..7facee78fb3d6 100644
--- a/Detectors/TPC/workflow/include/TPCWorkflow/TPCFourierTransformAggregatorSpec.h
+++ b/Detectors/TPC/workflow/include/TPCWorkflow/TPCFourierTransformAggregatorSpec.h
@@ -64,6 +64,8 @@ class TPCFourierTransformAggregatorSpec : public o2::framework::Task
     mLengthIDCScalerSeconds = ic.options().get<float>("tpcScalerLengthS");
     mDisableScaler = ic.options().get<bool>("disable-scaler");
     mEnableFFTCCDB = ic.options().get<bool>("enable-fft-CCDB");
+    int nthreads = ic.options().get<int>("nthreads");
+    TPCFourierTransformAggregatorSpec::IDCFType::setNThreads(nthreads);
     resizeBuffer(mInputLanes);
   }
 
@@ -448,7 +450,8 @@ DataProcessorSpec getTPCFourierTransformAggregatorSpec(const unsigned int rangeI
             {"dump-coefficients-agg", VariantType::Bool, false, {"Dump fourier coefficients to file"}},
             {"tpcScalerLengthS", VariantType::Float, 300.f, {"Length of the TPC scalers in seconds"}},
             {"disable-scaler", VariantType::Bool, false, {"Disable creation of IDC scaler"}},
-            {"enable-fft-CCDB", VariantType::Bool, false, {"Enable writing of FFT coefficients to CCDB"}}}};
+            {"enable-fft-CCDB", VariantType::Bool, false, {"Enable writing of FFT coefficients to CCDB"}},
+            {"nthreads", VariantType::Int, 1, {"Number of threads which will be used during the calculation of the fourier coefficients."}}}};
 }
 
 } // namespace o2::tpc
diff --git a/Detectors/TPC/workflow/src/tpc-fouriertransform-aggregator.cxx b/Detectors/TPC/workflow/src/tpc-fouriertransform-aggregator.cxx
index b0f09e02e627b..2f66a144251f1 100644
--- a/Detectors/TPC/workflow/src/tpc-fouriertransform-aggregator.cxx
+++ b/Detectors/TPC/workflow/src/tpc-fouriertransform-aggregator.cxx
@@ -26,7 +26,6 @@ void customize(std::vector<ConfigParamSpec>& workflowOptions)
   std::vector<ConfigParamSpec> options{
     {"rangeIDC", VariantType::Int, 200, {"Number of 1D-IDCs which will be used for the calculation of the fourier coefficients. TODO ALREADY SET IN ABERAGEGROUP"}},
     {"nFourierCoeff", VariantType::Int, 60, {"Number of fourier coefficients (real+imag) which will be stored in the CCDB. The maximum can be 'rangeIDC + 2'."}},
-    {"nthreads", VariantType::Int, 1, {"Number of threads which will be used during the calculation of the fourier coefficients."}},
     {"inputLanes", VariantType::Int, 2, {"Number of expected input lanes."}},
     {"sendOutput", VariantType::Bool, false, {"send fourier coefficients"}},
     {"use-naive-fft", VariantType::Bool, false, {"using naive fourier transform (true) or FFTW (false)"}},
@@ -51,8 +50,6 @@ WorkflowSpec defineDataProcessing(ConfigContext const& config)
   const bool processSACs = config.options().get<bool>("process-SACs");
   const auto rangeIDC = static_cast<unsigned int>(config.options().get<int>("rangeIDC"));
   const auto nFourierCoeff = std::clamp(static_cast<unsigned int>(config.options().get<int>("nFourierCoeff")), static_cast<unsigned int>(0), rangeIDC + 2);
-  const auto nthreadsFourier = static_cast<unsigned long>(config.options().get<int>("nthreads"));
-  TPCFourierTransformAggregatorSpec::IDCFType::setNThreads(nthreadsFourier);
   TPCFourierTransformAggregatorSpec::IDCFType::setFFT(!fft);
   const auto inputLanes = config.options().get<int>("inputLanes");
   WorkflowSpec workflow{getTPCFourierTransformAggregatorSpec(rangeIDC, nFourierCoeff, sendOutput, processSACs, inputLanes)};

From 5b4b32f9f0a761711c4e9448877f62f6cd0c2f78 Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Sun, 24 May 2026 00:04:36 +0200
Subject: [PATCH 14/17] DPL: add debug information for rate limiting

---
 Framework/Core/src/ControlWebSocketHandler.cxx |  9 +++++++++
 Framework/Core/src/DevicesManager.cxx          | 11 +++++++++++
 2 files changed, 20 insertions(+)

diff --git a/Framework/Core/src/ControlWebSocketHandler.cxx b/Framework/Core/src/ControlWebSocketHandler.cxx
index 35528a1d6dfec..8be91c0e22fc3 100644
--- a/Framework/Core/src/ControlWebSocketHandler.cxx
+++ b/Framework/Core/src/ControlWebSocketHandler.cxx
@@ -14,10 +14,13 @@
 #include "StatusWebSocketHandler.h"
 #include "Framework/DeviceMetricsHelper.h"
 #include "Framework/ServiceMetricsInfo.h"
+#include "Framework/Signpost.h"
 #include <regex>
 #include "Framework/Logger.h"
 #include "Framework/DeviceConfigInfo.h"
 
+O2_DECLARE_DYNAMIC_LOG(rate_limiting);
+
 namespace o2::framework
 {
 void ControlWebSocketHandler::frame(char const* frame, size_t s)
@@ -74,6 +77,10 @@ void ControlWebSocketHandler::endChunk()
   if (!didProcessMetric) {
     return;
   }
+  O2_SIGNPOST_ID_GENERATE(sid, rate_limiting);
+  O2_SIGNPOST_START(rate_limiting, sid, "endChunk",
+                    "Processing metrics from device %d (had new metric: %d)",
+                    mIndex, (int)didHaveNewMetric);
   size_t timestamp = (uv_hrtime() - mContext.driver->startTime) / 1000000 + mContext.driver->startTimeMsFromEpoch;
   assert(mContext.metrics);
   assert(mContext.infos);
@@ -91,6 +98,8 @@ void ControlWebSocketHandler::endChunk()
   for (auto& metricsInfo : *mContext.metrics) {
     std::fill(metricsInfo.changed.begin(), metricsInfo.changed.end(), false);
   }
+  O2_SIGNPOST_END(rate_limiting, sid, "endChunk",
+                  "Done processing metrics from device %d", mIndex);
 }
 
 void ControlWebSocketHandler::headers(std::map<std::string, std::string> const& headers)
diff --git a/Framework/Core/src/DevicesManager.cxx b/Framework/Core/src/DevicesManager.cxx
index e6fa2c2c61ae6..b427e72ca781d 100644
--- a/Framework/Core/src/DevicesManager.cxx
+++ b/Framework/Core/src/DevicesManager.cxx
@@ -13,12 +13,19 @@
 #include "Framework/RuntimeError.h"
 #include "Framework/Logger.h"
 #include "Framework/DeviceController.h"
+#include "Framework/Signpost.h"
+
+O2_DECLARE_DYNAMIC_LOG(devices_manager);
 
 namespace o2::framework
 {
 
 void DevicesManager::queueMessage(char const* target, char const* message)
 {
+  O2_SIGNPOST_ID_GENERATE(sid, devices_manager);
+  O2_SIGNPOST_EVENT_EMIT(devices_manager, sid, "queue",
+                         "Queuing message for %{public}s: %{public}s",
+                         target, message);
   for (int di = 0; di < specs.size(); ++di) {
     if (specs[di].id == target) {
       messages.push_back({di, message});
@@ -44,6 +51,10 @@ void DevicesManager::flush()
       LOGP(info, "Controller for {} now available.", specs[handle.ref.index].id);
       notifiedAvailable = true;
     }
+    O2_SIGNPOST_ID_GENERATE(sid, devices_manager);
+    O2_SIGNPOST_EVENT_EMIT(devices_manager, sid, "flush",
+                           "Flushing message to %{public}s: %{public}s",
+                           specs[handle.ref.index].id.c_str(), handle.message.c_str());
     controller->write(handle.message.c_str(), handle.message.size());
   }
 

From 19b82ba383c8000a3d1be23cc15d4cb0c36f83b3 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Nicol=C3=B2=20Jacazio?= <njacazio@users.noreply.github.com>
Date: Mon, 25 May 2026 09:14:46 +0200
Subject: [PATCH 15/17] [ALICE3] TRK: fix ACTS clusterer compilation (#15352)

Comment out the processing of digMC2ROFs if condition.
---
 .../ALICE3/TRK/reconstruction/src/ClustererACTS.cxx  | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/Detectors/Upgrades/ALICE3/TRK/reconstruction/src/ClustererACTS.cxx b/Detectors/Upgrades/ALICE3/TRK/reconstruction/src/ClustererACTS.cxx
index 2dbf56ae610e3..30ab503b7e250 100644
--- a/Detectors/Upgrades/ALICE3/TRK/reconstruction/src/ClustererACTS.cxx
+++ b/Detectors/Upgrades/ALICE3/TRK/reconstruction/src/ClustererACTS.cxx
@@ -387,10 +387,10 @@ void ClustererACTS::process(gsl::span<const Digit> digits,
                              outFirst, static_cast<int>(clusters.size()) - outFirst);
   }
 
-  if (clusterMC2ROFs && !digMC2ROFs.empty()) {
-    clusterMC2ROFs->reserve(clusterMC2ROFs->size() + digMC2ROFs.size());
-    for (const auto& in : digMC2ROFs) {
-      clusterMC2ROFs->emplace_back(in.eventRecordID, in.rofRecordID, in.minROF, in.maxROF);
-    }
-  }
+  // if (clusterMC2ROFs && !digMC2ROFs.empty()) {
+  //   clusterMC2ROFs->reserve(clusterMC2ROFs->size() + digMC2ROFs.size());
+  //   for (const auto& in : digMC2ROFs) {
+  //     clusterMC2ROFs->emplace_back(in.eventRecordID, in.rofRecordID, in.minROF, in.maxROF);
+  //   }
+  // }
 }

From 8936d29ea8fdee5085990361442be87c809ab99c Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Mon, 25 May 2026 10:01:47 +0200
Subject: [PATCH 16/17] DPL: fix mismatched type in signpost

---
 Framework/Core/src/ControlWebSocketHandler.cxx | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/Framework/Core/src/ControlWebSocketHandler.cxx b/Framework/Core/src/ControlWebSocketHandler.cxx
index 8be91c0e22fc3..8d2f85b034364 100644
--- a/Framework/Core/src/ControlWebSocketHandler.cxx
+++ b/Framework/Core/src/ControlWebSocketHandler.cxx
@@ -79,7 +79,7 @@ void ControlWebSocketHandler::endChunk()
   }
   O2_SIGNPOST_ID_GENERATE(sid, rate_limiting);
   O2_SIGNPOST_START(rate_limiting, sid, "endChunk",
-                    "Processing metrics from device %d (had new metric: %d)",
+                    "Processing metrics from device %zu (had new metric: %d)",
                     mIndex, (int)didHaveNewMetric);
   size_t timestamp = (uv_hrtime() - mContext.driver->startTime) / 1000000 + mContext.driver->startTimeMsFromEpoch;
   assert(mContext.metrics);
@@ -99,7 +99,7 @@ void ControlWebSocketHandler::endChunk()
     std::fill(metricsInfo.changed.begin(), metricsInfo.changed.end(), false);
   }
   O2_SIGNPOST_END(rate_limiting, sid, "endChunk",
-                  "Done processing metrics from device %d", mIndex);
+                  "Done processing metrics from device %zu", mIndex);
 }
 
 void ControlWebSocketHandler::headers(std::map<std::string, std::string> const& headers)

From 440a899dee1673762ec344cad73735ecf44ad57d Mon Sep 17 00:00:00 2001
From: Giulio Eulisse <10544+ktf@users.noreply.github.com>
Date: Mon, 25 May 2026 09:33:52 +0200
Subject: [PATCH 17/17] DPL: reduce number of spurious warnings

---
 Framework/Core/src/DataProcessingDevice.cxx | 22 +++++++++++++++------
 1 file changed, 16 insertions(+), 6 deletions(-)

diff --git a/Framework/Core/src/DataProcessingDevice.cxx b/Framework/Core/src/DataProcessingDevice.cxx
index be25133158072..b45a48c28f691 100644
--- a/Framework/Core/src/DataProcessingDevice.cxx
+++ b/Framework/Core/src/DataProcessingDevice.cxx
@@ -1420,14 +1420,24 @@ void DataProcessingDevice::Run()
           }
           return missingInfo.empty() ? std::string(" (policy: ") + spec.resourcePolicy.name + ")" : " -" + missingInfo;
         };
+        auto const timeSinceLastScheduled = lastSched ? uv_now(state.loop) - lastSched : 0;
         if (schedulingStats.numberOfUnscheduledSinceLastScheduled >= schedulingStats.nextWarnAt) {
           auto const missingStr = buildMissingInfo();
-          O2_SIGNPOST_EVENT_EMIT_WARN(scheduling, sid, "Run",
-                                      "Not enough resources to schedule computation on stream %d. %zu consecutive skips%s. Missing:%s. Data is not lost and it will be scheduled again.",
-                                      streamRef.index,
-                                      schedulingStats.numberOfUnscheduledSinceLastScheduled.load(),
-                                      schedInfo.c_str(),
-                                      missingStr.c_str());
+          if (timeSinceLastScheduled >= 50) {
+            O2_SIGNPOST_EVENT_EMIT_WARN(scheduling, sid, "Run",
+                                        "Not enough resources to schedule computation on stream %d. %zu consecutive skips%s. Missing:%s. Data is not lost and it will be scheduled again.",
+                                        streamRef.index,
+                                        schedulingStats.numberOfUnscheduledSinceLastScheduled.load(),
+                                        schedInfo.c_str(),
+                                        missingStr.c_str());
+          } else {
+            O2_SIGNPOST_EVENT_EMIT(scheduling, sid, "Run",
+                                   "Not enough resources to schedule computation on stream %d. %zu consecutive skips%s. Missing:%s. Data is not lost and it will be scheduled again.",
+                                   streamRef.index,
+                                   schedulingStats.numberOfUnscheduledSinceLastScheduled.load(),
+                                   schedInfo.c_str(),
+                                   missingStr.c_str());
+          }
           schedulingStats.nextWarnAt = schedulingStats.nextWarnAt * 2;
         } else {
           auto const missingStr = buildMissingInfo();