Skip to content

Commit f4e51af

Browse files
f3schf3sch
committed
ITS: simplify constants
1 parent 363a646 commit f4e51af

File tree

16 files changed

+98
-252
lines changed

16 files changed

+98
-252
lines changed

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,11 @@ namespace o2::its
2121
class CellSeed;
2222
namespace gpu
2323
{
24+
2425
#ifdef GPUCA_GPUCODE // GPUg() global kernels must only when compiled by GPU compiler
26+
27+
GPUdi() int4 getEmptyBinsRect() { return int4{0, 0, 0, 0}; }
28+
2529
GPUd() bool fitTrack(TrackITSExt& track,
2630
int start,
2731
int end,

‎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: 10 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -53,14 +53,8 @@ using namespace o2::track;
5353

5454
namespace o2::its
5555
{
56-
using namespace constants::its2;
5756
using Vertex = o2::dataformats::Vertex<o2::dataformats::TimeStamp<int>>;
5857

59-
GPUdii() float Sq(float v)
60-
{
61-
return v * v;
62-
}
63-
6458
namespace gpu
6559
{
6660

@@ -69,9 +63,9 @@ GPUd() const int4 getBinsRect(const Cluster& currentCluster, const int layerInde
6963
const float z1, const float z2, float maxdeltaz, float maxdeltaphi)
7064
{
7165
const float zRangeMin = o2::gpu::CAMath::Min(z1, z2) - maxdeltaz;
72-
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;
7367
const float zRangeMax = o2::gpu::CAMath::Max(z1, z2) + maxdeltaz;
74-
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;
7569

7670
if (zRangeMax < -utils.getLayerZ(layerIndex) ||
7771
zRangeMin > utils.getLayerZ(layerIndex) || zRangeMin > zRangeMax) {
@@ -99,7 +93,7 @@ GPUd() bool fitTrack(TrackITSExt& track,
9993
o2::base::PropagatorF::MatCorrType matCorrType)
10094
{
10195
for (int iLayer{start}; iLayer != end; iLayer += step) {
102-
if (track.getClusterIndex(iLayer) == constants::its::UnusedIndex) {
96+
if (track.getClusterIndex(iLayer) == constants::UnusedIndex) {
10397
continue;
10498
}
10599
const TrackingFrameInfo& trackingHit = tfInfos[iLayer][track.getClusterIndex(iLayer)];
@@ -286,7 +280,7 @@ GPUg() void fitTrackSeedsKernel(
286280
temporaryTrack.setChi2(0);
287281
int* clusters = seed.getClusters();
288282
for (int iL{0}; iL < 7; ++iL) {
289-
temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::its::UnusedIndex);
283+
temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::UnusedIndex);
290284
}
291285
bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track,
292286
0, // int lastLayer,
@@ -392,8 +386,6 @@ GPUg() void computeLayerCellsKernel(
392386
const float cellDeltaTanLambdaSigma,
393387
const float nSigmaCut)
394388
{
395-
constexpr float radl = 9.36f; // Radiation length of Si [cm].
396-
constexpr float rho = 2.33f; // Density of Si [g/cm^3].
397389
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.
398390
for (int iCurrentTrackletIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentTrackletIndex < nTrackletsCurrent; iCurrentTrackletIndex += blockDim.x * gridDim.x) {
399391
const Tracklet& currentTracklet = tracklets[layer][iCurrentTrackletIndex];
@@ -432,7 +424,7 @@ GPUg() void computeLayerCellsKernel(
432424
break;
433425
}
434426

435-
if (!track.correctForMaterial(layerxX0[layer + iC], layerxX0[layer] * radl * rho, true)) {
427+
if (!track.correctForMaterial(layerxX0[layer + iC], layerxX0[layer] * constants::Radl * constants::Rho, true)) {
436428
break;
437429
}
438430

@@ -518,12 +510,12 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
518510
if (primaryVertex.isFlagSet(2) && iteration != 3) {
519511
continue;
520512
}
521-
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));
522514
const float tanLambda{(currentCluster.zCoordinate - primaryVertex.getZ()) * inverseR0};
523515
const float zAtRmin{tanLambda * (minR - currentCluster.radius) + currentCluster.zCoordinate};
524516
const float zAtRmax{tanLambda * (maxR - currentCluster.radius) + currentCluster.zCoordinate};
525-
const float sqInverseDeltaZ0{1.f / (Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + 2.e-8f)}; /// protecting from overflows adding the detector resolution
526-
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))};
527519
const int4 selectedBinsRect{getBinsRect(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)};
528520
if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) {
529521
continue;
@@ -557,7 +549,7 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
557549
const float deltaPhi{o2::gpu::CAMath::Abs(currentCluster.phi - nextCluster.phi)};
558550
const float deltaZ{o2::gpu::CAMath::Abs(tanLambda * (nextCluster.radius - currentCluster.radius) + currentCluster.zCoordinate - nextCluster.zCoordinate)};
559551
const int nextSortedIndex{ROFClusters[layerIndex + 1][rof1] + nextClusterIndex};
560-
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)) {
561553
if constexpr (initRun) {
562554
trackletsLUT[layerIndex][currentSortedIndex]++; // we need l0 as well for usual exclusive sums.
563555
} else {
@@ -604,8 +596,6 @@ GPUg() void processNeighboursKernel(const int layer,
604596
const o2::base::Propagator* propagator,
605597
const o2::base::PropagatorF::MatCorrType matCorrType)
606598
{
607-
constexpr float radl = 9.36f; // Radiation length of Si [cm].
608-
constexpr float rho = 2.33f; // Density of Si [g/cm^3].
609599
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.
610600
for (unsigned int iCurrentCell = blockIdx.x * blockDim.x + threadIdx.x; iCurrentCell < nCurrentCells; iCurrentCell += blockDim.x * gridDim.x) {
611601
int foundSeeds{0};
@@ -648,7 +638,7 @@ GPUg() void processNeighboursKernel(const int layer,
648638
}
649639

650640
if (matCorrType == o2::base::PropagatorF::MatCorrType::USEMatCorrNONE) {
651-
if (!seed.correctForMaterial(layerxX0[layer - 1], layerxX0[layer - 1] * radl * rho, true)) {
641+
if (!seed.correctForMaterial(layerxX0[layer - 1], layerxX0[layer - 1] * constants::Radl * constants::Rho, true)) {
652642
continue;
653643
}
654644
}

‎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/Configuration.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -60,9 +60,9 @@ class Configuration : public Param
6060
};
6161

6262
struct TrackingParameters {
63-
int CellMinimumLevel() { return MinTrackLength - constants::its::ClustersPerCell + 1; }
64-
int CellsPerRoad() const { return NLayers - 2; }
65-
int TrackletsPerRoad() const { return NLayers - 1; }
63+
int CellMinimumLevel() const noexcept { return MinTrackLength - constants::ClustersPerCell + 1; }
64+
int CellsPerRoad() const noexcept { return NLayers - 2; }
65+
int TrackletsPerRoad() const noexcept { return NLayers - 1; }
6666
std::string asString() const;
6767

6868
int NLayers = 7;

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

Lines changed: 7 additions & 95 deletions
Original file line numberDiff line numberDiff line change
@@ -17,112 +17,24 @@
1717
#define TRACKINGITSU_INCLUDE_CONSTANTS_H_
1818

1919
#include "ITStracking/Definitions.h"
20-
#include "CommonConstants/MathConstants.h"
2120

22-
#include "GPUCommonMath.h"
23-
#include "GPUCommonDef.h"
24-
25-
#ifndef GPUCA_GPUCODE_DEVICE
26-
#include <climits>
27-
#include <vector>
28-
#include <array>
29-
#endif
30-
31-
namespace o2
32-
{
33-
namespace its
34-
{
35-
36-
namespace constants
21+
namespace o2::its::constants
3722
{
3823
constexpr float MB = 1024.f * 1024.f;
3924
constexpr float GB = 1024.f * 1024.f * 1024.f;
4025
constexpr bool DoTimeBenchmarks = true;
4126
constexpr bool SaveTimeBenchmarks = false;
4227

43-
namespace math
44-
{
45-
constexpr float Pi{3.14159265359f};
46-
constexpr float TwoPi{2.0f * Pi};
47-
constexpr float FloatMinThreshold{1e-20f};
48-
} // namespace math
49-
50-
namespace its
51-
{
52-
constexpr int LayersNumberVertexer{3};
5328
constexpr int ClustersPerCell{3};
5429
constexpr int UnusedIndex{-1};
5530
constexpr float Resolution{0.0005f};
56-
57-
GPUhdi() constexpr std::array<float, 3> VertexerHistogramVolume()
58-
{
59-
return std::array<float, 3>{{1.98, 1.98, 40.f}};
60-
}
61-
} // namespace its
62-
63-
namespace its2
64-
{
65-
constexpr int LayersNumber{7};
66-
constexpr int TrackletsPerRoad{LayersNumber - 1};
67-
constexpr int CellsPerRoad{LayersNumber - 2};
68-
69-
GPUhdi() constexpr std::array<float, LayersNumber> LayersZCoordinate()
70-
{
71-
constexpr double s = 1.; // safety margin
72-
return std::array<float, LayersNumber>{16.333f + s, 16.333f + s, 16.333f + s, 42.140f + s, 42.140f + s, 73.745f + s, 73.745f + s};
73-
}
74-
75-
GPUhdi() constexpr std::array<float, LayersNumber> LayersRCoordinate()
31+
constexpr float Radl = 9.36f; // Radiation length of Si [cm]
32+
constexpr float Rho = 2.33f; // Density of Si [g/cm^3]
33+
namespace its // to be removed
7634
{
77-
return std::array<float, LayersNumber>{{2.33959f, 3.14076f, 3.91924f, 19.6213f, 24.5597f, 34.388f, 39.3329f}};
78-
}
79-
80-
constexpr int ZBins{256};
81-
constexpr int PhiBins{128};
82-
constexpr float InversePhiBinSize{PhiBins / constants::math::TwoPi};
83-
GPUhdi() constexpr std::array<float, LayersNumber> InverseZBinSize()
84-
{
85-
constexpr auto zSize = LayersZCoordinate();
86-
return std::array<float, LayersNumber>{0.5f * ZBins / (zSize[0]), 0.5f * ZBins / (zSize[1]), 0.5f * ZBins / (zSize[2]),
87-
0.5f * ZBins / (zSize[3]), 0.5f * ZBins / (zSize[4]), 0.5f * ZBins / (zSize[5]),
88-
0.5f * ZBins / (zSize[6])};
89-
}
90-
91-
GPUhdi() constexpr float getInverseZCoordinate(const int layerIndex)
92-
{
93-
return 0.5f * ZBins / LayersZCoordinate()[layerIndex];
94-
}
95-
96-
GPUhdi() int getZBinIndex(const int layerIndex, const float zCoordinate)
97-
{
98-
return (zCoordinate + LayersZCoordinate()[layerIndex]) *
99-
InverseZBinSize()[layerIndex];
100-
}
101-
102-
GPUhdi() int getPhiBinIndex(const float currentPhi)
103-
{
104-
return (currentPhi * InversePhiBinSize);
105-
}
106-
107-
GPUhdi() int getBinIndex(const int zIndex, const int phiIndex)
108-
{
109-
return o2::gpu::GPUCommonMath::Min(phiIndex * ZBins + zIndex,
110-
ZBins * PhiBins - 1);
111-
}
112-
113-
GPUhdi() constexpr int4 getEmptyBinsRect() { return int4{0, 0, 0, 0}; }
114-
115-
} // namespace its2
116-
117-
namespace pdgcodes
118-
{
119-
constexpr int PionCode{211};
120-
}
121-
} // namespace constants
122-
#ifndef GPUCA_GPUCODE_DEVICE
123-
typedef std::vector<std::vector<int>> index_table_t;
124-
#endif
35+
constexpr int UnusedIndex{-1};
36+
constexpr float Resolution{0.0005f};
12537
} // namespace its
126-
} // namespace o2
38+
} // namespace o2::its::constants
12739

12840
#endif /* TRACKINGITSU_INCLUDE_CONSTANTS_H_ */

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16,9 +16,9 @@
1616
#ifndef TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_
1717
#define TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_
1818

19-
#include "ITStracking/Constants.h"
2019
#include "ITStracking/Configuration.h"
2120
#include "ITStracking/Definitions.h"
21+
#include "CommonConstants/MathConstants.h"
2222
#include "GPUCommonMath.h"
2323
#include "GPUCommonDef.h"
2424

@@ -55,7 +55,7 @@ class IndexTableUtils
5555
template <class T>
5656
inline void IndexTableUtils::setTrackingParameters(const T& params)
5757
{
58-
mInversePhiBinSize = params.PhiBins / constants::math::TwoPi;
58+
mInversePhiBinSize = params.PhiBins / o2::constants::math::TwoPI;
5959
mNzBins = params.ZBins;
6060
mNphiBins = params.PhiBins;
6161
for (int iLayer{0}; iLayer < params.LayerZ.size(); ++iLayer) {

0 commit comments

Comments
 (0)