diff --git a/Detectors/GlobalTrackingWorkflow/study/src/SVStudy.cxx b/Detectors/GlobalTrackingWorkflow/study/src/SVStudy.cxx
index c28048a1f9503..1e141a29d3f55 100644
--- a/Detectors/GlobalTrackingWorkflow/study/src/SVStudy.cxx
+++ b/Detectors/GlobalTrackingWorkflow/study/src/SVStudy.cxx
@@ -48,6 +48,7 @@
 // #include "GPUSettingsO2.h"
 #include "GPUParam.h"
 #include "GPUParam.inc"
+#include "GPUTPCGeometry.h"
 #include "GPUO2InterfaceRefit.h"
 #include "GPUO2InterfaceUtils.h"
 
@@ -254,7 +255,7 @@ o2::dataformats::V0Ext SVStudySpec::processV0(int iv, o2::globaltracking::RecoCo
         tpcTr.getClusterReference(clRefs, tpcTr.getNClusterReferences() - 1, clSect, clRow, clIdx);
         const auto& clus = recoData.getTPCClusters().clusters[clSect][clRow][clIdx];
         prInfo.lowestRow = clRow;
-        int npads = mParam->tpcGeometry.NPads(clRow);
+        int npads = o2::gpu::GPUTPCGeometry::NPads(clRow);
         prInfo.padFromEdge = uint8_t(clus.getPad());
         if (prInfo.padFromEdge > npads / 2) {
           prInfo.padFromEdge = npads - 1 - prInfo.padFromEdge;
diff --git a/Detectors/GlobalTrackingWorkflow/study/src/TrackingStudy.cxx b/Detectors/GlobalTrackingWorkflow/study/src/TrackingStudy.cxx
index db57ad5f8a7eb..f206c43f7f57a 100644
--- a/Detectors/GlobalTrackingWorkflow/study/src/TrackingStudy.cxx
+++ b/Detectors/GlobalTrackingWorkflow/study/src/TrackingStudy.cxx
@@ -47,6 +47,7 @@
 #include "GPUO2Interface.h" // Needed for propper settings in GPUParam.h
 #include "GPUParam.h"
 #include "GPUParam.inc"
+#include "GPUTPCGeometry.h"
 #include "Steer/MCKinematicsReader.h"
 #include "MathUtils/fit.h"
 #include <TF1.h>
@@ -301,7 +302,7 @@ void TrackingStudySpec::process(o2::globaltracking::RecoContainer& recoData)
       trExt.rowMinTPC = clRow;
       const auto& clus = tpcClusAcc.clusters[clSect][clRow][clIdx];
       trExt.padFromEdge = uint8_t(clus.getPad());
-      int npads = mTPCRefitter->getParam()->tpcGeometry.NPads(clRow);
+      int npads = o2::gpu::GPUTPCGeometry::NPads(clRow);
       if (trExt.padFromEdge > npads / 2) {
         trExt.padFromEdge = npads - 1 - trExt.padFromEdge;
       }
diff --git a/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx b/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx
index cc964ade0d87c..2efa7077be125 100644
--- a/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx
+++ b/Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx
@@ -230,9 +230,9 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
       int myThread = 0;
 #endif
       unsigned int count = 0;
-      const float x = mParam->tpcGeometry.Row2X(j);
+      const float x = GPUTPCGeometry::Row2X(j);
       auto checker = [i, j, firstIR, totalT, x, this, &preCl, &count, &outBuffer = tmpBuffer[myThread], &rejectHits, &clustersFiltered](const o2::tpc::ClusterNative& cl, unsigned int k) {
-        const float y = mParam->tpcGeometry.LinearPad2Y(i, j, cl.getPad());
+        const float y = GPUTPCGeometry::LinearPad2Y(i, j, cl.getPad());
         const float r = sqrtf(x * x + y * y);
         const float maxz = r * mEtaFactor + mMaxZ;
         const unsigned int deltaBC = std::max<float>(0.f, totalT - mFastTransform->convDeltaZtoDeltaTimeInTimeFrameAbs(maxz)) * constants::LHCBCPERTIMEBIN;
diff --git a/GPU/GPUTracking/Base/GPUParam.cxx b/GPU/GPUTracking/Base/GPUParam.cxx
index 192e46c36dc68..bbca150df405a 100644
--- a/GPU/GPUTracking/Base/GPUParam.cxx
+++ b/GPU/GPUTracking/Base/GPUParam.cxx
@@ -21,6 +21,7 @@
 #include "GPUDataTypes.h"
 #include "GPUConstantMem.h"
 #include "DetectorsBase/Propagator.h"
+#include "GPUTPCGeometry.h"
 
 using namespace o2::gpu;
 
@@ -32,7 +33,6 @@ using namespace o2::gpu;
 void GPUParam::SetDefaults(float solenoidBz)
 {
   memset((void*)this, 0, sizeof(*this));
-  new (&tpcGeometry) GPUTPCGeometry;
   new (&rec) GPUSettingsRec;
   occupancyMap = nullptr;
   occupancyTotal = 0;
@@ -178,8 +178,8 @@ void GPUParam::UpdateRun3ClusterErrors(const float* yErrorParam, const float* zE
     for (int32_t rowType = 0; rowType < 4; rowType++) {
       constexpr int32_t regionMap[4] = {0, 4, 6, 8};
       ParamErrors[yz][rowType][0] = param[0] * param[0];
-      ParamErrors[yz][rowType][1] = param[1] * param[1] * tpcGeometry.PadHeightByRegion(regionMap[rowType]);
-      ParamErrors[yz][rowType][2] = param[2] * param[2] / tpcGeometry.TPCLength() / tpcGeometry.PadHeightByRegion(regionMap[rowType]);
+      ParamErrors[yz][rowType][1] = param[1] * param[1] * GPUTPCGeometry::PadHeightByRegion(regionMap[rowType]);
+      ParamErrors[yz][rowType][2] = param[2] * param[2] / GPUTPCGeometry::TPCLength() / GPUTPCGeometry::PadHeightByRegion(regionMap[rowType]);
       ParamErrors[yz][rowType][3] = param[3] * param[3] * rec.tpc.clusterErrorOccupancyScaler * rec.tpc.clusterErrorOccupancyScaler;
     }
   }
diff --git a/GPU/GPUTracking/Base/GPUParam.h b/GPU/GPUTracking/Base/GPUParam.h
index fbce6246de112..4b77628c88775 100644
--- a/GPU/GPUTracking/Base/GPUParam.h
+++ b/GPU/GPUTracking/Base/GPUParam.h
@@ -19,7 +19,6 @@
 #include "GPUCommonMath.h"
 #include "GPUDef.h"
 #include "GPUSettings.h"
-#include "GPUTPCGeometry.h"
 #include "GPUTPCGMPolynomialField.h"
 
 #if !defined(GPUCA_GPUCODE)
@@ -59,7 +58,6 @@ struct GPUParam_t {
   int32_t continuousMaxTimeBin;
   int32_t tpcCutTimeBin;
 
-  GPUTPCGeometry tpcGeometry;              // TPC Geometry
   GPUTPCGMPolynomialField polynomialField; // Polynomial approx. of magnetic field for TPC GM
   const uint32_t* occupancyMap;            // Ptr to TPC occupancy map
   uint32_t occupancyTotal;                 // Total occupancy in the TPC (nCl / nHbf)
diff --git a/GPU/GPUTracking/Base/GPUParam.inc b/GPU/GPUTracking/Base/GPUParam.inc
index 19dc1fc4a3578..a118a8f639fe9 100644
--- a/GPU/GPUTracking/Base/GPUParam.inc
+++ b/GPU/GPUTracking/Base/GPUParam.inc
@@ -18,6 +18,7 @@
 #include "GPUParam.h"
 #include "GPUTPCGMMergedTrackHit.h"
 #include "GPUTPCClusterOccupancyMap.h"
+#include "GPUTPCGeometry.h"
 
 namespace o2::gpu
 {
@@ -42,14 +43,14 @@ GPUdi() void GPUParam::Global2Sector(int32_t iSector, float X, float Y, float Z,
 
 GPUdi() void GPUParam::GetClusterErrorsSeeding2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const
 {
-  const int32_t rowType = tpcGeometry.GetROC(iRow);
-  z = CAMath::Abs(tpcGeometry.TPCLength() - CAMath::Abs(z));
+  const int32_t rowType = GPUTPCGeometry::GetROC(iRow);
+  z = CAMath::Abs(GPUTPCGeometry::TPCLength() - CAMath::Abs(z));
   const float s2 = CAMath::Min(sinPhi * sinPhi, 0.95f * 0.95f);
   const float sec2 = 1.f / (1.f - s2);
   const float angleY2 = s2 * sec2;          // dy/dx
   const float angleZ2 = DzDs * DzDs * sec2; // dz/dx
 
-  const float unscaledMult = time >= 0.f ? GetUnscaledMult(time) / tpcGeometry.Row2X(iRow) : 0.f;
+  const float unscaledMult = time >= 0.f ? GetUnscaledMult(time) / GPUTPCGeometry::Row2X(iRow) : 0.f;
 
   ErrY2 = GetClusterErrorSeeding(0, rowType, z, angleY2, unscaledMult); // Returns Err2
   ErrZ2 = GetClusterErrorSeeding(1, rowType, z, angleZ2, unscaledMult); // Returns Err2
@@ -132,8 +133,8 @@ GPUdi() float GPUParam::GetClusterErrorSeeding(int32_t yz, int32_t type, float z
 
 GPUdi() void GPUParam::GetClusterErrorsSeeding2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const
 {
-  int32_t rowType = tpcGeometry.GetROC(iRow);
-  z = CAMath::Abs(tpcGeometry.TPCLength() - CAMath::Abs(z));
+  int32_t rowType = GPUTPCGeometry::GetROC(iRow);
+  z = CAMath::Abs(GPUTPCGeometry::TPCLength() - CAMath::Abs(z));
   const float s2 = CAMath::Min(sinPhi * sinPhi, 0.95f * 0.95f);
   float sec2 = 1.f / (1.f - s2);
   float angleY2 = s2 * sec2;          // dy/dx
@@ -172,14 +173,14 @@ GPUdi() float GPUParam::GetSystematicClusterErrorC122(float trackX, float trackY
 
 GPUdi() void GPUParam::GetClusterErrors2(uint8_t sector, int32_t iRow, float z, float sinPhi, float DzDs, float time, float avgInvCharge, float invCharge, float& ErrY2, float& ErrZ2) const
 {
-  const int32_t rowType = tpcGeometry.GetROC(iRow);
-  z = CAMath::Abs(tpcGeometry.TPCLength() - CAMath::Abs(z));
+  const int32_t rowType = GPUTPCGeometry::GetROC(iRow);
+  z = CAMath::Abs(GPUTPCGeometry::TPCLength() - CAMath::Abs(z));
   const float s2 = CAMath::Min(sinPhi * sinPhi, 0.95f * 0.95f);
   const float sec2 = 1.f / (1.f - s2);
   const float angleY2 = s2 * sec2;          // dy/dx
   const float angleZ2 = DzDs * DzDs * sec2; // dz/dx
 
-  const float unscaledMult = time >= 0.f ? GetUnscaledMult(time) / tpcGeometry.Row2X(iRow) : 0.f;
+  const float unscaledMult = time >= 0.f ? GetUnscaledMult(time) / GPUTPCGeometry::Row2X(iRow) : 0.f;
   const float scaledInvAvgCharge = avgInvCharge * rec.tpc.clusterErrorChargeScaler > 0.f ? avgInvCharge * rec.tpc.clusterErrorChargeScaler : 1.f;
   const float scaledInvCharge = invCharge * rec.tpc.clusterErrorChargeScaler > 0.f ? invCharge * rec.tpc.clusterErrorChargeScaler : 1.f;
 
@@ -218,7 +219,7 @@ GPUdi() float GPUParam::GetUnscaledMult(float time) const
 
 GPUdi() bool GPUParam::rejectEdgeClusterByY(float uncorrectedY, int32_t iRow, float trackSigmaY) const
 {
-  return CAMath::Abs(uncorrectedY) > (tpcGeometry.NPads(iRow) - 1) * 0.5f * tpcGeometry.PadWidth(iRow) + rec.tpc.rejectEdgeClustersMargin + trackSigmaY * rec.tpc.rejectEdgeClustersSigmaMargin;
+  return CAMath::Abs(uncorrectedY) > (GPUTPCGeometry::NPads(iRow) - 1) * 0.5f * GPUTPCGeometry::PadWidth(iRow) + rec.tpc.rejectEdgeClustersMargin + trackSigmaY * rec.tpc.rejectEdgeClustersSigmaMargin;
 }
 
 } // namespace o2::gpu
diff --git a/GPU/GPUTracking/Base/GPUReconstructionConvert.cxx b/GPU/GPUTracking/Base/GPUReconstructionConvert.cxx
index bc760f6188caa..e12ca7ec601ad 100644
--- a/GPU/GPUTracking/Base/GPUReconstructionConvert.cxx
+++ b/GPU/GPUTracking/Base/GPUReconstructionConvert.cxx
@@ -24,6 +24,7 @@
 #include "GPUTPCClusterData.h"
 #include "GPUO2DataTypes.h"
 #include "GPUDataTypes.h"
+#include "GPUTPCGeometry.h"
 #include "AliHLTTPCRawCluster.h"
 #include "GPUParam.h"
 #include "GPULogging.h"
@@ -275,10 +276,10 @@ struct zsEncoderRow : public zsEncoder {
 
 inline bool zsEncoderRow::sort(const o2::tpc::Digit a, const o2::tpc::Digit b)
 {
-  int32_t endpointa = param->tpcGeometry.GetRegion(a.getRow());
-  int32_t endpointb = param->tpcGeometry.GetRegion(b.getRow());
-  endpointa = 2 * endpointa + (a.getRow() >= param->tpcGeometry.GetRegionStart(endpointa) + param->tpcGeometry.GetRegionRows(endpointa) / 2);
-  endpointb = 2 * endpointb + (b.getRow() >= param->tpcGeometry.GetRegionStart(endpointb) + param->tpcGeometry.GetRegionRows(endpointb) / 2);
+  int32_t endpointa = GPUTPCGeometry::GetRegion(a.getRow());
+  int32_t endpointb = GPUTPCGeometry::GetRegion(b.getRow());
+  endpointa = 2 * endpointa + (a.getRow() >= GPUTPCGeometry::GetRegionStart(endpointa) + GPUTPCGeometry::GetRegionRows(endpointa) / 2);
+  endpointb = 2 * endpointb + (b.getRow() >= GPUTPCGeometry::GetRegionStart(endpointb) + GPUTPCGeometry::GetRegionRows(endpointb) / 2);
   if (endpointa != endpointb) {
     return endpointa <= endpointb;
   }
@@ -295,11 +296,11 @@ bool zsEncoderRow::checkInput(std::vector<o2::tpc::Digit>& tmpBuffer, uint32_t k
 {
   seqLen = 1;
   if (lastRow != tmpBuffer[k].getRow()) {
-    endpointStart = param->tpcGeometry.GetRegionStart(curRegion);
+    endpointStart = GPUTPCGeometry::GetRegionStart(curRegion);
     endpoint = curRegion * 2;
-    if (tmpBuffer[k].getRow() >= endpointStart + param->tpcGeometry.GetRegionRows(curRegion) / 2) {
+    if (tmpBuffer[k].getRow() >= endpointStart + GPUTPCGeometry::GetRegionRows(curRegion) / 2) {
       endpoint++;
-      endpointStart += param->tpcGeometry.GetRegionRows(curRegion) / 2;
+      endpointStart += GPUTPCGeometry::GetRegionRows(curRegion) / 2;
     }
   }
   for (uint32_t l = k + 1; l < tmpBuffer.size(); l++) {
@@ -408,7 +409,7 @@ void zsEncoderRow::decodePage(std::vector<o2::tpc::Digit>& outputBuffer, const z
   if ((uint32_t)region != decEndpoint / 2) {
     throw std::runtime_error("CRU ID / endpoint mismatch");
   }
-  int32_t nRowsRegion = param->tpcGeometry.GetRegionRows(region);
+  int32_t nRowsRegion = GPUTPCGeometry::GetRegionRows(region);
 
   int32_t timeBin = (decHDR->timeOffset + (uint64_t)(o2::raw::RDHUtils::getHeartBeatOrbit(*rdh) - firstOrbit) * o2::constants::lhc::LHCMaxBunches) / LHCBCPERTIMEBIN;
   for (int32_t l = 0; l < decHDR->nTimeBinSpan; l++) {
@@ -420,7 +421,7 @@ void zsEncoderRow::decodePage(std::vector<o2::tpc::Digit>& outputBuffer, const z
     if (tbHdr->rowMask != 0 && ((upperRows) ^ ((decEndpoint & 1) != 0))) {
       throw std::runtime_error("invalid endpoint");
     }
-    const int32_t rowOffset = param->tpcGeometry.GetRegionStart(region) + (upperRows ? (nRowsRegion / 2) : 0);
+    const int32_t rowOffset = GPUTPCGeometry::GetRegionStart(region) + (upperRows ? (nRowsRegion / 2) : 0);
     const int32_t nRows = upperRows ? (nRowsRegion - nRowsRegion / 2) : (nRowsRegion / 2);
     const int32_t nRowsUsed = __builtin_popcount((uint32_t)(tbHdr->rowMask & 0x7FFF));
     decPagePtr += nRowsUsed ? (2 * nRowsUsed) : 2;
@@ -513,7 +514,7 @@ void zsEncoderLinkBased::createBitmask(std::vector<o2::tpc::Digit>& tmpBuffer, u
   uint32_t l;
   for (l = k; l < tmpBuffer.size(); l++) {
     const auto& a = tmpBuffer[l];
-    int32_t cruinsector = param->tpcGeometry.GetRegion(a.getRow());
+    int32_t cruinsector = GPUTPCGeometry::GetRegion(a.getRow());
     o2::tpc::GlobalPadNumber pad = mapper.globalPadNumber(o2::tpc::PadPos(a.getRow(), a.getPad()));
     o2::tpc::FECInfo fec = mapper.fecInfo(pad);
     o2::tpc::CRU cru = cruinsector;
@@ -535,8 +536,8 @@ void zsEncoderLinkBased::createBitmask(std::vector<o2::tpc::Digit>& tmpBuffer, u
 bool zsEncoderLinkBased::sort(const o2::tpc::Digit a, const o2::tpc::Digit b)
 {
   // Fixme: this is blasphemy... one shoult precompute all values and sort an index array
-  int32_t cruinsectora = param->tpcGeometry.GetRegion(a.getRow());
-  int32_t cruinsectorb = param->tpcGeometry.GetRegion(b.getRow());
+  int32_t cruinsectora = GPUTPCGeometry::GetRegion(a.getRow());
+  int32_t cruinsectorb = GPUTPCGeometry::GetRegion(b.getRow());
   if (cruinsectora != cruinsectorb) {
     return cruinsectora < cruinsectorb;
   }
@@ -1124,7 +1125,7 @@ inline uint32_t zsEncoderRun<T>::run(std::vector<zsPage>* buffer, std::vector<o2
           }
         }
         if (lastRow != tmpBuffer[k].getRow()) {
-          curRegion = param->tpcGeometry.GetRegion(tmpBuffer[k].getRow());
+          curRegion = GPUTPCGeometry::GetRegion(tmpBuffer[k].getRow());
         }
         mustWriteSubPage = checkInput(tmpBuffer, k);
       } else {
diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt
index d5a90dbd65ea3..5ce96d450f765 100644
--- a/GPU/GPUTracking/CMakeLists.txt
+++ b/GPU/GPUTracking/CMakeLists.txt
@@ -10,7 +10,6 @@
 # or submit itself to any jurisdiction.
 
 set(MODULE GPUTracking)
-cmake_minimum_required(VERSION 3.27 FATAL_ERROR)
 
 # set(CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER} "${CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -O0") # to uncomment if needed, tired of typing this...
 # set(GPUCA_BUILD_DEBUG 1)
diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx
index 4831be9b12bcc..445c03113cd39 100644
--- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx
+++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx
@@ -18,7 +18,6 @@
 #include "GPUParam.h"
 #include "GPUCommonAlgorithm.h"
 #include "GPUTPCCompressionTrackModel.h"
-#include "GPUTPCGeometry.h"
 #include "GPUTPCClusterRejection.h"
 #include "GPUTPCCompressionKernels.inc"
 
@@ -68,9 +67,10 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
         continue; // No track model compression
       }
       const ClusterNative& GPUrestrict() orgCl = clusters -> clusters[hit.sector][hit.row][hit.num - clusters->clusterOffset[hit.sector][hit.row]];
-      float x = param.tpcGeometry.Row2X(hit.row);
-      float y = track.LinearPad2Y(hit.sector, orgCl.getPad(), param.tpcGeometry.PadWidth(hit.row), param.tpcGeometry.NPads(hit.row));
-      float z = param.tpcGeometry.LinearTime2Z(hit.sector, orgCl.getTime());
+      constexpr GPUTPCGeometry geo;
+      float x = geo.Row2X(hit.row);
+      float y = track.LinearPad2Y(hit.sector, orgCl.getPad(), geo.PadWidth(hit.row), geo.NPads(hit.row));
+      float z = geo.LinearTime2Z(hit.sector, orgCl.getTime());
       if (nClustersStored) {
         if ((hit.sector < GPUCA_NSECTORS) ^ (lastSector < GPUCA_NSECTORS)) {
           break;
@@ -78,7 +78,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
         if (lastLeg != hit.leg && track.Mirror()) {
           break;
         }
-        if (track.Propagate(param.tpcGeometry.Row2X(hit.row), param.SectorParam[hit.sector].Alpha)) {
+        if (track.Propagate(geo.Row2X(hit.row), param.SectorParam[hit.sector].Alpha)) {
           break;
         }
       }
@@ -115,9 +115,9 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step0at
         }
         c.rowDiffA[cidx] = row;
         c.sliceLegDiffA[cidx] = (hit.leg == lastLeg ? 0 : compressor.NSECTORS) + sector;
-        float pad = CAMath::Max(0.f, CAMath::Min((float)param.tpcGeometry.NPads(GPUCA_ROW_COUNT - 1), track.LinearY2Pad(hit.sector, track.Y(), param.tpcGeometry.PadWidth(hit.row), param.tpcGeometry.NPads(hit.row))));
+        float pad = CAMath::Max(0.f, CAMath::Min((float)geo.NPads(GPUCA_ROW_COUNT - 1), track.LinearY2Pad(hit.sector, track.Y(), geo.PadWidth(hit.row), geo.NPads(hit.row))));
         c.padResA[cidx] = orgCl.padPacked - orgCl.packPad(pad);
-        float time = CAMath::Max(0.f, param.tpcGeometry.LinearZ2Time(hit.sector, track.Z() + zOffset));
+        float time = CAMath::Max(0.f, geo.LinearZ2Time(hit.sector, track.Z() + zOffset));
         c.timeResA[cidx] = (orgCl.getTimePacked() - orgCl.packTime(time)) & 0xFFFFFF;
         lastLeg = hit.leg;
       }
diff --git a/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc b/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc
index 6c4f70d7c6884..e4de5a503e52d 100644
--- a/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc
+++ b/GPU/GPUTracking/DataCompression/TPCClusterDecompressionCore.inc
@@ -20,6 +20,7 @@
 #include "GPUTPCCompressionTrackModel.h"
 #include "GPUCommonAlgorithm.h"
 #include "GPUO2DataTypes.h"
+#include "GPUTPCGeometry.h"
 
 #ifndef GPUCA_GPUCODE
 #include <functional>
@@ -80,6 +81,7 @@ class TPCClusterDecompressionCore
     uint32_t slice = cmprClusters.sliceA[trackIndex];
     uint32_t row = cmprClusters.rowA[trackIndex];
     GPUTPCCompressionTrackModel track;
+    constexpr GPUTPCGeometry geo;
     uint32_t clusterIndex;
     for (clusterIndex = 0; clusterIndex < cmprClusters.nTrackClusters[trackIndex]; clusterIndex++) {
       uint32_t pad = 0, time = 0;
@@ -105,23 +107,23 @@ class TPCClusterDecompressionCore
         if (changeLeg && track.Mirror()) {
           break;
         }
-        if (track.Propagate(param.tpcGeometry.Row2X(row), param.SectorParam[slice].Alpha)) {
+        if (track.Propagate(geo.Row2X(row), param.SectorParam[slice].Alpha)) {
           break;
         }
         uint32_t timeTmp = cmprClusters.timeResA[clusterOffset - trackIndex - 1];
         if (timeTmp & 800000) {
           timeTmp |= 0xFF000000;
         }
-        time = timeTmp + ClusterNative::packTime(CAMath::Max(0.f, param.tpcGeometry.LinearZ2Time(slice, track.Z() + zOffset)));
-        float tmpPad = CAMath::Max(0.f, CAMath::Min((float)param.tpcGeometry.NPads(GPUCA_ROW_COUNT - 1), track.LinearY2Pad(slice, track.Y(), param.tpcGeometry.PadWidth(row), param.tpcGeometry.NPads(row))));
+        time = timeTmp + ClusterNative::packTime(CAMath::Max(0.f, geo.LinearZ2Time(slice, track.Z() + zOffset)));
+        float tmpPad = CAMath::Max(0.f, CAMath::Min((float)geo.NPads(GPUCA_ROW_COUNT - 1), track.LinearY2Pad(slice, track.Y(), geo.PadWidth(row), geo.NPads(row))));
         pad = cmprClusters.padResA[clusterOffset - trackIndex - 1] + ClusterNative::packPad(tmpPad);
         time = time & 0xFFFFFF;
         pad = (uint16_t)pad;
-        if (pad >= param.tpcGeometry.NPads(row) * ClusterNative::scalePadPacked) {
+        if (pad >= geo.NPads(row) * ClusterNative::scalePadPacked) {
           if (pad >= 0xFFFF - 11968) { // Constant 11968 = (2^15 - MAX_PADS(138) * scalePadPacked(64)) / 2
             pad = 0;
           } else {
-            pad = param.tpcGeometry.NPads(row) * ClusterNative::scalePadPacked - 1;
+            pad = geo.NPads(row) * ClusterNative::scalePadPacked - 1;
           }
         }
         if (param.continuousMaxTimeBin > 0 && time >= maxTime) {
@@ -136,11 +138,11 @@ class TPCClusterDecompressionCore
         pad = cmprClusters.padA[trackIndex];
       }
       const auto cluster = decompressTrackStore(cmprClusters, clusterOffset, slice, row, pad, time, args...);
-      float y = track.LinearPad2Y(slice, cluster.getPad(), param.tpcGeometry.PadWidth(row), param.tpcGeometry.NPads(row));
-      float z = param.tpcGeometry.LinearTime2Z(slice, cluster.getTime());
+      float y = track.LinearPad2Y(slice, cluster.getPad(), geo.PadWidth(row), geo.NPads(row));
+      float z = geo.LinearTime2Z(slice, cluster.getTime());
       if (clusterIndex == 0) {
         zOffset = z;
-        track.Init(param.tpcGeometry.Row2X(row), y, z - zOffset, param.SectorParam[slice].Alpha, cmprClusters.qPtA[trackIndex], param);
+        track.Init(geo.Row2X(row), y, z - zOffset, param.SectorParam[slice].Alpha, cmprClusters.qPtA[trackIndex], param);
       }
       if (clusterIndex + 1 < cmprClusters.nTrackClusters[trackIndex] && track.Filter(y, z - zOffset, row)) {
         break;
diff --git a/GPU/GPUTracking/DataTypes/TPCPadBitMap.cxx b/GPU/GPUTracking/DataTypes/TPCPadBitMap.cxx
index 40ce8c0ccda81..0b8e67fbe495e 100644
--- a/GPU/GPUTracking/DataTypes/TPCPadBitMap.cxx
+++ b/GPU/GPUTracking/DataTypes/TPCPadBitMap.cxx
@@ -21,7 +21,7 @@ using namespace o2::gpu;
 
 TPCPadBitMap::TPCPadBitMap()
 {
-  GPUTPCGeometry geo{};
+  constexpr GPUTPCGeometry geo;
   int32_t offset = 0;
   for (int32_t r = 0; r < GPUCA_ROW_COUNT; r++) {
     mPadOffsetPerRow[r] = offset;
diff --git a/GPU/GPUTracking/DataTypes/TPCPadGainCalib.cxx b/GPU/GPUTracking/DataTypes/TPCPadGainCalib.cxx
index 41f0ad819d1b6..a20f3dc8aac1d 100644
--- a/GPU/GPUTracking/DataTypes/TPCPadGainCalib.cxx
+++ b/GPU/GPUTracking/DataTypes/TPCPadGainCalib.cxx
@@ -21,7 +21,7 @@ using namespace o2::gpu;
 
 TPCPadGainCalib::TPCPadGainCalib()
 {
-  GPUTPCGeometry geo{};
+  constexpr GPUTPCGeometry geo{};
   int32_t offset = 0;
   for (int32_t r = 0; r < GPUCA_ROW_COUNT; r++) {
     mPadOffsetPerRow[r] = offset;
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
index 1c2a8e2b29a9c..f373d56ea0395 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx
@@ -36,6 +36,7 @@
 #include "GPUO2DataTypes.h"
 #include "TPCFastTransform.h"
 #include "GPUTPCConvertImpl.h"
+#include "GPUTPCGeometry.h"
 
 #include "GPUCommonMath.h"
 #include "GPUCommonAlgorithm.h"
@@ -601,13 +602,13 @@ GPUd() void GPUTPCGMMerger::MergeSectorsPrepareStep2(int32_t nBlocks, int32_t nT
   } else if (iBorder == 1) { // transport to the right edge of the sector and rotate horizontally
     dAlpha = -dAlpha - CAMath::Pi() / 2;
   } else if (iBorder == 2) { // transport to the middle of the sector and rotate vertically to the border on the left
-    x0 = Param().tpcGeometry.Row2X(63);
+    x0 = GPUTPCGeometry::Row2X(63);
   } else if (iBorder == 3) { // transport to the middle of the sector and rotate vertically to the border on the right
     dAlpha = -dAlpha;
-    x0 = Param().tpcGeometry.Row2X(63);
+    x0 = GPUTPCGeometry::Row2X(63);
   } else if (iBorder == 4) { // transport to the middle of the sßector, w/o rotation
     dAlpha = 0;
-    x0 = Param().tpcGeometry.Row2X(63);
+    x0 = GPUTPCGeometry::Row2X(63);
   }
 
   const float maxSin = CAMath::Sin(60.f / 180.f * CAMath::Pi());
@@ -955,7 +956,7 @@ template GPUdni() void GPUTPCGMMerger::MergeBorderTracks<2>(int32_t nBlocks, int
 
 GPUd() void GPUTPCGMMerger::MergeWithinSectorsPrepare(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread)
 {
-  float x0 = Param().tpcGeometry.Row2X(63);
+  float x0 = GPUTPCGeometry::Row2X(63);
   const float maxSin = CAMath::Sin(60.f / 180.f * CAMath::Pi());
 
   for (int32_t itr = iBlock * nThreads + iThread; itr < SectorTrackInfoLocalTotal(); itr += nThreads * nBlocks) {
@@ -1295,7 +1296,7 @@ GPUd() void GPUTPCGMMerger::MergeCEFill(const GPUTPCGMSectorTrack* track, const
   int32_t sector = track->Sector();
   for (int32_t attempt = 0; attempt < 2; attempt++) {
     GPUTPCGMBorderTrack b;
-    const float x0 = Param().tpcGeometry.Row2X(attempt == 0 ? 63 : cls.row);
+    const float x0 = GPUTPCGeometry::Row2X(attempt == 0 ? 63 : cls.row);
     if (track->TransportToX(this, x0, Param().bzCLight, b, GPUCA_MAX_SIN_PHI_LOW)) {
       b.SetTrackID(itr);
       b.SetNClusters(mOutputTracks[itr].NClusters());
@@ -1759,7 +1760,7 @@ GPUd() void GPUTPCGMMerger::CollectMergedTracks(int32_t nBlocks, int32_t nThread
     mergedTrack.SetCSide(p2.CSide());
 
     GPUTPCGMBorderTrack b;
-    const float toX = Param().par.earlyTpcTransform ? clXYZ[0].x : Param().tpcGeometry.Row2X(cl[0].row);
+    const float toX = Param().par.earlyTpcTransform ? clXYZ[0].x : GPUTPCGeometry::Row2X(cl[0].row);
     if (p2.TransportToX(this, toX, Param().bzCLight, b, GPUCA_MAX_SIN_PHI, false)) {
       p1.X() = toX;
       p1.Y() = b.Par()[0];
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx
index ae413aaa98648..1e4cc633eb4ca 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerDump.cxx
@@ -353,7 +353,7 @@ void GPUTPCGMMerger::DebugStreamerUpdate(int32_t iTrk, int32_t ihit, float xx, f
   auto uncorrectedYZ = StreamerUncorrectedZY(cluster.sector, cluster.row, track, prop);
   float invCharge = 1.f / clusterNative.qMax;
   int32_t iRow = cluster.row;
-  float unscaledMult = (time >= 0.f ? Param().GetUnscaledMult(time) / Param().tpcGeometry.Row2X(iRow) : 0.f);
+  float unscaledMult = (time >= 0.f ? Param().GetUnscaledMult(time) / GPUTPCGeometry::Row2X(iRow) : 0.f);
   const float clAlpha = Param().Alpha(cluster.sector);
   uint32_t occupancyTotal = Param().occupancyTotal;
   o2::utils::DebugStreamer::instance()->getStreamer("debug_update_track", "UPDATE") << o2::utils::DebugStreamer::instance()->getUniqueTreeName("tree_update_track").data()
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
index a439e6e653039..11b153c7f0d8b 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMSectorTrack.cxx
@@ -95,7 +95,7 @@ GPUd() void GPUTPCGMSectorTrack::SetParam2(const GPUTPCGMTrackParam& trk)
 GPUd() bool GPUTPCGMSectorTrack::FilterErrors(const GPUTPCGMMerger* merger, int32_t iSector, float maxSinPhi, float sinPhiMargin)
 {
   float lastX;
-  // float lastX = merger->Param().tpcGeometry.Row2X(mOrigTrack->Cluster(mOrigTrack->NClusters() - 1).GetRow()); // TODO: Why is this needed to be set below, Row2X should work, but looses some tracks
+  // float lastX = GPUTPCGeometry::Row2X(mOrigTrack->Cluster(mOrigTrack->NClusters() - 1).GetRow()); // TODO: Why is this needed to be set below, Row2X should work, but looses some tracks
   float y, z;
   int32_t row, index;
   const GPUTPCTracker& trk = merger->GetConstantMem()->tpcTrackers[iSector];
diff --git a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx
index be1d3803312fe..3bd2257d02e01 100644
--- a/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx
+++ b/GPU/GPUTracking/Merger/GPUTPCGMTrackParam.cxx
@@ -219,7 +219,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_
         bool dodEdx = param.par.dodEdx && param.dodEdxDownscaled && param.rec.tpc.adddEdxSubThresholdClusters && iWay == nWays - 1 && CAMath::Abs(cluster.row - lastRow) == 2 && cluster.leg == clusters[maxN - 1].leg;
         dodEdx = AttachClustersPropagate(merger, cluster.sector, lastRow, cluster.row, iTrk, cluster.leg == clusters[maxN - 1].leg, prop, inFlyDirection, GPUCA_MAX_SIN_PHI, dodEdx);
         if (dodEdx) {
-          dEdx.fillSubThreshold(lastRow - wayDirection, param);
+          dEdx.fillSubThreshold(lastRow - wayDirection);
         }
       }
 
@@ -384,7 +384,7 @@ GPUd() bool GPUTPCGMTrackParam::Fit(GPUTPCGMMerger* GPUrestrict() merger, int32_
           pad /= clusterCount;
           relTime /= clusterCount;
           relTime = relTime - CAMath::Round(relTime);
-          dEdx.fillCluster(qtot, qmax, cluster.row, cluster.sector, mP[2], mP[3], param, merger->GetConstantMem()->calibObjects, zz, pad, relTime);
+          dEdx.fillCluster(qtot, qmax, cluster.row, cluster.sector, mP[2], mP[3], merger->GetConstantMem()->calibObjects, zz, pad, relTime);
         }
       } else if (retVal >= GPUTPCGMPropagator::updateErrorClusterRejected) { // cluster far away form the track
         if (allowModification) {
@@ -650,7 +650,7 @@ GPUd() bool GPUTPCGMTrackParam::AttachClustersPropagate(const GPUTPCGMMerger* GP
     return dodEdx;
   }
   int32_t step = toRow > lastRow ? 1 : -1;
-  float xx = mX - Merger->Param().tpcGeometry.Row2X(lastRow);
+  float xx = mX - GPUTPCGeometry::Row2X(lastRow);
   for (int32_t iRow = lastRow + step; iRow != toRow; iRow += step) {
     if (CAMath::Abs(mP[2]) > maxSinPhi) {
       return dodEdx;
@@ -658,15 +658,15 @@ GPUd() bool GPUTPCGMTrackParam::AttachClustersPropagate(const GPUTPCGMMerger* GP
     if (CAMath::Abs(mP[0]) > CAMath::Abs(mX) * CAMath::Tan(kSectAngle / 2.f)) {
       return dodEdx;
     }
-    int32_t err = prop.PropagateToXAlpha(xx + Merger->Param().tpcGeometry.Row2X(iRow), prop.GetAlpha(), inFlyDirection);
+    int32_t err = prop.PropagateToXAlpha(xx + GPUTPCGeometry::Row2X(iRow), prop.GetAlpha(), inFlyDirection);
     if (err) {
       return dodEdx;
     }
     if (dodEdx && iRow + step == toRow) {
       float yUncorrected, zUncorrected;
       Merger->GetConstantMem()->calibObjects.fastTransformHelper->InverseTransformYZtoNominalYZ(sector, iRow, mP[0], mP[1], yUncorrected, zUncorrected);
-      uint32_t pad = CAMath::Float2UIntRn(Merger->Param().tpcGeometry.LinearY2Pad(sector, iRow, yUncorrected));
-      if (pad >= Merger->Param().tpcGeometry.NPads(iRow) || (Merger->GetConstantMem()->calibObjects.dEdxCalibContainer && Merger->GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(sector, iRow, pad))) {
+      uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(sector, iRow, yUncorrected));
+      if (pad >= GPUTPCGeometry::NPads(iRow) || (Merger->GetConstantMem()->calibObjects.dEdxCalibContainer && Merger->GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(sector, iRow, pad))) {
         dodEdx = false;
       }
     }
@@ -782,7 +782,7 @@ GPUdic(0, 1) int32_t GPUTPCGMTrackParam::FollowCircle(const GPUTPCGMMerger* GPUr
       }
       CADEBUG(printf("\tPropagated to y = %f: X %f Z %f SinPhi %f\n", mX, mP[0], mP[1], mP[2]));
       for (int32_t j = 0; j < GPUCA_ROW_COUNT; j++) {
-        float rowX = Merger->Param().tpcGeometry.Row2X(j);
+        float rowX = GPUTPCGeometry::Row2X(j);
         if (CAMath::Abs(rowX - (-mP[0] * lrFactor)) < 1.5f) {
           CADEBUG(printf("\t\tAttempt row %d (Y %f Z %f)\n", j, mX * lrFactor, mP[1]));
           AttachClusters(Merger, sector, j, iTrack, false, mX * lrFactor, mP[1]);
@@ -823,18 +823,18 @@ GPUdic(0, 1) int32_t GPUTPCGMTrackParam::FollowCircle(const GPUTPCGMMerger* GPUr
   prop.Rotate180();
   CADEBUG(printf("\tMirrored position: Alpha %f X %f Y %f Z %f SinPhi %f DzDs %f\n", prop.GetAlpha(), mX, mP[0], mP[1], mP[2], mP[3]));
   iRow = toRow;
-  float dx = toX - Merger->Param().tpcGeometry.Row2X(toRow);
+  float dx = toX - GPUTPCGeometry::Row2X(toRow);
   if (up ^ (toX > mX)) {
     if (up) {
-      while (iRow < GPUCA_ROW_COUNT - 2 && Merger->Param().tpcGeometry.Row2X(iRow + 1) + dx <= mX) {
+      while (iRow < GPUCA_ROW_COUNT - 2 && GPUTPCGeometry::Row2X(iRow + 1) + dx <= mX) {
         iRow++;
       }
     } else {
-      while (iRow > 1 && Merger->Param().tpcGeometry.Row2X(iRow - 1) + dx >= mX) {
+      while (iRow > 1 && GPUTPCGeometry::Row2X(iRow - 1) + dx >= mX) {
         iRow--;
       }
     }
-    prop.PropagateToXAlpha(Merger->Param().tpcGeometry.Row2X(iRow) + dx, prop.GetAlpha(), inFlyDirection);
+    prop.PropagateToXAlpha(GPUTPCGeometry::Row2X(iRow) + dx, prop.GetAlpha(), inFlyDirection);
     AttachClustersPropagate(Merger, sector, iRow, toRow, iTrack, false, prop, inFlyDirection);
   }
   if (prop.PropagateToXAlpha(toX, prop.GetAlpha(), inFlyDirection)) {
@@ -875,7 +875,7 @@ GPUdni() void GPUTPCGMTrackParam::AttachClustersMirror(const GPUTPCGMMerger* GPU
     return;
   }
   float dx = (toX - X) / count;
-  const float myRowX = Merger->Param().tpcGeometry.Row2X(iRow);
+  const float myRowX = GPUTPCGeometry::Row2X(iRow);
   // printf("AttachMirror\n");
   // printf("X %f Y %f Z %f SinPhi %f toY %f -->\n", mX, mP[0], mP[1], mP[2], toY);
   // printf("X %f Y %f Z %f SinPhi %f, count %d dx %f (to: %f)\n", X, Y, Z, SinPhi, count, dx, X + count * dx);
@@ -905,7 +905,7 @@ GPUdni() void GPUTPCGMTrackParam::AttachClustersMirror(const GPUTPCGMMerger* GPU
     int32_t step = paramX >= mX ? 1 : -1;
     int32_t found = 0;
     for (int32_t j = iRow; j >= 0 && j < GPUCA_ROW_COUNT && found < 3; j += step) {
-      float rowX = mX + Merger->Param().tpcGeometry.Row2X(j) - myRowX;
+      float rowX = mX + GPUTPCGeometry::Row2X(j) - myRowX;
       if (CAMath::Abs(rowX - paramX) < 1.5f) {
         // printf("Attempt row %d\n", j);
         AttachClusters(Merger, sector, j, iTrack, false, mP[2] > 0 ? X : -X, Z);
@@ -930,8 +930,8 @@ GPUd() void GPUTPCGMTrackParam::ShiftZ2(const GPUTPCGMMergedTrackHit* clusters,
     const auto& GPUrestrict() cls = merger->GetConstantMem()->ioPtrs.clustersNative->clustersLinear;
     tzInner = cls[clusters[N - 1].num].getTime();
     tzOuter = cls[clusters[0].num].getTime();
-    xInner = merger->Param().tpcGeometry.Row2X(clusters[N - 1].row);
-    xOuter = merger->Param().tpcGeometry.Row2X(clusters[0].row);
+    xInner = GPUTPCGeometry::Row2X(clusters[N - 1].row);
+    xOuter = GPUTPCGeometry::Row2X(clusters[0].row);
   }
   ShiftZ(merger, clusters[0].sector, tzInner, tzOuter, xInner, xOuter);
 }
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackingData.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackingData.cxx
index a3e73c377ed44..9a4d2eebcb953 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackingData.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackingData.cxx
@@ -19,6 +19,7 @@
 #include "GPUProcessor.h"
 #include "GPUO2DataTypes.h"
 #include "GPUTPCConvertImpl.h"
+#include "GPUTPCGeometry.h"
 #include "GPUCommonMath.h"
 
 #ifndef GPUCA_GPUCODE_DEVICE
@@ -39,7 +40,7 @@ void GPUTPCTrackingData::InitializeRows(const GPUParam& p)
     new (&mRows[i]) GPUTPCRow;
   }
   for (int32_t i = 0; i < GPUCA_ROW_COUNT; i++) {
-    mRows[i].mX = p.tpcGeometry.Row2X(i);
+    mRows[i].mX = GPUTPCGeometry::Row2X(i);
     mRows[i].mMaxY = CAMath::Tan(p.par.dAlpha / 2.f) * mRows[i].mX;
   }
 }
@@ -101,7 +102,7 @@ void* GPUTPCTrackingData::SetPointersRows(void* mem)
 GPUd() void GPUTPCTrackingData::GetMaxNBins(GPUconstantref() const GPUConstantMem* mem, GPUTPCRow* GPUrestrict() row, int32_t& maxY, int32_t& maxZ)
 {
   maxY = row->mMaxY * 2.f / GPUCA_MIN_BIN_SIZE + 1;
-  maxZ = (mem->param.continuousMaxTimeBin > 0 ? (mem->calibObjects.fastTransformHelper->getCorrMap()->convTimeToZinTimeFrame(0, 0, mem->param.continuousMaxTimeBin)) : mem->param.tpcGeometry.TPCLength()) + 50;
+  maxZ = (mem->param.continuousMaxTimeBin > 0 ? (mem->calibObjects.fastTransformHelper->getCorrMap()->convTimeToZinTimeFrame(0, 0, mem->param.continuousMaxTimeBin)) : GPUTPCGeometry::TPCLength()) + 50;
   maxZ = maxZ / GPUCA_MIN_BIN_SIZE + 1;
 }
 
diff --git a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
index 3aac31c87498c..5a7df0ba8b874 100644
--- a/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
+++ b/GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx
@@ -377,8 +377,8 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/,
       } while (false);
       (void)found;
       if (!found && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer) {
-        uint32_t pad = CAMath::Float2UIntRn(tracker.Param().tpcGeometry.LinearY2Pad(tracker.ISector(), iRow, yUncorrected));
-        if (pad < tracker.Param().tpcGeometry.NPads(iRow) && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(tracker.ISector(), iRow, pad)) {
+        uint32_t pad = CAMath::Float2UIntRn(GPUTPCGeometry::LinearY2Pad(tracker.ISector(), iRow, yUncorrected));
+        if (pad < GPUTPCGeometry::NPads(iRow) && tracker.GetConstantMem()->calibObjects.dEdxCalibContainer->isDead(tracker.ISector(), iRow, pad)) {
           r.mNMissed--;
           rowHit = CALINK_DEAD_CHANNEL;
         }
@@ -395,7 +395,7 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::UpdateTracklet(int32_t /*nBlocks*/,
     const float z1 = row1.Grid().ZMin() + hh1.y * row1.HstepZ();
     const float z2 = row2.Grid().ZMin() + hh2.y * row2.HstepZ();
     float oldOffset = tParam.ZOffset();
-    tParam.ShiftZ(z1, z2, tracker.Param().tpcGeometry.Row2X(r.mFirstRow), tracker.Param().tpcGeometry.Row2X(r.mLastRow), tracker.Param().bzCLight, tracker.Param().rec.tpc.defaultZOffsetOverR);
+    tParam.ShiftZ(z1, z2, GPUTPCGeometry::Row2X(r.mFirstRow), GPUTPCGeometry::Row2X(r.mLastRow), tracker.Param().bzCLight, tracker.Param().rec.tpc.defaultZOffsetOverR);
     r.mLastZ -= tParam.ZOffset() - oldOffset;
     CADEBUG(printf("Shifted z from %f to %f\n", oldOffset, tParam.ZOffset()));
   }
diff --git a/GPU/GPUTracking/TPCClusterFinder/ClusterAccumulator.cxx b/GPU/GPUTracking/TPCClusterFinder/ClusterAccumulator.cxx
index b3b3c64095017..622da856af805 100644
--- a/GPU/GPUTracking/TPCClusterFinder/ClusterAccumulator.cxx
+++ b/GPU/GPUTracking/TPCClusterFinder/ClusterAccumulator.cxx
@@ -13,9 +13,9 @@
 /// \author Felix Weiglhofer
 
 #include "ClusterAccumulator.h"
-#include "GPUTPCGeometry.h"
 #include "CfUtils.h"
 #include "GPUParam.h"
+#include "GPUTPCGeometry.h"
 #include "DataFormatsTPC/ClusterNative.h"
 
 using namespace o2::gpu;
@@ -81,27 +81,24 @@ GPUd() bool ClusterAccumulator::toNative(const ChargePos& pos, const Charge q, t
 
   bool isEdgeCluster;
   if (param.rec.tpc.cfEdgeTwoPads) {
-    isEdgeCluster = pad < 2 || pad >= param.tpcGeometry.NPads(pos.row()) - 2; // Geometrical edge check, peak within 2 pads of sector edge
+    isEdgeCluster = pad < 2 || pad >= GPUTPCGeometry::NPads(pos.row()) - 2; // Geometrical edge check, peak within 2 pads of sector edge
     if (isEdgeCluster) {
       bool leftEdge = (pad < 2);
-      if (leftEdge ? (pad == 1 && chargeMap[pos.delta({-1, 0})].unpack() < 1) : (pad == (param.tpcGeometry.NPads(pos.row()) - 2) && chargeMap[pos.delta({1, 0})].unpack() < 1)) {
+      if (leftEdge ? (pad == 1 && chargeMap[pos.delta({-1, 0})].unpack() < 1) : (pad == (GPUTPCGeometry::NPads(pos.row()) - 2) && chargeMap[pos.delta({1, 0})].unpack() < 1)) {
         isEdgeCluster = false; // No edge cluster if peak is close to edge but no charge at the edge.
       } else if (leftEdge ? (pad < mPadMean) : (pad > mPadMean)) {
         mPadMean = pad; // Correct to peak position if COG is close to middle of pad than peak
       }
     }
   } else {
-    isEdgeCluster = pad == 0 || pad == param.tpcGeometry.NPads(pos.row()) - 1;
+    isEdgeCluster = pad == 0 || pad == GPUTPCGeometry::NPads(pos.row()) - 1;
   }
 
   cn.qTot = CAMath::Float2UIntRn(mQtot);
   if (cn.qTot <= param.rec.tpc.cfQTotCutoff) {
     return false;
   }
-  cn.qMax = q;
-  if (cn.qMax <= param.rec.tpc.cfQMaxCutoff) {
-    return false;
-  }
+  cn.qMax = q; // cfQMaxCutoff check already done at PeakFinder level
   if (mTimeMean < param.rec.tpc.clustersShiftTimebinsClusterizer) {
     return false;
   }
diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx
index 4a167b7d53890..1e76860331de6 100644
--- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx
+++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx
@@ -15,6 +15,7 @@
 #include "GPUTPCCFCheckPadBaseline.h"
 #include "Array2D.h"
 #include "PackedCharge.h"
+#include "GPUTPCGeometry.h"
 #include "clusterFinderDefs.h"
 
 #ifndef GPUCA_GPUCODE
@@ -151,7 +152,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::Thread<0>(int32_t nBlocks, int32_t nThread
 
 GPUd() ChargePos GPUTPCCFCheckPadBaseline::padToChargePos(int32_t& pad, const GPUTPCClusterFinder& clusterer)
 {
-  const GPUTPCGeometry& geo = clusterer.Param().tpcGeometry;
+  constexpr GPUTPCGeometry geo;
 
   int32_t padOffset = 0;
   for (Row r = 0; r < GPUCA_ROW_COUNT; r++) {
diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx
index f1fd95d696f5d..6662b93eccb78 100644
--- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx
+++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx
@@ -22,6 +22,7 @@
 #include "GPUCommonAlgorithm.h"
 #include "TPCPadGainCalib.h"
 #include "TPCZSLinkMapping.h"
+#include "GPUTPCGeometry.h"
 
 using namespace o2::gpu;
 using namespace o2::gpu::tpccf;
@@ -57,8 +58,8 @@ GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUShared
   const size_t nDigits = clusterer.mPzsOffsets[iBlock].offset;
   if (iThread == 0) {
     const int32_t region = endpoint / 2;
-    s.nRowsRegion = clusterer.Param().tpcGeometry.GetRegionRows(region);
-    s.regionStartRow = clusterer.Param().tpcGeometry.GetRegionStart(region);
+    s.nRowsRegion = GPUTPCGeometry::GetRegionRows(region);
+    s.regionStartRow = GPUTPCGeometry::GetRegionStart(region);
     s.nThreadsPerRow = CAMath::Max(1u, nThreads / ((s.nRowsRegion + (endpoint & 1)) / 2));
     s.rowStride = nThreads / s.nThreadsPerRow;
     s.rowOffsetCounter = 0;
@@ -524,7 +525,7 @@ GPUd() o2::tpc::PadPos GPUTPCCFDecodeZSLinkBase::GetPadAndRowFromFEC(processorTy
 {
 #ifdef GPUCA_TPC_GEOMETRY_O2
   // Ported from tpc::Mapper (Not available on GPU...)
-  const GPUTPCGeometry& geo = clusterer.Param().tpcGeometry;
+  constexpr GPUTPCGeometry geo;
 
   const int32_t regionIter = cru % 2;
   const int32_t istreamm = ((rawFECChannel % 10) / 2);
diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx
index 30fdac92e8607..1de922f716c14 100644
--- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx
+++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx
@@ -42,7 +42,7 @@ GPUdii() bool GPUTPCCFPeakFinder::isPeak(
 {
   uint16_t ll = get_local_id(0);
 
-  bool belowThreshold = (q <= calib.tpc.cfQMaxCutoff);
+  bool belowThreshold = (uint32_t)q <= calib.tpc.cfQMaxCutoff;
 
   uint16_t lookForPeaks;
   uint16_t partId = CfUtils::partition<SCRATCH_PAD_WORK_GROUP_SIZE>(
diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx
index 25cd2497fbf62..379ea27443fea 100644
--- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx
+++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.cxx
@@ -14,6 +14,7 @@
 
 #include "GPUTPCNNClusterizerKernels.h"
 #include "GPUTPCCFClusterizer.h"
+#include "GPUTPCGeometry.h"
 
 using namespace o2::gpu;
 using namespace o2::gpu::tpccf;
@@ -102,9 +103,9 @@ GPUdii() void GPUTPCNNClusterizerKernels::Thread<GPUTPCNNClusterizerKernels::pub
 }
 
 // THe following arithmetic is done because the network is trained with a split between IROC and OROC boundary
-GPUd() int GPUTPCNNClusterizerKernels::padOffset(int row_ref, int row_current, const GPUTPCGeometry& geo)
+GPUd() int GPUTPCNNClusterizerKernels::padOffset(int row_ref, int row_current)
 {
-  return (int)((geo.NPads(row_current) - geo.NPads(row_ref)) / 2);
+  return (int)((GPUTPCGeometry::NPads(row_current) - GPUTPCGeometry::NPads(row_ref)) / 2);
 }
 
 GPUd() int GPUTPCNNClusterizerKernels::rowOffset(int row, int global_shift)
@@ -112,16 +113,16 @@ GPUd() int GPUTPCNNClusterizerKernels::rowOffset(int row, int global_shift)
   return (row > 62 ? global_shift : 0);
 }
 
-GPUd() bool GPUTPCNNClusterizerKernels::isBoundary(int row, int pad, int global_shift, const GPUTPCGeometry& geo)
+GPUd() bool GPUTPCNNClusterizerKernels::isBoundary(int row, int pad, int global_shift)
 {
   if (pad < 0 || row < 0) { // Faster short-circuit
     return true;
   } else if (row < 63) {
-    return (pad >= static_cast<int>(geo.NPads(row)));
+    return (pad >= static_cast<int>(GPUTPCGeometry::NPads(row)));
   } else if (row < (63 + global_shift)) { // to account for the gap between IROC and OROC. Charge will be set to -1 in order to signal boundary to the neural network
     return true;
   } else if (row < (o2::tpc::constants::MAXGLOBALPADROW + global_shift)) {
-    return (pad >= static_cast<int>(geo.NPads(row - global_shift)));
+    return (pad >= static_cast<int>(GPUTPCGeometry::NPads(row - global_shift)));
   } else {
     return true;
   }
@@ -152,9 +153,9 @@ GPUd() void GPUTPCNNClusterizerKernels::fillInputData(int32_t nBlocks, int32_t n
 #endif
   for (int r = -clustererNN.nnClusterizerSizeInputRow; r <= clustererNN.nnClusterizerSizeInputRow; r++) {
     bool is_row_boundary = ((row + r) > (o2::tpc::constants::MAXGLOBALPADROW - 1)) || ((row + r) < 0);
-    int pad_offset = is_row_boundary ? 0 : GPUTPCNNClusterizerKernels::padOffset(row, row + r, clusterer.Param().tpcGeometry);
+    int pad_offset = is_row_boundary ? 0 : GPUTPCNNClusterizerKernels::padOffset(row, row + r);
     for (int p = -clustererNN.nnClusterizerSizeInputPad + pad_offset; p <= clustererNN.nnClusterizerSizeInputPad + pad_offset; p++) {
-      bool is_boundary = is_row_boundary || GPUTPCNNClusterizerKernels::isBoundary(row + r + row_offset, pad + p, clustererNN.nnClusterizerSizeInputRow, clusterer.Param().tpcGeometry);
+      bool is_boundary = is_row_boundary || GPUTPCNNClusterizerKernels::isBoundary(row + r + row_offset, pad + p, clustererNN.nnClusterizerSizeInputRow);
       for (int t = -clustererNN.nnClusterizerSizeInputTime; t <= clustererNN.nnClusterizerSizeInputTime; t++) {
         if (!is_boundary) {
           ChargePos tmp_pos(row + r, pad + p, time + t);
@@ -183,11 +184,11 @@ GPUd() void GPUTPCNNClusterizerKernels::fillInputData(int32_t nBlocks, int32_t n
     if (dtype == 0) {
       clustererNN.inputData16[write_idx] = (OrtDataType::Float16_t)(clusterer.mISector / 36.f);
       clustererNN.inputData16[write_idx + 1] = (OrtDataType::Float16_t)(row / 152.f);
-      clustererNN.inputData16[write_idx + 2] = (OrtDataType::Float16_t)(static_cast<float>(pad) / clusterer.Param().tpcGeometry.NPads(row));
+      clustererNN.inputData16[write_idx + 2] = (OrtDataType::Float16_t)(static_cast<float>(pad) / GPUTPCGeometry::NPads(row));
     } else {
       clustererNN.inputData32[write_idx] = clusterer.mISector / 36.f;
       clustererNN.inputData32[write_idx + 1] = row / 152.f;
-      clustererNN.inputData32[write_idx + 2] = static_cast<float>(pad) / clusterer.Param().tpcGeometry.NPads(row);
+      clustererNN.inputData32[write_idx + 2] = static_cast<float>(pad) / GPUTPCGeometry::NPads(row);
     }
   }
 }
diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h
index c7bd18115d61f..e6c1dc508d6e4 100644
--- a/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h
+++ b/GPU/GPUTracking/TPCClusterFinder/GPUTPCNNClusterizerKernels.h
@@ -67,9 +67,9 @@ class GPUTPCNNClusterizerKernels : public GPUKernelTemplate
   static GPUd() void publishClustersReg1(uint, GPUSharedMemory&, processorType&, uint8_t, int8_t, int8_t, uint);
   static GPUd() void publishClustersReg2(uint, GPUSharedMemory&, processorType&, uint8_t, int8_t, int8_t, uint);
 
-  static GPUd() int padOffset(int, int, const GPUTPCGeometry&);
+  static GPUd() int padOffset(int, int);
   static GPUd() int rowOffset(int, int);
-  static GPUd() bool isBoundary(int, int, int, const GPUTPCGeometry&);
+  static GPUd() bool isBoundary(int, int, int);
 };
 
 } // namespace o2::gpu
diff --git a/GPU/GPUTracking/dEdx/GPUdEdx.cxx b/GPU/GPUTracking/dEdx/GPUdEdx.cxx
index 2e67ddda7c99c..b7da0de4c0e29 100644
--- a/GPU/GPUTracking/dEdx/GPUdEdx.cxx
+++ b/GPU/GPUTracking/dEdx/GPUdEdx.cxx
@@ -13,7 +13,6 @@
 /// \author David Rohr
 
 #include "GPUdEdx.h"
-#include "GPUTPCGeometry.h"
 #include "GPUdEdxInfo.h"
 #include "GPUCommonAlgorithm.h"
 #include "GPUParam.h"
diff --git a/GPU/GPUTracking/dEdx/GPUdEdx.h b/GPU/GPUTracking/dEdx/GPUdEdx.h
index 6c0a96d3adb75..bcd75af468c28 100644
--- a/GPU/GPUTracking/dEdx/GPUdEdx.h
+++ b/GPU/GPUTracking/dEdx/GPUdEdx.h
@@ -16,12 +16,12 @@
 #define GPUDEDX_H
 
 #include "GPUDef.h"
-#include "GPUTPCGeometry.h"
 #include "GPUCommonMath.h"
 #include "GPUParam.h"
 #include "GPUdEdxInfo.h"
 #include "DataFormatsTPC/Defs.h"
 #include "CalibdEdxContainer.h"
+#include "GPUTPCGeometry.h"
 #include "GPUDebugStreamer.h"
 
 namespace o2::gpu
@@ -32,8 +32,8 @@ class GPUdEdx
  public:
   // The driver must call clear(), fill clusters row by row outside-in, then run computedEdx() to get the result
   GPUd() void clear();
-  GPUd() void fillCluster(float qtot, float qmax, int32_t padRow, uint8_t sector, float trackSnp, float trackTgl, const GPUParam& param, const GPUCalibObjectsConst& calib, float z, float pad, float relTime);
-  GPUd() void fillSubThreshold(int32_t padRow, const GPUParam& param);
+  GPUd() void fillCluster(float qtot, float qmax, int32_t padRow, uint8_t sector, float trackSnp, float trackTgl, const GPUCalibObjectsConst& calib, float z, float pad, float relTime);
+  GPUd() void fillSubThreshold(int32_t padRow);
   GPUd() void computedEdx(GPUdEdxInfo& output, const GPUParam& param);
 
  private:
@@ -92,7 +92,7 @@ GPUdi() void GPUdEdx::checkSubThresh(int32_t roc)
   mLastROC = roc;
 }
 
-GPUdnii() void GPUdEdx::fillCluster(float qtot, float qmax, int32_t padRow, uint8_t sector, float trackSnp, float trackTgl, const GPUParam& GPUrestrict() param, const GPUCalibObjectsConst& calib, float z, float pad, float relTime)
+GPUdnii() void GPUdEdx::fillCluster(float qtot, float qmax, int32_t padRow, uint8_t sector, float trackSnp, float trackTgl, const GPUCalibObjectsConst& calib, float z, float pad, float relTime)
 {
   if (mCount >= MAX_NCL) {
     return;
@@ -100,8 +100,9 @@ GPUdnii() void GPUdEdx::fillCluster(float qtot, float qmax, int32_t padRow, uint
 
   // container containing all the dE/dx corrections
   auto calibContainer = calib.dEdxCalibContainer;
+  constexpr GPUTPCGeometry geo;
 
-  const int32_t roc = param.tpcGeometry.GetROC(padRow);
+  const int32_t roc = geo.GetROC(padRow);
   checkSubThresh(roc);
   float snp2 = trackSnp * trackSnp;
   if (snp2 > GPUCA_MAX_SIN_PHI_LOW) {
@@ -119,7 +120,7 @@ GPUdnii() void GPUdEdx::fillCluster(float qtot, float qmax, int32_t padRow, uint
   // getting the topology correction
   const uint32_t padPos = CAMath::Float2UIntRn(pad); // position of the pad is shifted half a pad ( pad=3 -> centre position of third pad)
   const float absRelPad = CAMath::Abs(pad - padPos);
-  const int32_t region = param.tpcGeometry.GetRegion(padRow);
+  const int32_t region = geo.GetRegion(padRow);
   z = CAMath::Abs(z);
   const float threshold = calibContainer->getZeroSupressionThreshold(sector, padRow, padPos); // TODO: Use the mean zero supresion threshold of all pads in the cluster?
   const bool useFullGainMap = calibContainer->isUsageOfFullGainMap();
@@ -161,8 +162,8 @@ GPUdnii() void GPUdEdx::fillCluster(float qtot, float qmax, int32_t padRow, uint
   }
 
   GPUCA_DEBUG_STREAMER_CHECK(if (o2::utils::DebugStreamer::checkStream(o2::utils::StreamFlags::streamdEdx)) {
-    float padlx = param.tpcGeometry.Row2X(padRow);
-    float padly = param.tpcGeometry.LinearPad2Y(sector, padRow, padPos);
+    float padlx = geo.Row2X(padRow);
+    float padly = geo.LinearPad2Y(sector, padRow, padPos);
     o2::utils::DebugStreamer::instance()->getStreamer("debug_dedx", "UPDATE") << o2::utils::DebugStreamer::instance()->getUniqueTreeName("tree_dedx").data()
                                                                               << "qTot=" << mChargeTot[mCount - 1]
                                                                               << "qMax=" << mChargeMax[mCount - 1]
@@ -189,9 +190,9 @@ GPUdnii() void GPUdEdx::fillCluster(float qtot, float qmax, int32_t padRow, uint
   })
 }
 
-GPUdi() void GPUdEdx::fillSubThreshold(int32_t padRow, const GPUParam& GPUrestrict() param)
+GPUdi() void GPUdEdx::fillSubThreshold(int32_t padRow)
 {
-  const int32_t roc = param.tpcGeometry.GetROC(padRow);
+  const int32_t roc = GPUTPCGeometry::GetROC(padRow);
   checkSubThresh(roc);
   mNSubThresh++;
 }
diff --git a/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx b/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx
index 6fd70354c9486..0a780732273db 100644
--- a/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx
+++ b/GPU/GPUTracking/display/render/GPUDisplayImportEvent.cxx
@@ -150,13 +150,13 @@ void GPUDisplay::DrawGLScene_updateEventData()
         float4* ptr = &mGlobalPos[cid];
         if (mParam->par.earlyTpcTransform) {
           const auto& cl = mIOPtrs->clusterData[iSector][i];
-          mParam->Sector2Global(iSector, (mCfgH.clustersOnNominalRow ? mParam->tpcGeometry.Row2X(row) : cl.x) + mCfgH.xAdd, cl.y, cl.z, &ptr->x, &ptr->y, &ptr->z);
+          mParam->Sector2Global(iSector, (mCfgH.clustersOnNominalRow ? GPUTPCGeometry::Row2X(row) : cl.x) + mCfgH.xAdd, cl.y, cl.z, &ptr->x, &ptr->y, &ptr->z);
         } else {
           float x, y, z;
           const auto& cln = mIOPtrs->clustersNative->clusters[iSector][0][i];
           GPUTPCConvertImpl::convert(*mCalib->fastTransform, *mParam, iSector, row, cln.getPad(), cln.getTime(), x, y, z);
           if (mCfgH.clustersOnNominalRow) {
-            x = mParam->tpcGeometry.Row2X(row);
+            x = GPUTPCGeometry::Row2X(row);
           }
           mParam->Sector2Global(iSector, x + mCfgH.xAdd, y, z, &ptr->x, &ptr->y, &ptr->z);
         }
diff --git a/GPU/GPUTracking/qa/genEvents.cxx b/GPU/GPUTracking/qa/genEvents.cxx
index 627cfc5f9909a..2e1bc1c5c64b2 100644
--- a/GPU/GPUTracking/qa/genEvents.cxx
+++ b/GPU/GPUTracking/qa/genEvents.cxx
@@ -222,7 +222,7 @@ int32_t genEvents::GenerateEvent(const GPUParam& param, char* filename)
 
     for (int32_t iRow = 0; iRow < GPUCA_ROW_COUNT; iRow++) {
       // if( iRow>=50 ) break; //SG!!!
-      float xRow = param.tpcGeometry.Row2X(iRow);
+      float xRow = GPUTPCGeometry::Row2X(iRow);
       // transport to row
       int32_t err = 0;
       for (int32_t itry = 0; itry < 1; itry++) {
diff --git a/Steer/CMakeLists.txt b/Steer/CMakeLists.txt
index 70f50f4ab8823..8e2706d31bb0a 100644
--- a/Steer/CMakeLists.txt
+++ b/Steer/CMakeLists.txt
@@ -34,11 +34,13 @@ o2_target_root_dictionary(Steer
                                   include/Steer/MCKinematicsReader.h
                                   include/Steer/MaterialBudgetMap.h)
 o2_add_test(InteractionSampler
+            COMPONENT_NAME steer
             PUBLIC_LINK_LIBRARIES O2::SimulationDataFormat
             SOURCES test/testInteractionSampler.cxx
             LABELS steer)
 
 o2_add_test(HitProcessingManager
+            COMPONENT_NAME steer
             PUBLIC_LINK_LIBRARIES O2::Steer
             SOURCES test/testHitProcessingManager.cxx
             LABELS steer)
diff --git a/run/CMakeLists.txt b/run/CMakeLists.txt
index 662716901ed0a..fd43207f92d1e 100644
--- a/run/CMakeLists.txt
+++ b/run/CMakeLists.txt
@@ -229,6 +229,7 @@ set_property(TEST o2sim_G4 APPEND PROPERTY ENVIRONMENT ${G4ENV})
 
 
 o2_add_test(CheckStackG4
+  COMPONENT_NAME sim
   SOURCES checkStack.cxx
   NAME o2sim_checksimkinematics_G4
   WORKING_DIRECTORY ${SIMTESTDIR}
@@ -276,6 +277,7 @@ set_tests_properties(o2sim_G3
                                 G3)
 
 o2_add_test(CheckStackG3
+  COMPONENT_NAME sim
   SOURCES checkStack.cxx
   NAME o2sim_checksimkinematics_G3
   WORKING_DIRECTORY ${SIMTESTDIR}