Skip to content

Commit 18abb6a

Browse files
mconcasdavidrohr
mconcas
authored and
davidrohr
committed
Finish vertexer code porting
1 parent 231a552 commit 18abb6a

10 files changed

Lines changed: 112 additions & 67 deletions

File tree

‎Common/MathUtils/include/MathUtils/CartesianGPU.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -26,8 +26,8 @@ template <typename T, int I>
2626
struct GPUPoint2D {
2727
GPUdDefault() GPUPoint2D() = default;
2828
GPUd() GPUPoint2D(T a, T b) : xx(a), yy(b) {}
29-
GPUd() float X() const { return xx; }
30-
GPUd() float Y() const { return yy; }
29+
GPUhd() float X() const { return xx; }
30+
GPUhd() float Y() const { return yy; }
3131
GPUd() float R() const { return o2::gpu::CAMath::Sqrt(xx * xx + yy * yy); }
3232
GPUd() void SetX(float v) { xx = v; }
3333
GPUd() void SetY(float v) { yy = v; }
@@ -39,7 +39,7 @@ template <typename T, int I>
3939
struct GPUPoint3D : public GPUPoint2D<T, I> {
4040
GPUdDefault() GPUPoint3D() = default;
4141
GPUd() GPUPoint3D(T a, T b, T c) : GPUPoint2D<T, I>(a, b), zz(c) {}
42-
GPUd() float Z() const { return zz; }
42+
GPUhd() float Z() const { return zz; }
4343
GPUd() float R() const { return o2::gpu::CAMath::Sqrt(GPUPoint2D<T, I>::xx * GPUPoint2D<T, I>::xx + GPUPoint2D<T, I>::yy * GPUPoint2D<T, I>::yy + zz * zz); }
4444
GPUd() void SetZ(float v) { zz = v; }
4545
T zz;

‎DataFormats/Reconstruction/include/ReconstructionDataFormats/Vertex.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -52,9 +52,9 @@ class VertexBase
5252
#endif
5353

5454
// getting the cartesian coordinates and errors
55-
GPUd() float getX() const { return mPos.X(); }
56-
GPUd() float getY() const { return mPos.Y(); }
57-
GPUd() float getZ() const { return mPos.Z(); }
55+
GPUhd() float getX() const { return mPos.X(); }
56+
GPUhd() float getY() const { return mPos.Y(); }
57+
GPUhd() float getZ() const { return mPos.Z(); }
5858
GPUd() float getSigmaX2() const { return mCov[kCovXX]; }
5959
GPUd() float getSigmaY2() const { return mCov[kCovYY]; }
6060
GPUd() float getSigmaZ2() const { return mCov[kCovZZ]; }

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,7 @@ struct GPUVertex final {
5151
float errorY;
5252
float errorZ;
5353
int contributors;
54+
int timeStamp;
5455
unsigned char realVertex;
5556
};
5657

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

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,8 @@
2020
#include "ITStracking/TimeFrame.h"
2121
#include "ITStracking/Configuration.h"
2222

23+
#include "ITStrackingGPU/ClusterLinesGPU.h"
24+
2325
#include "Array.h"
2426
#include "Vector.h"
2527

@@ -73,6 +75,7 @@ class TimeFrameGPU : public TimeFrame
7375
int* getDeviceZHistograms(const int rofId);
7476
cub::KeyValuePair<int, int>* getTmpVertexPositionBins(const int rofId);
7577
float* getDeviceBeamPosition(const int rofId);
78+
Vertex* getDeviceVertices(const int rofId);
7679

7780
private:
7881
TimeFrameGPUConfig mConfig;
@@ -99,6 +102,7 @@ class TimeFrameGPU : public TimeFrame
99102
std::array<Vector<int>, 3> mXYZHistograms;
100103
Vector<cub::KeyValuePair<int, int>> mTmpVertexPositionBins;
101104
Vector<float> mBeamPosition;
105+
Vector<Vertex> mGPUVertices;
102106
};
103107

104108
template <int NLayers>
@@ -250,6 +254,16 @@ inline float* TimeFrameGPU<NLayers>::getDeviceBeamPosition(const int rofId)
250254
return mBeamPosition.get() + 2 * rofId;
251255
}
252256

