Skip to content

Commit a6a09cb

Browse files
fweigdavidrohr
fweig
authored and
davidrohr
committed
TPCClusterFinder: Fix illegal write in noise suppression. Sort cluster before dumping them to file.
1 parent cf5c253 commit a6a09cb

5 files changed

Lines changed: 43 additions & 13 deletions

File tree

‎GPU/GPUTracking/TPCClusterFinder/ClusterAccumulator.cxx‎

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -24,11 +24,14 @@ using namespace GPUCA_NAMESPACE::gpu;
2424

2525
GPUd() void ClusterAccumulator::toNative(const deprecated::Digit& d, deprecated::ClusterNative& cn) const
2626
{
27-
uchar isEdgeCluster = CfUtils::isAtEdge(&d);
28-
uchar wasSplitInTime = mSplitInTime >= MIN_SPLIT_NUM;
29-
uchar wasSplitInPad = mSplitInPad >= MIN_SPLIT_NUM;
30-
uchar flags =
31-
(isEdgeCluster << deprecated::CN_FLAG_POS_IS_EDGE_CLUSTER) | (wasSplitInTime << deprecated::CN_FLAG_POS_SPLIT_IN_TIME) | (wasSplitInPad << deprecated::CN_FLAG_POS_SPLIT_IN_PAD);
27+
bool isEdgeCluster = CfUtils::isAtEdge(&d);
28+
bool wasSplitInTime = mSplitInTime >= MIN_SPLIT_NUM;
29+
bool wasSplitInPad = mSplitInPad >= MIN_SPLIT_NUM;
30+
31+
uchar flags = 0;
32+
flags |= (isEdgeCluster) ? deprecated::CN_FLAG_IS_EDGE_CLUSTER : 0;
33+
flags |= (wasSplitInTime) ? deprecated::CN_FLAG_SPLIT_IN_TIME : 0;
34+
flags |= (wasSplitInPad) ? deprecated::CN_FLAG_SPLIT_IN_PAD : 0;
3235

3336
cn.qmax = d.charge;
3437
cn.qtot = mQtot;

‎GPU/GPUTracking/TPCClusterFinder/Deconvolution.cxx‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,8 @@ GPUd() void Deconvolution::countPeaksImpl(int nBlocks, int nThreads, int iBlock,
3838

3939
char peakCount = (iamPeak) ? 1 : 0;
4040

41-
#if defined(BUILD_CLUSTER_SCRATCH_PAD) && defined(GPUCA_GPUCODE)
41+
#if defined(BUILD_CLUSTER_SCRATCH_PAD)
42+
/* #if defined(BUILD_CLUSTER_SCRATCH_PAD) && defined(GPUCA_GPUCODE) */
4243
/* #if 0 */
4344
ushort ll = get_local_id(0);
4445
ushort partId = ll;

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinderDump.cxx‎

Lines changed: 25 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -74,7 +74,7 @@ void GPUTPCClusterFinder::DumpSuppressedPeaks(std::ostream& out)
7474

7575
void GPUTPCClusterFinder::DumpSuppressedPeaksCompacted(std::ostream& out)
7676
{
77-
out << "Clusterer - Noise Suppression Peaks Compacted - Slice " << mISlice << ": " << mPmemory->nClusters <<"\n";
77+
out << "Clusterer - Noise Suppression Peaks Compacted - Slice " << mISlice << ": " << mPmemory->nClusters << "\n";
7878
for (size_t i = 0; i < mPmemory->nClusters; i++) {
7979
out << i << ": " << mPfilteredPeaks[i].charge << ", " << mPfilteredPeaks[i].time << ", " << (int)mPfilteredPeaks[i].pad << ", " << (int)mPfilteredPeaks[i].row << "\n";
8080
}
@@ -91,11 +91,31 @@ void GPUTPCClusterFinder::DumpCountedPeaks(std::ostream& out)
9191
void GPUTPCClusterFinder::DumpClusters(std::ostream& out)
9292
{
9393
out << "Clusterer - Clusters - Slice " << mISlice << "\n";
94+
9495
for (int i = 0; i < GPUCA_ROW_COUNT; i++) {
95-
out << "Row: " << i << ": " << mPclusterInRow[i] << "\n";
96-
for (unsigned int j = 0; j < mPclusterInRow[i]; j++) {
97-
const auto& cl = mPclusterByRow[i * mNMaxClusterPerRow + j];
98-
out << cl.timeFlagsPacked << ", " << cl.padPacked << ", " << (int)cl.sigmaTimePacked << ", " << (int)cl.sigmaPadPacked << ", " << cl.qmax << ", " << cl.qtot << "\n";
96+
size_t N = mPclusterInRow[i];
97+
out << "Row: " << i << ": " << N << "\n";
98+
std::vector<deprecated::ClusterNative> sortedCluster(N);
99+
deprecated::ClusterNative* row = &mPclusterByRow[i * mNMaxClusterPerRow];
100+
std::copy(row, &row[N], sortedCluster.begin());
101+
102+
std::sort(sortedCluster.begin(), sortedCluster.end(), [](const auto& c1, const auto& c2) {
103+
float t1 = deprecated::cnGetTime(&c1);
104+
float t2 = deprecated::cnGetTime(&c2);
105+
float p1 = deprecated::cnGetPad(&c1);
106+
float p2 = deprecated::cnGetPad(&c2);
107+
108+
if (t1 < t2) {
109+
return true;
110+
} else if (t1 == t2) {
111+
return p1 < p2;
112+
} else {
113+
return false;
114+
}
115+
});
116+
117+
for (const auto& cl : sortedCluster) {
118+
out << std::hex << cl.timeFlagsPacked << std::dec << ", " << cl.padPacked << ", " << (int)cl.sigmaTimePacked << ", " << (int)cl.sigmaPadPacked << ", " << cl.qmax << ", " << cl.qtot << "\n";
99119
}
100120
}
101121
}

‎GPU/GPUTracking/TPCClusterFinder/NoiseSuppression.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -248,9 +248,9 @@ GPUd() void NoiseSuppression::findMinimaAndPeaksScratchpad(
248248
SCRATCH_PAD_WORK_GROUP_SIZE,
249249
SCRATCH_PAD_WORK_GROUP_SIZE,
250250
ll,
251-
0,
251+
16,
252252
2,
253-
CfConsts::NoiseSuppressionNeighbors + 16,
253+
CfConsts::NoiseSuppressionNeighbors,
254254
posBcast,
255255
buf);
256256

‎GPU/GPUTracking/TPCClusterFinder/clusterFinderDefs.h‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -58,8 +58,14 @@ typedef unsigned long ulong;
5858

5959
#define SCRATCH_PAD_SEARCH_N 8
6060
#define SCRATCH_PAD_COUNT_N 16
61+
#if defined(GPUCA_GPUCODE)
6162
#define SCRATCH_PAD_BUILD_N 8
6263
#define SCRATCH_PAD_NOISE_N 8
64+
#else
65+
// Double shared memory on cpu as we can't reuse the memory from other threads
66+
#define SCRATCH_PAD_BUILD_N 16
67+
#define SCRATCH_PAD_NOISE_N 16
68+
#endif
6369

6470
#define PADDING_PAD 2
6571
#define PADDING_TIME 3

0 commit comments

Comments
 (0)