diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 54bdae302e643..b0db34199df09 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -40,6 +40,7 @@ GPUg() void fitTrackSeedsKernel( CellSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, o2::its::TrackITSExt* tracks, + const float* minPts, const unsigned int nSeeds, const float Bz, const int startLevel, @@ -182,6 +183,7 @@ void filterCellNeighboursHandler(std::vector&, void trackSeedHandler(CellSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, o2::its::TrackITSExt* tracks, + std::vector& minPtsHost, const unsigned int nSeeds, const float Bz, const int startLevel, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index ae86507e46325..4e34600562ae7 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -278,6 +278,9 @@ void TrackerTraitsGPU::findRoads(const int iteration) const int minimumLayer{startLevel - 1}; std::vector trackSeeds; for (int startLayer{mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) { + if ((mTrkParams[iteration].StartLayerMask & (1 << (startLayer + 2))) == 0) { + continue; + } std::vector lastCellId, updatedCellId; std::vector lastCellSeed, updatedCellSeed; @@ -308,6 +311,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* trackSeeds, mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo, mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* tracks, + mTrkParams[iteration].MinPt, // std::vector& minPtsHost, trackSeeds.size(), // const size_t nSeeds, mBz, // const float Bz, startLevel, // const int startLevel, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 229827611c077..a723f36c17dd0 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -267,6 +267,7 @@ GPUg() void fitTrackSeedsKernel( CellSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, o2::its::TrackITSExt* tracks, + const float* minPts, const unsigned int nSeeds, const float Bz, const int startLevel, @@ -317,7 +318,7 @@ GPUg() void fitTrackSeedsKernel( foundTrackingFrameInfo, // TrackingFrameInfo** trackingFrameInfo, propagator, // const o2::base::Propagator* propagator, matCorrType); // o2::base::PropagatorF::MatCorrType matCorrType - if (!fitSuccess) { + if (!fitSuccess || temporaryTrack.getPt() < minPts[nLayers - temporaryTrack.getNClusters()]) { continue; } tracks[iCurrentTrackSeedIndex] = temporaryTrack; @@ -1089,6 +1090,7 @@ void filterCellNeighboursHandler(std::vector& neighHost, void trackSeedHandler(CellSeed* trackSeeds, const TrackingFrameInfo** foundTrackingFrameInfo, o2::its::TrackITSExt* tracks, + std::vector& minPtsHost, const unsigned int nSeeds, const float Bz, const int startLevel, @@ -1099,17 +1101,19 @@ void trackSeedHandler(CellSeed* trackSeeds, const int nBlocks, const int nThreads) { + thrust::device_vector minPts(minPtsHost); gpu::fitTrackSeedsKernel<<>>( - trackSeeds, // CellSeed* - foundTrackingFrameInfo, // TrackingFrameInfo** - tracks, // TrackITSExt* - nSeeds, // const unsigned int - Bz, // const float - startLevel, // const int - maxChi2ClusterAttachment, // float - maxChi2NDF, // float - propagator, // const o2::base::Propagator* - matCorrType); // o2::base::PropagatorF::MatCorrType + trackSeeds, // CellSeed* + foundTrackingFrameInfo, // TrackingFrameInfo** + tracks, // TrackITSExt* + thrust::raw_pointer_cast(&minPts[0]), // const float* minPts, + nSeeds, // const unsigned int + Bz, // const float + startLevel, // const int + maxChi2ClusterAttachment, // float + maxChi2NDF, // float + propagator, // const o2::base::Propagator* + matCorrType); // o2::base::PropagatorF::MatCorrType gpuCheckError(cudaPeekAtLastError()); gpuCheckError(cudaDeviceSynchronize());