GPUTracking: rename the cluster finder's fragment to frag

fragment is a reserved word in MSL, where it qualifies a shader stage, so it
cannot name a member, a parameter or a local. The cluster finder used it for all
three, which was 38 of the errors left in the Metal compile.

Purely a rename, in code positions only: the option descriptions in
GPUSettingsList.h that mention fragments are strings and are untouched, as are
the comments.

Author: ktf

GPU/GPUTracking/Global/GPUChainTracking.h

@@ -300,11 +300,11 @@ class GPUChainTracking : public GPUChain
   int32_t RunTPCTrackingSectors_internal();
   int32_t RunTPCClusterizer_prepare(bool restorePointers, const GPUTPCExtraADC& extraADCs);
 #ifndef GPUCA_RUN2
-  std::pair<uint32_t, uint32_t> RunTPCClusterizer_transferZS(int32_t iSector, const CfFragment& fragment, int32_t lane, const GPUTPCExtraADC& extraADCs);
+  std::pair<uint32_t, uint32_t> RunTPCClusterizer_transferZS(int32_t iSector, const CfFragment& frag, int32_t lane, const GPUTPCExtraADC& extraADCs);
   void RunTPCClusterizer_compactPeaks(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int32_t stage, bool doGPU, int32_t lane);
-  std::pair<uint32_t, uint32_t> TPCClusterizerDecodeZSCount(uint32_t iSector, const CfFragment& fragment);
-  std::pair<uint32_t, uint32_t> TPCClusterizerDecodeZSCountUpdate(uint32_t iSector, const CfFragment& fragment);
-  void TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfFragment& fragment);
+  std::pair<uint32_t, uint32_t> TPCClusterizerDecodeZSCount(uint32_t iSector, const CfFragment& frag);
+  std::pair<uint32_t, uint32_t> TPCClusterizerDecodeZSCountUpdate(uint32_t iSector, const CfFragment& frag);
+  void TPCClusterizerEnsureZSOffsets(uint32_t iSector, const CfFragment& frag);
   void TPCClusterizerTransferExtraADC(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int lane, const GPUTPCExtraADC& extraADCs);
   void TPCClusterizerCheckExtraADCZeros(GPUTPCClusterFinder& clusterer, GPUTPCClusterFinder& clustererShadow, int lane, const GPUTPCExtraADC& extraADCs);
 #endif

GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx

@@ -24,11 +24,11 @@ template <>
 GPUdii() void GPUTPCCFChargeMapFiller::Thread<GPUTPCCFChargeMapFiller::fillIndexMap>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer)
 {
   CfArray2D<uint32_t> indexMap(clusterer.mPindexMap);
-  fillIndexMapImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer.mPmemory->fragment, clusterer.mPdigits, indexMap, clusterer.mPmemory->counters.nDigitsInFragment);
+  fillIndexMapImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer.mPmemory->frag, clusterer.mPdigits, indexMap, clusterer.mPmemory->counters.nDigitsInFragment);
 }
 
 GPUd() void GPUTPCCFChargeMapFiller::fillIndexMapImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread,
-                                                      const CfFragment& fragment,
+                                                      const CfFragment& frag,
                                                       const tpc::Digit* digits,
                                                       CfArray2D<uint32_t>& indexMap,
                                                       size_t maxDigit)
@@ -37,20 +37,20 @@ GPUd() void GPUTPCCFChargeMapFiller::fillIndexMapImpl(int32_t nBlocks, int32_t n
   if (idx >= maxDigit) {
     return;
   }
-  CPU_ONLY(idx += fragment.digitsStart);
+  CPU_ONLY(idx += frag.digitsStart);
   CPU_ONLY(tpc::Digit digit = digits[idx]);
-  CPU_ONLY(CfChargePos pos(digit.getRow(), digit.getPad(), fragment.toLocal(digit.getTimeStamp())));
+  CPU_ONLY(CfChargePos pos(digit.getRow(), digit.getPad(), frag.toLocal(digit.getTimeStamp())));
   CPU_ONLY(indexMap.safeWrite(pos, idx));
 }
 
 template <>
 GPUdii() void GPUTPCCFChargeMapFiller::Thread<GPUTPCCFChargeMapFiller::fillFromDigits>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer)
 {
   CfArray2D<PackedCharge> chargeMap(reinterpret_cast<PackedCharge*>(clusterer.mPchargeMap));
-  fillFromDigitsImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, clusterer.mPmemory->counters.nPositions, clusterer.mPdigits, clusterer.mPpositions, chargeMap);
+  fillFromDigitsImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->frag, clusterer.mPmemory->counters.nPositions, clusterer.mPdigits, clusterer.mPpositions, chargeMap);
 }
 
