Skip to content

Commit c383536

Browse files
davidrohrdavidrohr
committed
TPC: Add CTF skimming of attached clusters (full drift time for the moment)
1 parent 3c86833 commit c383536

4 files changed

Lines changed: 53 additions & 29 deletions

File tree

‎Detectors/TPC/reconstruction/include/TPCReconstruction/CTFCoder.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -273,8 +273,8 @@ o2::ctf::CTFIOSize CTFCoder::encode(VEC& buff, const CompressedClusters& ccl, co
273273
encodeTPC(ccl.qMaxU, ccl.qMaxU + (mCombineColumns ? 0 : ccl.nUnattachedClusters), CTF::BLCqMaxU, 0, rejectHits);
274274

275275
encodeTPC(ccl.flagsU, ccl.flagsU + ccl.nUnattachedClusters, CTF::BLCflagsU, 0, rejectHits);
276-
encodeTPC(ccl.padDiffU, ccl.padDiffU + ccl.nUnattachedClusters, CTF::BLCpadDiffU, 0, rejectHits);
277-
encodeTPC(ccl.timeDiffU, ccl.timeDiffU + ccl.nUnattachedClusters, CTF::BLCtimeDiffU, 0, rejectHits);
276+
encodeTPC(cclFiltered.padDiffU, cclFiltered.padDiffU + cclFiltered.nUnattachedClusters, CTF::BLCpadDiffU, 0);
277+
encodeTPC(cclFiltered.timeDiffU, cclFiltered.timeDiffU + cclFiltered.nUnattachedClusters, CTF::BLCtimeDiffU, 0);
278278

279279
if (mCombineColumns) {
280280
const auto [begin, end] = makeInputIterators(ccl.sigmaPadU, ccl.sigmaTimeU, ccl.nUnattachedClusters,

‎Detectors/TPC/workflow/src/EntropyEncoderSpec.cxx‎

Lines changed: 44 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@
2424
#include "GPUParam.h"
2525
#include "DataFormatsTPC/ClusterNative.h"
2626
#include "TPCClusterDecompressor.inc"
27+
#include "GPUTPCCompressionKernels.inc"
2728
#include "TPCCalibration/VDriftHelper.h"
2829

2930
using namespace o2::framework;
@@ -103,6 +104,7 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
103104
auto& buffer = pc.outputs().make<std::vector<o2::ctf::BufferType>>(Output{"TPC", "CTFDATA", 0, Lifetime::Timeframe});
104105
std::vector<bool> rejectHits, rejectTracks, rejectTrackHits, rejectTrackHitsReduced;
105106
CompressedClusters clustersFiltered = clusters;
107+
std::vector<std::pair<std::vector<unsigned int>, std::vector<unsigned short>>> tmpBuffer(std::max<int>(mNThreads, 1));
106108
if (mSelIR) {
107109
mCTFCoder.setSelectedIRFrames(pc.inputs().get<gsl::span<o2::dataformats::IRFrame>>("selIRFrames"));
108110
rejectHits.resize(clusters.nUnattachedClusters);
@@ -169,29 +171,52 @@ void EntropyEncoderSpec::run(ProcessingContext& pc)
169171
offset += (i * GPUCA_ROW_COUNT + j >= clusters.nSliceRows) ? 0 : clusters.nSliceRowClusters[i * GPUCA_ROW_COUNT + j];
170172
}
171173
}
174+
172175
#ifdef WITH_OPENMP
173-
#pragma omp parallel for num_threads(mNThreads)
176+
#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
174177
#endif
175-
for (unsigned int i = 0; i < GPUCA_NSLICES; i++) {
176-
for (unsigned int j = 0; j < GPUCA_ROW_COUNT; j++) {
177-
if (i * GPUCA_ROW_COUNT + j >= clusters.nSliceRows) {
178-
break;
179-
}
180-
auto checker = [i, j, &rejectHits, &clustersFiltered](const o2::tpc::ClusterNative& cl, unsigned int k) {
181-
if (false) {
182-
rejectHits[k] = true;
183-
static std::atomic_flag lock = ATOMIC_FLAG_INIT;
184-
while (lock.test_and_set(std::memory_order_acquire)) {
185-
}
186-
clustersFiltered.nSliceRowClusters[i * GPUCA_ROW_COUNT + j]--;
187-
clustersFiltered.nUnattachedClusters--;
188-
lock.clear(std::memory_order_release);
178+
for (unsigned int ii = 0; ii < clusters.nSliceRows; ii++) {
179+
unsigned int i = ii / GPUCA_ROW_COUNT;
180+
unsigned int j = ii % GPUCA_ROW_COUNT;
181+
o2::tpc::ClusterNative preCl;
182+
#ifdef WITH_OPENMP
183+
int myThread = omp_get_thread_num();
184+
#else
185+
int myThread = 0;
186+
#endif
187+
unsigned int count = 0;
188+
auto checker = [i, j, firstIR, totalT, this, &preCl, &count, &outBuffer = tmpBuffer[myThread], &rejectHits, &clustersFiltered](const o2::tpc::ClusterNative& cl, unsigned int k) {
189+
const auto chkVal = firstIR + (cl.getTime() * constants::LHCBCPERTIMEBIN);
190+
const auto chkExt = totalT * constants::LHCBCPERTIMEBIN;
191+
const bool reject = mCTFCoder.getIRFramesSelector().check(o2::dataformats::IRFrame(chkVal, chkVal + 1), chkExt, 0) < 0;
192+
if (reject) {
193+
rejectHits[k] = true;
194+
clustersFiltered.nSliceRowClusters[i * GPUCA_ROW_COUNT + j]--;
195+
static std::atomic_flag lock = ATOMIC_FLAG_INIT;
196+
while (lock.test_and_set(std::memory_order_acquire)) {
189197
}
190-
};
191-
unsigned int end = offsets[i][j] + clusters.nSliceRowClusters[i * GPUCA_ROW_COUNT + j];
192-
o2::gpu::TPCClusterDecompressor::decompressHits(&clusters, offsets[i][j], end, checker);
193-
}
198+
clustersFiltered.nUnattachedClusters--;
199+
lock.clear(std::memory_order_release);
200+
} else {
201+
outBuffer.first.emplace_back(0);
202+
outBuffer.second.emplace_back(0);
203+
GPUTPCCompression_EncodeUnattached(clustersFiltered.nComppressionModes, cl, outBuffer.first.back(), outBuffer.second.back(), count++ ? &preCl : nullptr);
204+
preCl = cl;
205+
}
206+
};
207+
unsigned int end = offsets[i][j] + clusters.nSliceRowClusters[i * GPUCA_ROW_COUNT + j];
208+
o2::gpu::TPCClusterDecompressor::decompressHits(&clusters, offsets[i][j], end, checker);
209+
}
210+
tmpBuffer[0].first.reserve(clustersFiltered.nUnattachedClusters);
211+
tmpBuffer[0].second.reserve(clustersFiltered.nUnattachedClusters);
212+
for (int i = 1; i < mNThreads; i++) {
213+
tmpBuffer[0].first.insert(tmpBuffer[0].first.end(), tmpBuffer[i].first.begin(), tmpBuffer[i].first.end());
214+
tmpBuffer[i].first.clear();
215+
tmpBuffer[0].second.insert(tmpBuffer[0].second.end(), tmpBuffer[i].second.begin(), tmpBuffer[i].second.end());
216+
tmpBuffer[i].second.clear();
194217
}
218+
clustersFiltered.timeDiffU = tmpBuffer[0].first.data();
219+
clustersFiltered.padDiffU = tmpBuffer[0].second.data();
195220
}
196221
auto iosize = mCTFCoder.encode(buffer, clusters, clustersFiltered, mSelIR ? &rejectHits : nullptr, mSelIR ? &rejectTracks : nullptr, mSelIR ? &rejectTrackHits : nullptr, mSelIR ? &rejectTrackHitsReduced : nullptr);
197222
pc.outputs().snapshot({"ctfrep", 0}, iosize);

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -275,7 +275,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
275275
const ClusterNative& GPUrestrict() orgCl = clusters->clusters[iSlice][iRow][sortBuffer[j]];
276276

277277
int preId = j != 0 ? (int)sortBuffer[j - 1] : (totalCount != 0 ? (int)smem.lastIndex : -1);
278-
GPUTPCCompression_EncodeUnattached(param, orgCl, c, outidx, preId == -1 ? nullptr : &clusters->clusters[iSlice][iRow][preId]);
278+
GPUTPCCompression_EncodeUnattached(param.rec.tpc.compressionTypeMask, orgCl, c.timeDiffU[outidx], c.padDiffU[outidx], preId == -1 ? nullptr : &clusters->clusters[iSlice][iRow][preId]);
279279

280280
unsigned short qtot = orgCl.qTot, qmax = orgCl.qMax;
281281
unsigned char sigmapad = orgCl.sigmaPadPacked, sigmatime = orgCl.sigmaTimePacked;

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.inc‎

Lines changed: 6 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -20,21 +20,20 @@
2020
namespace o2::gpu
2121
{
2222

23-
template <class T>
24-
GPUdii() void GPUTPCCompression_EncodeUnattached(const GPUParam& param, const o2::tpc::ClusterNative& orgCl, T& c, int outidx, const o2::tpc::ClusterNative* orgClPre = nullptr)
23+
GPUdii() void GPUTPCCompression_EncodeUnattached(unsigned char nComppressionModes, const o2::tpc::ClusterNative& orgCl, unsigned int& outTime, unsigned short& outPad, const o2::tpc::ClusterNative* orgClPre = nullptr)
2524
{
26-
if (param.rec.tpc.compressionTypeMask & GPUSettings::CompressionDifferences) {
25+
if (nComppressionModes & GPUSettings::CompressionDifferences) {
2726
unsigned int lastTime = 0, lastPad = 0;
2827
if (orgClPre) {
2928
lastPad = orgClPre->padPacked;
3029
lastTime = orgClPre->getTimePacked();
3130
}
3231

33-
c.padDiffU[outidx] = orgCl.padPacked - lastPad;
34-
c.timeDiffU[outidx] = (orgCl.getTimePacked() - lastTime) & 0xFFFFFF;
32+
outPad = orgCl.padPacked - lastPad;
33+
outTime = (orgCl.getTimePacked() - lastTime) & 0xFFFFFF;
3534
} else {
36-
c.padDiffU[outidx] = orgCl.padPacked;
37-
c.timeDiffU[outidx] = orgCl.getTimePacked();
35+
outPad = orgCl.padPacked;
36+
outTime = orgCl.getTimePacked();
3837
}
3938
}
4039

0 commit comments

Comments
 (0)