257+
template <int NLayers>
258+
inline Vertex* TimeFrameGPU<NLayers>::getDeviceVertices(const int rofId)
259+
{
260+
if (rofId < 0 || rofId >= mNrof) {
261+
LOG(error) << "Invalid rofId: " << rofId << "/" << mNrof << ", returning nullptr";
262+
return nullptr;
263+
}
264+
return mGPUVertices.get() + rofId * mConfig.maxVerticesCapacity;
265+
}
266+
253267
} // namespace gpu
254268
} // namespace its
255269
} // namespace o2

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -70,6 +70,7 @@ TimeFrameGPU<NLayers>::TimeFrameGPU()
7070
}
7171
mTmpVertexPositionBins = Vector<cub::KeyValuePair<int, int>>{3 * mConfig.nMaxROFs, 3 * mConfig.nMaxROFs};
7272
mBeamPosition = Vector<float>{2 * mConfig.nMaxROFs, 2 * mConfig.nMaxROFs};
73+
mGPUVertices = Vector<Vertex>{mConfig.nMaxROFs * mConfig.maxVerticesCapacity, mConfig.nMaxROFs * mConfig.maxVerticesCapacity};
7374
}
7475

7576
template <int NLayers>
@@ -95,6 +96,7 @@ float TimeFrameGPU<NLayers>::getDeviceMemory()
9596
}
9697
totalMemory += 3 * mConfig.nMaxROFs * sizeof(cub::KeyValuePair<int, int>);
9798
totalMemory += 2 * mConfig.nMaxROFs * sizeof(float);
99+
totalMemory += mConfig.nMaxROFs * mConfig.maxVerticesCapacity * sizeof(Vertex);
98100

99101
LOGP(info, "Total requested memory for GPU: {:.2f} MB", totalMemory / MB);
100102
LOGP(info, "\t- Clusters: {:.2f} MB", NLayers * mConfig.clustersPerLayerCapacity * sizeof(Cluster) / MB);
@@ -115,6 +117,7 @@ float TimeFrameGPU<NLayers>::getDeviceMemory()
115117
LOGP(info, "\t- Z histograms: {:.2f} MB", mConfig.nMaxROFs * mConfig.histConf.nBinsXYZ[2] * sizeof(int) / MB);
116118
LOGP(info, "\t- TMP Vertex position bins: {:.2f} MB", 3 * mConfig.nMaxROFs * sizeof(cub::KeyValuePair<int, int>) / MB);
117119
LOGP(info, "\t- Beam positions: {:.2f} MB", 2 * mConfig.nMaxROFs * sizeof(float) / MB);
120+
LOGP(info, "\t- Vertices: {:.2f} MB", mConfig.nMaxROFs * mConfig.maxVerticesCapacity * sizeof(Vertex) / MB);
118121

119122
return totalMemory;
120123
}

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

Lines changed: 71 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -30,11 +30,12 @@
3030
#include "ITStrackingGPU/ClusterLinesGPU.h"
3131
#include "ITStrackingGPU/VertexerTraitsGPU.h"
3232

