Skip to content

Commit 51fc357

Browse files
fweigfweig
committed
Forward MC for saturated clusters
1 parent 711eba5 commit 51fc357

11 files changed

Lines changed: 99 additions & 43 deletions

‎GPU/GPUTracking/Global/GPUChainTrackingClusterizer.cxx‎

Lines changed: 15 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1059,8 +1059,12 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
10591059
return;
10601060
}
10611061

1062-
if (propagateMCLabels && fragment.index == 0) {
1063-
clusterer.PrepareMC();
1062+
if (propagateMCLabels) {
1063+
if (fragment.index == 0) {
1064+
// Must be only called on the first fragment as some buffers are used across the whole timeframe
1065+
clusterer.AllocMCBuffers();
1066+
}
1067+
clusterer.InitMCBuffersForFragment();
10641068
clusterer.mPinputLabels = digitsMC->v[iSector];
10651069
if (clusterer.mPinputLabels == nullptr) {
10661070
GPUFatal("MC label container missing, sector %d", iSector);
@@ -1358,9 +1362,7 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
13581362

13591363
if (doGPU && propagateMCLabels) {
13601364
TransferMemoryResourceLinkToHost(RecoStep::TPCClusterFinding, clusterer.mScratchId, lane);
1361-
if (doGPU) {
1362-
SynchronizeStream(lane);
1363-
}
1365+
SynchronizeStream(lane);
13641366
runKernel<GPUTPCCFClusterizer>({GetGrid(clusterer.mPmemory->counters.nClusters, lane, GPUReconstruction::krnlDeviceType::CPU), {iSector}}, 1); // Computes MC labels
13651367
}
13661368
} // if (nRegularClusters != 0) {
@@ -1369,11 +1371,16 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
13691371
// TODO: Move this right after CheckPadBaseline once tail zeroing is moved into this kernel.
13701372
if (rec()->GetParam().rec.tpc.hipTailFilter) {
13711373
runKernel<GPUTPCCFHIPTailConnector>({GetGridBlk(GPUTPCGeometry::NROWS, lane), {iSector}});
1372-
runKernel<GPUTPCCFHIPClusterizer>({GetGridBlk(GPUTPCGeometry::NROWS, lane), {iSector}});
1374+
runKernel<GPUTPCCFHIPClusterizer>({GetGridBlk(GPUTPCGeometry::NROWS, lane), {iSector}}, 0);
13731375
if (doGPU && (nRegularClusters == 0 || GetProcessingSettings().debugLevel >= 3)) {
13741376
TransferMemoryResourceLinkToHost(RecoStep::TPCClusterFinding, clusterer.mMemoryId, lane);
13751377
SynchronizeStream(lane);
13761378
}
1379+
if (doGPU && propagateMCLabels) {
1380+
TransferMemoryResourceLinkToHost(RecoStep::TPCClusterFinding, clusterer.mScratchId, lane);
1381+
SynchronizeStream(lane);
1382+
runKernel<GPUTPCCFHIPClusterizer>({GetGrid(GPUTPCGeometry::NROWS, lane, GPUReconstruction::krnlDeviceType::CPU), {iSector}}, 1); // Computes MC labels
1383+
}
13771384
}
13781385

13791386
bool hasClusters = nRegularClusters != 0;
@@ -1445,20 +1452,16 @@ int32_t GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
14451452
}
14461453

14471454
if (not propagateMCLabels || not laneHasData[lane]) {
1448-
assert(propagateMCLabels ? mcLinearLabels.header.size() == nClsTotal : true);
14491455
continue;
14501456
}
14511457

14521458
runKernel<GPUTPCCFMCLabelFlattener, GPUTPCCFMCLabelFlattener::setRowOffsets>({GetGrid(GPUTPCGeometry::NROWS, lane, GPUReconstruction::krnlDeviceType::CPU), {iSector}});
14531459
GPUTPCCFMCLabelFlattener::setGlobalOffsetsAndAllocate(clusterer, mcLinearLabels);
14541460
runKernel<GPUTPCCFMCLabelFlattener, GPUTPCCFMCLabelFlattener::flatten>({GetGrid(GPUTPCGeometry::NROWS, lane, GPUReconstruction::krnlDeviceType::CPU), {iSector}}, &mcLinearLabels);
1455-
clusterer.clearMCMemory();
14561461
assert(propagateMCLabels ? mcLinearLabels.header.size() == nClsTotal : true);
14571462
}
1458-
if (propagateMCLabels) {
1459-
for (int32_t lane = 0; lane < maxLane; lane++) {
1460-
processors()->tpcClusterer[iSectorBase + lane].clearMCMemory();
1461-
}
1463+
for (int32_t lane = 0; lane < maxLane; lane++) {
1464+
processors()->tpcClusterer[iSectorBase + lane].FreeMCBuffers();
14621465
}
14631466
if (buildNativeHost && buildNativeGPU && anyLaneHasData) {
14641467
if (GetProcessingSettings().delayedOutput) {

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx‎

Lines changed: 31 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,7 @@
2020
#include "GPUCommonAlgorithm.h"
2121

2222
#ifndef GPUCA_GPUCODE
23+
#include "MCLabelAccumulator.h"
2324
#include "utils/VcShim.h"
2425
#endif
2526

@@ -504,7 +505,7 @@ GPUd() void GPUTPCCFHIPTailConnector::Thread<0>(int32_t nBlocks, int32_t nThread
504505
// ======== HIP Clusterizer Kernel ========
505506

506507
template <>
507-
GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer)
508+
GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, uint8_t onlyMC)
508509
{
509510
if (iBlock >= (int32_t)GPUTPCGeometry::NROWS) {
510511
return;
@@ -514,16 +515,22 @@ GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads,
514515
uint32_t nTails = clusterer.mPnHIPTails[row];
515516
nTails = CAMath::Min(nTails, (uint32_t)MaxHIPTailsPerRow - 1);
516517

517-
HIPTailDescriptor* tails = GetHIPTails(clusterer, row);
518+
const auto* tails = GetHIPTails(clusterer, row);
518519
const auto& fragment = clusterer.mPmemory->fragment;
519520

521+
auto* clusterPosInRow = clusterer.mPhipClusterPosInRow
522+
? clusterer.mPhipClusterPosInRow + row * MaxHIPTailsPerRow
523+
: nullptr;
524+
520525
for (uint32_t iTail = iThread + 1; iTail <= nTails; iTail += nThreads) {
521526

522527
const auto* tail = &tails[iTail];
523528
if (tail->iPrev != 0) {
524529
continue;
525530
}
526531

532+
CPU_ONLY(auto labelAcc = MCLabelAccumulator{clusterer});
533+
527534
float qTot = 0;
528535
float qMax = 0;
529536
float padSum = 0;
@@ -544,6 +551,8 @@ GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads,
544551
timeSum += tailWeight * tailTime;
545552
tailStart = CAMath::Min<uint32_t>(tailStart, tail->tailStart);
546553
tailEnd = CAMath::Max<uint32_t>(tailEnd, tail->tailEnd);
554+
555+
CPU_ONLY(labelAcc.collectTail(row, tail->pad, tail->tailStart, tail->tailEnd));
547556
}
548557

549558
const float weightSum = CAMath::Max(qTot, 1.f);
@@ -561,13 +570,26 @@ GPUd() void GPUTPCCFHIPClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads,
561570
cn.setSigmaPad(padSigma);
562571

563572
if (cn.qMax >= 1023) {
564-
// Cut off clusters where the tail connection failed for some reason
565-
// TODO: Deduplicate with GPUTPCCFClusterizer::sortIntoBuckets (can't call cross-kernel).
566-
// TODO: Add error reporting for row cluster overflow.
567-
uint32_t index = CAMath::AtomicAdd(&clusterer.mPclusterInRow[row], 1u);
568-
if (index < clusterer.mNMaxClusterPerRow) {
569-
clusterer.mPclusterByRow[clusterer.mNMaxClusterPerRow * row + index] = cn;
573+
574+
uint32_t index;
575+
576+
if (!onlyMC) {
577+
// Cut off clusters where the tail connection failed for some reason
578+
// TODO: Deduplicate with GPUTPCCFClusterizer::sortIntoBuckets (can't call cross-kernel).
579+
// TODO: Add error reporting for row cluster overflow.
580+
index = CAMath::AtomicAdd(&clusterer.mPclusterInRow[row], 1u);
581+
if (index < clusterer.mNMaxClusterPerRow) {
582+
clusterer.mPclusterByRow[clusterer.mNMaxClusterPerRow * row + index] = cn;
583+
}
584+
if (clusterPosInRow) {
585+
clusterPosInRow[iTail] = index;
586+
}
587+
} else {
588+
index = clusterPosInRow[iTail];
570589
}
590+
591+
CPU_ONLY(labelAcc.commit(row, index, clusterer.mNMaxClusterPerRow));
571592
}
572-
}
593+
594+
} // for (uint32_t iTail = iThread + 1; iTail <= nTails; iTail += nThreads)
573595
}

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -193,7 +193,7 @@ class GPUTPCCFHIPClusterizer : public GPUKernelTemplate
193193
}
194194

195195
template <int32_t iKernel = defaultKernel>
196-
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer);
196+
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUSharedMemory& smem, processorType& clusterer, uint8_t onlyMC);
197197
};
198198

