TPCClusterFinder: Fix buffer overflow in compaction step.

Author: fweig

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFStreamCompaction.cxx

@@ -25,7 +25,8 @@ template <>
 GPUdii() void GPUTPCCFStreamCompaction::Thread<GPUTPCCFStreamCompaction::nativeScanUpStart>(int nBlocks, int nThreads, int iBlock, int iThread, GPUSharedMemory& smem, processorType& clusterer, int iBuf, int stage)
 {
   int nElems = compactionElems(clusterer, stage);
-  nativeScanUpStartImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, clusterer.mPisPeak, clusterer.mPbuf + (iBuf - 1) * clusterer.mBufSize, clusterer.mPbuf + iBuf * clusterer.mBufSize, nElems);
+  size_t bufferSize = (stage) ? clusterer.mNMaxClusters : clusterer.mNMaxPeaks;
+  nativeScanUpStartImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, clusterer.mPisPeak, clusterer.mPbuf + (iBuf - 1) * clusterer.mBufSize, clusterer.mPbuf + iBuf * clusterer.mBufSize, nElems, bufferSize);
 }
 
 GPUd() void GPUTPCCFStreamCompaction::nativeScanUpStartImpl(int nBlocks, int nThreads, int iBlock, int iThread, GPUSharedMemory& smem,
@@ -122,6 +123,7 @@ template <>
 GPUdii() void GPUTPCCFStreamCompaction::Thread<GPUTPCCFStreamCompaction::compactDigit>(int nBlocks, int nThreads, int iBlock, int iThread, GPUSharedMemory& smem, processorType& clusterer, int iBuf, int stage, deprecated::PackedDigit* in, deprecated::PackedDigit* out)
 {
   unsigned int nElems = compactionElems(clusterer, stage);
+
   compactDigitImpl(get_num_groups(0), get_local_size(0), get_group_id(0), get_local_id(0), smem, in, out, clusterer.mPisPeak, clusterer.mPbuf + (iBuf - 1) * clusterer.mBufSize, clusterer.mPbuf + iBuf * clusterer.mBufSize, nElems);
   unsigned int lastId = get_global_size(0) - 1;
   if ((unsigned int)get_global_id(0) == lastId) {
@@ -151,17 +153,18 @@ GPUd() void GPUTPCCFStreamCompaction::compactDigitImpl(int nBlocks, int nThreads
   int pred = (iAmDummy) ? 0 : predicate[idx];
   int scanRes = work_group_scan_inclusive_add(pred);
 
-  int compIdx = scanRes;
+  size_t compIdx = scanRes;
   if (gid) {
     compIdx += incr[gid - 1];
   }
 
-  if (pred) {
-    out[compIdx - 1] = in[idx];
+  size_t tgtIdx = compIdx - 1;
+  if (pred && tgtIdx < bufferSize) {
+    out[tgtIdx] = in[idx];
   }
 
   if (idx == lastItem) {
-    newIdx[idx] = compIdx; // TODO: Eventually, we can just return the last value, no need to store to memory
+    newIdx[idx] = CAMath::Min(compIdx, bufferSize); // TODO: Eventually, we can just return the last value, no need to store to memory
   }
 }
 

GPU/GPUTracking/TPCClusterFinder/GPUTPCCFStreamCompaction.h

@@ -55,7 +55,7 @@ class GPUTPCCFStreamCompaction
   static GPUd() void compactDigitImpl(int, int, int, int, GPUSharedMemory&,
                                       const deprecated::Digit*, deprecated::Digit*,
                                       const uchar*, int*, const int*,
-                                      int);
+                                      int, size_t)
 
 #ifdef HAVE_O2HEADERS
   typedef GPUTPCClusterFinder processorType;

GPU/GPUTracking/TPCClusterFinder/GPUTPCClusterFinder.cxx

@@ -68,7 +68,7 @@ void GPUTPCClusterFinder::RegisterMemoryAllocation()
 void GPUTPCClusterFinder::SetMaxData(const GPUTrackingInOutPointers& io)
 {
   mNMaxPeaks = 0.5f * mNMaxDigits;
-  mNMaxClusters = 0.2f * mNMaxPeaks;
+  mNMaxClusters = mNMaxPeaks;
   mNMaxClusterPerRow = 0.01f * mNMaxDigits;
   mBufSize = nextMultipleOf<std::max<int>(GPUCA_MEMALIGN, mScanWorkGroupSize)>(mNMaxDigits);
   mNBufs = getNSteps(mBufSize);