Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly

Appearance settings
Original file line number Diff line number Diff line change
Expand Up @@ -141,7 +141,7 @@ class Vertex : public VertexBase
GPUd() ushort getFlags() const { return mBits; }
GPUd() bool isFlagSet(uint f) const { return mBits & (FlagsMask & f); }
GPUd() void setFlags(ushort f) { mBits |= FlagsMask & f; }
GPUd() void resetFrags(ushort f = FlagsMask) { mBits &= ~(FlagsMask & f); }
GPUd() void resetFlags(ushort f = FlagsMask) { mBits &= ~(FlagsMask & f); }

GPUd() void setChi2(float v) { mChi2 = v; }
GPUd() float getChi2() const { return mChi2; }
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -42,11 +42,11 @@ class VertexerTraitsGPU : public VertexerTraits
public:
VertexerTraitsGPU();
~VertexerTraitsGPU() override;
void initialise(const TrackingParameters&) override;
void initialise(const TrackingParameters&, const int iteration = 0) override;
void adoptTimeFrame(TimeFrame*) override;
void computeTracklets() override;
void computeTrackletMatching() override;
void computeVertices() override;
void computeTracklets(const int iteration = 0) override;
void computeTrackletMatching(const int iteration = 0) override;
void computeVertices(const int iteration = 0) override;

// Hybrid
void initialiseHybrid(const TrackingParameters& pars) override { VertexerTraits::initialise(pars); }
Expand All @@ -55,7 +55,7 @@ class VertexerTraitsGPU : public VertexerTraits
void computeTrackletMatchingHybrid() override { VertexerTraits::computeTrackletMatching(); }
void computeVerticesHybrid() override { VertexerTraits::computeVertices(); }

void updateVertexingParameters(const VertexingParameters&, const TimeFrameGPUParameters&) override;
void updateVertexingParameters(const std::vector<VertexingParameters>&, const TimeFrameGPUParameters&) override;

void computeVerticesHist();

Expand Down
44 changes: 23 additions & 21 deletions 44 Detectors/ITSMFT/ITS/tracking/GPU/cuda/VertexerTraitsGPU.cu
Original file line number Diff line number Diff line change
Expand Up @@ -95,7 +95,7 @@ VertexerTraitsGPU::~VertexerTraitsGPU()
{
}

void VertexerTraitsGPU::initialise(const TrackingParameters& trackingParams)
void VertexerTraitsGPU::initialise(const TrackingParameters& trackingParams, const int iteration)
{
mTimeFrameGPU->initialise(0, trackingParams, 3, &mIndexTableUtils, &mTfGPUParams);
}
Expand Down Expand Up @@ -595,17 +595,19 @@ GPUg() void computeVertexKernel(
}
} // namespace gpu

void VertexerTraitsGPU::updateVertexingParameters(const VertexingParameters& vrtPar, const TimeFrameGPUParameters& tfPar)
void VertexerTraitsGPU::updateVertexingParameters(const std::vector<VertexingParameters>& vrtPar, const TimeFrameGPUParameters& tfPar)
{
mVrtParams = vrtPar;
mTfGPUParams = tfPar;
mIndexTableUtils.setTrackingParameters(vrtPar);
mVrtParams.phiSpan = static_cast<int>(std::ceil(mIndexTableUtils.getNphiBins() * mVrtParams.phiCut /
constants::math::TwoPi));
mVrtParams.zSpan = static_cast<int>(std::ceil(mVrtParams.zCut * mIndexTableUtils.getInverseZCoordinate(0)));
mVrtParams = vrtPar;
mIndexTableUtils.setTrackingParameters(vrtPar[0]);
for (auto& par : mVrtParams) {
par.phiSpan = static_cast<int>(std::ceil(mIndexTableUtils.getNphiBins() * par.phiCut / constants::math::TwoPi));
par.zSpan = static_cast<int>(std::ceil(par.zCut * mIndexTableUtils.getInverseZCoordinate(0)));
}
}

void VertexerTraitsGPU::computeTracklets()
void VertexerTraitsGPU::computeTracklets(const int iteration)
{
if (!mTimeFrameGPU->getClusters().size()) {
return;
Expand Down Expand Up @@ -635,8 +637,8 @@ void VertexerTraitsGPU::computeTracklets()
mTimeFrameGPU->getDeviceIndexTableUtils(), // const IndexTableUtils* utils,
offset, // const unsigned int startRofId,
rofs, // const unsigned int rofSize,
mVrtParams.phiCut, // const float phiCut,
mVrtParams.maxTrackletsPerCluster); // const size_t maxTrackletsPerCluster = 1e2
mVrtParams[iteration].phiCut, // const float phiCut,
mVrtParams[iteration].maxTrackletsPerCluster); // const size_t maxTrackletsPerCluster = 1e2

