Skip to content

Commit 60b41a1

Browse files
fweigdavidrohr
fweig
and
davidrohr
authored
TPCClusterFinder: Add optional filter for noisy pads. (#4850)
* TPCClusterFinder: Add optional step to check if pads have lost baseline. * fixfelix Co-authored-by: David Rohr <drohr@jwdt.org>
1 parent 949d94f commit 60b41a1

15 files changed

Lines changed: 209 additions & 9 deletions

‎GPU/Common/GPUDefGPUParameters.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -352,6 +352,7 @@
352352

353353
#define GPUCA_THREAD_COUNT_SCAN 512 // TODO: WARNING!!! Must not be GPUTYPE-dependent right now! // TODO: Fix!
354354

355+
#define GPUCA_LB_GPUTPCCFCheckPadBaseline GPUCA_WARP_SIZE
355356
#define GPUCA_LB_GPUTPCCFChargeMapFiller_fillIndexMap GPUCA_LB_CLUSTER_FINDER
356357
#define GPUCA_LB_GPUTPCCFChargeMapFiller_fillFromDigits GPUCA_LB_CLUSTER_FINDER
357358
#define GPUCA_LB_GPUTPCCFChargeMapFiller_findFragmentStart GPUCA_LB_CLUSTER_FINDER

‎GPU/GPUTracking/Base/GPUReconstructionIncludesDevice.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -77,6 +77,7 @@ using namespace GPUCA_NAMESPACE::gpu;
7777
#include "GPUTPCCFClusterizer.cxx"
7878
#include "GPUTPCCFDeconvolution.cxx"
7979
#include "GPUTPCCFMCLabelFlattener.cxx"
80+
#include "GPUTPCCFCheckPadBaseline.cxx"
8081
#include "GPUTPCCFDecodeZS.cxx"
8182
#include "GPUTPCCFGather.cxx"
8283

‎GPU/GPUTracking/Base/GPUReconstructionKernels.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -70,6 +70,7 @@ GPUCA_KRNL_LB((GPUTPCCompressionGatherKernels, buffered64 ), (simple), (), ())
7070
GPUCA_KRNL_LB((GPUTPCCompressionGatherKernels, buffered128 ), (simple), (), ())
7171
GPUCA_KRNL_LB((GPUTPCCompressionGatherKernels, multiBlock ), (simple), (), ())
7272

73+
GPUCA_KRNL_LB((GPUTPCCFCheckPadBaseline ), (single), (), ())
7374
GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, fillIndexMap ), (single), (), ())
7475
GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, fillFromDigits ), (single), (), ())
7576
GPUCA_KRNL_LB((GPUTPCCFChargeMapFiller, findFragmentStart ), (single), (, char setPositions), (, setPositions))

