Skip to content

ITS-GPU: finalise processNeighbours porting #13899

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
Jan 25, 2025
Merged

ITS-GPU: finalise processNeighbours porting #13899

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
4 changes: 2 additions & 2 deletions DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -160,7 +160,7 @@ class TrackParametrization
GPUd() value_t getZ() const;
GPUd() value_t getSnp() const;
GPUd() value_t getTgl() const;
GPUd() value_t getQ2Pt() const;
GPUhd() value_t getQ2Pt() const;
GPUd() value_t getCharge2Pt() const;
GPUd() int getAbsCharge() const;
GPUd() PID getPID() const;
Expand Down Expand Up @@ -357,7 +357,7 @@ GPUdi() auto TrackParametrization<value_T>::getTgl() const -> value_t

//____________________________________________________________
template <typename value_T>
GPUdi() auto TrackParametrization<value_T>::getQ2Pt() const -> value_t
GPUhdi() auto TrackParametrization<value_T>::getQ2Pt() const -> value_t
{
return mP[kQ2Pt];
}
Expand Down
2 changes: 2 additions & 0 deletions Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -116,6 +116,7 @@ class TimeFrameGPU : public TimeFrame
int* getDeviceNeighboursLUT(const int layer) { return mNeighboursLUTDevice[layer]; }
gsl::span<int*> getDeviceNeighboursLUTs() { return mNeighboursLUTDevice; }
gpuPair<int, int>* getDeviceNeighbourPairs(const int layer) { return mNeighbourPairsDevice[layer]; }
std::array<int*, nLayers - 2>& getDeviceNeighboursAll() { return mNeighboursDevice; }
int* getDeviceNeighbours(const int layer) { return mNeighboursDevice[layer]; }
int** getDeviceNeighboursArray() { return mNeighboursDeviceArray; }
TrackingFrameInfo* getDeviceTrackingFrameInfo(const int);
Expand All @@ -142,6 +143,7 @@ class TimeFrameGPU : public TimeFrame
// Host-specific getters
gsl::span<int, nLayers - 1> getNTracklets() { return mNTracklets; }
gsl::span<int, nLayers - 2> getNCells() { return mNCells; }
std::array<int, nLayers - 2>& getArrayNCells() { return mNCells; }

// Host-available device getters
gsl::span<int*> getDeviceTrackletsLUTs() { return mTrackletsLUTDevice; }
Expand Down
10 changes: 4 additions & 6 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 @@ -186,19 +186,17 @@ void processNeighboursHandler(const int startLayer,
const int startLevel,
CellSeed** allCellSeeds,
CellSeed* currentCellSeeds,
const unsigned int nCurrentCells,
std::array<int, nLayers - 2>& nCells,
const unsigned char** usedClusters,
int* neighbours,
std::array<int*, nLayers - 2>& neighbours,
gsl::span<int*> neighboursDeviceLUTs,
const TrackingFrameInfo** foundTrackingFrameInfo,
std::vector<CellSeed>& seedsHost,
const float bz,
const float MaxChi2ClusterAttachment,
const float maxChi2NDF,
const o2::base::Propagator* propagator,
const o2::base::PropagatorF::MatCorrType matCorrType,
const std::vector<int>& lastCellIdHost, // temporary host vector
const std::vector<CellSeed>& lastCellSeedHost, // temporary host vector
std::vector<int>& updatedCellIdHost, // temporary host vector
std::vector<CellSeed>& updatedCellSeedHost, // temporary host vector
const int nBlocks,
const int nThreads);

Expand Down
2 changes: 1 addition & 1 deletion Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@
if(CUDA_ENABLED)
find_package(CUDAToolkit)
message(STATUS "Building ITS CUDA tracker")
# add_compile_options(-O0 -g -lineinfo -fPIC)
add_compile_options(-O0 -g -lineinfo -fPIC)
# add_compile_definitions(ITS_MEASURE_GPU_TIME)
o2_add_library(ITStrackingCUDA
SOURCES ClusterLinesGPU.cu
Expand Down
63 changes: 11 additions & 52 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 @@ -205,9 +205,6 @@ void TrackerTraitsGPU<nLayers>::computeCellsHybrid(const int iteration)
conf.nBlocks,
conf.nThreads);
}
// Needed for processNeighbours() which is still on CPU.
mTimeFrameGPU->downloadCellsDevice();
mTimeFrameGPU->downloadCellsLUTDevice();
}

