Skip to content

Commit 4613f73

Browse files
mconcasmconcas
committed
Add tracklet counting
1 parent 2bdeefe commit 4613f73

File tree

6 files changed

+215
-74
lines changed

6 files changed

+215
-74
lines changed

‎Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,7 @@ class TimeFrameGPU : public TimeFrame
7272
void loadTrackSeedsChi2Device();
7373
void loadRoadsDevice();
7474
void loadTrackSeedsDevice(std::vector<CellSeed>&);
75+
void createTrackletsBuffers();
7576
void createCellsBuffers(const int);
7677
void createCellsDevice();
7778
void createCellsLUTDevice();
@@ -139,6 +140,7 @@ class TimeFrameGPU : public TimeFrame
139140
gsl::span<int> getHostNCells(const int chunkId);
140141

141142
// Host-available device getters
143+
gsl::span<int*> getDeviceTrackletsLUTs() { return mTrackletsLUTDevice; }
142144
gsl::span<int*> getDeviceCellLUTs() { return mCellsLUTDevice; }
143145
gsl::span<CellSeed*> getDeviceCells() { return mCellsDevice; }
144146
gsl::span<int, nLayers - 2> getNCellsDevice() { return mNCells; }
@@ -151,6 +153,7 @@ class TimeFrameGPU : public TimeFrame
151153
StaticTrackingParameters<nLayers> mStaticTrackingParams;
152154

153155
// Host-available device buffer sizes
156+
std::array<int, nLayers - 1> mNTracklets;
154157
std::array<int, nLayers - 2> mNCells;
155158

156159
// Device pointers

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

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,36 @@ GPUg() void fitTrackSeedsKernel(
5050
#endif
5151
} // namespace gpu
5252

53+
template <int nLayers = 7>
54+
void countTrackletsInROFsHandler(const IndexTableUtils* utils,
55+
const uint8_t* multMask,
56+
const int startROF,
57+
const int endROF,
58+
const int maxROF,
59+
const int deltaROF,
60+
const int vertexId,
61+
const Vertex* vertices,
62+
const int* rofPV,
63+
const int nVertices,
64+
const Cluster** clusters,
65+
std::vector<unsigned int> nClusters,
66+
const int** ROFClusters,
67+
const unsigned char** usedClusters,
68+
const int** clustersIndexTables,
69+
int** trackletsLUTs,
70+
gsl::span<int*> trackletsLUTsHost,
71+
const int iteration,
72+
const float NSigmaCut,
73+
std::vector<float>& phiCuts,
74+
const float resolutionPV,
75+
std::vector<float>& minR,
76+
std::vector<float>& maxR,
77+
std::vector<float>& resolutions,
78+
std::vector<float>& radii,
79+
std::vector<float>& mulScatAng,
80+
const int nBlocks,
81+
const int nThreads);
82+
5383
template <int nLayers = 7>
5484
void computeTrackletsInROFsHandler(const IndexTableUtils* utils,
5585
const uint8_t* multMask,
@@ -66,6 +96,7 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils,
6696
const int** ROFClusters,
6797
const unsigned char** usedClusters,
6898
const int** clustersIndexTables,
99+
Tracklet* tracklets,
69100
int** trackletsLUTs,
70101
const int iteration,
71102
const float NSigmaCut,

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

Lines changed: 11 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -259,17 +259,17 @@ void TimeFrameGPU<nLayers>::createTrackletsLUTDevice()
259259
STOP_GPU_STREAM_TIMER(mGpuStreams[0].get());
260260
}
261261

262-
// template<int nLayers> void TimeFrameGPU<nLayers>::createTrackletsBuffers()
263-
// {
264-
// START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "creating cells buffers");
265-
// for (auto iLayer{0}; iLayer < nLayers - 1; ++iLayer) {
266-
// mNTracklets[iLayer] = 0;
267-
// checkGPUError(cudaMemcpyAsync(&mNTracklets[iLayer], mTrackletsLUTDevice[iLayer] + mClusters[iLayer].size(), sizeof(int), cudaMemcpyDeviceToHost));
268-
// LOGP(debug, "gpu-transfer: creating tracklets buffer for {} elements on layer {}, for {} MB.", mNTracklets[layer], iLayer, mNTracklets[iLayer] * sizeof(CellSeed) / MB);
269-
// allocMemAsync(reinterpret_cast<void**>(&mTrackletsDevice[iLayer]), mNTracklets[iLayer] * sizeof(Tracklet), nullptr, getExtAllocator());
270-
// }
271-
// STOP_GPU_STREAM_TIMER(mGpuStreams[0].get());
272-
// }
262+
template<int nLayers> void TimeFrameGPU<nLayers>::createTrackletsBuffers()
263+
{
264+
START_GPU_STREAM_TIMER(mGpuStreams[0].get(), "creating cells buffers");
265+
for (auto iLayer{0}; iLayer < nLayers - 1; ++iLayer) {
266+
mNTracklets[iLayer] = 0;
267+
checkGPUError(cudaMemcpyAsync(&mNTracklets[iLayer], mTrackletsLUTDevice[iLayer] + mClusters[iLayer].size(), sizeof(int), cudaMemcpyDeviceToHost));
268+
LOGP(info, "gpu-transfer: creating tracklets buffer for {} elements on layer {}, for {} MB.", mNTracklets[iLayer], iLayer, mNTracklets[iLayer] * sizeof(Tracklet) / MB);
269+
allocMemAsync(reinterpret_cast<void**>(&mTrackletsDevice[iLayer]), mNTracklets[iLayer] * sizeof(Tracklet), nullptr, getExtAllocator());
270+
}
271+
STOP_GPU_STREAM_TIMER(mGpuStreams[0].get());
272+
}
273273