‎GPU/GPUTracking/Base/GPUSettingsList.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,9 @@ AddOptionRTC(tpcTubeMaxSize2, float, 2.5f * 2.5f, "", 0, "Square of max tube siz
5050
AddOptionRTC(trdMinTrackPt, float, .5f, "", 0, "Min Pt for tracks to be propagated through the TRD")
5151
AddOptionRTC(trdMaxChi2, float, 15.f, "", 0, "Max chi2 for TRD tracklets to be matched to a track")
5252
AddOptionRTC(trdPenaltyChi2, float, 12.f, "", 0, "Chi2 penalty for no available TRD tracklet (effective chi2 cut value)")
53+
AddOptionRTC(noisyPadsQuickCheck, unsigned char, 1, "", 0, "Only check first fragment for noisy pads instead of all fragments (when test is enabled).")
54+
AddOptionRTC(maxTimeBinAboveThresholdIn1000Bin, unsigned short, 500, "", 0, "Except pad from cluster finding if total number of charges in a fragment is above this baseline (disable = 0)")
55+
AddOptionRTC(maxConsecTimeBinAboveThreshold, unsigned short, 200, "", 0, "Except pad from cluster finding if number of consecutive charges in a fragment is above this baseline (disable = 0)")
5356
AddOptionRTC(tpcCFqmaxCutoff, unsigned char, 3, "", 0, "Cluster Finder rejects cluster with qmax below this threshold")
5457
AddOptionRTC(tpcCFqtotCutoff, unsigned char, 0, "", 0, "Cluster Finder rejects cluster with qtot below this threshold")
5558
AddOptionRTC(tpcCFinnerThreshold, unsigned char, 0, "", 0, "Cluster Finder extends cluster if inner charge above this threshold")

‎GPU/GPUTracking/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -163,6 +163,7 @@ if(ALIGPU_BUILD_TYPE STREQUAL "O2" OR CONFIG_O2_EXTENSIONS)
163163
TPCClusterFinder/GPUTPCClusterFinder.cxx
164164
TPCClusterFinder/ClusterAccumulator.cxx
165165
TPCClusterFinder/MCLabelAccumulator.cxx
166+
TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx
166167
TPCClusterFinder/GPUTPCCFStreamCompaction.cxx
167168
TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx
168169
TPCClusterFinder/GPUTPCCFPeakFinder.cxx

‎GPU/GPUTracking/Global/GPUChainTracking.cxx‎

Lines changed: 13 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1162,6 +1162,10 @@ int GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
11621162
using PeakMapType = decltype(*clustererShadow.mPpeakMap);
11631163
runKernel<GPUMemClean16>(GetGridAutoStep(lane, RecoStep::TPCClusterFinding), krnlRunRangeNone, {}, clustererShadow.mPchargeMap, TPCMapMemoryLayout<ChargeMapType>::items() * sizeof(ChargeMapType));
11641164
runKernel<GPUMemClean16>(GetGridAutoStep(lane, RecoStep::TPCClusterFinding), krnlRunRangeNone, {}, clustererShadow.mPpeakMap, TPCMapMemoryLayout<PeakMapType>::items() * sizeof(PeakMapType));
1165+
if (fragment.index == 0) {
1166+
using HasLostBaselineType = decltype(*clustererShadow.mPpadHasLostBaseline);
1167+
runKernel<GPUMemClean16>(GetGridAutoStep(lane, RecoStep::TPCClusterFinding), krnlRunRangeNone, {}, clustererShadow.mPpadHasLostBaseline, TPC_PADS_IN_SECTOR * sizeof(HasLostBaselineType));
1168+
}
11651169
DoDebugAndDump(RecoStep::TPCClusterFinding, 0, clusterer, &GPUTPCClusterFinder::DumpChargeMap, *mDebugFile, "Zeroed Charges");
11661170

11671171
if (mIOPtrs.tpcZS && mCFContext->nPagesSector[iSlice]) {
@@ -1221,6 +1225,14 @@ int GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
12211225
runKernel<GPUTPCCFChargeMapFiller, GPUTPCCFChargeMapFiller::fillIndexMap>(GetGrid(clusterer.mPmemory->counters.nDigitsInFragment, lane, GPUReconstruction::krnlDeviceType::CPU), {iSlice}, {});
12221226
}
12231227

1228+
bool checkForNoisyPads = (rec()->GetParam().rec.maxTimeBinAboveThresholdIn1000Bin > 0) || (rec()->GetParam().rec.maxConsecTimeBinAboveThreshold > 0);
1229+
checkForNoisyPads &= (rec()->GetParam().rec.noisyPadsQuickCheck ? fragment.index == 0 : true);
1230+
1231+
if (checkForNoisyPads) {
1232+
int nBlocks = TPC_PADS_IN_SECTOR / GPUTPCCFCheckPadBaseline::getPadsPerBlock(doGPU);
1233+
runKernel<GPUTPCCFCheckPadBaseline>(GetGridBlk(nBlocks, lane), {iSlice}, {});
1234+
}
1235+
12241236
runKernel<GPUTPCCFPeakFinder>(GetGrid(clusterer.mPmemory->counters.nPositions, lane), {iSlice}, {});
12251237
DoDebugAndDump(RecoStep::TPCClusterFinding, 0, clusterer, &GPUTPCClusterFinder::DumpPeaks, *mDebugFile);
12261238

@@ -1310,7 +1322,7 @@ int GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
13101322
transferRunning[lane] = 1;
13111323
}
13121324