gpu::trackleterKernelMultipleRof<TrackletMode::Layer1Layer2><<<rofs, 1024, 0, mTimeFrameGPU->getStream(chunkId).get()>>>(
mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(2), // const Cluster* clustersNextLayer, // 0 2
Expand All @@ -649,8 +651,8 @@ void VertexerTraitsGPU::computeTracklets()
mTimeFrameGPU->getDeviceIndexTableUtils(), // const IndexTableUtils* utils,
offset, // const unsigned int startRofId,
rofs, // const unsigned int rofSize,
mVrtParams.phiCut, // const float phiCut,
mVrtParams.maxTrackletsPerCluster); // const size_t maxTrackletsPerCluster = 1e2
mVrtParams[iteration].phiCut, // const float phiCut,
mVrtParams[iteration].maxTrackletsPerCluster); // const size_t maxTrackletsPerCluster = 1e2

gpu::trackletSelectionKernelMultipleRof<true><<<rofs, 1024, 0, mTimeFrameGPU->getStream(chunkId).get()>>>(
mTimeFrameGPU->getChunk(chunkId).getDeviceClusters(0), // const Cluster* clusters0, // Clusters on layer 0
Expand All @@ -667,9 +669,9 @@ void VertexerTraitsGPU::computeTracklets()
mTimeFrameGPU->getChunk(chunkId).getDeviceNExclusiveFoundLines(), // int* nExclusiveFoundLines, // Number of found lines exclusive scan
offset, // const unsigned int startRofId, // Starting ROF ID
rofs, // const unsigned int rofSize, // Number of ROFs to consider
mVrtParams.maxTrackletsPerCluster, // const int maxTrackletsPerCluster = 1e2, // Maximum number of tracklets per cluster
mVrtParams.tanLambdaCut, // const float tanLambdaCut = 0.025f, // Cut on tan lambda
mVrtParams.phiCut); // const float phiCut = 0.002f) // Cut on phi
mVrtParams[iteration].maxTrackletsPerCluster, // const int maxTrackletsPerCluster = 1e2, // Maximum number of tracklets per cluster
mVrtParams[iteration].tanLambdaCut, // const float tanLambdaCut = 0.025f, // Cut on tan lambda
mVrtParams[iteration].phiCut); // const float phiCut = 0.002f) // Cut on phi

discardResult(cub::DeviceScan::ExclusiveSum(mTimeFrameGPU->getChunk(chunkId).getDeviceCUBTmpBuffer(),
mTimeFrameGPU->getChunk(chunkId).getTimeFrameGPUParameters()->tmpCUBBufferSize,
Expand All @@ -681,7 +683,7 @@ void VertexerTraitsGPU::computeTracklets()
// Reset used tracklets
checkGPUError(cudaMemsetAsync(mTimeFrameGPU->getChunk(chunkId).getDeviceUsedTracklets(),
false,
sizeof(unsigned char) * mVrtParams.maxTrackletsPerCluster * mTimeFrameGPU->getTotalClustersPerROFrange(offset, rofs, 1),
sizeof(unsigned char) * mVrtParams[iteration].maxTrackletsPerCluster * mTimeFrameGPU->getTotalClustersPerROFrange(offset, rofs, 1),
mTimeFrameGPU->getStream(chunkId).get()),
__FILE__, __LINE__);

Expand All @@ -700,9 +702,9 @@ void VertexerTraitsGPU::computeTracklets()
mTimeFrameGPU->getChunk(chunkId).getDeviceNExclusiveFoundLines(), // int* nExclusiveFoundLines, // Number of found lines exclusive scan
offset, // const unsigned int startRofId, // Starting ROF ID
rofs, // const unsigned int rofSize, // Number of ROFs to consider
mVrtParams.maxTrackletsPerCluster, // const int maxTrackletsPerCluster = 1e2, // Maximum number of tracklets per cluster
mVrtParams.tanLambdaCut, // const float tanLambdaCut = 0.025f, // Cut on tan lambda
mVrtParams.phiCut); // const float phiCut = 0.002f) // Cut on phi
mVrtParams[iteration].maxTrackletsPerCluster, // const int maxTrackletsPerCluster = 1e2, // Maximum number of tracklets per cluster
mVrtParams[iteration].tanLambdaCut, // const float tanLambdaCut = 0.025f, // Cut on tan lambda
mVrtParams[iteration].phiCut); // const float phiCut = 0.002f) // Cut on phi