199199
} // namespace o2::gpu

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ GPUdii() void GPUTPCCFClusterizer::Thread<0>(int32_t nBlocks, int32_t nThreads,
3333
CfArray2D<PackedCharge> chargeMap(reinterpret_cast<PackedCharge*>(clusterer.mPchargeMap));
3434
CPU_ONLY(MCLabelAccumulator labelAcc(clusterer));
3535

36-
tpc::ClusterNative* clusterOut = (onlyMC) ? nullptr : clusterer.mPclusterByRow;
36+
tpc::ClusterNative* clusterOut = onlyMC ? nullptr : clusterer.mPclusterByRow;
3737

3838
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);
3939
}

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.inc‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -112,7 +112,6 @@ GPUdii() void GPUTPCCFClusterizer::updateClusterInner(
112112
PackedCharge p = buf[N * lid + i];
113113

114114
Charge q = cluster->updateInner(p, d);
115-
116115
CPU_ONLY(labelAcc->collect(pos.delta(d), q));
117116

118117
aboveThreshold |= (uint8_t(q > calib.tpc.cfInnerThreshold) << i);
@@ -139,9 +138,7 @@ GPUdii() void GPUTPCCFClusterizer::updateClusterOuter(
139138

140139
Delta2 d = cfconsts::OuterNeighbors[i];
141140

142-
Charge q = cluster->updateOuter(p, d);
143-
static_cast<void>(q); // Avoid unused varible warning on GPU.
144-
141+
[[maybe_unused]] Charge q = cluster->updateOuter(p, d);
145142
CPU_ONLY(labelAcc->collect(pos.delta(d), q));
146143
}
147144
}

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFMCLabelFlattener.cxx‎

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -49,12 +49,7 @@ GPUd() void GPUTPCCFMCLabelFlattener::Thread<GPUTPCCFMCLabelFlattener::setRowOff
4949
const Row row = get_global_id(0);
5050
const size_t clusterInRow = clusterer.mPclusterInRow[row];
5151

