Skip to content

ITS::gpu: Update track selection logics to the state of the art #13816

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
mconcas merged 1 commit into from
Dec 19, 2024
Merged

ITS::gpu: Update track selection logics to the state of the art #13816

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -182,6 +183,7 @@ void filterCellNeighboursHandler(std::vector<int>&,
void trackSeedHandler(CellSeed* trackSeeds,
const TrackingFrameInfo** foundTrackingFrameInfo,
o2::its::TrackITSExt* tracks,
std::vector<float>& minPtsHost,
const unsigned int nSeeds,
const float Bz,
const int startLevel,
Expand Down
4 changes: 4 additions & 0 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -278,6 +278,9 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
const int minimumLayer{startLevel - 1};
std::vector<CellSeed> trackSeeds;
for (int startLayer{mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) {
if ((mTrkParams[iteration].StartLayerMask & (1 << (startLayer + 2))) == 0) {
continue;
}
std::vector<int> lastCellId, updatedCellId;
std::vector<CellSeed> lastCellSeed, updatedCellSeed;

Expand Down Expand Up @@ -308,6 +311,7 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* trackSeeds,
mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo,
mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* tracks,
mTrkParams[iteration].MinPt, // std::vector<float>& minPtsHost,
trackSeeds.size(), // const size_t nSeeds,
mBz, // const float Bz,
startLevel, // const int startLevel,
Expand Down
26 changes: 15 additions & 11 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -1089,6 +1090,7 @@ void filterCellNeighboursHandler(std::vector<int>& neighHost,
void trackSeedHandler(CellSeed* trackSeeds,
const TrackingFrameInfo** foundTrackingFrameInfo,
o2::its::TrackITSExt* tracks,
std::vector<float>& minPtsHost,
const unsigned int nSeeds,
const float Bz,
const int startLevel,
Expand All @@ -1099,17 +1101,19 @@ void trackSeedHandler(CellSeed* trackSeeds,
const int nBlocks,
const int nThreads)
{
thrust::device_vector<float> minPts(minPtsHost);
gpu::fitTrackSeedsKernel<<<nBlocks, nThreads>>>(
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());
Expand Down
Loading