template <int nLayers>
Expand All @@ -221,11 +218,11 @@ void TrackerTraitsGPU<nLayers>::findCellsNeighboursHybrid(const int iteration)
mTimeFrameGPU->getCellsNeighboursLUT()[iLayer].clear();
mTimeFrameGPU->getCellsNeighboursLUT()[iLayer].resize(nextLayerCellsNum, 0);

if (mTimeFrameGPU->getCells()[iLayer + 1].empty() ||
mTimeFrameGPU->getCellsLookupTable()[iLayer].empty()) {
mTimeFrameGPU->getCellsNeighbours()[iLayer].clear();
continue;
}
// if (mTimeFrameGPU->getCells()[iLayer + 1].empty() ||
// mTimeFrameGPU->getCellsLookupTable()[iLayer].empty()) {
// mTimeFrameGPU->getCellsNeighbours()[iLayer].clear();
// continue;
// }

mTimeFrameGPU->createNeighboursLUTDevice(iLayer, nextLayerCellsNum);
countCellNeighboursHandler(mTimeFrameGPU->getDeviceArrayCells(),
Expand Down Expand Up @@ -267,7 +264,6 @@ void TrackerTraitsGPU<nLayers>::findCellsNeighboursHybrid(const int iteration)
cellsNeighboursLayer[iLayer].size());
}
mTimeFrameGPU->createNeighboursDeviceArray();
mTimeFrameGPU->downloadCellsDevice();
mTimeFrameGPU->unregisterRest();
};

Expand All @@ -289,55 +285,21 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
startLevel,
mTimeFrameGPU->getDeviceArrayCells(),
mTimeFrameGPU->getDeviceCells()[startLayer],
mTimeFrameGPU->getNCells()[startLayer],
mTimeFrameGPU->getArrayNCells(),
mTimeFrameGPU->getDeviceArrayUsedClusters(),
mTimeFrameGPU->getDeviceNeighbours(startLayer - 1),
mTimeFrameGPU->getDeviceNeighboursAll(),
mTimeFrameGPU->getDeviceNeighboursLUTs(),
mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
trackSeeds,
mBz,
mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment
mTrkParams[0].MaxChi2ClusterAttachment,
mTrkParams[0].MaxChi2NDF,
mTimeFrameGPU->getDevicePropagator(),
mCorrType,
lastCellId, // temporary host vector
lastCellSeed, // temporary host vector
updatedCellId, // temporary host vectors
updatedCellSeed, // temporary host vectors
conf.nBlocks,
conf.nThreads);

int level = startLevel;
for (int iLayer{startLayer - 1}; iLayer > 0 && level > 2; --iLayer) {
lastCellSeed.swap(updatedCellSeed);
lastCellId.swap(updatedCellId);
std::vector<CellSeed>().swap(updatedCellSeed); /// tame the memory peaks
updatedCellId.clear();
processNeighboursHandler<nLayers>(iLayer,
--level,
mTimeFrameGPU->getDeviceArrayCells(),
mTimeFrameGPU->getDeviceCells()[iLayer],
mTimeFrameGPU->getNCells()[iLayer],
mTimeFrameGPU->getDeviceArrayUsedClusters(),
mTimeFrameGPU->getDeviceNeighbours(iLayer - 1),
mTimeFrameGPU->getDeviceNeighboursLUTs(),
mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
mBz,
mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment
mTimeFrameGPU->getDevicePropagator(),
mCorrType,
lastCellId, // temporary host vector
lastCellSeed, // temporary host vector
updatedCellId, // temporary host vectors
updatedCellSeed, // temporary host vectors
conf.nBlocks,
conf.nThreads);
}
for (auto& seed : updatedCellSeed) {
if (seed.getQ2Pt() > 1.e3 || seed.getChi2() > mTrkParams[0].MaxChi2NDF * ((startLevel + 2) * 2 - 5)) {
continue;
}
trackSeeds.push_back(seed);
}
}
// fixme: I don't want to move tracks back and forth, but I need a way to use a thrust::allocator that is aware of our managed memory.
if (!trackSeeds.size()) {
LOGP(info, "No track seeds found, skipping track finding");
continue;
Expand All @@ -362,9 +324,6 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
mTimeFrameGPU->downloadTrackITSExtDevice(trackSeeds);

auto& tracks = mTimeFrameGPU->getTrackITSExt();
std::sort(tracks.begin(), tracks.end(), [](const TrackITSExt& a, const TrackITSExt& b) {
return a.getChi2() < b.getChi2();
});

for (auto& track : tracks) {
if (!track.getChi2()) {
Expand Down
Loading
Loading