1313-
if (not propagateMCLabels) {
1325+
if (not propagateMCLabels || clusterer.mPmemory->counters.nClusters == 0) {
13141326
continue;
13151327
}
13161328

‎GPU/GPUTracking/TPCClusterFinder/CfFragment.h‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,21 @@ struct CfFragment {
7272
return (hasBacklog ? t < OverlapTimebins : false) || (hasFuture ? t >= (length - OverlapTimebins) : false);
7373
}
7474

75+
GPUdi() tpccf::TPCFragmentTime lengthWithoutOverlap() const
76+
{
77+
return length - (hasBacklog ? OverlapTimebins : 0) - (hasFuture ? OverlapTimebins : 0);
78+
}
79+
80+
GPUdi() tpccf::TPCFragmentTime firstNonOverlapTimeBin() const
81+
{
82+
return (hasBacklog ? OverlapTimebins : 0);
83+
}
84+
85+
GPUdi() tpccf::TPCFragmentTime lastNonOverlapTimeBin() const
86+
{
87+
return length - (hasFuture ? OverlapTimebins : 0);
88+
}
89+
7590
GPUdi() tpccf::TPCFragmentTime toLocal(tpccf::TPCTime t) const
7691
{
7792
return t - first();

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.cxx‎

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,80 @@
1+
// Copyright CERN and copyright holders of ALICE O2. This software is
2+
// distributed under the terms of the GNU General Public License v3 (GPL
3+
// Version 3), copied verbatim in the file "COPYING".
4+
//
5+
// See http://alice-o2.web.cern.ch/license for full licensing information.
6+
//
7+
// In applying this license CERN does not waive the privileges and immunities
8+
// granted to it by virtue of its status as an Intergovernmental Organization
9+
// or submit itself to any jurisdiction.
10+
11+
/// \file GPUTPCCFCheckPadBaseline.h
12+
/// \author Felix Weiglhofer
13+
14+
#include "GPUTPCCFCheckPadBaseline.h"
15+
#include "Array2D.h"
16+
#include "PackedCharge.h"
17+
#include "clusterFinderDefs.h"
18+
19+
using namespace GPUCA_NAMESPACE::gpu;
20+
using namespace GPUCA_NAMESPACE::gpu::tpccf;
21+
22+
template <>
23+
GPUd() void GPUTPCCFCheckPadBaseline::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUSharedMemory& smem, processorType& clusterer)
24+
{
25+
static_assert(TPC_MAX_FRAGMENT_LEN % NumOfCachedTimebins == 0);
26+
27+
Array2D<PackedCharge> chargeMap(reinterpret_cast<PackedCharge*>(clusterer.mPchargeMap));
28+
29+
int totalCharges = 0;
30+
int consecCharges = 0;
31+
int maxConsecCharges = 0;
32+
33+
int localPadId = iThread / NumOfCachedTimebins;
34+
int localTimeBin = iThread % NumOfCachedTimebins;
35+
bool handlePad = localTimeBin == 0;
36+
int basePad = iBlock * PadsPerBlock;
37+
38+
CfFragment& fragment = clusterer.mPmemory->fragment;
39+
40+
ChargePos basePos = padToChargePos(basePad + localPadId, clusterer);
41+
42+
for (tpccf::TPCFragmentTime t = localTimeBin + fragment.firstNonOverlapTimeBin(); t < fragment.lastNonOverlapTimeBin(); t += NumOfCachedTimebins) {
43+
ChargePos pos = basePos.delta({0, t});
44+
smem.charges[localPadId][localTimeBin] = (pos.valid()) ? chargeMap[pos].unpack() : 0;
45+
GPUbarrierWarp();
46+
if (handlePad) {
47+
for (int i = 0; i < NumOfCachedTimebins; i++) {
48+
Charge q = smem.charges[localPadId][i];
49+
totalCharges += (q > 0);
50+
consecCharges = (q > 0) ? consecCharges + 1 : 0;
51+
maxConsecCharges = CAMath::Max(consecCharges, maxConsecCharges);
52+
}
53+
}
54+
}
55+
56+
GPUbarrierWarp();
57+
58+
if (handlePad) {
59+
int totalChargesBaseline = clusterer.Param().rec.maxTimeBinAboveThresholdIn1000Bin * fragment.lengthWithoutOverlap() / 1000;
60+
int consecChargesBaseline = clusterer.Param().rec.maxConsecTimeBinAboveThreshold;
61+
bool hasLostBaseline = (totalChargesBaseline > 0 && totalCharges >= totalChargesBaseline) || (consecChargesBaseline > 0 && maxConsecCharges >= consecChargesBaseline);
62+
clusterer.mPpadHasLostBaseline[basePad + localPadId] |= hasLostBaseline;
63+
}
64+
}
65+
66+
GPUd() ChargePos GPUTPCCFCheckPadBaseline::padToChargePos(int pad, const GPUTPCClusterFinder& clusterer)
67+
{
68+
const GPUTPCGeometry& geo = clusterer.Param().tpcGeometry;
69+
70+
int padOffset = 0;
71+
for (Row r = 0; r < TPC_NUM_OF_ROWS; r++) {
72+
int padInRow = pad - padOffset;
73+
if (0 <= padInRow && padInRow < geo.NPads(r)) {
74+
return ChargePos{r, Pad(padInRow), 0};
75+
}
76+
padOffset += geo.NPads(r);
77+
}
78+
79+
return ChargePos{0, 0, INVALID_TIME_BIN};
80+
}

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFCheckPadBaseline.h‎

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
// Copyright CERN and copyright holders of ALICE O2. This software is
2+
// distributed under the terms of the GNU General Public License v3 (GPL
3+
// Version 3), copied verbatim in the file "COPYING".
4+
//
5+
// See http://alice-o2.web.cern.ch/license for full licensing information.
6+
//
7+
// In applying this license CERN does not waive the privileges and immunities
8+
// granted to it by virtue of its status as an Intergovernmental Organization
9+
// or submit itself to any jurisdiction.
10+
11+
/// \file GPUTPCCFCheckPadBaseline.h
12+
/// \author Felix Weiglhofer
13+
14+
#ifndef O2_GPU_GPU_TPC_CF_CHECK_PAD_BASELINE_H
15+
#define O2_GPU_GPU_TPC_CF_CHECK_PAD_BASELINE_H
16+
17+
#include "GPUGeneralKernels.h"
18+
#include "GPUConstantMem.h"
19+
20+
#include "clusterFinderDefs.h"
21+
22+
namespace GPUCA_NAMESPACE::gpu
23+
{
24+
25+
class GPUTPCCFCheckPadBaseline : public GPUKernelTemplate
26+
{
27+
28+
private:
29+
// Only use these constants on device side...
30+
// Use getPadsPerBlock() for host side
31+
enum {
32+
PadsPerBlockGPU = 4, // Number of pads in a single cache line
33+
PadsPerBlockCPU = 1,
34+
#ifdef GPUCA_GPUCODE
35+
PadsPerBlock = PadsPerBlockGPU,
36+
#else
37+
PadsPerBlock = PadsPerBlockCPU,
38+
#endif
39+
NumOfCachedTimebins = GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCCFCheckPadBaseline) / PadsPerBlock,
40+
};
41+
42+
public:
43+
struct GPUSharedMemory {
44+
tpccf::Charge charges[PadsPerBlock][NumOfCachedTimebins];
45+
};
46+
47+
#ifdef HAVE_O2HEADERS
48+
typedef GPUTPCClusterFinder processorType;
49+
GPUhdi() static processorType* Processor(GPUConstantMem& processors)
50+
{
51+
return processors.tpcClusterer;
52+
}
53+
#endif
54+
55+
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep()
56+
{
57+
return GPUDataTypes::RecoStep::TPCClusterFinding;
58+
}
59+
60+
// Use this to get num of pads per block on host side. Can't use constant there.
61+
static int getPadsPerBlock(bool isGPU)
62+
{
63+
return (isGPU) ? PadsPerBlockGPU : PadsPerBlockCPU;
64+
}
65+
66+
template <int iKernel = defaultKernel>
67+
GPUd() static void Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUSharedMemory& smem, processorType& clusterer);
68+
69+
private:
70+
GPUd() static ChargePos padToChargePos(int pad, const GPUTPCClusterFinder&);
71+
};
72+
73+
} // namespace GPUCA_NAMESPACE::gpu
74+
75+
#endif

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFPeakFinder.cxx‎

Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ GPUdii() void GPUTPCCFPeakFinder::Thread<0>(int nBlocks, int nThreads, int iBloc
2626
{
2727
Array2D<PackedCharge> chargeMap(reinterpret_cast<PackedCharge*>(clusterer.mPchargeMap));
2828
Array2D<uchar> isPeakMap(clusterer.mPpeakMap);
29-
findPeaksImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, chargeMap, clusterer.mPpositions, clusterer.mPmemory->counters.nPositions, clusterer.Param().rec, clusterer.mPisPeak, isPeakMap);
29+
findPeaksImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, chargeMap, clusterer.mPpadHasLostBaseline, clusterer.mPpositions, clusterer.mPmemory->counters.nPositions, clusterer.Param().rec, *clusterer.GetConstantMem()->calibObjects.tpcPadGain, clusterer.mPisPeak, isPeakMap);
3030
}
3131

