Skip to content

Commit 6d863f4

Browse files
f3schFelix Schlepper
f3sch
and
Felix Schlepper
authored
ITS: template Vertexer&Traits, IndexTableUtils (#14606)
Co-authored-by: Felix Schlepper <fschlepp@aliceml.cern.ch>
1 parent 8735573 commit 6d863f4

31 files changed

Lines changed: 188 additions & 161 deletions

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

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -28,6 +28,7 @@ template <int nLayers = 7>
2828
class TimeFrameGPU : public TimeFrame<nLayers>
2929
{
3030
using typename TimeFrame<nLayers>::CellSeedN;
31+
using typename TimeFrame<nLayers>::IndexTableUtilsN;
3132

3233
public:
3334
TimeFrameGPU();
@@ -36,8 +37,8 @@ class TimeFrameGPU : public TimeFrame<nLayers>
3637
/// Most relevant operations
3738
void registerHostMemory(const int);
3839
void unregisterHostMemory(const int);
39-
void initialise(const int, const TrackingParameters&, const int, IndexTableUtils* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr);
40-
void initDevice(IndexTableUtils*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int);
40+
void initialise(const int, const TrackingParameters&, const int, IndexTableUtilsN* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr);
41+
void initDevice(IndexTableUtilsN*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int);
4142
void initDeviceSAFitting();
4243
void loadIndexTableUtils(const int);
4344
void loadTrackingFrameInfoDevice(const int, const int);
@@ -98,7 +99,7 @@ class TimeFrameGPU : public TimeFrame<nLayers>
9899

99100
/// interface
100101
int getNClustersInRofSpan(const int, const int, const int) const;
101-
IndexTableUtils* getDeviceIndexTableUtils() { return mIndexTableUtilsDevice; }
102+
IndexTableUtilsN* getDeviceIndexTableUtils() { return mIndexTableUtilsDevice; }
102103
int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; }
103104
auto& getTrackITSExt() { return mTrackITSExt; }
104105
Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; }
@@ -165,7 +166,7 @@ class TimeFrameGPU : public TimeFrame<nLayers>
165166
std::array<int, nLayers - 3> mNNeighbours;
166167

167168
// Device pointers
168-
IndexTableUtils* mIndexTableUtilsDevice;
169+
IndexTableUtilsN* mIndexTableUtilsDevice;
169170

170171
// Hybrid pref
171172
uint8_t* mMultMaskDevice;

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

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,8 @@ namespace o2::its
2222
template <int nLayers = 7>
2323
class TrackerTraitsGPU final : public TrackerTraits<nLayers>
2424
{
25+
using typename TrackerTraits<nLayers>::IndexTableUtilsN;
26+
2527
public:
2628
TrackerTraitsGPU() = default;
2729
~TrackerTraitsGPU() final = default;
@@ -48,7 +50,7 @@ class TrackerTraitsGPU final : public TrackerTraits<nLayers>
4850
int getTFNumberOfCells() const override;
4951

5052
private:
51-
IndexTableUtils* mDeviceIndexTableUtils;
53+
IndexTableUtilsN* mDeviceIndexTableUtils;
5254
gpu::TimeFrameGPU<nLayers>* mTimeFrameGPU;
5355
};
5456

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

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,13 +27,14 @@ template <int>
2727
class CellSeed;
2828
class TrackingFrameInfo;
2929
class Tracklet;
30+
template <int>
3031
class IndexTableUtils;
3132
class Cluster;
3233
class TrackITSExt;
3334
class ExternalAllocator;
3435

3536
template <int nLayers = 7>
36-
void countTrackletsInROFsHandler(const IndexTableUtils* utils,
37+
void countTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
3738
const uint8_t* multMask,
3839
const int layer,
3940
const int startROF,
@@ -66,7 +67,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils,
6667
gpu::Streams& streams);
6768