int nClusters = mTimeFrameGPU->getTotalClustersPerROFrange(offset, rofs, 1);
int lastFoundLines;
Expand Down Expand Up @@ -758,7 +760,7 @@ void VertexerTraitsGPU::computeTracklets()
gsl::span<const Vertex> rofVerts{mTimeFrameGPU->getVerticesInChunks()[chunkId].data() + start, static_cast<gsl::span<Vertex>::size_type>(mTimeFrameGPU->getNVerticesInChunks()[chunkId][rofId])};
mTimeFrameGPU->addPrimaryVertices(rofVerts);
if (mTimeFrameGPU->hasMCinformation()) {
mTimeFrameGPU->getVerticesLabels().emplace_back();
// mTimeFrameGPU->getVerticesLabels().emplace_back();
// TODO: add MC labels
}
start += mTimeFrameGPU->getNVerticesInChunks()[chunkId][rofId];
Expand All @@ -767,11 +769,11 @@ void VertexerTraitsGPU::computeTracklets()
mTimeFrameGPU->wipe(3);
}

void VertexerTraitsGPU::computeTrackletMatching()
void VertexerTraitsGPU::computeTrackletMatching(const int iteration)
{
}

void VertexerTraitsGPU::computeVertices()
void VertexerTraitsGPU::computeVertices(const int iteration)
{
}

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -98,6 +98,7 @@ struct TrackingParameters {
bool FindShortTracks = false;
bool PerPrimaryVertexProcessing = false;
bool SaveTimeBenchmarks = false;
bool DoUPCIteration = false;
};

inline int TrackingParameters::CellMinimumLevel()
Expand All @@ -106,6 +107,8 @@ inline int TrackingParameters::CellMinimumLevel()
}

struct VertexingParameters {
int nIterations = 1; // Number of vertexing passes to perform
int vertPerRofThreshold = 0; // Maximum number of vertices per ROF to trigger second a round
bool allowSingleContribClusters = false;
std::vector<float> LayerZ = {16.333f + 1, 16.333f + 1, 16.333f + 1, 42.140f + 1, 42.140f + 1, 73.745f + 1, 73.745f + 1};
std::vector<float> LayerRadii = {2.33959f, 3.14076f, 3.91924f, 19.6213f, 24.5597f, 34.388f, 39.3329f};
Expand Down
30 changes: 23 additions & 7 deletions 30 Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h
Original file line number Diff line number Diff line change
Expand Up @@ -80,13 +80,16 @@ class TimeFrame
const Vertex& getPrimaryVertex(const int) const;
gsl::span<const Vertex> getPrimaryVertices(int tf) const;
gsl::span<const Vertex> getPrimaryVertices(int romin, int romax) const;
gsl::span<const std::vector<MCCompLabel>> getPrimaryVerticesLabels(const int rof) const;
gsl::span<const std::pair<MCCompLabel, float>> getPrimaryVerticesMCRecInfo(const int rof) const;
gsl::span<const std::array<float, 2>> getPrimaryVerticesXAlpha(int tf) const;
void fillPrimaryVerticesXandAlpha();
int getPrimaryVerticesNum(int rofID = -1) const;
void addPrimaryVertices(const std::vector<Vertex>& vertices);
void addPrimaryVerticesLabels(std::vector<std::pair<MCCompLabel, float>>& labels);
void addPrimaryVertices(const gsl::span<const Vertex>& vertices);
void addPrimaryVertices(const std::vector<lightVertex>&);
void addPrimaryVerticesInROF(const std::vector<Vertex>& vertices, const int rofId);
void addPrimaryVerticesLabelsInROF(const std::vector<std::pair<MCCompLabel, float>>& labels, const int rofId);
void removePrimaryVerticesInROf(const int rofId);
int loadROFrameData(const o2::itsmft::ROFRecord& rof, gsl::span<const itsmft::Cluster> clusters,
const dataformats::MCTruthContainer<MCCompLabel>* mcLabels = nullptr);
Expand All @@ -98,6 +101,7 @@ class TimeFrame
const dataformats::MCTruthContainer<MCCompLabel>* mcLabels = nullptr);

