Skip to content

ITS-GPU: Fixup for running on PbPb in async mode #12489

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 2 commits into from
Jan 8, 2024
Merged

ITS-GPU: Fixup for running on PbPb in async mode #12489

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
eccb9c0
Fix tracking mode not correctly forwarded
mconcas Nov 15, 2023
145d9f8
Make ITS GPU running on PbPb in async
mconcas Dec 21, 2023
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 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 @@ -190,7 +190,7 @@ class TimeFrameGPU : public TimeFrame
void initialiseHybrid(const int, const TrackingParameters&, const int, IndexTableUtils* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr);
void initDevice(const int, IndexTableUtils*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int);
void initDeviceSAFitting();
void loadTrackingFrameInfoDevice();
void loadTrackingFrameInfoDevice(const int);
void loadUnsortedClustersDevice();
void loadClustersDevice();
void loadTrackletsDevice();
Expand All @@ -199,7 +199,7 @@ class TimeFrameGPU : public TimeFrame
void loadTrackSeedsChi2Device();
void loadRoadsDevice();
void loadTrackSeedsDevice(std::vector<CellSeed>&);
void createTrackITSExtDevice(const std::vector<CellSeed>& seeds);
void createTrackITSExtDevice(const std::vector<CellSeed>&);
void createTrackITSExtDevice(); // deprecated
void downloadTrackITSExtDevice();
void initDeviceChunks(const int, const int);
Expand Down
37 changes: 17 additions & 20 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -345,7 +345,7 @@ void TimeFrameGPU<nLayers>::allocMemAsync(void** ptr, size_t size, Stream* strPt
if (extAllocator) {
*ptr = mAllocator->allocate(size);
} else {
LOGP(debug, "Calling default CUDA allocator");
LOGP(info, "Calling default CUDA allocator");
checkGPUError(cudaMallocAsync(reinterpret_cast<void**>(ptr), size, strPtr->get()));
}
}
Expand Down Expand Up @@ -376,17 +376,10 @@ void TimeFrameGPU<nLayers>::registerHostMemory(const int maxLayers)
template <int nLayers>
void TimeFrameGPU<nLayers>::unregisterHostMemory(const int maxLayers)
{
if (!mHostRegistered) {
return;
}
for (auto iLayer{0}; iLayer < maxLayers; ++iLayer) {
checkGPUError(cudaHostUnregister(mClusters[iLayer].data()));
checkGPUError(cudaHostUnregister(mNClustersPerROF[iLayer].data()));
checkGPUError(cudaHostUnregister(mIndexTables[iLayer].data()));
for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {
checkGPUError(cudaHostUnregister(mTrackingFrameInfo[iLayer].data()));
}
checkGPUError(cudaHostUnregister(mHostNTracklets.data()));
checkGPUError(cudaHostUnregister(mHostNCells.data()));
mHostRegistered = false;
checkGPUError(cudaHostUnregister(mTrackingFrameInfoDevice.data()));
}

template <int nLayers>
Expand Down Expand Up @@ -525,17 +518,21 @@ void TimeFrameGPU<nLayers>::loadClustersDevice()
}