274274
template <int nLayers>
275275
void TimeFrameGPU<nLayers>::loadTrackletsDevice()

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

Lines changed: 29 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -99,33 +99,35 @@ void TrackerTraitsGPU<nLayers>::computeTrackletsHybrid(const int iteration, int
9999
int startROF{mTrkParams[iteration].nROFsPerIterations > 0 ? iROFslice * mTrkParams[iteration].nROFsPerIterations : 0};
100100
int endROF{mTrkParams[iteration].nROFsPerIterations > 0 ? (iROFslice + 1) * mTrkParams[iteration].nROFsPerIterations + mTrkParams[iteration].DeltaROF : mTimeFrameGPU->getNrof()};
101101

102-
computeTrackletsInROFsHandler<nLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
103-
mTimeFrameGPU->getDeviceMultCutMask(),
104-
startROF,
105-
endROF,
106-
mTimeFrameGPU->getNrof(),
107-
mTrkParams[iteration].DeltaROF,
108-
iVertex,
109-
mTimeFrameGPU->getDeviceVertices(),
110-
mTimeFrameGPU->getDeviceROFramesPV(),
111-
mTimeFrameGPU->getPrimaryVerticesNum(),
112-
mTimeFrameGPU->getDeviceArrayClusters(),
113-
mTimeFrameGPU->getClusterSizes(),
114-
mTimeFrameGPU->getDeviceROframeClusters(),
115-
mTimeFrameGPU->getDeviceArrayUsedClusters(),
116-
mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
117-
mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
118-
iteration,
119-
mTrkParams[iteration].NSigmaCut,
120-
mTimeFrameGPU->getPhiCuts(),
121-
mTrkParams[iteration].PVres,
122-
mTimeFrameGPU->getMinRs(),
123-
mTimeFrameGPU->getMaxRs(),
124-
mTimeFrameGPU->getPositionResolutions(),
125-
mTrkParams[iteration].LayerRadii,
126-
mTimeFrameGPU->getMSangles(),
127-
conf.nBlocks,
128-
conf.nThreads);
102+
countTrackletsInROFsHandler<nLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
103+
mTimeFrameGPU->getDeviceMultCutMask(),
104+
startROF,
105+
endROF,
106+
mTimeFrameGPU->getNrof(),
107+
mTrkParams[iteration].DeltaROF,
108+
iVertex,
109+
mTimeFrameGPU->getDeviceVertices(),
110+
mTimeFrameGPU->getDeviceROFramesPV(),
111+
mTimeFrameGPU->getPrimaryVerticesNum(),
112+
mTimeFrameGPU->getDeviceArrayClusters(),
113+
mTimeFrameGPU->getClusterSizes(),
114+
mTimeFrameGPU->getDeviceROframeClusters(),
115+
mTimeFrameGPU->getDeviceArrayUsedClusters(),
116+
mTimeFrameGPU->getDeviceArrayClustersIndexTables(),
117+
mTimeFrameGPU->getDeviceArrayTrackletsLUT(),
118+
mTimeFrameGPU->getDeviceTrackletsLUTs(),
119+
iteration,
120+
mTrkParams[iteration].NSigmaCut,
121+
mTimeFrameGPU->getPhiCuts(),
122+
mTrkParams[iteration].PVres,
123+
mTimeFrameGPU->getMinRs(),
124+
mTimeFrameGPU->getMaxRs(),
125+
mTimeFrameGPU->getPositionResolutions(),
126+
mTrkParams[iteration].LayerRadii,
127+
mTimeFrameGPU->getMSangles(),
128+
conf.nBlocks,
129+
conf.nThreads);
130+
mTimeFrameGPU->createTrackletsBuffers();
129131
}
130132

131133
template <int nLayers>

0 commit comments

Comments
 (0)