int getTotalClusters() const;
std::vector<int>& getTotVertIteration() { return mTotVertPerIteration; }
bool empty() const;
bool isGPU() const { return mIsGPU; }
int getSortedIndex(int rof, int layer, int i) const;
Expand Down Expand Up @@ -143,7 +147,7 @@ class TimeFrame
std::vector<MCCompLabel>& getCellsLabel(int layer) { return mCellLabels[layer]; }

bool hasMCinformation() const;
void initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers = 7);
void initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers = 7, bool resetVertices = true);
void resetRofPV()
{
mPrimaryVertices.clear();
Expand All @@ -168,7 +172,7 @@ class TimeFrame
std::vector<TrackITSExt>& getTracks(int rof) { return mTracks[rof]; }
std::vector<MCCompLabel>& getTracksLabel(const int rof) { return mTracksLabel[rof]; }
std::vector<MCCompLabel>& getLinesLabel(const int rof) { return mLinesLabels[rof]; }
std::vector<std::vector<MCCompLabel>>& getVerticesLabels() { return mVerticesLabels; }
std::vector<std::pair<MCCompLabel, float>>& getVerticesMCRecInfo() { return mVerticesMCRecInfo; }

int getNumberOfClusters() const;
int getNumberOfCells() const;
Expand All @@ -195,6 +199,7 @@ class TimeFrame
uint32_t getTotalTrackletsTF(const int iLayer) { return mTotalTracklets[iLayer]; }
int getTotalClustersPerROFrange(int rofMin, int range, int layerId) const;
std::array<float, 2>& getBeamXY() { return mBeamPos; }
unsigned int& getNoVertexROF() { return mNoVertexROF; }
// \Vertexer

void initialiseRoadLabels();
Expand All @@ -203,6 +208,8 @@ class TimeFrame
bool isRoadFake(int i) const;

void setMultiplicityCutMask(const std::vector<bool>& cutMask) { mMultiplicityCutMask = cutMask; }
void setROFMask(const std::vector<bool>& rofMask) { mROFMask = rofMask; }
void swapMasks() { mMultiplicityCutMask.swap(mROFMask); }

int hasBogusClusters() const { return std::accumulate(mBogusClusters.begin(), mBogusClusters.end(), 0); }

Expand All @@ -220,7 +227,10 @@ class TimeFrame
}
}

virtual void setDevicePropagator(const o2::base::PropagatorImpl<float>*){};
virtual void setDevicePropagator(const o2::base::PropagatorImpl<float>*)
{
return;
};
const o2::base::PropagatorImpl<float>* getDevicePropagator() const { return mPropagatorDevice; }

template <typename... T>
Expand All @@ -233,6 +243,7 @@ class TimeFrame

void setExtAllocator(bool ext) { mExtAllocator = ext; }
bool getExtAllocator() const { return mExtAllocator; }

/// Debug and printing
void checkTrackletLUTs();
void printROFoffsets();
Expand Down Expand Up @@ -282,7 +293,9 @@ class TimeFrame
}

private:
void prepareClusters(const TrackingParameters& trkParam, const int maxLayers);
float mBz = 5.;
unsigned int mNTotalLowPtVertices = 0;
int mBeamPosWeight = 0;
std::array<float, 2> mBeamPos = {0.f, 0.f};
bool isBeamPositionOverridden = false;
Expand All @@ -293,6 +306,7 @@ class TimeFrame
std::vector<float> mPositionResolution;
std::vector<uint8_t> mClusterSize;
std::vector<bool> mMultiplicityCutMask;
std::vector<bool> mROFMask;
std::vector<std::array<float, 2>> mPValphaX; /// PV x and alpha for track propagation
std::vector<std::vector<MCCompLabel>> mTrackletLabels;
std::vector<std::vector<MCCompLabel>> mCellLabels;
Expand All @@ -311,8 +325,10 @@ class TimeFrame
std::vector<std::vector<ClusterLines>> mTrackletClusters;
std::vector<std::vector<int>> mTrackletsIndexROf;
std::vector<std::vector<MCCompLabel>> mLinesLabels;
std::vector<std::vector<MCCompLabel>> mVerticesLabels;
std::vector<std::pair<MCCompLabel, float>> mVerticesMCRecInfo;
std::array<uint32_t, 2> mTotalTracklets = {0, 0};
unsigned int mNoVertexROF = 0;
std::vector<int> mTotVertPerIteration;
// \Vertexer
};

Expand All @@ -326,12 +342,12 @@ inline gsl::span<const Vertex> TimeFrame::getPrimaryVertices(int rof) const
return {&mPrimaryVertices[start], static_cast<gsl::span<const Vertex>::size_type>(delta)};
}