52-
// Label Flattener assumes 1 label container per cluster,
53-
// but HIP clusters don't support MC labels yet and containers are missing for those clusters.
54-
// So append empty label container for each HIP cluster.
55-
// Note: This assumes that HIP cluster are store behind regular clusters!
5652
auto& labels = clusterer.mPlabelsByRow[row].data;
57-
labels.resize(std::max(labels.size(), clusterInRow));
5853

5954
uint32_t labelCount = 0;
6055

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.cxx‎

Lines changed: 16 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,8 @@
2727
#include "CfArray2D.h"
2828
#include "GPUTPCCFCheckPadBaseline.h"
2929

30+
#include <algorithm>
31+
3032
using namespace o2::gpu;
3133
using namespace o2::tpc;
3234

@@ -38,7 +40,7 @@ void GPUTPCClusterFinder::InitializeProcessor()
3840
GPUTPCClusterFinder::~GPUTPCClusterFinder()
3941
{
4042
delete[] mMinMaxCN;
41-
clearMCMemory();
43+
FreeMCBuffers();
4244
}
4345

4446
void* GPUTPCClusterFinder::SetPointersMemory(void* mem)
@@ -86,8 +88,10 @@ void* GPUTPCClusterFinder::SetPointersScratch(void* mem)
8688
computePointerWithAlignment(mem, mPfilteredPeakPositions, mNMaxClusters);
8789
if (mRec->GetProcessingSettings().runMC) {
8890
computePointerWithAlignment(mem, mPclusterPosInRow, mNMaxClusters);
91+
computePointerWithAlignment(mem, mPhipClusterPosInRow, GPUTPCGeometry::NROWS * GPUTPCCFHIPClusterizer::MaxHIPTailsPerRow);
8992
} else {
9093
mPclusterPosInRow = nullptr;
94+
mPhipClusterPosInRow = nullptr;
9195
}
9296
computePointerWithAlignment(mem, mPisPeak, mNMaxDigitsFragment);
9397
computePointerWithAlignment(mem, mPchargeMap, TPCMapMemoryLayout<decltype(*mPchargeMap)>::items(mRec->GetProcessingSettings().overrideClusterizerFragmentLen));
@@ -165,17 +169,24 @@ uint32_t GPUTPCClusterFinder::getNSteps(size_t items) const
165169
return c;
166170
}
167171