-GPUd() void GPUTPCCFChargeMapFiller::fillFromDigitsImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, processorType& clusterer, const CfFragment& fragment, size_t digitNum,
+GPUd() void GPUTPCCFChargeMapFiller::fillFromDigitsImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, processorType& clusterer, const CfFragment& frag, size_t digitNum,
                                                         const tpc::Digit* digits,
                                                         CfChargePos* positions,
                                                         CfArray2D<PackedCharge>& chargeMap)
@@ -59,9 +59,9 @@ GPUd() void GPUTPCCFChargeMapFiller::fillFromDigitsImpl(int32_t nBlocks, int32_t
   if (idx >= digitNum) {
     return;
   }
-  tpc::Digit digit = digits[fragment.digitsStart + idx];
+  tpc::Digit digit = digits[frag.digitsStart + idx];
 
-  CfChargePos pos(digit.getRow(), digit.getPad(), fragment.toLocal(digit.getTimeStamp()));
+  CfChargePos pos(digit.getRow(), digit.getPad(), frag.toLocal(digit.getTimeStamp()));
   positions[idx] = pos;
   float q = digit.getChargeFloat();
   q *= clusterer.GetConstantMem()->calibObjects.tpcPadGain->getGainCorrection(clusterer.mISector, digit.getRow(), digit.getPad());
@@ -77,10 +77,10 @@ GPUdii() void GPUTPCCFChargeMapFiller::Thread<GPUTPCCFChargeMapFiller::findFragm
 
   size_t nDigits = clusterer.mPmemory->counters.nDigits;
   const tpc::Digit* digits = clusterer.mPdigits;
-  size_t st = findTransition(clusterer.mPmemory->fragment.first(), digits, nDigits, 0);
-  size_t end = findTransition(clusterer.mPmemory->fragment.last(), digits, nDigits, st);
+  size_t st = findTransition(clusterer.mPmemory->frag.first(), digits, nDigits, 0);
+  size_t end = findTransition(clusterer.mPmemory->frag.last(), digits, nDigits, st);
 
-  clusterer.mPmemory->fragment.digitsStart = st;
+  clusterer.mPmemory->frag.digitsStart = st;
 
   size_t elems = end - st;
 

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx

@@ -236,7 +236,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineGPU(int32_t nBlocks, int32_t
     return;
   }
 
-  const CfFragment& fragment = clusterer.mPmemory->fragment;
+  const CfFragment& frag = clusterer.mPmemory->frag;
   const bool hipFilterOn = clusterer.Param().rec.tpc.hipTailFilter;
   const Charge hipTailThreshold = clusterer.Param().rec.tpc.hipTailFilterThreshold;
   const Charge hipTailFilterAlpha = clusterer.Param().rec.tpc.hipTailFilterAlpha;
@@ -265,7 +265,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineGPU(int32_t nBlocks, int32_t
   // saturated signal in overlap region can create tails in the next fragment
   // even when cleared in current fragment as they're decoded twice
   const TPCFragmentTime firstTB = 0;
-  const TPCFragmentTime lastTB = fragment.length;
+  const TPCFragmentTime lastTB = frag.length;
 
   for (uint16_t t = firstTB; t < lastTB; t += NumOfCachedTBs) {
 
@@ -418,7 +418,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t
   const int32_t nPads = geo.NPads(row);
   const int32_t nVecPads = (nPads + PadsPerCacheline - 1) / PadsPerCacheline;
 
-  const CfFragment& fragment = clusterer.mPmemory->fragment;
+  const CfFragment& frag = clusterer.mPmemory->frag;
   const bool hipFilterOn = clusterer.Param().rec.tpc.hipTailFilter;
   const Charge hipTailThreshold = clusterer.Param().rec.tpc.hipTailFilterThreshold;
   const Charge hipTailFilterAlpha = clusterer.Param().rec.tpc.hipTailFilterAlpha;
@@ -449,7 +449,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t
   std::vector<Short8> activeHIPTailSatEndV(nVecPads, -1);   // end of that plateau, extended while the plateau continues into later chunks
   std::vector<Charge8> tailFilterChargeV(nVecPads, Charge8{Vc::Zero});
 
-  for (int16_t t = 0; t < fragment.length; t += NumOfCachedTBs) {
+  for (int16_t t = 0; t < frag.length; t += NumOfCachedTBs) {
 
     bool hasAnyTrigger = false;
 
@@ -478,7 +478,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t
       for (tpccf::TPCFragmentTime localtime = 0; localtime < NumOfCachedTBs; localtime++) {
 
         const uint16_t* packedChargeStart = reinterpret_cast<uint16_t*>(&chargeMap[basePos.delta({0, localtime})]);
-        const UShort8 packedCharges = t + localtime < fragment.length
+        const UShort8 packedCharges = t + localtime < frag.length
                                         ? UShort8{packedChargeStart, Vc::Aligned}
                                         : UShort8{Vc::Zero};
         const auto isCharge = packedCharges != 0;
@@ -661,7 +661,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t
       tailFilterChargeV[iVecPad] = tailFilterCharge;
 
     } // for (int32_t iVecPad = 0; iVecPad < nVecPads; iVecPad++)
-  } // for (auto t = 0; t < fragment.length; t += TimebinsPerCacheline)
+  } // for (auto t = 0; t < frag.length; t += TimebinsPerCacheline)
 
   // Close old tails for all pads, open new tails in case of overlap
   for (int16_t iVecPad = 0; iVecPad < nVecPads; iVecPad++) {
@@ -672,7 +672,7 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t
     const auto activeHIPTailSatEnd = activeHIPTailSatEndV[iVecPad];
 
     const auto shouldCloseTail = activeHIPTailStart > -1;
-    activeHIPTailEnd(shouldCloseTail && activeHIPTailEnd < 0) = fragment.length;
+    activeHIPTailEnd(shouldCloseTail && activeHIPTailEnd < 0) = frag.length;
 
     if (hipFilterOn && shouldCloseTail.isNotEmpty()) {
       for (int16_t p = 0; p < PadsPerCacheline; p++) {
@@ -728,8 +728,8 @@ GPUd() void GPUTPCCFCheckPadBaseline::CheckBaselineCPU(int32_t nBlocks, int32_t
 
 GPUd() void GPUTPCCFCheckPadBaseline::updatePadBaseline(int32_t pad, const GPUTPCClusterFinder& clusterer, int32_t totalCharges, int32_t consecCharges, Charge maxCharge)
 {
-  const CfFragment& fragment = clusterer.mPmemory->fragment;
-  const int32_t totalChargesBaseline = clusterer.Param().rec.tpc.maxTimeBinAboveThresholdIn1000Bin * fragment.lengthWithoutOverlap() / 1000;
+  const CfFragment& frag = clusterer.mPmemory->frag;
+  const int32_t totalChargesBaseline = clusterer.Param().rec.tpc.maxTimeBinAboveThresholdIn1000Bin * frag.lengthWithoutOverlap() / 1000;
   const int32_t consecChargesBaseline = clusterer.Param().rec.tpc.maxConsecTimeBinAboveThreshold;
   const uint16_t saturationThreshold = clusterer.Param().rec.tpc.noisyPadSaturationThreshold;
   const bool isNoisy = (!saturationThreshold || maxCharge < saturationThreshold) && ((totalChargesBaseline > 0 && totalCharges >= totalChargesBaseline) || (consecChargesBaseline > 0 && consecCharges >= consecChargesBaseline));
@@ -821,7 +821,7 @@ GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads,
   nTails = CAMath::Min(nTails, (uint32_t)MaxHIPTailsPerRow - 1);
 
   const auto* tails = GetHIPTails(clusterer, row);
-  const auto& fragment = clusterer.mPmemory->fragment;
+  const auto& frag = clusterer.mPmemory->frag;
 
   auto* clusterPosInRow = clusterer.mPhipClusterPosInRow
                             ? clusterer.mPhipClusterPosInRow + row * MaxHIPTailsPerRow
@@ -879,7 +879,7 @@ GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads,
       // Use the middle of the saturated plateau, averaged over all tails of the cluster that were triggered by saturation on their own pad.
       // Computed only here: chains consisting only of tails inherited from neighboring pads have no saturated plateau,
       // but these never contain a saturated charge and are dropped by the qMax cut.
-      const float clusterTime = fragment.start + satTimeSum / nSatTails - clusterer.Param().rec.tpc.clustersShiftTimebinsClusterizer;
+      const float clusterTime = frag.start + satTimeSum / nSatTails - clusterer.Param().rec.tpc.clustersShiftTimebinsClusterizer;
       assert(!CAMath::IsNaN(clusterTime));
       cn.setTimeFlags(clusterTime, 0);
 

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx

@@ -35,5 +35,5 @@ GPUdii() void GPUTPCCFClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads,
 
   tpc::ClusterNative* clusterOut = onlyMC ? nullptr : clusterer.mPclusterByRow;
 
-  GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->fragment, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, true);
+  GPUTPCCFClusterizer::computeClustersImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), clusterer, clusterer.mPmemory->frag, smem, chargeMap, clusterer.mPfilteredPeakPositions, clusterer.Param().rec, CPU_PTR(&labelAcc), clusterer.mPmemory->counters.nClusters, clusterer.mNMaxClusterPerRow, clusterer.mPclusterInRow, clusterOut, clusterer.mPclusterPosInRow, true);
 }

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx

@@ -17,7 +17,7 @@
 
 GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread,
                                                        processorType& clusterer,
-                                                       const CfFragment& fragment,
+                                                       const CfFragment& frag,
                                                        GPUSharedMemory& smem,
                                                        const CfArray2D<PackedCharge>& chargeMap,
                                                        const CfChargePos* filteredPeakPositions,
@@ -54,14 +54,14 @@ GPUdii() void GPUTPCCFClusterizer::computeClustersImpl(int32_t nBlocks, int32_t
   if (idx >= clusternum) {
     return;
   }
-  if (fragment.isOverlap(pos.time())) {
+  if (frag.isOverlap(pos.time())) {
     if (clusterPosInRow) {
       clusterPosInRow[idx] = maxClusterPerRow;
     }
     return;
   }
   tpc::ClusterNative myCluster;
-  pc.finalize(pos, charge, fragment.start);
+  pc.finalize(pos, charge, frag.start);
   bool rejectCluster = !pc.toNative(pos, charge, myCluster, clusterer.Param(), chargeMap);
   if (!isAccepted) {
     rejectCluster = true;

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc

@@ -172,13 +172,13 @@ GPUdii() void GPUTPCCFDecodeZS::decode(GPUTPCClusterFinder& clusterer, GPUShared
                     seqLen = rowData[(nSeq + 1) * 2] - rowData[nSeq * 2];
                     pad = rowData[nSeq++ * 2 + 1];
                   }
-                  const CfFragment& fragment = clusterer.mPmemory->fragment;
+                  const CfFragment& frag = clusterer.mPmemory->frag;
                   TPCTime globalTime = timeBin + l;
-                  bool discardTimeBin = not fragment.contains(globalTime);
+                  bool discardTimeBin = not frag.contains(globalTime);
                   discardTimeBin |= (tpcTimeBinCut > 0 && globalTime > tpcTimeBinCut);
 
                   Row row = rowOffset + m;
-                  CfChargePos pos(row, Pad(pad), discardTimeBin ? INVALID_TIME_BIN : fragment.toLocal(globalTime));
+                  CfChargePos pos(row, Pad(pad), discardTimeBin ? INVALID_TIME_BIN : frag.toLocal(globalTime));
                   positions[nDigitsTmp++] = pos;
 
                   if (!discardTimeBin) {
@@ -219,7 +219,7 @@ GPUdii() void GPUTPCCFDecodeZSLink::Thread<0>(int32_t nBlocks, int32_t nThreads,
 
 GPUd() size_t GPUTPCCFDecodeZSLink::DecodePage(GPUSharedMemory& smem, DecodeCtx& ctx)
 {
-  const CfFragment& fragment = ctx.clusterer.mPmemory->fragment;
+  const CfFragment& frag = ctx.clusterer.mPmemory->frag;
 
   const auto* rdHdr = ConsumeHeader<header::RAWDataHeader>(ctx.page);
 
@@ -246,7 +246,7 @@ GPUd() size_t GPUTPCCFDecodeZSLink::DecodePage(GPUSharedMemory& smem, DecodeCtx&
 
     nDecoded += nAdc;
 
-    bool discardTimeBin = not fragment.contains(timeBin);
+    bool discardTimeBin = not frag.contains(timeBin);
     discardTimeBin |= (ctx.tpcTimeBinCut > 0 && timeBin > ctx.tpcTimeBinCut);
 
     if (discardTimeBin) {
@@ -331,9 +331,9 @@ GPUd() void GPUTPCCFDecodeZSLink::DecodeTB(
     }
 
     o2::tpc::PadPos padAndRow = GetPadAndRowFromFEC(ctx.clusterer, cru, rawFECChannel, fecInPartition);
-    const CfFragment& fragment = ctx.clusterer.mPmemory->fragment;
+    const CfFragment& frag = ctx.clusterer.mPmemory->frag;
     float charge = ADCToFloat(adc, DECODE_MASK, DECODE_BITS_FACTOR);
-    WriteCharge(ctx.clusterer, charge, padAndRow, fragment.toLocal(timeBin), ctx.pageDigitOffset + myOffset);
+    WriteCharge(ctx.clusterer, charge, padAndRow, frag.toLocal(timeBin), ctx.pageDigitOffset + myOffset);
 
   } // for (uint8_t i = iThread; blockOffset < nAdc; i += NThreads)
 }
@@ -651,7 +651,7 @@ GPUd() int16_t GPUTPCCFDecodeZSDenseLink::DecodeTB(
   constexpr int32_t NTHREADS = GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCFDecodeZSDenseLink);
   static_assert(NTHREADS == GPUCA_WARP_SIZE, "Decoding TB Headers in parallel assumes block size is a single warp.");
 
-  const CfFragment& fragment = ctx.clusterer.mPmemory->fragment;
+  const CfFragment& frag = ctx.clusterer.mPmemory->frag;
 
   // Read timebin block header
   uint16_t tbbHdr = ConsumeByte(ctx.page);
@@ -721,7 +721,7 @@ GPUd() int16_t GPUTPCCFDecodeZSDenseLink::DecodeTB(
   const uint8_t* adcData = ConsumeBytes(ctx.page, (nSamplesInTB * DECODE_BITS + 7) / 8);
   MAYBE_PAGE_OVERFLOW(ctx.page);
 
-  bool discardTimeBin = not fragment.contains(timeBin);
+  bool discardTimeBin = not frag.contains(timeBin);
   discardTimeBin |= (ctx.tpcTimeBinCut > 0 && timeBin > ctx.tpcTimeBinCut);
 
   if (discardTimeBin) {
@@ -754,7 +754,7 @@ GPUd() int16_t GPUTPCCFDecodeZSDenseLink::DecodeTB(
     o2::tpc::PadPos padAndRow = GetPadAndRowFromFEC(ctx.clusterer, cru, rawFECChannelLink, smem.linkIds[iLink]);
 
     float charge = ADCToFloat(adc, DECODE_MASK, DECODE_BITS_FACTOR);
-    WriteCharge(ctx.clusterer, charge, padAndRow, fragment.toLocal(timeBin), ctx.pageDigitOffset + sample);
+    WriteCharge(ctx.clusterer, charge, padAndRow, frag.toLocal(timeBin), ctx.pageDigitOffset + sample);
 
   } // for (uint16_t sample = iThread; sample < nSamplesInTB; sample += NTHREADS)
 

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx

@@ -63,7 +63,7 @@ class GPUTPCClusterFinder : public GPUProcessor
       uint32_t maxTimeBin = 0;
       uint32_t nPagesSubsector = 0;
     } counters;
-    CfFragment fragment;
+    CfFragment frag;
   };
 
   struct ZSOffset {
@@ -105,9 +105,9 @@ class GPUTPCClusterFinder : public GPUProcessor
   CfChargePos* mPpeakPositions = nullptr;
   CfChargePos* mPfilteredPeakPositions = nullptr;
   uint8_t* mPisPeak = nullptr;
-  uint32_t* mPclusterPosInRow = nullptr; // store the index where the corresponding cluster is stored in a bucket.
-                                         // Required when MC are enabled to write the mc data to the correct position.
-                                         // Set to >= mNMaxClusterPerRow if cluster was discarded.
+  uint32_t* mPclusterPosInRow = nullptr;    // store the index where the corresponding cluster is stored in a bucket.
+                                            // Required when MC are enabled to write the mc data to the correct position.
+                                            // Set to >= mNMaxClusterPerRow if cluster was discarded.
   uint32_t* mPhipClusterPosInRow = nullptr; // Identical to mPclusterPosInRow. Need a seperate array for HIP cluster because tail index is used to identify clusters across GPU and CPU
   uint16_t* mPchargeMap = nullptr;
   uint8_t* mPpeakMap = nullptr;