33+
#include "GPUCommonArray.h"
34+
3335
namespace o2
3436
{
3537
namespace its
3638
{
37-
3839
using constants::its::VertexerHistogramVolume;
3940
using constants::math::TwoPi;
4041
using gpu::utils::host::checkGPUError;
@@ -119,9 +120,9 @@ GPUd() void printVectorOnThread(const char* name, Vector<int>& vector, size_t si
119120
// {
120121
// if (blockIdx.x * blockDim.x + threadIdx.x == threadId) {
121122
// printf("XmaxBin: %d at index: %d | YmaxBin: %d at index: %d | ZmaxBin: %d at index: %d\n",
122-
// store.getTmpVertexPositionBins()[0].value, store.getTmpVertexPositionBins()[0].key,
123-
// store.getTmpVertexPositionBins()[1].value, store.getTmpVertexPositionBins()[1].key,
124-
// store.getTmpVertexPositionBins()[2].value, store.getTmpVertexPositionBins()[2].key);
123+
// tmpVertexBins[0].value, tmpVertexBins[0].key,
124+
// tmpVertexBins[1].value, tmpVertexBins[1].key,
125+
// tmpVertexBins[2].value, tmpVertexBins[2].key);
125126
// }
126127
// }
127128

@@ -313,35 +314,44 @@ GPUg() void computeZCentroidsKernel(const int nLines,
313314
}
314315
}
315316

316-
// GPUg() void computeVertexKernel(DeviceStoreVertexerGPU& store, const int vertIndex, const int minContributors, const int binOpeningZ)
317-
// {
318-
// for (size_t currentThreadIndex = blockIdx.x * blockDim.x + threadIdx.x; currentThreadIndex < binOpeningZ; currentThreadIndex += blockDim.x * gridDim.x) {
319-
// if (currentThreadIndex == 0) {
320-
// if (store.getTmpVertexPositionBins()[2].value > 1 && (store.getTmpVertexPositionBins()[0].value || store.getTmpVertexPositionBins()[1].value)) {
321-
// float z{store.getConfig().histConf.lowHistBoundariesXYZ[2] + store.getTmpVertexPositionBins()[2].key * store.getConfig().histConf.binSizeHistZ + store.getConfig().histConf.binSizeHistZ / 2};
322-
// float ex{0.f};
323-
// float ey{0.f};
324-
// float ez{0.f};
325-
// int sumWZ{store.getTmpVertexPositionBins()[2].value};
326-
// float wZ{z * store.getTmpVertexPositionBins()[2].value};
327-
// for (int iBin{o2::gpu::GPUCommonMath::Max(0, store.getTmpVertexPositionBins()[2].key - binOpeningZ)}; iBin < o2::gpu::GPUCommonMath::Min(store.getTmpVertexPositionBins()[2].key + binOpeningZ + 1, store.getConfig().histConf.nBinsXYZ[2] - 1); ++iBin) {
328-
// if (iBin != store.getTmpVertexPositionBins()[2].key) {
329-
// wZ += (store.getConfig().histConf.lowHistBoundariesXYZ[2] + iBin * store.getConfig().histConf.binSizeHistZ + store.getConfig().histConf.binSizeHistZ / 2) * store.getHistogramXYZ()[2].get()[iBin];
330-
// sumWZ += store.getHistogramXYZ()[2].get()[iBin];
331-
// }
332-
// store.getHistogramXYZ()[2].get()[iBin] = 0;
333-
// }
334-
// if (sumWZ > minContributors || vertIndex == 0) {
335-
// store.getVertices().emplace(vertIndex, store.getBeamPosition()[0], store.getBeamPosition()[1], wZ / sumWZ, ex, ey, ez, sumWZ);
336-
// } else {
337-
// store.getVertices().emplace(vertIndex);
338-
// }
339-
// } else {
340-
// store.getVertices().emplace(vertIndex);
341-
// }
342-
// }
343-
// }
344-
// }
317+
GPUg() void computeVertexKernel(cub::KeyValuePair<int, int>* tmpVertexBins,
318+
int* histZ, // Z
319+
const float lowHistZ,
320+
const float binSizeHistZ,
321+
const int nBinsHistZ,
322+
Vertex* vertices,
323+
float* beamPosition,
324+
const int vertIndex,
325+
const int minContributors,
326+
const int binOpeningZ)
327+
{
328+
for (size_t currentThreadIndex = blockIdx.x * blockDim.x + threadIdx.x; currentThreadIndex < binOpeningZ; currentThreadIndex += blockDim.x * gridDim.x) {
329+
if (currentThreadIndex == 0) {
330+
if (tmpVertexBins[2].value > 1 && (tmpVertexBins[0].value || tmpVertexBins[1].value)) {
331+
float z{lowHistZ + tmpVertexBins[2].key * binSizeHistZ + binSizeHistZ / 2};
332+
float ex{0.f};
333+
float ey{0.f};
334+
float ez{0.f};
335+
int sumWZ{tmpVertexBins[2].value};
336+
float wZ{z * tmpVertexBins[2].value};
337+
for (int iBin{o2::gpu::GPUCommonMath::Max(0, tmpVertexBins[2].key - binOpeningZ)}; iBin < o2::gpu::GPUCommonMath::Min(tmpVertexBins[2].key + binOpeningZ + 1, nBinsHistZ - 1); ++iBin) {
338+
if (iBin != tmpVertexBins[2].key) {
339+
wZ += (lowHistZ + iBin * binSizeHistZ + binSizeHistZ / 2) * histZ[iBin];
340+
sumWZ += histZ[iBin];
341+
}
342+
histZ[iBin] = 0;
343+
}
344+
if (sumWZ > minContributors || vertIndex == 0) {
345+
new (vertices + vertIndex) Vertex{o2::math_utils::Point3D<float>(beamPosition[0], beamPosition[1], wZ / sumWZ), o2::gpu::gpustd::array<float, 6>{ex, 0, ey, 0, 0, ez}, sumWZ, 0};
346+
} else {
347+
new (vertices + vertIndex) Vertex{};
348+
}
349+
} else {
350+
new (vertices + vertIndex) Vertex{};
351+
}
352+
}
353+
}
354+
}
345355
} // namespace gpu
346356

347357
void VertexerTraitsGPU::computeTracklets()
@@ -538,23 +548,34 @@ void VertexerTraitsGPU::computeVertices()
538548
mTimeFrameGPU->getConfig().histConf.lowHistBoundariesXYZ[2], // lower_level
539549
mTimeFrameGPU->getConfig().histConf.highHistBoundariesXYZ[2], // fupper_level
540550
nLines)); // num_row_pixels
541-
// for (int iVertex{0}; iVertex < mStoreVertexerGPU.getConfig().nMaxVertices; ++iVertex) {
542-
// discardResult(cub::DeviceReduce::ArgMax(reinterpret_cast<void*>(mStoreVertexerGPU.getCUBTmpBuffer().get()),
543-
// bufferSize,
544-
// mStoreVertexerGPU.getHistogramXYZ()[2].get(),
545-
// mStoreVertexerGPU.getTmpVertexPositionBins().get() + 2,
546-
// mStoreVertexerGPU.getConfig().histConf.nBinsXYZ[2]));
547-
// gpu::computeVertexKernel<<<blocksGrid, 5>>>(getDeviceContext(), iVertex, mVrtParams.clusterContributorsCut, mStoreVertexerGPU.getConfig().histConf.binSpanXYZ[2]);
548-
// }
549-
// std::vector<gpu::GPUVertex> vertices;
550-
// vertices.resize(mStoreVertexerGPU.getConfig().nMaxVertices);
551-
// mStoreVertexerGPU.getVertices().copyIntoSizedVector(vertices);
552-
553-
// for (auto& vertex : vertices) {
554-
// if (vertex.realVertex) {
555-
// mVertices.emplace_back(vertex.xCoord, vertex.yCoord, vertex.zCoord, std::array<float, 6>{0.f, 0.f, 0.f, 0.f, 0.f, 0.f}, vertex.contributors, 0.f, -9);
556-
// }
557-
// }
551+
for (int iVertex{0}; iVertex < mTimeFrameGPU->getConfig().maxVerticesCapacity; ++iVertex) {
552+
discardResult(cub::DeviceReduce::ArgMax(reinterpret_cast<void*>(mTimeFrameGPU->getDeviceCUBBuffer(rofId)),
553+
bufferSize,
554+
mTimeFrameGPU->getDeviceZHistograms(rofId),
555+
mTimeFrameGPU->getTmpVertexPositionBins(rofId) + 2,
556+
mTimeFrameGPU->getConfig().histConf.nBinsXYZ[2]));
557+
558+
gpu::computeVertexKernel<<<blocksGrid, 5>>>(mTimeFrameGPU->getTmpVertexPositionBins(rofId),
559+
mTimeFrameGPU->getDeviceZHistograms(rofId),
560+
mTimeFrameGPU->getConfig().histConf.lowHistBoundariesXYZ[2],
561+
mTimeFrameGPU->getConfig().histConf.binSizeHistZ,
562+
mTimeFrameGPU->getConfig().histConf.nBinsXYZ[2],
563+
mTimeFrameGPU->getDeviceVertices(rofId),
564+
mTimeFrameGPU->getDeviceBeamPosition(rofId),
565+
iVertex,
566+
mVrtParams.clusterContributorsCut,
567+
mTimeFrameGPU->getConfig().histConf.binSpanXYZ[2]);
568+
}
569+
std::vector<Vertex> GPUvertices;
570+
GPUvertices.resize(mTimeFrameGPU->getConfig().maxVerticesCapacity);
571+
checkGPUError(cudaMemcpy(GPUvertices.data(), mTimeFrameGPU->getDeviceVertices(rofId), sizeof(gpu::GPUVertex) * mTimeFrameGPU->getConfig().maxVerticesCapacity, cudaMemcpyDeviceToHost), __FILE__, __LINE__);
572+
int nRealVertices{0};
573+
for (auto& vertex : GPUvertices) {
574+
if (vertex.getX() || vertex.getY() || vertex.getZ()) {
575+
++nRealVertices;
576+
}
577+
}
578+
mTimeFrameGPU->addPrimaryVertices(gsl::span<const Vertex>{GPUvertices.data(), static_cast<gsl::span<const Vertex>::size_type>(nRealVertices)});
558579
}
559580
gpuThrowOnError();
560581
}