168-
void GPUTPCClusterFinder::PrepareMC()
172+
void GPUTPCClusterFinder::AllocMCBuffers()
169173
{
170174
assert(mNMaxClusterPerRow > 0);
171175

172-
clearMCMemory();
173-
mPindexMap = new uint32_t[TPCMapMemoryLayout<decltype(*mPindexMap)>::items(mRec->GetProcessingSettings().overrideClusterizerFragmentLen)];
176+
FreeMCBuffers();
177+
const size_t nItems = TPCMapMemoryLayout<decltype(*mPindexMap)>::items(mRec->GetProcessingSettings().overrideClusterizerFragmentLen);
178+
mPindexMap = new uint32_t[nItems];
174179
mPlabelsByRow = new GPUTPCClusterMCInterimArray[GPUTPCGeometry::NROWS];
175180
mPlabelsInRow = new uint32_t[GPUTPCGeometry::NROWS];
176181
}
177182

178-
void GPUTPCClusterFinder::clearMCMemory()
183+
void GPUTPCClusterFinder::InitMCBuffersForFragment()
184+
{
185+
const size_t nItems = TPCMapMemoryLayout<decltype(*mPindexMap)>::items(mRec->GetProcessingSettings().overrideClusterizerFragmentLen);
186+
std::fill_n(mPindexMap, nItems, uint32_t(-1));
187+
}
188+
189+
void GPUTPCClusterFinder::FreeMCBuffers()
179190
{
180191
delete[] mPindexMap;
181192
mPindexMap = nullptr;

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.h‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -92,8 +92,9 @@ class GPUTPCClusterFinder : public GPUProcessor
9292
uint32_t getNSteps(size_t items) const;
9393
void SetNMaxDigits(size_t nDigits, size_t nPages, size_t nDigitsFragment, size_t nDigitsEndpointMax);
9494

95-
void PrepareMC();
96-
void clearMCMemory();
95+
void AllocMCBuffers();
96+
void InitMCBuffersForFragment();
97+
void FreeMCBuffers();
9798
#endif
9899
uint8_t* mPzs = nullptr;
99100
ZSOffset* mPzsOffsets = nullptr;
@@ -107,6 +108,7 @@ class GPUTPCClusterFinder : public GPUProcessor
107108
uint32_t* mPclusterPosInRow = nullptr; // store the index where the corresponding cluster is stored in a bucket.
108109
// Required when MC are enabled to write the mc data to the correct position.
109110
// Set to >= mNMaxClusterPerRow if cluster was discarded.
111+
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
110112
uint16_t* mPchargeMap = nullptr;
111113
uint8_t* mPpeakMap = nullptr;
112114
uint32_t* mPindexMap = nullptr;

‎GPU/GPUTracking/TPCClusterFinder/MCLabelAccumulator.cxx‎

Lines changed: 24 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,13 +26,20 @@ MCLabelAccumulator::MCLabelAccumulator(GPUTPCClusterFinder& clusterer)
2626
{
2727
}
2828

29-
void MCLabelAccumulator::collect(const CfChargePos& pos, Charge q)
29+
MCLabelAccumulator::~MCLabelAccumulator() {}
30+
31+
void MCLabelAccumulator::collect(const CfChargePos& pos, float q)
3032
{
3133
if (q == 0 || !engaged()) {
3234
return;
3335
}
3436

37+
// Use -1 as sentinel to indicate a missing label.
38+
// Can't use zero charge, as HIP filter will zero existing digits.
3539
uint32_t index = mIndexMap[pos];
40+
if (index == uint32_t(-1)) {
41+
return;
42+
}
3643

3744
const auto& labels = mLabels->getLabels(index);
3845

@@ -51,6 +58,22 @@ void MCLabelAccumulator::collect(const CfChargePos& pos, Charge q)
5158
}
5259
}
5360

