Please consider the following formatting changes to #13999

Detectors/TPC/monitor/src/SimpleEventDisplayGUI.cxx

@@ -1227,7 +1227,7 @@ void SimpleEventDisplayGUI::showClusters(int roc, int row)
         }
         if (fillSingleTB && std::abs(cl.getTime() - timeBin) < 2) {
           const auto ly = gpuGeom.LinearPad2Y(sector, irow, cl.getPad() + 0.5);
-          mClustersRowPad->SetNextPoint(gpuGeom.Row2X(irow), (sector >= GPUCA_NSLICES / 2) ? -ly : ly);
+          mClustersRowPad->SetNextPoint(gpuGeom.Row2X(irow), (sector >= GPUCA_NSECTORS / 2) ? -ly : ly);
         }
       }
       // fmt::print("\n");

Detectors/TPC/reconstruction/test/testGPUCATracking.cxx

@@ -70,7 +70,7 @@ BOOST_AUTO_TEST_CASE(CATracking_test1)
   config.configReconstruction.tpc.searchWindowDZDR = 2.5f; //Should always be 2.5 for looper-finding and/or continuous tracking
   config.configReconstruction.tpc.trackReferenceX = refX;
 
-  config.configWorkflow.steps.set(GPUDataTypes::RecoStep::TPCConversion, GPUDataTypes::RecoStep::TPCSliceTracking,
+  config.configWorkflow.steps.set(GPUDataTypes::RecoStep::TPCConversion, GPUDataTypes::RecoStep::TPCSectorTracking,
                                   GPUDataTypes::RecoStep::TPCMerging, GPUDataTypes::RecoStep::TPCCompression, GPUDataTypes::RecoStep::TPCdEdx);
   config.configWorkflow.inputs.set(GPUDataTypes::InOutType::TPCClusters);
   config.configWorkflow.outputs.set(GPUDataTypes::InOutType::TPCMergedTracks);

Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx

@@ -159,7 +159,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
 
     const auto& tinfo = pc.services().get<o2::framework::TimingInfo>();
     const auto firstIR = o2::InteractionRecord(0, tinfo.firstTForbit);
-    const float totalT = std::max(mFastTransform->getMaxDriftTime(0), mFastTransform->getMaxDriftTime(GPUCA_NSLICES / 2));
+    const float totalT = std::max(mFastTransform->getMaxDriftTime(0), mFastTransform->getMaxDriftTime(GPUCA_NSECTORS / 2));
 
     unsigned int offset = 0, lasti = 0;
     const unsigned int maxTime = (mParam->continuousMaxTimeBin + 1) * o2::tpc::ClusterNative::scaleTimePacked - 1;
@@ -206,8 +206,8 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
       }
     }
     offset = 0;