inline gsl::span<const std::vector<MCCompLabel>> TimeFrame::getPrimaryVerticesLabels(const int rof) const
inline gsl::span<const std::pair<MCCompLabel, float>> TimeFrame::getPrimaryVerticesMCRecInfo(const int rof) const
{
const int start = mROframesPV[rof];
const int stop_idx = rof >= mNrof - 1 ? mNrof : rof + 1;
int delta = mMultiplicityCutMask[rof] ? mROframesPV[stop_idx] - start : 0; // return empty span if Rof is excluded
return {&mVerticesLabels[start], static_cast<gsl::span<const Vertex>::size_type>(delta)};
return {&(mVerticesMCRecInfo[start]), static_cast<gsl::span<const std::pair<MCCompLabel, float>>::size_type>(delta)};
}

inline gsl::span<const Vertex> TimeFrame::getPrimaryVertices(int romin, int romax) const
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,8 @@ namespace its

struct VertexerParamConfig : public o2::conf::ConfigurableParamHelper<VertexerParamConfig> {

int nIterations = 1; // Number of vertexing passes to perform
int vertPerRofThreshold = 0; // Maximum number of vertices per ROF to trigger second a round
bool allowSingleContribClusters = false;

// geometrical cuts
Expand Down Expand Up @@ -80,6 +82,7 @@ struct TrackerParamConfig : public o2::conf::ConfigurableParamHelper<TrackerPara
bool saveTimeBenchmarks = false;
bool overrideBeamEstimation = false; // used by gpuwf only
int trackingMode = -1; // -1: unset, 0=sync, 1=async, 2=cosmics used by gpuwf only
bool doUPCIteration = false;

O2ParamDef(TrackerParamConfig, "ITSCATrackerParam");
};
Expand Down
30 changes: 21 additions & 9 deletions 30 Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,8 @@ class Vertexer
Vertexer& operator=(const Vertexer&) = delete;

void adoptTimeFrame(TimeFrame& tf);
VertexingParameters& getVertParameters() const;
std::vector<VertexingParameters>& getVertParameters() const;
void setParameters(std::vector<VertexingParameters>& vertParams);
void getGlobalConfiguration();

std::vector<Vertex> exportVertices();
Expand All @@ -62,16 +63,18 @@ class Vertexer
float clustersToVertices(std::function<void(std::string s)> = [](std::string s) { std::cout << s << std::endl; });
float clustersToVerticesHybrid(std::function<void(std::string s)> = [](std::string s) { std::cout << s << std::endl; });
void filterMCTracklets();
void validateTracklets();
void validateTrackletsHybrid();

template <typename... T>
void findTracklets(T&&... args);
template <typename... T>
void findTrackletsHybrid(T&&... args);

void findTrivialMCTracklets();
void findVertices();
template <typename... T>
void validateTracklets(T&&... args);
void validateTrackletsHybrid();
template <typename... T>
void findVertices(T&&... args);
void findVerticesHybrid();
void findHistVertices();

Expand All @@ -95,6 +98,8 @@ class Vertexer

VertexerTraits* mTraits = nullptr; /// Observer pointer, not owned by this class
TimeFrame* mTimeFrame = nullptr; /// Observer pointer, not owned by this class

std::vector<VertexingParameters> mVertParams;
};

template <typename... T>
Expand All @@ -109,24 +114,31 @@ void Vertexer::findTracklets(T&&... args)
mTraits->computeTracklets(std::forward<T>(args)...);
}

inline VertexingParameters& Vertexer::getVertParameters() const
inline std::vector<VertexingParameters>& Vertexer::getVertParameters() const
{
return mTraits->getVertexingParameters();
}

inline void Vertexer::setParameters(std::vector<VertexingParameters>& vertParams)
{
mVertParams = vertParams;
}

inline void Vertexer::dumpTraits()
{
mTraits->dumpVertexerTraits();
}

inline void Vertexer::validateTracklets()
template <typename... T>
inline void Vertexer::validateTracklets(T&&... args)
{
mTraits->computeTrackletMatching();
mTraits->computeTrackletMatching(std::forward<T>(args)...);
}

inline void Vertexer::findVertices()
template <typename... T>
inline void Vertexer::findVertices(T&&... args)
{
mTraits->computeVertices();
mTraits->computeVertices(std::forward<T>(args)...);
}

template <typename... T>
Expand Down
Loading
Morty Proxy This is a proxified and sanitized view of the page, visit original site.