Skip to content

Commit 765e6d7

Browse files
f3schf3sch
committed
ITS: GPU vetexer
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent ada64c3 commit 765e6d7

2 files changed

Lines changed: 11 additions & 178 deletions

File tree

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,6 @@ class VertexerTraitsGPU final : public VertexerTraits
4040
void computeTrackletMatching(const int iteration = 0) final;
4141
void computeVertices(const int iteration = 0) final;
4242
void updateVertexingParameters(const std::vector<VertexingParameters>&, const TimeFrameGPUParameters&) final;
43-
void computeVerticesHist();
4443

4544
bool isGPU() const noexcept final { return true; }
4645
const char* getName() const noexcept final { return "GPU"; }

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

Lines changed: 11 additions & 177 deletions
Original file line numberDiff line numberDiff line change
@@ -11,18 +11,6 @@
1111
//
1212
/// \author matteo.concas@cern.ch
1313

14-
#include <iostream>
15-
#include <sstream>
16-
#include <fstream>
17-
#include <array>
18-
#include <cassert>
19-
#include <thread>
20-
21-
#ifdef VTX_DEBUG
22-
#include "TTree.h"
23-
#include "TFile.h"
24-
#endif
25-
2614
#include "ITStrackingGPU/VertexingKernels.h"
2715
#include "ITStrackingGPU/VertexerTraitsGPU.h"
2816

@@ -32,6 +20,10 @@ namespace o2::its
3220
void VertexerTraitsGPU::initialise(const TrackingParameters& trackingParams, const int iteration)
3321
{
3422
mTimeFrameGPU->initialise(0, trackingParams, 3, &mIndexTableUtils, &mTfGPUParams);
23+
// mTimeFrameGPU->loadClustersDevice(iteration);
24+
// mTimeFrameGPU->loadUnsortedClustersDevice(iteration);
25+
// mTimeFrameGPU->loadClustersIndexTables(iteration);
26+
// mTimeFrameGPU->loadROframeClustersDevice(iteration);
3527
}
3628