3232
GPUdii() bool GPUTPCCFPeakFinder::isPeak(
@@ -91,9 +91,11 @@ GPUdii() bool GPUTPCCFPeakFinder::isPeak(
9191

9292
GPUd() void GPUTPCCFPeakFinder::findPeaksImpl(int nBlocks, int nThreads, int iBlock, int iThread, GPUSharedMemory& smem,
9393
const Array2D<PackedCharge>& chargeMap,
94+
const uchar* padHasLostBaseline,
9495
const ChargePos* positions,
9596
SizeT digitnum,
9697
const GPUSettingsRec& calib,
98+
const TPCPadGainCalib& gainCorrection, // Only used for globalPad() function
9799
uchar* isPeakPredicate,
98100
Array2D<uchar>& peakMap)
99101
{
@@ -105,8 +107,10 @@ GPUd() void GPUTPCCFPeakFinder::findPeaksImpl(int nBlocks, int nThreads, int iBl
105107
ChargePos pos = positions[CAMath::Min(idx, (SizeT)(digitnum - 1))];
106108
Charge charge = pos.valid() ? chargeMap[pos].unpack() : Charge(0);
107109

108-
uchar peak;
109-
peak = isPeak(smem, charge, pos, SCRATCH_PAD_SEARCH_N, chargeMap, calib, smem.posBcast, smem.buf);
110+
bool hasLostBaseline = padHasLostBaseline[gainCorrection.globalPad(pos.row(), pos.pad())];
111+
charge = (hasLostBaseline) ? 0.f : charge;
112+
113+
uchar peak = isPeak(smem, charge, pos, SCRATCH_PAD_SEARCH_N, chargeMap, calib, smem.posBcast, smem.buf);
110114

111115
// Exit early if dummy. See comment above.
112116
bool iamDummy = (idx >= digitnum);

0 commit comments

Comments
 (0)