6869
template <int nLayers = 7>
69-
void computeTrackletsInROFsHandler(const IndexTableUtils* utils,
70+
void computeTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
7071
const uint8_t* multMask,
7172
const int layer,
7273
const int startROF,

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,8 @@ if(CUDA_ENABLED)
2222
TimeFrameGPU.cu
2323
TracerGPU.cu
2424
TrackingKernels.cu
25-
VertexingKernels.cu
26-
VertexerTraitsGPU.cxx
25+
# VertexingKernels.cu
26+
# VertexerTraitsGPU.cxx
2727
PUBLIC_INCLUDE_DIRECTORIES ../
2828
PUBLIC_LINK_LIBRARIES O2::ITStracking
2929
O2::SimConfig

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

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -61,11 +61,11 @@ void TimeFrameGPU<nLayers>::loadIndexTableUtils(const int iteration)
6161
{
6262
GPUTimer timer("loading indextable utils");
6363
if (!iteration) {
64-
GPULog("gpu-allocation: allocating IndexTableUtils buffer, for {:.2f} MB.", sizeof(IndexTableUtils) / constants::MB);
65-
allocMem(reinterpret_cast<void**>(&mIndexTableUtilsDevice), sizeof(IndexTableUtils), this->getExtAllocator());
64+
GPULog("gpu-allocation: allocating IndexTableUtils buffer, for {:.2f} MB.", sizeof(IndexTableUtilsN) / constants::MB);
65+
allocMem(reinterpret_cast<void**>(&mIndexTableUtilsDevice), sizeof(IndexTableUtilsN), this->getExtAllocator());
6666
}
67-
GPULog("gpu-transfer: loading IndexTableUtils object, for {:.2f} MB.", sizeof(IndexTableUtils) / constants::MB);
68-
GPUChkErrS(cudaMemcpy(mIndexTableUtilsDevice, &(this->mIndexTableUtils), sizeof(IndexTableUtils), cudaMemcpyHostToDevice));
67+
GPULog("gpu-transfer: loading IndexTableUtils object, for {:.2f} MB.", sizeof(IndexTableUtilsN) / constants::MB);
68+
GPUChkErrS(cudaMemcpy(mIndexTableUtilsDevice, &(this->mIndexTableUtils), sizeof(IndexTableUtilsN), cudaMemcpyHostToDevice));
6969
}
7070

7171
template <int nLayers>
@@ -547,7 +547,7 @@ template <int nLayers>
547547
void TimeFrameGPU<nLayers>::initialise(const int iteration,
548548
const TrackingParameters& trkParam,
549549
const int maxLayers,
550-
IndexTableUtils* utils,
550+
IndexTableUtilsN* utils,
551551
const TimeFrameGPUParameters* gpuParam)
552552
{
553553
mGpuStreams.resize(nLayers);

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

Lines changed: 10 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -95,8 +95,9 @@ GPUdii() int4 getEmptyBinsRect()
9595
return int4{0, 0, 0, 0};
9696
}
9797

98+
template <int nLayers>
9899
GPUdii() const int4 getBinsRect(const Cluster& currentCluster, const int layerIndex,
99-
const o2::its::IndexTableUtils& utils,
100+
const IndexTableUtils<nLayers>& utils,
100101
const float z1, const float z2, float maxdeltaz, float maxdeltaphi)
101102
{
102103
const float zRangeMin = o2::gpu::CAMath::Min(z1, z2) - maxdeltaz;
@@ -331,7 +332,7 @@ GPUg() void fitTrackSeedsKernel(
331332
temporaryTrack.resetCovariance();
332333
temporaryTrack.setChi2(0);
333334
auto& clusters = seed.getClusters();
334-
for (int iL{0}; iL < 7; ++iL) {
335+
for (int iL{0}; iL < nLayers; ++iL) {
335336
temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::UnusedIndex);
336337
}
337338
bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track,
@@ -523,9 +524,9 @@ GPUg() void computeLayerCellsKernel(
523524
}
524525
}
525526

526-
template <bool initRun>
527+
template <bool initRun, int nLayers>
527528
GPUg() void computeLayerTrackletsMultiROFKernel(
528-
const IndexTableUtils* utils,
529+
const IndexTableUtils<nLayers>* utils,
529530
const uint8_t* multMask,
530531
const int layerIndex,
531532
const int startROF,
@@ -601,7 +602,7 @@ GPUg() void computeLayerTrackletsMultiROFKernel(
601602
const float zAtRmax{tanLambda * (maxR - currentCluster.radius) + currentCluster.zCoordinate};
602603
const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + constants::Tolerance)}; /// protecting from overflows adding the detector resolution
603604
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))};
604-
const int4 selectedBinsRect{getBinsRect(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)};
605+
const int4 selectedBinsRect{getBinsRect<nLayers>(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)};
605606
if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) {
606607
continue;
607608
}
@@ -769,7 +770,7 @@ GPUhi() void deallocateMemory(void* p, size_t bytes, cudaStream_t stream = nullp
769770
} // namespace gpu
770771