61+
void MCLabelAccumulator::collectTail(tpccf::Row row, tpccf::Pad pad, uint16_t tailStart, uint16_t tailEnd)
62+
{
63+
if (!engaged()) {
64+
return;
65+
}
66+
67+
const auto basePos = CfChargePos{row, pad, 0};
68+
69+
for (uint16_t t = tailStart; t < tailEnd; t++) {
70+
const auto pos = basePos.delta({0, (int16_t)t});
71+
// Charge passed to collect() doesn't matter, collect() skips zero charges
72+
// But we know there's an interesting value, but it was zeroed in chargeMap by tail filter
73+
collect(pos, 1023.f);
74+
}
75+
}
76+
5477
void MCLabelAccumulator::commit(Row row, uint32_t indexInRow, uint32_t maxElemsPerBucket)
5578
{
5679
if (indexInRow >= maxElemsPerBucket || !engaged()) {

‎GPU/GPUTracking/TPCClusterFinder/MCLabelAccumulator.h‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -43,12 +43,15 @@ class MCLabelAccumulator
4343

4444
public:
4545
MCLabelAccumulator(GPUTPCClusterFinder&);
46+
~MCLabelAccumulator(); // Explicit destructor to allow forward declaring MCCompLabel with std::vector
4647

47-
void collect(const CfChargePos&, tpccf::Charge);
48+
void collect(const CfChargePos& pos, float q);
49+
50+
void collectTail(tpccf::Row row, tpccf::Pad pad, uint16_t tailStart, uint16_t tailEnd);
4851

4952
bool engaged() const { return mLabels != nullptr && mOutput != nullptr; }
5053

51-
void commit(tpccf::Row, uint32_t, uint32_t);
54+
void commit(tpccf::Row row, uint32_t indexInRow, uint32_t maxElemsPerBucket);
5255

5356
private:
5457
CfArray2D<const uint32_t> mIndexMap;

0 commit comments

Comments
 (0)