ITSGPU: further optimisations

Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>

Author: f3sch

DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TimeEstBC.h

@@ -22,9 +22,9 @@
 namespace o2::its
 {
 // Time estimates are given in BC
-// error needs to cover maximum 1 orbit
+// error needs to cover maximum 1 orbit (uint16_t), but increased due to 2 byte padding
 using TimeStampType = uint32_t;
-using TimeStampErrorType = uint16_t;
+using TimeStampErrorType = uint32_t;
 // this is an symmetric time error [t0-tE, t0+tE]
 using TimeStamp = o2::dataformats::TimeStampWithError<float, float>;
 // this is an asymmetric time interval [t0, t0+tE] used for internal calculations
@@ -95,7 +95,7 @@ class TimeEstBC : public o2::dataformats::TimeStampWithError<TimeStampType, Time
     this->setTimeStampError(static_cast<TimeStampErrorType>(hi - lo));
   }
 
-  ClassDefNV(TimeEstBC, 1);
+  ClassDefNV(TimeEstBC, 2);
 };
 
 } // namespace o2::its

DataFormats/common/src/CommonDataFormatLinkDef.h

@@ -31,7 +31,7 @@
 #pragma link C++ class o2::dataformats::TimeStampWithError < float, float> + ;
 #pragma link C++ class o2::dataformats::TimeStampWithError < double, double> + ;
 #pragma link C++ class o2::dataformats::TimeStampWithError < int, int> + ;
-#pragma link C++ class o2::dataformats::TimeStampWithError < uint32_t, uint16_t> + ;
+#pragma link C++ class o2::dataformats::TimeStampWithError < uint32_t, uint32_t> + ;
 
 #pragma link C++ class o2::dataformats::EvIndex < int, int> + ;
 #pragma link C++ class o2::dataformats::RangeReference < int, int> + ;

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

@@ -101,6 +101,7 @@ struct TrackingKernels {
                                            const float maxChi2ClusterAttachment,
                                            const float bz,
                                            const unsigned int nCells,
+                                           o2::its::ExternalAllocator* alloc,
                                            gpu::Stream& stream);
 
   static void processNeighboursHandler(const int startLevel,

Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt

@@ -32,6 +32,7 @@ if(CUDA_ENABLED)
     set_property(TARGET ${targetName} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
     target_compile_options(${targetName} PRIVATE
         $<$<COMPILE_LANGUAGE:CUDA>:-diag-error=20014>
+        $<$<COMPILE_LANGUAGE:CUDA>:-lineinfo>
     #     $<$<COMPILE_LANGUAGE:CUDA>:-G;-O0;-Xptxas=-O0>
     #     $<$<COMPILE_LANGUAGE:CXX>:-O0;-g>
     )

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

@@ -263,12 +263,12 @@ void TimeFrameGPU<NLayers>::createUsedClustersDevice(const int layer)
 template <int NLayers>
 void TimeFrameGPU<NLayers>::loadUsedClustersDevice()
 {
+  GPUTimer timer("loading used clusters flags");
   for (auto iLayer{0}; iLayer < NLayers; ++iLayer) {
-    GPUTimer timer(mGpuStreams[iLayer], "loading used clusters flags", iLayer);
     const auto& used = this->mUsedClusters[iLayer];
     GPULog("gpu-transfer: loading {} used clusters flags on layer {}, for {:.2f} MB.", used.size(), iLayer, used.size() * sizeof(unsigned char) / constants::MB);
     if (!used.empty()) {
-      GPUChkErrS(cudaMemcpyAsync(mUsedClustersDevice[iLayer], used.data(), used.size() * sizeof(unsigned char), cudaMemcpyHostToDevice, mGpuStreams[iLayer].get()));
+      GPUChkErrS(cudaMemcpyAsync(mUsedClustersDevice[iLayer], used.data(), used.size() * sizeof(unsigned char), cudaMemcpyHostToDevice, Stream::DefaultStream));
     }
   }
 }

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

@@ -242,6 +242,7 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
                                                                  this->mTrkParams[iteration].MaxChi2ClusterAttachment,
                                                                  this->mBz,
                                                                  mTimeFrameGPU->getNCells()[sourceCellTopologyId],
+                                                                 mTimeFrameGPU->getFrameworkAllocator(),
                                                                  stream);
         }
         return finalizeCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbours(targetCellTopologyId),

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

@@ -29,13 +29,14 @@ enum SlabSite : uint8_t {
   Tracklets = 0,
   Cells,
   Neighbours,
+  RoadCandidates,
   Roads,
   TrackSeeds,
   TracksExtended,
   Tracks,
   NSlabSite,
 };
-constexpr const char* const SlabSiteNames[SlabSite::NSlabSite]{"Tracklets", "Cells", "Neighbours", "Roads", "TrackSeeds", "TracksExtended", "Tracks"};
+constexpr const char* const SlabSiteNames[SlabSite::NSlabSite]{"Tracklets", "Cells", "Neighbours", "RoadCandidates", "Roads", "TrackSeeds", "TracksExtended", "Tracks"};
 
 class CapacityEstimator
 {

Detectors/ITSMFT/ITS/tracking/include/ITStracking/CapacityEstimator.h

@@ -65,7 +65,7 @@ struct Tracklet final {
   int secondClusterIndex{constants::UnusedIndex};
   float tanLambda{constants::UnsetValue};
   float phi{constants::UnsetValue};
-  TimeEstBC mTime;
+  TimeEstBC mTime{};
 
   ClassDefNV(Tracklet, 1);
 };