Skip to content

Commit 4d37b74

Browse files
f3schf3sch
committed
more
1 parent 20c398d commit 4d37b74

File tree

8 files changed

+43
-72
lines changed

8 files changed

+43
-72
lines changed

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

Lines changed: 3 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,6 @@
1818
#define ITSTRACKINGGPU_VERTEXERTRAITSGPU_H_
1919

2020
#include <vector>
21-
#include <array>
2221

2322
#include "ITStracking/VertexerTraits.h"
2423
#include "ITStracking/Configuration.h"
@@ -29,13 +28,8 @@
2928

3029
#include "ITStrackingGPU/TimeFrameGPU.h"
3130

32-
namespace o2
31+
namespace o2::its
3332
{
34-
namespace its
35-
{
36-
class ROframe;
37-
38-
using constants::its2::InversePhiBinSize;
3933

4034
class VertexerTraitsGPU final : public VertexerTraits
4135
{
@@ -63,6 +57,6 @@ inline void VertexerTraitsGPU::adoptTimeFrame(TimeFrame<7>* tf) noexcept
6357
mTimeFrame = static_cast<TimeFrame<7>*>(tf);
6458
}
6559

66-
} // namespace its
67-
} // namespace o2
60+
} // namespace o2::its
61+
6862
#endif

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