771772
template <int nLayers>
772-
void countTrackletsInROFsHandler(const IndexTableUtils* utils,
773+
void countTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
773774
const uint8_t* multMask,
774775
const int layer,
775776
const int startROF,
@@ -833,7 +834,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils,
833834
}
834835

835836
template <int nLayers>
836-
void computeTrackletsInROFsHandler(const IndexTableUtils* utils,
837+
void computeTrackletsInROFsHandler(const IndexTableUtils<nLayers>* utils,
837838
const uint8_t* multMask,
838839
const int layer,
839840
const int startROF,
@@ -1241,7 +1242,7 @@ void trackSeedHandler(CellSeed<nLayers>* trackSeeds,
12411242
}
12421243

12431244
/// Explicit instantiation of ITS2 handlers
1244-
template void countTrackletsInROFsHandler<7>(const IndexTableUtils* utils,
1245+
template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
12451246
const uint8_t* multMask,
12461247
const int layer,
12471248
const int startROF,
@@ -1273,7 +1274,7 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils* utils,
12731274
const int nThreads,
12741275
gpu::Streams& streams);
12751276

1276-
template void computeTrackletsInROFsHandler<7>(const IndexTableUtils* utils,
1277+
template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
12771278
const uint8_t* multMask,
12781279
const int layer,
12791280
const int startROF,

‎Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +19,8 @@ if(HIP_ENABLED)
1919
../cuda/TrackerTraitsGPU.cxx
2020
../cuda/TracerGPU.cu
2121
../cuda/TrackingKernels.cu
22-
../cuda/VertexingKernels.cu
23-
../cuda/VertexerTraitsGPU.cxx
22+
# ../cuda/VertexingKernels.cu
23+
# ../cuda/VertexerTraitsGPU.cxx
2424
PUBLIC_INCLUDE_DIRECTORIES ../
2525
PUBLIC_LINK_LIBRARIES O2::ITStracking
2626
O2::GPUTracking

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

Lines changed: 6 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -22,32 +22,23 @@
2222
namespace o2::its
2323
{
2424

25+
template <int>
2526
class IndexTableUtils;
2627

2728
struct Cluster final {
2829
GPUhdDefault() Cluster() = default;
2930
GPUhd() Cluster(const float x, const float y, const float z, const int idx);
30-
GPUhd() Cluster(const int, const IndexTableUtils& utils, const Cluster&);
31-
GPUhd() Cluster(const int, const float3&, const IndexTableUtils& utils, const Cluster&);
31+
template <int nLayers>
32+
GPUhd() Cluster(const int, const IndexTableUtils<nLayers>& utils, const Cluster&);
33+
template <int nLayers>
34+
GPUhd() Cluster(const int, const float3&, const IndexTableUtils<nLayers>& utils, const Cluster&);
3235
GPUhdDefault() Cluster(const Cluster&) = default;
3336
GPUhdDefault() Cluster(Cluster&&) noexcept = default;
3437
GPUhdDefault() ~Cluster() = default;
3538

3639
GPUhdDefault() Cluster& operator=(const Cluster&) = default;
3740
GPUhdDefault() Cluster& operator=(Cluster&&) noexcept = default;
38-
39-
// TODO
40-
/*GPUhdDefault() bool operator==(const Cluster&) const = default;*/
41-
GPUhd() bool operator==(const Cluster& other) const
42-
{
43-
return xCoordinate == other.xCoordinate &&
44-
yCoordinate == other.yCoordinate &&
45-
zCoordinate == other.zCoordinate &&
46-
phi == other.phi &&
47-
radius == other.radius &&
48-
clusterId == other.clusterId &&
49-
indexTableBinIndex == other.indexTableBinIndex;
50-
}
41+
GPUhdDefault() bool operator==(const Cluster&) const = default;
5142

5243
GPUhd() void print() const;
5344

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

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

19+
#include <array>
20+
21+
#include "ITStracking/Constants.h"
1922
#include "ITStracking/Configuration.h"
2023
#include "ITStracking/Definitions.h"
2124
#include "CommonConstants/MathConstants.h"
2225
#include "GPUCommonMath.h"
2326
#include "GPUCommonDef.h"
2427

25-
namespace o2
26-
{
27-
namespace its
28+
namespace o2::its
2829
{
30+
31+
template <int nLayers>
2932
class IndexTableUtils
3033
{
3134
public:
@@ -48,12 +51,13 @@ class IndexTableUtils
4851
int mNzBins = 0;
4952
int mNphiBins = 0;
5053
float mInversePhiBinSize = 0.f;
51-
float mLayerZ[8] = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
52-
float mInverseZBinSize[8] = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
54+
std::array<float, nLayers> mLayerZ{};
55+
std::array<float, nLayers> mInverseZBinSize{};
5356
};
5457

58+
template <int nLayers>
5559
template <class T>
56-
inline void IndexTableUtils::setTrackingParameters(const T& params)
60+
inline void IndexTableUtils<nLayers>::setTrackingParameters(const T& params)
5761
{
5862
mInversePhiBinSize = params.PhiBins / o2::constants::math::TwoPI;
5963
mNzBins = params.ZBins;
@@ -66,43 +70,48 @@ inline void IndexTableUtils::setTrackingParameters(const T& params)
6670
}
6771
}
6872

69-
inline float IndexTableUtils::getInverseZCoordinate(const int layerIndex) const
73+
template <int nLayers>
74+
inline float IndexTableUtils<nLayers>::getInverseZCoordinate(const int layerIndex) const
7075
{
7176
return 0.5f * mNzBins / mLayerZ[layerIndex];
7277
}
7378

74-
GPUhdi() int IndexTableUtils::getZBinIndex(const int layerIndex, const float zCoordinate) const
79+
template <int nLayers>
80+
GPUhdi() int IndexTableUtils<nLayers>::getZBinIndex(const int layerIndex, const float zCoordinate) const
7581
{
7682
return (zCoordinate + mLayerZ[layerIndex]) * mInverseZBinSize[layerIndex];
7783
}
7884

79-
GPUhdi() int IndexTableUtils::getPhiBinIndex(const float currentPhi) const
85+
template <int nLayers>
86+
GPUhdi() int IndexTableUtils<nLayers>::getPhiBinIndex(const float currentPhi) const
8087
{
8188
return (currentPhi * mInversePhiBinSize);
8289
}
8390

84-
GPUhdi() int IndexTableUtils::getBinIndex(const int zIndex, const int phiIndex) const
91+
template <int nLayers>
92+
GPUhdi() int IndexTableUtils<nLayers>::getBinIndex(const int zIndex, const int phiIndex) const
8593
{
8694
return o2::gpu::GPUCommonMath::Min(phiIndex * mNzBins + zIndex, mNzBins * mNphiBins - 1);
8795
}
8896

89-
GPUhdi() int IndexTableUtils::countRowSelectedBins(const int* indexTable, const int phiBinIndex,
90-
const int minZBinIndex, const int maxZBinIndex) const
97+
template <int nLayers>
98+
GPUhdi() int IndexTableUtils<nLayers>::countRowSelectedBins(const int* indexTable, const int phiBinIndex,
99+
const int minZBinIndex, const int maxZBinIndex) const
91100
{
92101
const int firstBinIndex{getBinIndex(minZBinIndex, phiBinIndex)};
93102
const int maxBinIndex{firstBinIndex + maxZBinIndex - minZBinIndex + 1};
94103

95104
return indexTable[maxBinIndex] - indexTable[firstBinIndex];
96105
}
97106

98-
GPUhdi() void IndexTableUtils::print() const
107+
template <int nLayers>
108+
GPUhdi() void IndexTableUtils<nLayers>::print() const
99109
{
100110
printf("NzBins: %d, NphiBins: %d, InversePhiBinSize: %f\n", mNzBins, mNphiBins, mInversePhiBinSize);
101-
for (int iLayer{0}; iLayer < 7; ++iLayer) {
111+
for (int iLayer{0}; iLayer < nLayers; ++iLayer) {
102112
printf("Layer %d: Z: %f, InverseZBinSize: %f\n", iLayer, mLayerZ[iLayer], mInverseZBinSize[iLayer]);
103113
}
104114
}
105-
} // namespace its
106-
} // namespace o2
107115

116+
} // namespace o2::its
108117
#endif /* TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_ */

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -65,6 +65,7 @@ class TimeFrameGPU;
6565

6666
template <int nLayers = 7>
6767
struct TimeFrame {
68+
using IndexTableUtilsN = IndexTableUtils<nLayers>;
6869
using CellSeedN = CellSeed<nLayers>;
6970
friend class gpu::TimeFrameGPU<nLayers>;
7071

@@ -273,7 +274,7 @@ struct TimeFrame {
273274
void printCellLUTs();
274275
void printSliceInfo(const int, const int);
275276

276-
IndexTableUtils mIndexTableUtils;
277+
IndexTableUtilsN mIndexTableUtils;
277278

278279
bool mIsGPU = false;
279280

0 commit comments

Comments
 (0)