-    unsigned int offsets[GPUCA_NSLICES][GPUCA_ROW_COUNT];
-    for (unsigned int i = 0; i < GPUCA_NSLICES; i++) {
+    unsigned int offsets[GPUCA_NSECTORS][GPUCA_ROW_COUNT];
+    for (unsigned int i = 0; i < GPUCA_NSECTORS; i++) {
       for (unsigned int j = 0; j < GPUCA_ROW_COUNT; j++) {
         if (i * GPUCA_ROW_COUNT + j >= clusters.nSliceRows) {
           break;
@@ -218,7 +218,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
     }
 
 #ifdef WITH_OPENMP
-#pragma omp parallel for num_threads(mNThreads) schedule(static, (GPUCA_NSLICES + mNThreads - 1) / mNThreads) // Static round-robin scheduling with one chunk per thread to ensure correct order of the final vector
+#pragma omp parallel for num_threads(mNThreads) schedule(static, (GPUCA_NSECTORS + mNThreads - 1) / mNThreads) // Static round-robin scheduling with one chunk per thread to ensure correct order of the final vector
 #endif
     for (unsigned int ii = 0; ii < clusters.nSliceRows; ii++) {
       unsigned int i = ii / GPUCA_ROW_COUNT;

GPU/GPUTracking/Base/GPUConstantMem.h

@@ -34,21 +34,19 @@
 #include "GPUKernelDebugOutput.h"
 #endif
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 struct GPUConstantMem {
   GPUParam param;
   GPUTPCTracker
-    tpcTrackers[GPUCA_NSLICES];
+    tpcTrackers[GPUCA_NSECTORS];
   GPUTPCConvert tpcConverter;
   GPUTPCCompression tpcCompressor;
   GPUTPCDecompression tpcDecompressor;
   GPUTPCGMMerger tpcMerger;
   GPUTRDTrackerGPU trdTrackerGPU;
   GPUTRDTracker trdTrackerO2;
-  GPUTPCClusterFinder tpcClusterer[GPUCA_NSLICES];
+  GPUTPCClusterFinder tpcClusterer[GPUCA_NSECTORS];
   GPUITSFitter itsFitter;
   GPUTrackingRefitProcessor trackingRefit;
   GPUTrackingInOutPointers ioPtrs;
@@ -90,14 +88,11 @@ union GPUConstantMemCopyable {
 #if defined(GPUCA_GPUCODE)
 static constexpr size_t gGPUConstantMemBufferSize = (sizeof(GPUConstantMem) + sizeof(uint4) - 1);
 #endif
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 #if defined(GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM) && !defined(GPUCA_GPUCODE_HOSTONLY)
 GPUconstant() o2::gpu::GPUConstantMemCopyable gGPUConstantMemBuffer;
 #endif // GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 
 // Must be placed here, to avoid circular header dependency
@@ -120,7 +115,6 @@ GPUdi() void GPUProcessor::raiseError(uint32_t code, uint32_t param1, uint32_t p
   GetConstantMem()->errorCodes.raiseError(code, param1, param2, param3);
 }
 
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #endif

GPU/GPUTracking/Base/GPUGeneralKernels.h

@@ -32,9 +32,7 @@
 #define GPUCA_CUB cub
 #endif
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 struct GPUConstantMem;
 
@@ -110,8 +108,7 @@ class GPUitoa : public GPUKernelTemplate
   GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUglobalref() int32_t* ptr, uint64_t size);
 };
 
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #undef GPUCA_CUB
 

GPU/GPUTracking/Base/GPUKernelDebugOutput.h

@@ -19,9 +19,7 @@
 #include "GPUProcessor.h"
 #ifdef GPUCA_KERNEL_DEBUGGER_OUTPUT
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 
 class GPUKernelDebugOutput : public GPUProcessor
@@ -75,8 +73,7 @@ class GPUKernelDebugOutput : public GPUProcessor
   mutable int32_t* mDebugOutMemory;
 };
 
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #endif
 #endif

GPU/GPUTracking/Base/GPUMemoryResource.h

@@ -18,9 +18,7 @@
 #include "GPUCommonDef.h"
 #include "GPUProcessor.h"
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 
 struct GPUMemoryReuse {
@@ -103,7 +101,6 @@ class GPUMemoryResource
   int32_t mReuse;
   MemoryType mType;
 };
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #endif

GPU/GPUTracking/Base/GPUParam.cxx

@@ -91,22 +91,22 @@ void GPUParam::SetDefaults(float solenoidBz)
   constexpr float plusZmax = 249.778;
   constexpr float minusZmin = -249.645;
   constexpr float minusZmax = -0.0799937;
-  for (int32_t i = 0; i < GPUCA_NSLICES; i++) {
-    const bool zPlus = (i < GPUCA_NSLICES / 2);
-    SliceParam[i].ZMin = zPlus ? plusZmin : minusZmin;
-    SliceParam[i].ZMax = zPlus ? plusZmax : minusZmax;
+  for (int32_t i = 0; i < GPUCA_NSECTORS; i++) {
+    const bool zPlus = (i < GPUCA_NSECTORS / 2);
+    SectorParam[i].ZMin = zPlus ? plusZmin : minusZmin;
+    SectorParam[i].ZMax = zPlus ? plusZmax : minusZmax;
     int32_t tmp = i;
-    if (tmp >= GPUCA_NSLICES / 2) {
-      tmp -= GPUCA_NSLICES / 2;
+    if (tmp >= GPUCA_NSECTORS / 2) {
+      tmp -= GPUCA_NSECTORS / 2;
     }
-    if (tmp >= GPUCA_NSLICES / 4) {
-      tmp -= GPUCA_NSLICES / 2;
+    if (tmp >= GPUCA_NSECTORS / 4) {
+      tmp -= GPUCA_NSECTORS / 2;
     }
-    SliceParam[i].Alpha = 0.174533f + par.dAlpha * tmp;
-    SliceParam[i].CosAlpha = CAMath::Cos(SliceParam[i].Alpha);
-    SliceParam[i].SinAlpha = CAMath::Sin(SliceParam[i].Alpha);
-    SliceParam[i].AngleMin = SliceParam[i].Alpha - par.dAlpha / 2.f;
-    SliceParam[i].AngleMax = SliceParam[i].Alpha + par.dAlpha / 2.f;
+    SectorParam[i].Alpha = 0.174533f + par.dAlpha * tmp;
+    SectorParam[i].CosAlpha = CAMath::Cos(SectorParam[i].Alpha);
+    SectorParam[i].SinAlpha = CAMath::Sin(SectorParam[i].Alpha);
+    SectorParam[i].AngleMin = SectorParam[i].Alpha - par.dAlpha / 2.f;
+    SectorParam[i].AngleMax = SectorParam[i].Alpha + par.dAlpha / 2.f;
   }
 
   par.assumeConstantBz = false;

GPU/GPUTracking/Base/GPUParam.h

@@ -31,19 +31,17 @@ using Propagator = PropagatorImpl<float>;
 } // namespace o2::base
 #endif
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 struct GPUSettingsRec;
 struct GPUSettingsGTP;
 struct GPURecoStepConfiguration;
 
-struct GPUParamSlice {
-  float Alpha;              // slice angle
-  float CosAlpha, SinAlpha; // sign and cosine of the slice angle
+struct GPUParamSector {
+  float Alpha;              // sector angle
+  float CosAlpha, SinAlpha; // sign and cosine of the sector angle
   float AngleMin, AngleMax; // minimal and maximal angle
-  float ZMin, ZMax;         // slice Z range
+  float ZMin, ZMax;         // sector Z range
 };
 
 namespace internal
@@ -66,7 +64,7 @@ struct GPUParam_t {
   const uint32_t* occupancyMap;                     // Ptr to TPC occupancy map
   uint32_t occupancyTotal;                          // Total occupancy in the TPC (nCl / nHbf)
 
-  GPUParamSlice SliceParam[GPUCA_NSLICES];
+  GPUParamSector SectorParam[GPUCA_NSECTORS];
 
  protected:
 #ifdef GPUCA_TPC_GEOMETRY_O2
@@ -88,15 +86,15 @@ struct GPUParam : public internal::GPUParam_t<GPUSettingsRec, GPUSettingsParam>
   void UpdateRun3ClusterErrors(const float* yErrorParam, const float* zErrorParam);
 #endif
 
-  GPUd() float Alpha(int32_t iSlice) const
+  GPUd() float Alpha(int32_t iSector) const
   {
-    if (iSlice >= GPUCA_NSLICES / 2) {
-      iSlice -= GPUCA_NSLICES / 2;
+    if (iSector >= GPUCA_NSECTORS / 2) {
+      iSector -= GPUCA_NSECTORS / 2;
     }
-    if (iSlice >= GPUCA_NSLICES / 4) {
-      iSlice -= GPUCA_NSLICES / 2;
+    if (iSector >= GPUCA_NSECTORS / 4) {
+      iSector -= GPUCA_NSECTORS / 2;
     }
-    return 0.174533f + par.dAlpha * iSlice;
+    return 0.174533f + par.dAlpha * iSector;
   }
   GPUd() float GetClusterErrorSeeding(int32_t yz, int32_t type, float zDiff, float angle2, float unscaledMult) const;
   GPUd() void GetClusterErrorsSeeding2(uint8_t sector, int32_t row, float z, float sinPhi, float DzDs, float time, float& ErrY2, float& ErrZ2) const;
@@ -108,13 +106,12 @@ struct GPUParam : public internal::GPUParam_t<GPUSettingsRec, GPUSettingsParam>
   GPUd() void UpdateClusterError2ByState(int16_t clusterState, float& ErrY2, float& ErrZ2) const;
   GPUd() float GetUnscaledMult(float time) const;
 
-  GPUd() void Slice2Global(int32_t iSlice, float x, float y, float z, float* X, float* Y, float* Z) const;
-  GPUd() void Global2Slice(int32_t iSlice, float x, float y, float z, float* X, float* Y, float* Z) const;
+  GPUd() void Sector2Global(int32_t iSector, float x, float y, float z, float* X, float* Y, float* Z) const;
+  GPUd() void Global2Sector(int32_t iSector, float x, float y, float z, float* X, float* Y, float* Z) const;
 
   GPUd() bool rejectEdgeClusterByY(float uncorrectedY, int32_t iRow, float trackSigmaY) const;
 };
 
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #endif

GPU/GPUTracking/Base/GPUParam.inc

@@ -19,24 +19,22 @@
 #include "GPUTPCGMMergedTrackHit.h"
 #include "GPUTPCClusterOccupancyMap.h"
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 
-GPUdi() void GPUParam::Slice2Global(int32_t iSlice, float x, float y, float z, float* X, float* Y, float* Z) const
+GPUdi() void GPUParam::Sector2Global(int32_t iSector, float x, float y, float z, float* X, float* Y, float* Z) const
 {
   // conversion of coordinates sector->global
-  *X = x * SliceParam[iSlice].CosAlpha - y * SliceParam[iSlice].SinAlpha;
-  *Y = y * SliceParam[iSlice].CosAlpha + x * SliceParam[iSlice].SinAlpha;
+  *X = x * SectorParam[iSector].CosAlpha - y * SectorParam[iSector].SinAlpha;
+  *Y = y * SectorParam[iSector].CosAlpha + x * SectorParam[iSector].SinAlpha;
   *Z = z;
 }
 
-GPUdi() void GPUParam::Global2Slice(int32_t iSlice, float X, float Y, float Z, float* x, float* y, float* z) const
+GPUdi() void GPUParam::Global2Sector(int32_t iSector, float X, float Y, float Z, float* x, float* y, float* z) const
 {
   // conversion of coordinates global->sector
-  *x = X * SliceParam[iSlice].CosAlpha + Y * SliceParam[iSlice].SinAlpha;
-  *y = Y * SliceParam[iSlice].CosAlpha - X * SliceParam[iSlice].SinAlpha;
+  *x = X * SectorParam[iSector].CosAlpha + Y * SectorParam[iSector].SinAlpha;
+  *y = Y * SectorParam[iSector].CosAlpha - X * SectorParam[iSector].SinAlpha;
   *z = Z;
 }
 
@@ -117,7 +115,7 @@ GPUdi() float GPUParam::GetSystematicClusterErrorC122(float x, float y, uint8_t
     return 0.f;
   }
   constexpr float dEdgeInv = 18.f / CAMath::Pi();
-  const float dy = (sector == (GPUCA_NSLICES / 2 + 1) ? 0.5f : -0.5f) * (y / x) * dEdgeInv + 0.5f;
+  const float dy = (sector == (GPUCA_NSECTORS / 2 + 1) ? 0.5f : -0.5f) * (y / x) * dEdgeInv + 0.5f;
   const float errC12 = rec.tpc.sysClusErrorC12Norm * occupancyTotal * dy;
   return errC12 * errC12;
 }
@@ -223,7 +221,6 @@ GPUdi() bool GPUParam::rejectEdgeClusterByY(float uncorrectedY, int32_t iRow, fl
   return CAMath::Abs(uncorrectedY) > (tpcGeometry.NPads(iRow) - 1) * 0.5f * tpcGeometry.PadWidth(iRow) + rec.tpc.rejectEdgeClustersMargin + trackSigmaY * rec.tpc.rejectEdgeClustersSigmaMargin;
 }
 
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #endif

GPU/GPUTracking/Base/GPUParamRTC.h

@@ -18,9 +18,7 @@
 #include "GPUParam.h"
 #include <string>
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 namespace gpu_rtc
 {
@@ -38,7 +36,6 @@ struct GPUParamRTC : public internal::GPUParam_t<gpu_rtc::GPUSettingsRec, gpu_rt
   static std::string generateRTCCode(const GPUParam& param, bool useConstexpr);
 };
 
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #endif

GPU/GPUTracking/Base/GPUProcessor.h

@@ -23,9 +23,7 @@
 #include <algorithm>
 #endif
 
-namespace o2
-{
-namespace gpu
+namespace o2::gpu
 {
 struct GPUTrackingInOutPointers;
 class GPUReconstruction;
@@ -157,7 +155,6 @@ class GPUProcessor
 
   friend class GPUTPCNeighboursFinder;
 };
-} // namespace gpu
-} // namespace o2
+} // namespace o2::gpu
 
 #endif