Skip to content

Commit 9732e5c

Browse files
davidrohrdavidrohr
committed
GPU: Improve Clusterizer PEAK_OVERFLOW error reporting
1 parent 7cbbc7c commit 9732e5c

8 files changed

Lines changed: 17 additions & 17 deletions

File tree

‎GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -266,7 +266,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread<GPUTPCCompressionKernels::step1un
266266
for (unsigned int j = get_local_id(0); j < count; j += get_local_size(0)) {
267267
int outidx = idOffsetOut + totalCount + j;
268268
if (outidx >= idOffsetOutMax) {
269-
compressor.raiseError(GPUErrors::ERROR_COMPRESSION_ROW_HIT_OVERFLOW, outidx, idOffsetOutMax);
269+
compressor.raiseError(GPUErrors::ERROR_COMPRESSION_ROW_HIT_OVERFLOW, iSlice * 1000 + iRow, outidx, idOffsetOutMax);
270270
count = 0;
271271
break;
272272
}

‎GPU/GPUTracking/SliceTracker/GPUTPCGlobalTracking.cxx‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -75,7 +75,7 @@ GPUd() int GPUTPCGlobalTracking::PerformGlobalTrackingRun(GPUTPCTracker& tracker
7575
// GPUInfo("%d hits found", nHits);
7676
unsigned int hitId = CAMath::AtomicAdd(&tracker.CommonMemory()->nTrackHits, (unsigned int)nHits);
7777
if (hitId + nHits > tracker.NMaxTrackHits()) {
78-
tracker.raiseError(GPUErrors::ERROR_GLOBAL_TRACKING_TRACK_HIT_OVERFLOW, hitId + nHits, tracker.NMaxTrackHits());
78+
tracker.raiseError(GPUErrors::ERROR_GLOBAL_TRACKING_TRACK_HIT_OVERFLOW, tracker.ISlice(), hitId + nHits, tracker.NMaxTrackHits());
7979
CAMath::AtomicExch(&tracker.CommonMemory()->nTrackHits, tracker.NMaxTrackHits());
8080
return (0);
8181
}
@@ -124,7 +124,7 @@ GPUd() void GPUTPCGlobalTracking::PerformGlobalTracking(int nBlocks, int nThread
124124
const GPUTPCRow& GPUrestrict() row = tracker.Row(rowIndex);
125125
float Y = (float)tracker.Data().HitDataY(row, tracker.TrackHits()[tmpHit].HitIndex()) * row.HstepY() + row.Grid().YMin();
126126
if (sliceTarget.CommonMemory()->nTracks >= sliceTarget.NMaxTracks()) { // >= since will increase by 1
127-
sliceTarget.raiseError(GPUErrors::ERROR_GLOBAL_TRACKING_TRACK_OVERFLOW, sliceTarget.CommonMemory()->nTracks, sliceTarget.NMaxTracks());
127+
sliceTarget.raiseError(GPUErrors::ERROR_GLOBAL_TRACKING_TRACK_OVERFLOW, sliceTarget.ISlice(), sliceTarget.CommonMemory()->nTracks, sliceTarget.NMaxTracks());
128128
return;
129129
}
130130
if (!right && Y < -row.MaxY() * GPUCA_GLOBAL_TRACKING_Y_RANGE_LOWER) {
@@ -145,7 +145,7 @@ GPUd() void GPUTPCGlobalTracking::PerformGlobalTracking(int nBlocks, int nThread
145145
const GPUTPCRow& GPUrestrict() row = tracker.Row(rowIndex);
146146
float Y = (float)tracker.Data().HitDataY(row, tracker.TrackHits()[tmpHit].HitIndex()) * row.HstepY() + row.Grid().YMin();
147147
if (sliceTarget.CommonMemory()->nTracks >= sliceTarget.NMaxTracks()) { // >= since will increase by 1
148-
sliceTarget.raiseError(GPUErrors::ERROR_GLOBAL_TRACKING_TRACK_OVERFLOW, sliceTarget.CommonMemory()->nTracks, sliceTarget.NMaxTracks());
148+
sliceTarget.raiseError(GPUErrors::ERROR_GLOBAL_TRACKING_TRACK_OVERFLOW, sliceTarget.ISlice(), sliceTarget.CommonMemory()->nTracks, sliceTarget.NMaxTracks());
149149
return;
150150
}
151151
if (!right && Y < -row.MaxY() * GPUCA_GLOBAL_TRACKING_Y_RANGE_UPPER) {

‎GPU/GPUTracking/SliceTracker/GPUTPCSliceData.cxx‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -219,7 +219,7 @@ GPUdii() int GPUTPCSliceData::InitFromClusterData(int nBlocks, int nThreads, int
219219
CONSTEXPR unsigned int maxN = 1u << (sizeof(calink) < 3 ? (sizeof(calink) * 8) : 24);
220220
if (NumberOfClusters >= maxN) {
221221
if (iThread == 0) {
222-
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_HITINROW_OVERFLOW, rowIndex, NumberOfClusters, maxN);
222+
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_HITINROW_OVERFLOW, iSlice * 1000 + rowIndex, NumberOfClusters, maxN);
223223
}
224224
return 1;
225225
}
@@ -314,15 +314,15 @@ GPUdii() int GPUTPCSliceData::InitFromClusterData(int nBlocks, int nThreads, int
314314
CONSTEXPR int maxBins = sizeof(calink) < 4 ? (int)(1ul << (sizeof(calink) * 8)) : 0x7FFFFFFF; // NOLINT: false warning
315315
if (sizeof(calink) < 4 && numberOfBins >= maxBins) {
316316
if (iThread == 0) {
317-
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_BIN_OVERFLOW, rowIndex, numberOfBins, maxBins);
317+
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_BIN_OVERFLOW, iSlice * 1000 + rowIndex, numberOfBins, maxBins);
318318
}
319319
return 1;
320320
}
321321
const unsigned int nn = numberOfBins + grid.Ny() + 3;
322322
const unsigned int maxnn = GetGridSize(NumberOfClusters, 1);
323323
if (nn >= maxnn) {
324324
if (iThread == 0) {
325-
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_FIRSTHITINBIN_OVERFLOW, nn, maxnn);
325+
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_FIRSTHITINBIN_OVERFLOW, iSlice, nn, maxnn);
326326
}
327327
return 1;
328328
}
@@ -395,7 +395,7 @@ GPUdii() int GPUTPCSliceData::InitFromClusterData(int nBlocks, int nThreads, int
395395
if (iThread == 0) {
396396
const float maxAbsZ = CAMath::Max(CAMath::Abs(tmpMinMax[2]), CAMath::Abs(tmpMinMax[3]));
397397
if (maxAbsZ > 300 && !mem->param.par.continuousTracking) {
398-
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_Z_OVERFLOW, (unsigned int)maxAbsZ);
398+
mem->errorCodes.raiseError(GPUErrors::ERROR_SLICEDATA_Z_OVERFLOW, iSlice, (unsigned int)maxAbsZ);
399399
return 1;
400400
}
401401
}

‎GPU/GPUTracking/SliceTracker/GPUTPCStartHitsFinder.cxx‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -43,12 +43,12 @@ GPUdii() void GPUTPCStartHitsFinder::Thread<0>(int /*nBlocks*/, int nThreads, in
4343
GPUglobalref() GPUTPCHitId* const GPUrestrict() startHits = tracker.mTrackletTmpStartHits + s.mIRow * tracker.mNMaxRowStartHits;
4444
unsigned int nextRowStartHits = CAMath::AtomicAddShared(&s.mNRowStartHits, 1u);
4545
if (nextRowStartHits >= tracker.mNMaxRowStartHits) {
46-
tracker.raiseError(GPUErrors::ERROR_ROWSTARTHIT_OVERFLOW, nextRowStartHits, tracker.mNMaxRowStartHits);
46+
tracker.raiseError(GPUErrors::ERROR_ROWSTARTHIT_OVERFLOW, tracker.ISlice(), nextRowStartHits, tracker.mNMaxRowStartHits);
4747
#else
4848
GPUglobalref() GPUTPCHitId* const GPUrestrict() startHits = tracker.mTrackletStartHits;
4949
unsigned int nextRowStartHits = CAMath::AtomicAdd(&tracker.mCommonMem->nStartHits, 1u);
5050
if (nextRowStartHits >= tracker.mNMaxStartHits) {
51-
tracker.raiseError(GPUErrors::ERROR_STARTHIT_OVERFLOW, nextRowStartHits, tracker.mNMaxStartHits);
51+
tracker.raiseError(GPUErrors::ERROR_STARTHIT_OVERFLOW, tracker.ISlice(), nextRowStartHits, tracker.mNMaxStartHits);
5252
#endif
5353
CAMath::AtomicExch(&tracker.mCommonMem->nStartHits, 0u);
5454
break;
@@ -63,7 +63,7 @@ GPUdii() void GPUTPCStartHitsFinder::Thread<0>(int /*nBlocks*/, int nThreads, in
6363
unsigned int nOffset = CAMath::AtomicAdd(&tracker.mCommonMem->nStartHits, s.mNRowStartHits);
6464
tracker.mRowStartHitCountOffset[s.mIRow] = s.mNRowStartHits;
6565
if (nOffset + s.mNRowStartHits > tracker.mNMaxStartHits) {
66-
tracker.raiseError(GPUErrors::ERROR_STARTHIT_OVERFLOW, nOffset + s.mNRowStartHits, tracker.mNMaxStartHits);
66+
tracker.raiseError(GPUErrors::ERROR_STARTHIT_OVERFLOW, tracker.ISlice(), nOffset + s.mNRowStartHits, tracker.mNMaxStartHits);
6767
CAMath::AtomicExch(&tracker.mCommonMem->nStartHits, 0u);
6868
}
6969
}

‎GPU/GPUTracking/SliceTracker/GPUTPCTrackletConstructor.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -71,9 +71,9 @@ GPUd() void GPUTPCTrackletConstructor::StoreTracklet(int /*nBlocks*/, int /*nThr
7171
unsigned int hitout = CAMath::AtomicAdd(tracker.NRowHits(), nHits);
7272
if (itrout >= tracker.NMaxTracklets() || hitout + nHits > tracker.NMaxRowHits()) {
7373
if (itrout >= tracker.NMaxTracklets()) {
74-
tracker.raiseError(GPUErrors::ERROR_TRACKLET_OVERFLOW, itrout, tracker.NMaxTracklets());
74+
tracker.raiseError(GPUErrors::ERROR_TRACKLET_OVERFLOW, tracker.ISlice(), itrout, tracker.NMaxTracklets());
7575
} else {
76-
tracker.raiseError(GPUErrors::ERROR_TRACKLET_HIT_OVERFLOW, hitout + nHits, tracker.NMaxRowHits());
76+
tracker.raiseError(GPUErrors::ERROR_TRACKLET_HIT_OVERFLOW, tracker.ISlice(), hitout + nHits, tracker.NMaxRowHits());
7777
}
7878
CAMath::AtomicExch(tracker.NTracklets(), 0u);
7979
if (hitout + nHits > tracker.NMaxRowHits()) {

‎GPU/GPUTracking/SliceTracker/GPUTPCTrackletSelector.cxx‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -91,9 +91,9 @@ GPUdii() void GPUTPCTrackletSelector::Thread<0>(int nBlocks, int nThreads, int i
9191
unsigned int nFirstTrackHit = CAMath::AtomicAdd(tracker.NTrackHits(), (unsigned int)nHits);
9292
if (itrout >= tracker.NMaxTracks() || nFirstTrackHit + nHits > tracker.NMaxTrackHits()) {
9393
if (itrout >= tracker.NMaxTracks()) {
94-
tracker.raiseError(GPUErrors::ERROR_TRACK_OVERFLOW, itrout, tracker.NMaxTracks());
94+
tracker.raiseError(GPUErrors::ERROR_TRACK_OVERFLOW, tracker.ISlice(), itrout, tracker.NMaxTracks());
9595
} else {
96-
tracker.raiseError(GPUErrors::ERROR_TRACK_HIT_OVERFLOW, nFirstTrackHit + nHits, tracker.NMaxTrackHits());
96+
tracker.raiseError(GPUErrors::ERROR_TRACK_HIT_OVERFLOW, tracker.ISlice(), nFirstTrackHit + nHits, tracker.NMaxTrackHits());
9797
}
9898
CAMath::AtomicExch(tracker.NTracks(), 0u);
9999
if (nFirstTrackHit + nHits > tracker.NMaxTrackHits()) {

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFClusterizer.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -260,7 +260,7 @@ GPUd() uint GPUTPCCFClusterizer::sortIntoBuckets(processorType& clusterer, const
260260
if (index < maxElemsPerBucket) {
261261
buckets[maxElemsPerBucket * row + index] = cluster;
262262
} else {
263-
clusterer.raiseError(GPUErrors::ERROR_CF_ROW_CLUSTER_OVERFLOW, index, maxElemsPerBucket);
263+
clusterer.raiseError(GPUErrors::ERROR_CF_ROW_CLUSTER_OVERFLOW, clusterer.mISlice * 1000 + row, index, maxElemsPerBucket);
264264
CAMath::AtomicExch(&elemsInBucket[row], maxElemsPerBucket);
265265
}
266266
return index;

‎GPU/GPUTracking/TPCClusterFinder/GPUTPCCFStreamCompaction.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -127,7 +127,7 @@ GPUdii() void GPUTPCCFStreamCompaction::Thread<GPUTPCCFStreamCompaction::compact
127127
if ((unsigned int)get_global_id(0) == lastId) {
128128
SizeT nFinal = clusterer.mPbuf[lastId];
129129
if (nFinal > bufferSize) {
130-
clusterer.raiseError(stage ? GPUErrors::ERROR_CF_CLUSTER_OVERFLOW : GPUErrors::ERROR_CF_PEAK_OVERFLOW, nFinal, bufferSize);
130+
clusterer.raiseError(stage ? GPUErrors::ERROR_CF_CLUSTER_OVERFLOW : GPUErrors::ERROR_CF_PEAK_OVERFLOW, clusterer.mISlice, nFinal, bufferSize);
131131
nFinal = bufferSize;
132132
}
133133
if (stage) {

0 commit comments

Comments
 (0)