GPU TPC: Rename slice to sector

Author: davidrohr

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

@@ -41,14 +41,14 @@ namespace 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;

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

@@ -39,11 +39,11 @@ 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 +66,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 +88,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,8 +108,8 @@ 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;
 };

GPU/GPUTracking/Base/GPUParam.inc

@@ -24,19 +24,19 @@ namespace o2
 namespace 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 +117,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;
 }

GPU/GPUTracking/Base/GPUReconstruction.cxx

@@ -90,9 +90,9 @@ GPUReconstruction::GPUReconstruction(const GPUSettingsDeviceBackend& cfg) : mHos
   new (&mGRPSettings) GPUSettingsGRP;
   param().SetDefaults(&mGRPSettings);
   mMemoryScalers.reset(new GPUMemorySizeScalers);
-  for (uint32_t i = 0; i < NSLICES; i++) {
-    processors()->tpcTrackers[i].SetSlice(i); // TODO: Move to a better place
-    processors()->tpcClusterer[i].mISlice = i;
+  for (uint32_t i = 0; i < NSECTORS; i++) {
+    processors()->tpcTrackers[i].SetSector(i); // TODO: Move to a better place
+    processors()->tpcClusterer[i].mISector = i;
   }
 #ifndef GPUCA_NO_ROOT
   mROOTDump = GPUROOTDumpCore::getAndCreate();
@@ -273,8 +273,8 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice()
     if (mProcessingSettings.trackletSelectorInPipeline < 0) {
       mProcessingSettings.trackletSelectorInPipeline = 1;
     }
-    if (mProcessingSettings.trackletSelectorSlices < 0) {
-      mProcessingSettings.trackletSelectorSlices = 1;
+    if (mProcessingSettings.trackletSelectorSectors < 0) {
+      mProcessingSettings.trackletSelectorSectors = 1;
     }
   }
   if (mProcessingSettings.createO2Output > 1 && mProcessingSettings.runQA && mProcessingSettings.qcRunFraction == 100.f) {
@@ -334,14 +334,14 @@ int32_t GPUReconstruction::InitPhaseBeforeDevice()
   }
 
   if (mProcessingSettings.nTPCClustererLanes == -1) {
-    mProcessingSettings.nTPCClustererLanes = (GetRecoStepsGPU() & RecoStep::TPCClusterFinding) ? 3 : std::max<int32_t>(1, std::min<int32_t>(GPUCA_NSLICES, mProcessingSettings.inKernelParallel ? (mMaxHostThreads >= 4 ? std::min<int32_t>(mMaxHostThreads / 2, mMaxHostThreads >= 32 ? GPUCA_NSLICES : 4) : 1) : mMaxHostThreads));
+    mProcessingSettings.nTPCClustererLanes = (GetRecoStepsGPU() & RecoStep::TPCClusterFinding) ? 3 : std::max<int32_t>(1, std::min<int32_t>(GPUCA_NSECTORS, mProcessingSettings.inKernelParallel ? (mMaxHostThreads >= 4 ? std::min<int32_t>(mMaxHostThreads / 2, mMaxHostThreads >= 32 ? GPUCA_NSECTORS : 4) : 1) : mMaxHostThreads));
   }
   if (mProcessingSettings.overrideClusterizerFragmentLen == -1) {
     mProcessingSettings.overrideClusterizerFragmentLen = ((GetRecoStepsGPU() & RecoStep::TPCClusterFinding) || (mMaxHostThreads / mProcessingSettings.nTPCClustererLanes >= 3)) ? TPC_MAX_FRAGMENT_LEN_GPU : TPC_MAX_FRAGMENT_LEN_HOST;
   }
-  if (mProcessingSettings.nTPCClustererLanes > GPUCA_NSLICES) {
+  if (mProcessingSettings.nTPCClustererLanes > GPUCA_NSECTORS) {
     GPUError("Invalid value for nTPCClustererLanes: %d", mProcessingSettings.nTPCClustererLanes);
-    mProcessingSettings.nTPCClustererLanes = GPUCA_NSLICES;
+    mProcessingSettings.nTPCClustererLanes = GPUCA_NSECTORS;
   }
 
   if (mProcessingSettings.doublePipeline && (mChains.size() != 1 || mChains[0]->SupportsDoublePipeline() == false || !IsGPU() || mProcessingSettings.memoryAllocationStrategy != GPUMemoryResource::ALLOCATION_GLOBAL)) {

GPU/GPUTracking/Base/GPUReconstruction.h

@@ -31,7 +31,7 @@
 #include "GPUOutputControl.h"
 #include "GPUMemoryResource.h"
 #include "GPUConstantMem.h"
-#include "GPUTPCSliceOutput.h"
+#include "GPUTPCSectorOutput.h"
 #include "GPULogging.h"
 
 namespace o2
@@ -76,7 +76,7 @@ class GPUReconstruction
   GPUReconstruction& operator=(const GPUReconstruction&) = delete;
 
   // General definitions
-  constexpr static uint32_t NSLICES = GPUCA_NSLICES;
+  constexpr static uint32_t NSECTORS = GPUCA_NSECTORS;
 
   using GeometryType = GPUDataTypes::GeometryType;
   using DeviceType = GPUDataTypes::DeviceType;
@@ -94,8 +94,8 @@ class GPUReconstruction
 
   static DeviceType GetDeviceType(const char* type);
   enum InOutPointerType : uint32_t { CLUSTER_DATA = 0,
-                                     SLICE_OUT_TRACK = 1,
-                                     SLICE_OUT_CLUSTER = 2,
+                                     SECTOR_OUT_TRACK = 1,
+                                     SECTOR_OUT_CLUSTER = 2,
                                      MC_LABEL_TPC = 3,
                                      MC_INFO_TPC = 4,
                                      MERGED_TRACK = 5,
@@ -113,10 +113,10 @@ class GPUReconstruction
                                      TRD_SPACEPOINT = 17,
                                      TRD_TRIGGERRECORDS = 18,
                                      TF_SETTINGS = 19 };
-  static constexpr const char* const IOTYPENAMES[] = {"TPC HLT Clusters", "TPC Slice Tracks", "TPC Slice Track Clusters", "TPC Cluster MC Labels", "TPC Track MC Informations", "TPC Tracks", "TPC Track Clusters", "TRD Tracks", "TRD Tracklets",
+  static constexpr const char* const IOTYPENAMES[] = {"TPC HLT Clusters", "TPC Sector Tracks", "TPC Sector Track Clusters", "TPC Cluster MC Labels", "TPC Track MC Informations", "TPC Tracks", "TPC Track Clusters", "TRD Tracks", "TRD Tracklets",
                                                       "TPC Raw Clusters", "TPC Native Clusters", "TRD Tracklet MC Labels", "TPC Compressed Clusters", "TPC Digit", "TPC ZS Page", "TPC Native Clusters MC Labels", "TPC Digit MC Labeels",
                                                       "TRD Spacepoints", "TRD Triggerrecords", "TF Settings"};
-  static uint32_t getNIOTypeMultiplicity(InOutPointerType type) { return (type == CLUSTER_DATA || type == SLICE_OUT_TRACK || type == SLICE_OUT_CLUSTER || type == RAW_CLUSTERS || type == TPC_DIGIT || type == TPC_DIGIT_MC) ? NSLICES : 1; }
+  static uint32_t getNIOTypeMultiplicity(InOutPointerType type) { return (type == CLUSTER_DATA || type == SECTOR_OUT_TRACK || type == SECTOR_OUT_CLUSTER || type == RAW_CLUSTERS || type == TPC_DIGIT || type == TPC_DIGIT_MC) ? NSECTORS : 1; }
 
   // Functionality to create an instance of GPUReconstruction for the desired device
   static GPUReconstruction* CreateInstance(const GPUSettingsDeviceBackend& cfg);

GPU/GPUTracking/Base/GPUReconstructionCPU.cxx

@@ -18,8 +18,8 @@
 #include "GPUChain.h"
 
 #include "GPUTPCClusterData.h"
-#include "GPUTPCSliceOutput.h"
-#include "GPUTPCSliceOutCluster.h"
+#include "GPUTPCSectorOutput.h"
+#include "GPUTPCSectorOutCluster.h"
 #include "GPUTPCGMMergedTrack.h"
 #include "GPUTPCGMMergedTrackHit.h"
 #include "GPUTRDTrackletWord.h"