template <int nLayers>
void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice()
void TimeFrameGPU<nLayers>::loadTrackingFrameInfoDevice(const int iteration)
{
for (auto iLayer{0}; iLayer < nLayers; ++iLayer) {
LOGP(debug, "gpu-transfer: loading {} tfinfo on layer {}, for {} MB.", mTrackingFrameInfo[iLayer].size(), iLayer, mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo) / MB);
allocMemAsync(reinterpret_cast<void**>(&mTrackingFrameInfoDevice[iLayer]), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), nullptr, getExtAllocator());
// Register and move data
checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));
if (!iteration) {
checkGPUError(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable));
}
checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDevice[iLayer], mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
}
allocMemAsync(reinterpret_cast<void**>(&mTrackingFrameInfoDeviceArray), nLayers * sizeof(TrackingFrameInfo*), nullptr, getExtAllocator());
checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));
if (!iteration) {
checkGPUError(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable));
}
checkGPUError(cudaMemcpyAsync(mTrackingFrameInfoDeviceArray, mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
}

Expand Down Expand Up @@ -573,7 +570,7 @@ template <int nLayers>
void TimeFrameGPU<nLayers>::loadRoadsDevice()
{
LOGP(debug, "gpu-transfer: loading {} roads, for {} MB.", mRoads.size(), mRoads.size() * sizeof(Road<nLayers - 2>) / MB);
allocMemAsync(reinterpret_cast<void**>(&mRoadsDevice), mRoads.size() * sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), false);
allocMemAsync(reinterpret_cast<void**>(&mRoadsDevice), mRoads.size() * sizeof(Road<nLayers - 2>), &(mGpuStreams[0]), getExtAllocator());
checkGPUError(cudaHostRegister(mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaHostRegisterPortable));
checkGPUError(cudaMemcpyAsync(mRoadsDevice, mRoads.data(), mRoads.size() * sizeof(Road<nLayers - 2>), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
}
Expand All @@ -582,7 +579,7 @@ template <int nLayers>
void TimeFrameGPU<nLayers>::loadTrackSeedsDevice(std::vector<CellSeed>& seeds)
{
LOGP(debug, "gpu-transfer: loading {} track seeds, for {} MB.", seeds.size(), seeds.size() * sizeof(CellSeed) / MB);
allocMemAsync(reinterpret_cast<void**>(&mTrackSeedsDevice), seeds.size() * sizeof(CellSeed), &(mGpuStreams[0]), false);
allocMemAsync(reinterpret_cast<void**>(&mTrackSeedsDevice), seeds.size() * sizeof(CellSeed), &(mGpuStreams[0]), getExtAllocator());
checkGPUError(cudaHostRegister(seeds.data(), seeds.size() * sizeof(CellSeed), cudaHostRegisterPortable));
checkGPUError(cudaMemcpyAsync(mTrackSeedsDevice, seeds.data(), seeds.size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0].get()));
}
Expand All @@ -592,8 +589,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice(const std::vector<CellSeed>&
{
mTrackITSExt.clear();
mTrackITSExt.resize(seeds.size());
LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);
allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), seeds.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);
LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB);
allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), seeds.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());
checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));
checkGPUError(cudaHostRegister(mTrackITSExt.data(), seeds.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));
}
Expand All @@ -603,8 +600,8 @@ void TimeFrameGPU<nLayers>::createTrackITSExtDevice()
{
mTrackITSExt.clear();
mTrackITSExt.resize(mRoads.size());
LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);
allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), mRoads.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), false);
LOGP(info, "gpu-allocation: reserving {} tracks, for {} MB.", mRoads.size(), mRoads.size() * sizeof(o2::its::TrackITSExt) / MB);
allocMemAsync(reinterpret_cast<void**>(&mTrackITSExtDevice), mRoads.size() * sizeof(o2::its::TrackITSExt), &(mGpuStreams[0]), getExtAllocator());
checkGPUError(cudaMemsetAsync(mTrackITSExtDevice, 0, mRoads.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0].get()));
checkGPUError(cudaHostRegister(mTrackITSExt.data(), mRoads.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable));
}
Expand Down
16 changes: 9 additions & 7 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -870,9 +870,7 @@ template <int nLayers>
void TrackerTraitsGPU<nLayers>::initialiseTimeFrame(const int iteration)
{
mTimeFrameGPU->initialiseHybrid(iteration, mTrkParams[iteration], nLayers);
mTimeFrameGPU->loadClustersDevice();
mTimeFrameGPU->loadUnsortedClustersDevice();
mTimeFrameGPU->loadTrackingFrameInfoDevice();
mTimeFrameGPU->loadTrackingFrameInfoDevice(iteration);
}

