TPCClusterFinder: Improve performance for mc labels.

fweig · davidrohr · commit dc09cd0aee26 · 2020-08-25T17:00:30.000+02:00
diff --git a/GPU/GPUTracking/Global/GPUChainTracking.cxx b/GPU/GPUTracking/Global/GPUChainTracking.cxx
@@ -1111,7 +1111,7 @@ int GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)
     std::vector<bool> laneHasData(GetProcessingSettings().nTPCClustererLanes, false);
     for (CfFragment fragment = mCFContext->fragmentFirst; !fragment.isEnd(); fragment = fragment.next()) {
       if (GetProcessingSettings().debugLevel >= 3) {
-        GPUInfo("Processing time bins [%d, %d)", fragment.start, fragment.last());
+        GPUInfo("Processing time bins [%d, %d) for sectors %d to %d", fragment.start, fragment.last(), iSliceBase, iSliceBase + GetProcessingSettings().nTPCClustererLanes - 1);
       }
       for (int lane = 0; lane < GetProcessingSettings().nTPCClustererLanes && iSliceBase + lane < NSLICES; lane++) {
         if (fragment.index != 0) {
diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.cxx
@@ -67,21 +67,37 @@ GPUd() void GPUTPCCFChargeMapFiller::fillFromDigitsImpl(int nBlocks, int nThread
 template <>
 GPUdii() void GPUTPCCFChargeMapFiller::Thread<GPUTPCCFChargeMapFiller::findFragmentStart>(int nBlocks, int nThreads, int iBlock, int iThread, GPUSharedMemory& smem, processorType& clusterer)
 {
-  // TODO: use binary search
-  if (get_global_id(0) != 0) {
+  if (iThread != 0) {
     return;
   }
 
   size_t nDigits = clusterer.mPmemory->counters.nDigits;
   const tpc::Digit* digits = clusterer.mPdigits;
-  size_t st = 0;
-  for (; st < nDigits && digits[st].getTimeStamp() < clusterer.mPmemory->fragment.first(); st++) {
-  }
-
-  size_t end = st;
-  for (; end < nDigits && digits[end].getTimeStamp() < clusterer.mPmemory->fragment.last(); end++) {
-  }
+  size_t st = findTransition(clusterer.mPmemory->fragment.first(), digits, nDigits, 0);
+  size_t end = findTransition(clusterer.mPmemory->fragment.last(), digits, nDigits, st);
 
   clusterer.mPmemory->fragment.digitsStart = st;
   clusterer.mPmemory->counters.nPositions = end - st;
 }
+
+GPUd() size_t GPUTPCCFChargeMapFiller::findTransition(int time, const tpc::Digit* digits, size_t nDigits, size_t lower)
+{
+  size_t upper = nDigits - 1;
+
+  while (lower <= upper) {
+    size_t middle = (lower + upper) / 2;
+
+    if (middle == 0) {
+      return 0;
+    }
+
+    if (digits[middle].getTimeStamp() < time) {
+      lower = middle + 1;
+    } else if (digits[middle - 1].getTimeStamp() < time) {
+      return middle;
+    } else {
+      upper = middle - 1;
+    }
+  }
+  return lower;
+}
diff --git a/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.h b/GPU/GPUTracking/TPCClusterFinder/GPUTPCCFChargeMapFiller.h
@@ -64,6 +64,9 @@ class GPUTPCCFChargeMapFiller : public GPUKernelTemplate
   static GPUd() void fillIndexMapImpl(int, int, int, int, const CfFragment&, const tpc::Digit*, Array2D<uint>&, size_t);
 
   static GPUd() void fillFromDigitsImpl(int, int, int, int, const CfFragment&, size_t, const tpc::Digit*, ChargePos*, Array2D<PackedCharge>&);
+
+ private:
+  static GPUd() size_t findTransition(int, const tpc::Digit*, size_t, size_t);
 };
 
 } // namespace gpu

Original file line number	Diff line number	Diff line change
`@@ -1111,7 +1111,7 @@ int GPUChainTracking::RunTPCClusterizer(bool synchronizeOutput)`
`1111`	`1111`	`std::vector<bool> laneHasData(GetProcessingSettings().nTPCClustererLanes, false);`
`1112`	`1112`	`for (CfFragment fragment = mCFContext->fragmentFirst; !fragment.isEnd(); fragment = fragment.next()) {`
`1113`	`1113`	`if (GetProcessingSettings().debugLevel >= 3) {`
`1114`		`- GPUInfo("Processing time bins [%d, %d)", fragment.start, fragment.last());`
	`1114`	`+ GPUInfo("Processing time bins [%d, %d) for sectors %d to %d", fragment.start, fragment.last(), iSliceBase, iSliceBase + GetProcessingSettings().nTPCClustererLanes - 1);`
`1115`	`1115`	`}`
`1116`	`1116`	`for (int lane = 0; lane < GetProcessingSettings().nTPCClustererLanes && iSliceBase + lane < NSLICES; lane++) {`
`1117`	`1117`	`if (fragment.index != 0) {`