Fix indices for used clusters

Author: mconcas

Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h

@@ -63,7 +63,7 @@ class TimeFrameGPU : public TimeFrame
   void loadVertices(const int);
 
   ///
-  void createTrackletsLUTDevice();
+  void createTrackletsLUTDevice(const int);
   void loadTrackletsDevice();
   void loadTrackletsLUTDevice();
   void loadCellsDevice();

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu

@@ -176,7 +176,7 @@ void TimeFrameGPU<nLayers>::createUsedClustersDevice(const int iteration)
 template <int nLayers>
 void TimeFrameGPU<nLayers>::loadUsedClustersDevice()
 {
-  START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "creating used clusters flags");
+  START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "loading used clusters flags");
   for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {
     LOGP(debug, "gpu-transfer: loading {} used clusters flags on layer {}, for {} MB.", mUsedClusters[iLayer].size(), iLayer, mClusters[iLayer].size() * sizeof(unsigned char) / MB);
     checkGPUError(cudaMemcpyAsync(mUsedClustersDevice[iLayer], mUsedClusters[iLayer].data(), mUsedClusters[iLayer].size() * sizeof(unsigned char), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
@@ -246,20 +246,25 @@ void TimeFrameGPU<nLayers>::loadVertices(const int iteration)
 }
 
 template <int nLayers>
-void TimeFrameGPU<nLayers>::createTrackletsLUTDevice()
+void TimeFrameGPU<nLayers>::createTrackletsLUTDevice(const int iteration)
 {
-  START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "creating cells LUTs");
+  START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "creating tracklets LUTs");
   for (auto iLayer{0}; iLayer < nLayers - 1; ++iLayer) {
-    LOGP(debug, "gpu-transfer: creating tracklets LUT for {} elements on layer {}, for {} MB.", mClusters[iLayer].size() + 1, iLayer, (mClusters[iLayer].size() + 1) * sizeof(int) / MB);
-    allocMemAsync(reinterpret_cast<void**>(&mTrackletsLUTDevice[iLayer]), (mClusters[iLayer].size() + 1) * sizeof(int), nullptr, getExtAllocator());
+    if (!iteration) {
+      LOGP(debug, "gpu-transfer: creating tracklets LUT for {} elements on layer {}, for {} MB.", mClusters[iLayer].size() + 1, iLayer, (mClusters[iLayer].size() + 1) * sizeof(int) / MB);
+      allocMemAsync(reinterpret_cast<void**>(&mTrackletsLUTDevice[iLayer]), (mClusters[iLayer].size() + 1) * sizeof(int), nullptr, getExtAllocator());
+    }
     checkGPUError(cudaMemsetAsync(mTrackletsLUTDevice[iLayer], 0, (mClusters[iLayer].size() + 1) * sizeof(int), mGpuStreams[0].get()));
   }
-  allocMemAsync(reinterpret_cast<void**>(&mTrackletsLUTDeviceArray), (nLayers - 2) * sizeof(int*), nullptr, getExtAllocator());
-  checkGPUError(cudaMemcpyAsync(mTrackletsLUTDeviceArray, mTrackletsLUTDevice.data(), mTrackletsLUTDevice.size() * sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
+  if (!iteration) {
+    allocMemAsync(reinterpret_cast<void**>(&mTrackletsLUTDeviceArray), (nLayers - 1) * sizeof(int*), nullptr, getExtAllocator());
+    checkGPUError(cudaMemcpyAsync(mTrackletsLUTDeviceArray, mTrackletsLUTDevice.data(), mTrackletsLUTDevice.size() * sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
+  }
   STOP_GPU_STREAM_TIMER(mGpuStreams[0].get());
 }
 
-template<int nLayers> void TimeFrameGPU<nLayers>::createTrackletsBuffers()
+template <int nLayers>
+void TimeFrameGPU<nLayers>::createTrackletsBuffers()
 {
   START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "creating cells buffers");
   for (auto iLayer{0}; iLayer < nLayers - 1; ++iLayer) {

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx

@@ -92,7 +92,7 @@ void TrackerTraitsGPU<nLayers>::computeTrackletsHybrid(const int iteration, int
 {
   auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
   TrackerTraits::computeLayerTracklets(iteration, iROFslice, iVertex);
-  mTimeFrameGPU->createTrackletsLUTDevice();
+  mTimeFrameGPU->createTrackletsLUTDevice(iteration);
 
   const Vertex diamondVert({mTrkParams[iteration].Diamond[0], mTrkParams[iteration].Diamond[1], mTrkParams[iteration].Diamond[2]}, {25.e-6f, 0.f, 0.f, 25.e-6f, 0.f, 36.f}, 1, 1.f);
   gsl::span<const Vertex> diamondSpan(&diamondVert, 1);

Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu

@@ -472,8 +472,8 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
   const float NSigmaCut,
   const float phiCut,
   const float resolutionPV,
-  const float maxR,
   const float minR,
+  const float maxR,
   const float positionResolution,
   const float meanDeltaR = -666.f,
   const float MSAngle = -666.f)
@@ -496,7 +496,7 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
       unsigned int storedTracklets{0};
       auto currentCluster{clustersCurrentLayer[currentClusterIndex]};
       const int currentSortedIndex{ROFClusters[layerIndex][rof0] + currentClusterIndex};
-      if (usedClusters[layerIndex][currentSortedIndex]) {
+      if (usedClusters[layerIndex][currentCluster.clusterId]) {
         continue;
       }
 
@@ -774,8 +774,6 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils,
       resolutions[iLayer],
       radii[iLayer + 1] - radii[iLayer],
       mulScatAng[iLayer]);
-    // gpuCheckError(cudaPeekAtLastError());
-    // gpuCheckError(cudaDeviceSynchronize());
     void* d_temp_storage = nullptr;
     size_t temp_storage_bytes = 0;
     gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage,            // d_temp_storage

Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx

@@ -193,6 +193,18 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in
           }
         }
       }
+      // if (rof0 == 81) {
+      //   printf("CPU layer: %d -> %f %f %f %f %f %f %f %f\n",
+      //          iLayer,
+      //          mTrkParams[iteration].NSigmaCut,
+      //          tf->getPhiCut(iLayer),
+      //          mTrkParams[iteration].PVres,
+      //          tf->getMinR(iLayer + 1),
+      //          tf->getMaxR(iLayer + 1),
+      //          tf->getPositionResolution(iLayer),
+      //          meanDeltaR,
+      //          tf->getMSangle(iLayer));
+      // }
     }
   }
   if (!tf->checkMemory(mTrkParams[iteration].MaxMemory)) {