Lines changed: 8 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -55,11 +55,6 @@ namespace o2::its
5555
{
5656
using Vertex = o2::dataformats::Vertex<o2::dataformats::TimeStamp<int>>;
5757

58-
GPUdii() float Sq(float v)
59-
{
60-
return v * v;
61-
}
62-
6358
namespace gpu
6459
{
6560

@@ -68,9 +63,9 @@ GPUd() const int4 getBinsRect(const Cluster& currentCluster, const int layerInde
6863
const float z1, const float z2, float maxdeltaz, float maxdeltaphi)
6964
{
7065
const float zRangeMin = o2::gpu::CAMath::Min(z1, z2) - maxdeltaz;
71-
const float phiRangeMin = (maxdeltaphi > constants::math::Pi) ? 0.f : currentCluster.phi - maxdeltaphi;
66+
const float phiRangeMin = (maxdeltaphi > o2::constants::math::PI) ? 0.f : currentCluster.phi - maxdeltaphi;
7267
const float zRangeMax = o2::gpu::CAMath::Max(z1, z2) + maxdeltaz;
73-
const float phiRangeMax = (maxdeltaphi > constants::math::Pi) ? constants::math::TwoPi : currentCluster.phi + maxdeltaphi;
68+
const float phiRangeMax = (maxdeltaphi > o2::constants::math::PI) ? o2::constants::math::TwoPI : currentCluster.phi + maxdeltaphi;
7469

7570
if (zRangeMax < -utils.getLayerZ(layerIndex) ||
7671
zRangeMin > utils.getLayerZ(layerIndex) || zRangeMin > zRangeMax) {
@@ -98,7 +93,7 @@ GPUd() bool fitTrack(TrackITSExt& track,
9893
o2::base::PropagatorF::MatCorrType matCorrType)
9994
{
10095
for (int iLayer{start}; iLayer != end; iLayer += step) {
101-
if (track.getClusterIndex(iLayer) == constants::its::UnusedIndex) {
96+
if (track.getClusterIndex(iLayer) == constants::UnusedIndex) {
10297
continue;
10398
}
10499
const TrackingFrameInfo& trackingHit = tfInfos[iLayer][track.getClusterIndex(iLayer)];
@@ -285,7 +280,7 @@ GPUg() void fitTrackSeedsKernel(
285280
temporaryTrack.setChi2(0);
286281
int* clusters = seed.getClusters();
287282
for (int iL{0}; iL < 7; ++iL) {
288-
temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::its::UnusedIndex);
283+
temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::UnusedIndex);
289284
}
290285
bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track,
291286
0, // int lastLayer,
@@ -391,8 +386,6 @@ GPUg() void computeLayerCellsKernel(
391386
const float cellDeltaTanLambdaSigma,
392387
const float nSigmaCut)
393388
{
394-
constexpr float constants::Radl = 9.36f; // Radiation length of Si [cm].
395-
constexpr float constants::Rho = 2.33f; // Density of Si [g/cm^3].
396389
constexpr float layerxX0[7] = {5.e-3f, 5.e-3f, 5.e-3f, 1.e-2f, 1.e-2f, 1.e-2f, 1.e-2f}; // Hardcoded here for the moment.
397390
for (int iCurrentTrackletIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentTrackletIndex < nTrackletsCurrent; iCurrentTrackletIndex += blockDim.x * gridDim.x) {
398391
const Tracklet& currentTracklet = tracklets[layer][iCurrentTrackletIndex];
@@ -517,12 +510,12 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
517510
if (primaryVertex.isFlagSet(2) && iteration != 3) {
518511
continue;
519512
}
520-
const float resolution = o2::gpu::CAMath::Sqrt(Sq(resolutionPV) / primaryVertex.getNContributors() + Sq(positionResolution));
513+
const float resolution = o2::gpu::CAMath::Sqrt(math_utils::Sq(resolutionPV) / primaryVertex.getNContributors() + math_utils::Sq(positionResolution));
521514
const float tanLambda{(currentCluster.zCoordinate - primaryVertex.getZ()) * inverseR0};
522515
const float zAtRmin{tanLambda * (minR - currentCluster.radius) + currentCluster.zCoordinate};
523516
const float zAtRmax{tanLambda * (maxR - currentCluster.radius) + currentCluster.zCoordinate};
524-
const float sqInverseDeltaZ0{1.f / (Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + 2.e-8f)}; /// protecting from overflows adding the detector resolution
525-
const float sigmaZ{o2::gpu::CAMath::Sqrt(Sq(resolution) * Sq(tanLambda) * ((Sq(inverseR0) + sqInverseDeltaZ0) * Sq(meanDeltaR) + 1.f) + Sq(meanDeltaR * MSAngle))};
517+
const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + 2.e-8f)}; /// protecting from overflows adding the detector resolution
518+
const float sigmaZ{o2::gpu::CAMath::Sqrt(math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInverseDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * MSAngle))};
526519
const int4 selectedBinsRect{getBinsRect(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)};
527520
if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) {
528521
continue;
@@ -556,7 +549,7 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
556549
const float deltaPhi{o2::gpu::CAMath::Abs(currentCluster.phi - nextCluster.phi)};
557550
const float deltaZ{o2::gpu::CAMath::Abs(tanLambda * (nextCluster.radius - currentCluster.radius) + currentCluster.zCoordinate - nextCluster.zCoordinate)};
558551
const int nextSortedIndex{ROFClusters[layerIndex + 1][rof1] + nextClusterIndex};
559-
if (deltaZ / sigmaZ < NSigmaCut && (deltaPhi < phiCut || o2::gpu::CAMath::Abs(deltaPhi - constants::math::TwoPi) < phiCut)) {
552+
if (deltaZ / sigmaZ < NSigmaCut && (deltaPhi < phiCut || o2::gpu::CAMath::Abs(deltaPhi - o2::constants::math::TwoPI) < phiCut)) {
560553
if constexpr (initRun) {
561554
trackletsLUT[layerIndex][currentSortedIndex]++; // we need l0 as well for usual exclusive sums.
562555
} else {

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,7 @@ void VertexerTraitsGPU::updateVertexingParameters(const std::vector<VertexingPar
4040
mTfGPUParams = tfPar;
4141
mIndexTableUtils.setTrackingParameters(vrtPar[0]);
4242
for (auto& par : mVrtParams) {
43-
par.phiSpan = static_cast<int>(std::ceil(mIndexTableUtils.getNphiBins() * par.phiCut / constants::math::TwoPi));
43+
par.phiSpan = static_cast<int>(std::ceil(mIndexTableUtils.getNphiBins() * par.phiCut / o2::constants::math::TwoPI));
4444
par.zSpan = static_cast<int>(std::ceil(par.zCut * mIndexTableUtils.getInverseZCoordinate(0)));
4545
}
4646
}

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

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -20,10 +20,7 @@ namespace o2
2020
{
2121
namespace its
2222
{
23-
using constants::its::VertexerHistogramVolume;
24-
using constants::math::TwoPi;
2523
using math_utils::getNormalizedPhi;
26-
using namespace constants::its2;
2724

2825
namespace gpu
2926
{
@@ -58,11 +55,6 @@ void trackletFinderHandler(const Cluster* clustersNextLayer, // 0 2
5855
maxTrackletsPerCluster); // const unsigned int maxTrackletsPerCluster = 1e2
5956
}
6057
/*
61-
GPUd() float smallestAngleDifference(float a, float b)
62-
{
63-
float diff = fmod(b - a + constants::math::Pi, constants::math::TwoPi) - constants::math::Pi;
64-
return (diff < -constants::math::Pi) ? diff + constants::math::TwoPi : ((diff > constants::math::Pi) ? diff - constants::math::TwoPi : diff);
65-
}
6658
6759
GPUd() const int4 getBinsRect(const Cluster& currentCluster, const int layerIndex,
6860
const float z1, float maxdeltaz, float maxdeltaphi)

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

Lines changed: 20 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -13,8 +13,8 @@
1313
/// \brief
1414
///
1515

16-
#ifndef TRACKINGITSU_INCLUDE_CAUTILS_H_
17-
#define TRACKINGITSU_INCLUDE_CAUTILS_H_
16+
#ifndef O2_ITS_TRACKING_MATHUTILS_H_
17+
#define O2_ITS_TRACKING_MATHUTILS_H_
1818

1919
#include "CommonConstants/MathConstants.h"
2020
#include "MathUtils/Utils.h"
@@ -75,6 +75,23 @@ GPUhdi() float computeTanDipAngle(float x1, float y1, float x2, float y2, float
7575
return (z1 - z2) / o2::gpu::CAMath::Sqrt((x1 - x2) * (x1 - x2) + (y1 - y2) * (y1 - y2));
7676
}
7777

78+
GPUhdi() float smallestAngleDifference(float a, float b)
79+
{
80+
float diff = fmod(b - a + o2::constants::math::PI, o2::constants::math::TwoPI) - o2::constants::math::PI;
81+
return (diff < -o2::constants::math::PI) ? diff + o2::constants::math::TwoPI : ((diff > o2::constants::math::PI) ? diff - o2::constants::math::TwoPI : diff);
82+
}
83+
84+
GPUhdi() float Sq(float v)
85+
{
86+
return v * v;
87+
}
88+
89+
GPUhdi() float MSangle(float mass, float p, float xX0)
90+
{
91+
float beta = p / o2::gpu::CAMath::Hypot(mass, p);
92+
return 0.0136f * o2::gpu::CAMath::Sqrt(xX0) * (1.f + 0.038f * o2::gpu::CAMath::Log(xX0)) / (beta * p);
93+
}
94+
7895
} // namespace o2::its::math_utils
7996

80-
#endif /* TRACKINGITSU_INCLUDE_CAUTILS_H_ */
97+
#endif

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

Lines changed: 5 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -34,18 +34,6 @@ struct ClusterHelper {
3434
int bin;
3535
int ind;
3636
};
37-
38-
inline float MSangle(float mass, float p, float xX0)
39-
{
40-
float beta = p / o2::gpu::CAMath::Hypot(mass, p);
41-
return 0.0136f * o2::gpu::CAMath::Sqrt(xX0) * (1.f + 0.038f * o2::gpu::CAMath::Log(xX0)) / (beta * p);
42-
}
43-
44-
inline float Sq(float v)
45-
{
46-
return v * v;
47-
}
48-
4937
} // namespace
5038

5139
namespace o2::its
@@ -323,7 +311,7 @@ void TimeFrame<nLayers>::initialise(const int iteration, const TrackingParameter
323311
for (unsigned int iLayer{0}; iLayer < std::min((int)mClusters.size(), maxLayers); ++iLayer) {
324312
clearResizeBoundedVector(mClusters[iLayer], mUnsortedClusters[iLayer].size(), mMemoryPool.get());
325313
clearResizeBoundedVector(mUsedClusters[iLayer], mUnsortedClusters[iLayer].size(), mMemoryPool.get());
326-
mPositionResolution[iLayer] = o2::gpu::CAMath::Sqrt(0.5 * (trkParam.SystErrorZ2[iLayer] + trkParam.SystErrorY2[iLayer]) + trkParam.LayerResolution[iLayer] * trkParam.LayerResolution[iLayer]);
314+
mPositionResolution[iLayer] = o2::gpu::CAMath::Sqrt(0.5f * (trkParam.SystErrorZ2[iLayer] + trkParam.SystErrorY2[iLayer]) + trkParam.LayerResolution[iLayer] * trkParam.LayerResolution[iLayer]);
327315
}
328316
clearResizeBoundedArray(mIndexTables, mNrof * (trkParam.ZBins * trkParam.PhiBins + 1), mMemoryPool.get());
329317
clearResizeBoundedVector(mLines, mNrof, mMemoryPool.get());
@@ -364,17 +352,17 @@ void TimeFrame<nLayers>::initialise(const int iteration, const TrackingParameter
364352

365353
float oneOverR{0.001f * 0.3f * std::abs(mBz) / trkParam.TrackletMinPt};
366354
for (unsigned int iLayer{0}; iLayer < mClusters.size(); ++iLayer) {
367-
mMSangles[iLayer] = MSangle(0.14f, trkParam.TrackletMinPt, trkParam.LayerxX0[iLayer]);
355+
mMSangles[iLayer] = math_utils::MSangle(0.14f, trkParam.TrackletMinPt, trkParam.LayerxX0[iLayer]);
368356
mPositionResolution[iLayer] = o2::gpu::CAMath::Sqrt(0.5f * (trkParam.SystErrorZ2[iLayer] + trkParam.SystErrorY2[iLayer]) + trkParam.LayerResolution[iLayer] * trkParam.LayerResolution[iLayer]);
369357
if (iLayer < mClusters.size() - 1) {
370358
const float& r1 = trkParam.LayerRadii[iLayer];
371359
const float& r2 = trkParam.LayerRadii[iLayer + 1];
372360
const float res1 = o2::gpu::CAMath::Hypot(trkParam.PVres, mPositionResolution[iLayer]);
373361
const float res2 = o2::gpu::CAMath::Hypot(trkParam.PVres, mPositionResolution[iLayer + 1]);
374-
const float cosTheta1half = o2::gpu::CAMath::Sqrt(1.f - Sq(0.5f * r1 * oneOverR));
375-
const float cosTheta2half = o2::gpu::CAMath::Sqrt(1.f - Sq(0.5f * r2 * oneOverR));
362+
const float cosTheta1half = o2::gpu::CAMath::Sqrt(1.f - math_utils::Sq(0.5f * r1 * oneOverR));
363+
const float cosTheta2half = o2::gpu::CAMath::Sqrt(1.f - math_utils::Sq(0.5f * r2 * oneOverR));
376364
float x = r2 * cosTheta1half - r1 * cosTheta2half;
377-
float delta = o2::gpu::CAMath::Sqrt(1.f / (1.f - 0.25f * Sq(x * oneOverR)) * (Sq(0.25f * r1 * r2 * Sq(oneOverR) / cosTheta2half + cosTheta1half) * Sq(res1) + Sq(0.25f * r1 * r2 * Sq(oneOverR) / cosTheta1half + cosTheta2half) * Sq(res2)));
365+
float delta = o2::gpu::CAMath::Sqrt(1.f / (1.f - 0.25f * math_utils::Sq(x * oneOverR)) * (math_utils::Sq(0.25f * r1 * r2 * math_utils::Sq(oneOverR) / cosTheta2half + cosTheta1half) * math_utils::Sq(res1) + math_utils::Sq(0.25f * r1 * r2 * math_utils::Sq(oneOverR) / cosTheta1half + cosTheta2half) * math_utils::Sq(res2)));
378366
mPhiCuts[iLayer] = std::min(o2::gpu::CAMath::ASin(0.5f * x * oneOverR) + 2.f * mMSangles[iLayer] + delta, o2::constants::math::PI * 0.5f);
379367
}
380368
}

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

Lines changed: 4 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,7 @@
2222
#include <format>
2323
#endif
2424

25+
#include <oneapi/tbb/blocked_range.h>
2526
#include <oneapi/tbb/parallel_sort.h>
2627

2728
#include "CommonConstants/MathConstants.h"
@@ -37,14 +38,6 @@
3738

3839
using o2::base::PropagatorF;
3940

40-
namespace
41-
{
42-
inline float Sq(float q)
43-
{
44-
return q * q;
45-
}
46-
} // namespace
47-
4841
namespace o2::its
4942
{
5043

@@ -104,15 +97,15 @@ void TrackerTraits<nLayers>::computeLayerTracklets(const int iteration, int iROF
10497
if (primaryVertex.isFlagSet(2) && iteration != 3) {
10598
continue;
10699
}
107-
const float resolution = o2::gpu::CAMath::Sqrt(Sq(mTrkParams[iteration].PVres) / primaryVertex.getNContributors() + Sq(mTimeFrame->getPositionResolution(iLayer)));
100+
const float resolution = o2::gpu::CAMath::Sqrt(math_utils::Sq(mTrkParams[iteration].PVres) / primaryVertex.getNContributors() + math_utils::Sq(mTimeFrame->getPositionResolution(iLayer)));
108101

109102
const float tanLambda{(currentCluster.zCoordinate - primaryVertex.getZ()) * inverseR0};
110103

111104
const float zAtRmin{tanLambda * (mTimeFrame->getMinR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate};
112105
const float zAtRmax{tanLambda * (mTimeFrame->getMaxR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate};
113106

114-
const float sqInverseDeltaZ0{1.f / (Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + 2.e-8f)}; /// protecting from overflows adding the detector resolution
115-
const float sigmaZ{o2::gpu::CAMath::Sqrt(Sq(resolution) * Sq(tanLambda) * ((Sq(inverseR0) + sqInverseDeltaZ0) * Sq(meanDeltaR) + 1.f) + Sq(meanDeltaR * mTimeFrame->getMSangle(iLayer)))};
107+
const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + 2.e-8f)}; /// protecting from overflows adding the detector resolution
108+
const float sigmaZ{o2::gpu::CAMath::Sqrt(math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInverseDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * mTimeFrame->getMSangle(iLayer)))};
116109

117110
const int4 selectedBinsRect{getBinsRect(currentCluster, iLayer + 1, zAtRmin, zAtRmax, sigmaZ * mTrkParams[iteration].NSigmaCut, mTimeFrame->getPhiCut(iLayer))};
118111
if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) {

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

Lines changed: 2 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -31,12 +31,6 @@
3131

3232
using namespace o2::its;
3333

34-
float smallestAngleDifference(float a, float b)
35-
{
36-
float diff = fmod(b - a + o2::constants::math::PI, o2::constants::math::TwoPI) - o2::constants::math::PI;
37-
return (diff < -o2::constants::math::PI) ? diff + o2::constants::math::TwoPI : ((diff > o2::constants::math::PI) ? diff - o2::constants::math::TwoPI : diff);
38-
}
39-
4034
template <TrackletMode Mode, bool EvalRun>
4135
void trackleterKernelHost(
4236
const gsl::span<const Cluster>& clustersNextLayer, // 0 2
@@ -75,7 +69,7 @@ void trackleterKernelHost(
7569
continue;
7670
}
7771
const Cluster& nextCluster{clustersNextLayer[iNextLayerClusterIndex]};
78-
if (o2::gpu::GPUCommonMath::Abs(smallestAngleDifference(currentCluster.phi, nextCluster.phi)) < phiCut) {
72+
if (o2::gpu::GPUCommonMath::Abs(math_utils::smallestAngleDifference(currentCluster.phi, nextCluster.phi)) < phiCut) {
7973
if (storedTracklets < maxTrackletsPerCluster) {
8074
if constexpr (!EvalRun) {
8175
if constexpr (Mode == TrackletMode::Layer0Layer1) {
@@ -128,7 +122,7 @@ void trackletSelectionKernelHost(
128122
continue;
129123
}
130124
const float deltaTanLambda{o2::gpu::GPUCommonMath::Abs(tracklet01.tanLambda - tracklet12.tanLambda)};
131-
const float deltaPhi{o2::gpu::GPUCommonMath::Abs(smallestAngleDifference(tracklet01.phi, tracklet12.phi))};
125+
const float deltaPhi{o2::gpu::GPUCommonMath::Abs(math_utils::smallestAngleDifference(tracklet01.phi, tracklet12.phi))};
132126
if (!usedTracklets[iTracklet01] && deltaTanLambda < tanLambdaCut && deltaPhi < phiCut && validTracklets != maxTracklets) {
133127
usedClusters0[tracklet01.firstClusterIndex] = true;
134128
usedClusters2[tracklet12.secondClusterIndex] = true;

0 commit comments

Comments
 (0)