‎Detectors/ITSMFT/ITS/tracking/include/ITStracking/Configuration.h‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -149,7 +149,7 @@ struct TimeFrameGPUConfig {
149149
size_t cluLayCap,
150150
size_t cluROfCap,
151151
size_t maxTrkCap,
152-
size_t maxVert);
152+
size_t maxVertCap);
153153

154154
size_t tmpCUBBufferSize = 1e3; // In average in pp events there are required 767 bytes
155155
size_t maxTrackletsPerCluster = 50;
@@ -158,7 +158,7 @@ struct TimeFrameGPUConfig {
158158
size_t trackletsCapacity = maxTrackletsPerCluster * clustersPerLayerCapacity;
159159
size_t maxLinesCapacity = 1e2;
160160
size_t maxCentroidsXYCapacity = std::ceil(maxLinesCapacity * (maxLinesCapacity - 1) / (float)2);
161-
size_t nMaxVertices = 10;
161+
size_t maxVerticesCapacity = 10;
162162
size_t nMaxROFs = 1e3;
163163

164164
VertexerHistogramsConfiguration histConf; // <==== split into separate configs
@@ -169,12 +169,12 @@ inline TimeFrameGPUConfig::TimeFrameGPUConfig(size_t cubBufferSize,
169169
size_t cluLayCap,
170170
size_t cluROfCap,
171171
size_t maxTrkCap,
172-
size_t maxVert) : tmpCUBBufferSize{cubBufferSize},
172+
size_t maxVertCap) : tmpCUBBufferSize{cubBufferSize},
173173
maxTrackletsPerCluster{maxTrkClu},
174174
clustersPerLayerCapacity{cluLayCap},
175175
clustersPerROfCapacity{cluROfCap},
176176
maxLinesCapacity{maxTrkCap},
177-
nMaxVertices{maxVert}
177+
maxVerticesCapacity{maxVertCap}
178178
{
179179
maxCentroidsXYCapacity = std::ceil(maxLinesCapacity * (maxLinesCapacity - 1) / 2);
180180
trackletsCapacity = maxTrackletsPerCluster * clustersPerLayerCapacity;

‎Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -67,6 +67,7 @@ class TimeFrame
6767
gsl::span<const Vertex> getPrimaryVertices(int romin, int romax) const;
6868
int getPrimaryVerticesNum(int rofID = -1) const;
6969
void addPrimaryVertices(const std::vector<Vertex>& vertices);
70+
void addPrimaryVertices(const gsl::span<const Vertex>& vertices);
7071
void removePrimaryVerticesInROf(const int rofId);
7172
int loadROFrameData(const o2::itsmft::ROFRecord& rof, gsl::span<const itsmft::Cluster> clusters,
7273
const dataformats::MCTruthContainer<MCCompLabel>* mcLabels = nullptr);
@@ -226,7 +227,6 @@ inline gsl::span<const Vertex> TimeFrame::getPrimaryVertices(int rof) const
226227
const int start = mROframesPV[rof];
227228
const int stop_idx = rof >= mNrof - 1 ? mNrof : rof + 1;
228229
int delta = mMultiplicityCutMask[rof] ? mROframesPV[stop_idx] - start : 0; // return empty span if Rof is excluded
229-
LOGP(info, "delta: {}", delta);
230230
return {&mPrimaryVertices[start], static_cast<gsl::span<const Vertex>::size_type>(delta)};
231231
}
232232