template <int nLayers>
Expand Down Expand Up @@ -1165,14 +1163,12 @@ template <int nLayers>
void TrackerTraitsGPU<nLayers>::computeTrackletsHybrid(const int iteration)
{
TrackerTraits::computeLayerTracklets(iteration);
mTimeFrameGPU->loadTrackletsDevice();
}

template <int nLayers>
void TrackerTraitsGPU<nLayers>::computeCellsHybrid(const int iteration)
{
TrackerTraits::computeLayerCells(iteration);
mTimeFrameGPU->loadCellsDevice();
};

template <int nLayers>
Expand Down Expand Up @@ -1203,6 +1199,10 @@ void TrackerTraitsGPU<nLayers>::findRoadsHybrid(const int iteration)
}
trackSeeds.insert(trackSeeds.end(), updatedCellSeed.begin(), updatedCellSeed.end());
}
if (!trackSeeds.size()) {
LOGP(info, "No track seeds found, skipping track finding");
continue;
}
mTimeFrameGPU->createTrackITSExtDevice(trackSeeds);
mTimeFrameGPU->loadTrackSeedsDevice(trackSeeds);

Expand Down Expand Up @@ -1265,12 +1265,15 @@ void TrackerTraitsGPU<nLayers>::findRoadsHybrid(const int iteration)
mTimeFrame->getTracks(std::min(rofs[0], rofs[1])).emplace_back(track);
}
}
if (iteration == 2) {
mTimeFrameGPU->unregisterHostMemory(0); // FIXME this needs to work also with sync
}
};

