Skip to content

GPUTracking: rename the cluster finder's fragment to frag #15924

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
davidrohr merged 1 commit into from
Oct 10, 2026
Merged

GPUTracking: rename the cluster finder's fragment to frag #15924

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions GPU/GPUTracking/Global/GPUChainTracking.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
76 changes: 38 additions & 38 deletions GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx
View file
Open in desktop

Large diffs are not rendered by default.

22 changes: 11 additions & 11 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand All @@ -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)
Expand All @@ -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());
Expand All @@ -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;

Expand Down
22 changes: 11 additions & 11 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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) {

Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;

Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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++) {
Expand All @@ -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++) {
Expand Down Expand Up @@ -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));
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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);

Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
6 changes: 3 additions & 3 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -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;
Expand Down
20 changes: 10 additions & 10 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCCFDecodeZS.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down Expand Up @@ -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);

Expand All @@ -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) {
Expand Down Expand Up @@ -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)
}
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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) {
Expand Down Expand Up @@ -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)

Expand Down
8 changes: 4 additions & 4 deletions GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,7 +63,7 @@ class GPUTPCClusterFinder : public GPUProcessor
uint32_t maxTimeBin = 0;
uint32_t nPagesSubsector = 0;
} counters;
CfFragment fragment;
CfFragment frag;
};

struct ZSOffset {
Expand Down Expand Up @@ -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;
Expand Down
Loading
Loading