‎Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h‎

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -35,14 +35,8 @@ namespace o2
3535
{
3636
class MCCompLabel;
3737

38-
namespace utils
39-
{
40-
class TreeStreamRedirector;
41-
}
42-
4338
namespace its
4439
{
45-
class StandaloneDebugger;
4640
class ROframe;
4741

4842
using constants::its::LayersNumberVertexer;

‎Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -81,6 +81,18 @@ void TimeFrame::addPrimaryVertices(const std::vector<Vertex>& vertices)
8181
mROframesPV.push_back(mPrimaryVertices.size());
8282
}
8383

84+
void TimeFrame::addPrimaryVertices(const gsl::span<const Vertex>& vertices)
85+
{
86+
for (const auto& vertex : vertices) {
87+
mPrimaryVertices.emplace_back(vertex);
88+
const int w{vertex.getNContributors()};
89+
mBeamPos[0] = (mBeamPos[0] * mBeamPosWeight + vertex.getX() * w) / (mBeamPosWeight + w);
90+
mBeamPos[1] = (mBeamPos[1] * mBeamPosWeight + vertex.getY() * w) / (mBeamPosWeight + w);
91+
mBeamPosWeight += w;
92+
}
93+
mROframesPV.push_back(mPrimaryVertices.size());
94+
}
95+
8496
int TimeFrame::loadROFrameData(const o2::itsmft::ROFRecord& rof, gsl::span<const itsmft::Cluster> clusters,
8597
const dataformats::MCTruthContainer<MCCompLabel>* mcLabels)
8698
{

0 commit comments

Comments
 (0)