template <int nLayers>
void TrackerTraitsGPU<nLayers>::findTracksHybrid(const int iteration)
{
// LOGP(info, "propagator device pointer: {}", (void*)mTimeFrameGPU->getDevicePropagator());
LOGP(info, "propagator device pointer: {}", (void*)mTimeFrameGPU->getDevicePropagator());
mTimeFrameGPU->createTrackITSExtDevice();
gpu::fitTracksKernel<<<20, 512>>>(mTimeFrameGPU->getDeviceArrayClusters(), // Cluster** foundClusters,
mTimeFrameGPU->getDeviceArrayUnsortedClusters(), // Cluster** foundUnsortedClusters,
Expand All @@ -1287,7 +1290,6 @@ void TrackerTraitsGPU<nLayers>::findTracksHybrid(const int iteration)
mTrkParams[0].MaxChi2NDF, // float maxChi2NDF,
mTimeFrameGPU->getDevicePropagator()); // const o2::base::Propagator* propagator
mTimeFrameGPU->downloadTrackITSExtDevice();
discardResult(cudaDeviceSynchronize());
auto& tracks = mTimeFrameGPU->getTrackITSExt();
std::sort(tracks.begin(), tracks.end(),
[](TrackITSExt& track1, TrackITSExt& track2) { return track1.isBetter(track2, 1.e6f); });
Expand Down
5 changes: 3 additions & 2 deletions Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -106,6 +106,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,
double total{0.};
mTraits->UpdateTrackingParameters(mTrkParams);
for (int iteration = 0; iteration < (int)mTrkParams.size(); ++iteration) {
LOGP(info, "Iteration {}", iteration);
total += evaluateTask(&Tracker::initialiseTimeFrameHybrid, "Hybrid Timeframe initialisation", logger, iteration);
total += evaluateTask(&Tracker::computeTrackletsHybrid, "Hybrid Tracklet finding", logger, iteration);
logger(fmt::format("\t- Number of tracklets: {}", mTraits->getTFNumberOfTracklets()));
Expand All @@ -115,7 +116,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,
}
float trackletsPerCluster = mTraits->getTFNumberOfClusters() > 0 ? float(mTraits->getTFNumberOfTracklets()) / mTraits->getTFNumberOfClusters() : 0.f;
if (trackletsPerCluster > mTrkParams[iteration].TrackletsPerClusterLimit) {
error(fmt::format("Too many tracklets per cluster ({}), check the detector status and/or the selections.", trackletsPerCluster));
error(fmt::format("Too many tracklets per cluster ({}), check the detector status and/or the selections. Current limit is {}", trackletsPerCluster, mTrkParams[iteration].TrackletsPerClusterLimit));
break;
}

Expand All @@ -127,7 +128,7 @@ void Tracker::clustersToTracksHybrid(std::function<void(std::string s)> logger,
}
float cellsPerCluster = mTraits->getTFNumberOfClusters() > 0 ? float(mTraits->getTFNumberOfCells()) / mTraits->getTFNumberOfClusters() : 0.f;
if (cellsPerCluster > mTrkParams[iteration].CellsPerClusterLimit) {
error(fmt::format("Too many cells per cluster ({}), check the detector status and/or the selections.", cellsPerCluster));
error(fmt::format("Too many cells per cluster ({}), check the detector status and/or the selections. Current limit is {}", cellsPerCluster, mTrkParams[iteration].CellsPerClusterLimit));
break;
}
total += evaluateTask(&Tracker::findCellsNeighboursHybrid, "Hybrid Neighbour finding", logger, iteration);
Expand Down
1 change: 1 addition & 0 deletions Detectors/ITSMFT/ITS/workflow/src/RecoWorkflow.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,7 @@ framework::WorkflowSpec getWorkflow(bool useMC,
cfg.runITSTracking = true;
cfg.itsTriggerType = useTrig;
cfg.itsOverrBeamEst = overrideBeamPosition;
cfg.itsTrackingMode = trmode == "sync" ? 0 : (trmode == "async" ? 1 : 2);

Inputs ggInputs;
auto ggRequest = std::make_shared<o2::base::GRPGeomRequest>(false, true, false, true, true, o2::base::GRPGeomRequest::Aligned, ggInputs, true);
Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/Base/GPUReconstruction.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -676,7 +676,7 @@ void* GPUReconstruction::AllocateUnmanagedMemory(size_t size, int type)
char* retVal;
GPUProcessor::computePointerWithAlignment(pool, retVal, size);
if (pool > poolend) {
GPUError("Insufficient unmanaged memory: missing %lu", (size_t)((char*)pool - (char*)poolend));
GPUError("Insufficient unmanaged memory: missing %lu bytes", (size_t)((char*)pool - (char*)poolend));
throw std::bad_alloc();
}
UpdateMaxMemoryUsed();
Expand Down
2 changes: 1 addition & 1 deletion GPU/GPUTracking/Global/GPUChainITS.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,7 @@ o2::its::TimeFrame* GPUChainITS::GetITSTimeframe()
mFrameworkAllocator.reset(new o2::its::GPUFrameworkExternalAllocator);
mFrameworkAllocator->setReconstructionFramework(rec());
mITSTimeFrame->setExternalAllocator(mFrameworkAllocator.get());
LOGP(debug, "GPUChainITS is giving me ps: {} prop: {}", (void*)processorsShadow(), (void*)processorsShadow()->calibObjects.o2Propagator);
LOGP(info, "GPUChainITS is giving me ps: {} prop: {} allocator: {}", (void*)processorsShadow(), (void*)processorsShadow()->calibObjects.o2Propagator, (void*)mFrameworkAllocator.get());
mITSTimeFrame->setDevicePropagator(processorsShadow()->calibObjects.o2Propagator);
}
#endif
Expand Down
2 changes: 1 addition & 1 deletion GPU/Workflow/include/GPUWorkflow/GPUWorkflowSpec.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -126,6 +126,7 @@ class GPURecoWorkflowSpec : public o2::framework::Task
int lumiScaleType = 0; // 0=off, 1=CTP, 2=TPC scalers
bool outputErrorQA = false;
bool runITSTracking = false;
int itsTrackingMode = 0; // 0=sync, 1=async, 2=cosmics
bool itsOverrBeamEst = false;
bool tpcTriggerHandling = false;
};
Expand Down Expand Up @@ -232,7 +233,6 @@ class GPURecoWorkflowSpec : public o2::framework::Task
bool mPropagatorInstanceCreated = false;
bool mITSRunVertexer = false;
bool mITSCosmicsProcessing = false;
std::string mITSMode = "sync";
};

} // end namespace gpu
Expand Down
18 changes: 11 additions & 7 deletions GPU/Workflow/src/GPUWorkflowITS.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -184,6 +184,7 @@ int GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc)
pattIt = patterns.begin();
std::vector<int> savedROF;
auto logger = [&](std::string s) { LOG(info) << s; };
auto fatalLogger = [&](std::string s) { LOG(fatal) << s; };
auto errorLogger = [&](std::string s) { LOG(error) << s; };