3729
void VertexerTraitsGPU::updateVertexingParameters(const std::vector<VertexingParameters>& vrtPar, const TimeFrameGPUParameters& tfPar)
@@ -47,165 +39,15 @@ void VertexerTraitsGPU::updateVertexingParameters(const std::vector<VertexingPar
4739

4840
void VertexerTraitsGPU::computeTracklets(const int iteration)
4941
{
50-
if (!mTimeFrameGPU->getClusters().size()) {
42+
if (mTimeFrameGPU->getClusters().empty()) {
5143
return;
5244
}
53-
// std::vector<std::thread> threads(mTimeFrameGPU->getNChunks());
54-
// for (int chunkId{0}; chunkId < mTimeFrameGPU->getNChunks(); ++chunkId) {
55-
// int rofPerChunk{mTimeFrameGPU->mNrof / (int)mTimeFrameGPU->getNChunks()};
56-
// mTimeFrameGPU->getVerticesInChunks()[chunkId].clear();
57-
// mTimeFrameGPU->getNVerticesInChunks()[chunkId].clear();
58-
// mTimeFrameGPU->getLabelsInChunks()[chunkId].clear();
59-
// auto doVertexReconstruction = [&, chunkId, rofPerChunk]() -> void {
60-
// auto offset = chunkId * rofPerChunk;
61-
// auto maxROF = offset + rofPerChunk;
62-
// while (offset < maxROF) {
63-
// auto rofs = mTimeFrameGPU->loadChunkData<gpu::Task::Vertexer>(chunkId, offset, maxROF);
64-
// RANGE("chunk_gpu_vertexing", 1);
65-
// // gpu::GpuTimer timer{offset, mTimeFrameGPU->getStream(chunkId).get()};
66-
// // timer.Start("vtTrackletFinder");
67-
// gpu::trackleterKernelMultipleRof<TrackletMode::Layer0Layer1><<<rofs, 1024, 0, mTimeFrameGPU->getStream(chunkId).get()>>>(
68-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(0), // const Cluster* clustersNextLayer, // 0 2
69-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(1), // const Cluster* clustersCurrentLayer, // 1 1
70-
// mTimeFrameGPU->getDeviceROframesClusters(0), // const int* sizeNextLClusters,
71-
// mTimeFrameGPU->getDeviceROframesClusters(1), // const int* sizeCurrentLClusters,
72-
// mTimeFrameGPU->getChunk(chunkId).getDeviceIndexTables(0), // const int* nextIndexTables,
73-
// mTimeFrameGPU->getChunk(chunkId).getDeviceTracklets(0), // Tracklet* Tracklets,
74-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNTrackletCluster(0), // int* foundTracklets,
75-
// mTimeFrameGPU->getDeviceIndexTableUtils(), // const IndexTableUtils* utils,
76-
// offset, // const unsigned int startRofId,
77-
// rofs, // const unsigned int rofSize,
78-
// mVrtParams.phiCut, // const float phiCut,
79-
// mVrtParams.maxTrackletsPerCluster); // const size_t maxTrackletsPerCluster = 1e2
80-
81-
// gpu::trackleterKernelMultipleRof<TrackletMode::Layer1Layer2><<<rofs, 1024, 0, mTimeFrameGPU->getStream(chunkId).get()>>>(
82-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(2), // const Cluster* clustersNextLayer, // 0 2
83-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(1), // const Cluster* clustersCurrentLayer, // 1 1
84-
// mTimeFrameGPU->getDeviceROframesClusters(2), // const int* sizeNextLClusters,
85-
// mTimeFrameGPU->getDeviceROframesClusters(1), // const int* sizeCurrentLClusters,
86-
// mTimeFrameGPU->getChunk(chunkId).getDeviceIndexTables(2), // const int* nextIndexTables,
87-
// mTimeFrameGPU->getChunk(chunkId).getDeviceTracklets(1), // Tracklet* Tracklets,
88-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNTrackletCluster(1), // int* foundTracklets,
89-
// mTimeFrameGPU->getDeviceIndexTableUtils(), // const IndexTableUtils* utils,
90-
// offset, // const unsigned int startRofId,
91-
// rofs, // const unsigned int rofSize,
92-
// mVrtParams.phiCut, // const float phiCut,
93-
// mVrtParams.maxTrackletsPerCluster); // const size_t maxTrackletsPerCluster = 1e2
94-
95-
// gpu::trackletSelectionKernelMultipleRof<true><<<rofs, 1024, 0, mTimeFrameGPU->getStream(chunkId).get()>>>(
96-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(0), // const Cluster* clusters0, // Clusters on layer 0
97-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(1), // const Cluster* clusters1, // Clusters on layer 1
98-
// mTimeFrameGPU->getDeviceROframesClusters(0), // const int* sizeClustersL0, // Number of clusters on layer 0 per ROF
99-
// mTimeFrameGPU->getDeviceROframesClusters(1), // const int* sizeClustersL1, // Number of clusters on layer 1 per ROF
100-
// mTimeFrameGPU->getChunk(chunkId).getDeviceTracklets(0), // Tracklet* tracklets01, // Tracklets on layer 0-1
101-
// mTimeFrameGPU->getChunk(chunkId).getDeviceTracklets(1), // Tracklet* tracklets12, // Tracklets on layer 1-2
102-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNTrackletCluster(0), // const int* nFoundTracklets01, // Number of tracklets found on layers 0-1
103-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNTrackletCluster(1), // const int* nFoundTracklet12, // Number of tracklets found on layers 1-2
104-
// mTimeFrameGPU->getChunk(chunkId).getDeviceUsedTracklets(), // unsigned char* usedTracklets, // Used tracklets
105-
// mTimeFrameGPU->getChunk(chunkId).getDeviceLines(), // Line* lines, // Lines
106-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNFoundLines(), // int* nFoundLines, // Number of found lines
107-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNExclusiveFoundLines(), // int* nExclusiveFoundLines, // Number of found lines exclusive scan
108-
// offset, // const unsigned int startRofId, // Starting ROF ID
109-
// rofs, // const unsigned int rofSize, // Number of ROFs to consider
110-
// mVrtParams.maxTrackletsPerCluster, // const int maxTrackletsPerCluster = 1e2, // Maximum number of tracklets per cluster
111-
// mVrtParams.tanLambdaCut, // const float tanLambdaCut = 0.025f, // Cut on tan lambda
112-
// mVrtParams.phiCut); // const float phiCut = 0.002f) // Cut on phi
113-
114-
// discardResult(cub::DeviceScan::ExclusiveSum(mTimeFrameGPU->getChunk(chunkId).getDeviceCUBTmpBuffer(),
115-
// mTimeFrameGPU->getChunk(chunkId).getTimeFrameGPUParameters()->tmpCUBBufferSize,
116-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNFoundLines(),
117-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNExclusiveFoundLines(),
118-
// mTimeFrameGPU->getTotalClustersPerROFrange(offset, rofs, 1),
119-
// mTimeFrameGPU->getStream(chunkId).get()));
120-
121-
// // Reset used tracklets
122-
// checkGPUError(cudaMemsetAsync(mTimeFrameGPU->getChunk(chunkId).getDeviceUsedTracklets(),
123-
// false,
124-
// sizeof(unsigned char) * mVrtParams.maxTrackletsPerCluster * mTimeFrameGPU->getTotalClustersPerROFrange(offset, rofs, 1),
125-
// mTimeFrameGPU->getStream(chunkId).get()),
126-
// __FILE__, __LINE__);
127-
128-
// gpu::trackletSelectionKernelMultipleRof<false><<<rofs, 1024, 0, mTimeFrameGPU->getStream(chunkId).get()>>>(
129-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(0), // const Cluster* clusters0, // Clusters on layer 0
130-
// mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(1), // const Cluster* clusters1, // Clusters on layer 1
131-
// mTimeFrameGPU->getDeviceROframesClusters(0), // const int* sizeClustersL0, // Number of clusters on layer 0 per ROF
132-
// mTimeFrameGPU->getDeviceROframesClusters(1), // const int* sizeClustersL1, // Number of clusters on layer 1 per ROF
133-
// mTimeFrameGPU->getChunk(chunkId).getDeviceTracklets(0), // Tracklet* tracklets01, // Tracklets on layer 0-1
134-
// mTimeFrameGPU->getChunk(chunkId).getDeviceTracklets(1), // Tracklet* tracklets12, // Tracklets on layer 1-2
135-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNTrackletCluster(0), // const int* nFoundTracklets01, // Number of tracklets found on layers 0-1
136-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNTrackletCluster(1), // const int* nFoundTracklet12, // Number of tracklets found on layers 1-2
137-
// mTimeFrameGPU->getChunk(chunkId).getDeviceUsedTracklets(), // unsigned char* usedTracklets, // Used tracklets
138-
// mTimeFrameGPU->getChunk(chunkId).getDeviceLines(), // Line* lines, // Lines
139-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNFoundLines(), // int* nFoundLines, // Number of found lines
140-
// mTimeFrameGPU->getChunk(chunkId).getDeviceNExclusiveFoundLines(), // int* nExclusiveFoundLines, // Number of found lines exclusive scan
141-
// offset, // const unsigned int startRofId, // Starting ROF ID
142-
// rofs, // const unsigned int rofSize, // Number of ROFs to consider
143-
// mVrtParams.maxTrackletsPerCluster, // const int maxTrackletsPerCluster = 1e2, // Maximum number of tracklets per cluster
144-
// mVrtParams.tanLambdaCut, // const float tanLambdaCut = 0.025f, // Cut on tan lambda
145-
// mVrtParams.phiCut); // const float phiCut = 0.002f) // Cut on phi
146-
147-
// int nClusters = mTimeFrameGPU->getTotalClustersPerROFrange(offset, rofs, 1);
148-
// int lastFoundLines;
149-
// std::vector<int> exclusiveFoundLinesHost(nClusters + 1);
150-
151-
// // Obtain whole exclusive sum including nCluster+1 element (nCluster+1)th element is the total number of found lines.
152-
// checkGPUError(cudaMemcpyAsync(exclusiveFoundLinesHost.data(), mTimeFrameGPU->getChunk(chunkId).getDeviceNExclusiveFoundLines(), (nClusters) * sizeof(int), cudaMemcpyDeviceToHost, mTimeFrameGPU->getStream(chunkId).get()));
153-
// checkGPUError(cudaMemcpyAsync(&lastFoundLines, mTimeFrameGPU->getChunk(chunkId).getDeviceNFoundLines() + nClusters - 1, sizeof(int), cudaMemcpyDeviceToHost, mTimeFrameGPU->getStream(chunkId).get()));
154-
// exclusiveFoundLinesHost[nClusters] = exclusiveFoundLinesHost[nClusters - 1] + lastFoundLines;
155-
156-
// std::vector<Line> lines(exclusiveFoundLinesHost[nClusters]);
157-
158-
// checkGPUError(cudaMemcpyAsync(lines.data(), mTimeFrameGPU->getChunk(chunkId).getDeviceLines(), sizeof(Line) * lines.size(), cudaMemcpyDeviceToHost, mTimeFrameGPU->getStream(chunkId).get()));
159-
// checkGPUError(cudaStreamSynchronize(mTimeFrameGPU->getStream(chunkId).get()));
160-
161-
// // Compute vertices
162-
// std::vector<ClusterLines> clusterLines;
163-
// std::vector<bool> usedLines;
164-
// for (int rofId{0}; rofId < rofs; ++rofId) {
165-
// auto rof = offset + rofId;
166-
// auto clustersL1offsetRof = mTimeFrameGPU->getROframeClusters(1)[rof] - mTimeFrameGPU->getROframeClusters(1)[offset]; // starting cluster offset for this ROF
167-
// auto nClustersL1Rof = mTimeFrameGPU->getROframeClusters(1)[rof + 1] - mTimeFrameGPU->getROframeClusters(1)[rof]; // number of clusters for this ROF
168-
// auto linesOffsetRof = exclusiveFoundLinesHost[clustersL1offsetRof]; // starting line offset for this ROF
169-
// auto nLinesRof = exclusiveFoundLinesHost[clustersL1offsetRof + nClustersL1Rof] - linesOffsetRof;
170-
// gsl::span<const o2::its::Line> linesInRof(lines.data() + linesOffsetRof, static_cast<gsl::span<o2::its::Line>::size_type>(nLinesRof));
171-
172-
// usedLines.resize(linesInRof.size(), false);
173-
// usedLines.assign(linesInRof.size(), false);
174-
// clusterLines.clear();
175-
// clusterLines.reserve(nClustersL1Rof);
176-
// computeVerticesInRof(rof,
177-
// linesInRof,
178-
// usedLines,
179-
// clusterLines,
180-
// mTimeFrameGPU->getBeamXY(),
181-
// mTimeFrameGPU->getVerticesInChunks()[chunkId],
182-
// mTimeFrameGPU->getNVerticesInChunks()[chunkId],
183-
// mTimeFrameGPU,
184-
// mTimeFrameGPU->hasMCinformation() ? &mTimeFrameGPU->getLabelsInChunks()[chunkId] : nullptr);
185-
// }
186-
// offset += rofs;
187-
// }
188-
// };
189-
// // Do work
190-
// threads[chunkId] = std::thread(doVertexReconstruction);
191-
// }
192-
// for (auto& thread : threads) {
193-
// thread.join();
194-
// }
195-
// for (int chunkId{0}; chunkId < mTimeFrameGPU->getNChunks(); ++chunkId) {
196-
// int start{0};
197-
// for (int rofId{0}; rofId < mTimeFrameGPU->getNVerticesInChunks()[chunkId].size(); ++rofId) {
198-
// gsl::span<const Vertex> rofVerts{mTimeFrameGPU->getVerticesInChunks()[chunkId].data() + start, static_cast<gsl::span<Vertex>::size_type>(mTimeFrameGPU->getNVerticesInChunks()[chunkId][rofId])};
199-
// mTimeFrameGPU->addPrimaryVertices(rofVerts);
200-
// if (mTimeFrameGPU->hasMCinformation()) {
201-
// mTimeFrameGPU->getVerticesLabels().emplace_back();
202-
// // TODO: add MC labels
203-
// }
204-
// start += mTimeFrameGPU->getNVerticesInChunks()[chunkId][rofId];
205-
// }
206-
// }
207-
// mTimeFrameGPU->wipe(3);
208-
// }
45+
for (short pivotRofId{0}; pivotRofId < mTimeFrameGPU->getNrof(); ++pivotRofId) {
46+
bool skipROF = iteration && (int)mTimeFrame->getPrimaryVertices(pivotRofId).size() > mVrtParams[iteration].vertPerRofThreshold;
47+
short startROF{std::max((short)0, static_cast<short>(pivotRofId - mVrtParams[iteration].deltaRof))};
48+
short endROF{std::min(static_cast<short>(mTimeFrame->getNrof()), static_cast<short>(pivotRofId + mVrtParams[iteration].deltaRof + 1))};
49+
}
50+
// gpu::computeTrackletsInROFsHandler
20951
}
21052

21153
void VertexerTraitsGPU::computeTrackletMatching(const int iteration)
@@ -216,12 +58,4 @@ void VertexerTraitsGPU::computeVertices(const int iteration)
21658
{
21759
}
21860

219-
void VertexerTraitsGPU::computeVerticesHist()
220-
{
221-
}
222-
223-
VertexerTraits* createVertexerTraitsGPU()
224-
{
225-
return new VertexerTraitsGPU;
226-
}
22761
} // namespace o2::its

0 commit comments

Comments
 (0)