o2::its::FastMultEst multEst; // mult estimator
Expand Down Expand Up @@ -252,7 +253,11 @@ int GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc)

mITSTimeFrame->setMultiplicityCutMask(processingMask);
// Run CA tracker
mITSTracker->clustersToTracksHybrid(logger, errorLogger);
if (!mSpecConfig.itsTrackingMode) {
mITSTracker->clustersToTracksHybrid(logger, errorLogger);
} else {
mITSTracker->clustersToTracksHybrid(logger, fatalLogger);
}
size_t totTracks{mITSTimeFrame->getNumberOfTracks()}, totClusIDs{mITSTimeFrame->getNumberOfUsedClusters()};
allTracks.reserve(totTracks);
allClusIdx.reserve(totClusIDs);
Expand Down Expand Up @@ -303,7 +308,6 @@ int GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc)

void GPURecoWorkflowSpec::initFunctionITS(InitContext& ic)
{
std::transform(mITSMode.begin(), mITSMode.end(), mITSMode.begin(), [](unsigned char c) { return std::tolower(c); });
o2::its::VertexerTraits* vtxTraits = nullptr;
o2::its::TrackerTraits* trkTraits = nullptr;
mGPUReco->GetITSTraits(trkTraits, vtxTraits, mITSTimeFrame);
Expand All @@ -315,25 +319,27 @@ void GPURecoWorkflowSpec::initFunctionITS(InitContext& ic)
mITSCosmicsProcessing = false;
std::vector<o2::its::TrackingParameters> trackParams;

if (mITSMode == "async") {
if (mSpecConfig.itsTrackingMode == 1) {
trackParams.resize(3);
for (auto& param : trackParams) {
param.ZBins = 64;
param.PhiBins = 32;
param.CellsPerClusterLimit = 1.e3f;
param.TrackletsPerClusterLimit = 1.e3f;
}
trackParams[1].TrackletMinPt = 0.2f;
trackParams[1].CellDeltaTanLambdaSigma *= 2.;
trackParams[2].TrackletMinPt = 0.1f;
trackParams[2].CellDeltaTanLambdaSigma *= 4.;
trackParams[2].MinTrackLength = 4;
LOG(info) << "Initializing tracker in async. phase reconstruction with " << trackParams.size() << " passes";
} else if (mITSMode == "sync") {
} else if (mSpecConfig.itsTrackingMode == 0) {
trackParams.resize(1);
trackParams[0].ZBins = 64;
trackParams[0].PhiBins = 32;
trackParams[0].MinTrackLength = 4;
LOG(info) << "Initializing tracker in sync. phase reconstruction with " << trackParams.size() << " passes";
} else if (mITSMode == "cosmics") {
} else if (mSpecConfig.itsTrackingMode == 2) {
mITSCosmicsProcessing = true;
mITSRunVertexer = false;
trackParams.resize(1);
Expand All @@ -347,8 +353,6 @@ void GPURecoWorkflowSpec::initFunctionITS(InitContext& ic)
trackParams[0].TrackletsPerClusterLimit = 100.;
trackParams[0].CellsPerClusterLimit = 100.;
LOG(info) << "Initializing tracker in reconstruction for cosmics with " << trackParams.size() << " passes";
} else {
throw std::runtime_error(fmt::format("Unsupported ITS tracking mode {:s} ", mITSMode));
}

for (auto& params : trackParams) {
Expand Down