Skip to content

Commit ff9580b

Browse files
ktfktf
authored
Merge 6da9bd8 into sapling-pr-archive-ktf
2 parents 6e40a7c + 6da9bd8 commit ff9580b

40 files changed

+350
-754
lines changed

‎DataFormats/Detectors/ITSMFT/common/CMakeLists.txt‎

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -18,17 +18,18 @@ o2_add_library(DataFormatsITSMFT
1818
src/ClusterPattern.cxx
1919
src/ClusterTopology.cxx
2020
src/TopologyDictionary.cxx
21+
src/TimeDeadMap.cxx
2122
src/CTF.cxx
2223
PUBLIC_LINK_LIBRARIES O2::ITSMFTBase
2324
O2::ReconstructionDataFormats
2425
Microsoft.GSL::GSL)
2526

2627
o2_target_root_dictionary(DataFormatsITSMFT
2728
HEADERS include/DataFormatsITSMFT/ROFRecord.h
28-
include/DataFormatsITSMFT/Digit.h
29-
include/DataFormatsITSMFT/GBTCalibData.h
30-
include/DataFormatsITSMFT/NoiseMap.h
31-
include/DataFormatsITSMFT/TimeDeadMap.h
29+
include/DataFormatsITSMFT/Digit.h
30+
include/DataFormatsITSMFT/GBTCalibData.h
31+
include/DataFormatsITSMFT/NoiseMap.h
32+
include/DataFormatsITSMFT/TimeDeadMap.h
3233
include/DataFormatsITSMFT/Cluster.h
3334
include/DataFormatsITSMFT/CompCluster.h
3435
include/DataFormatsITSMFT/ClusterPattern.h

‎DataFormats/Detectors/ITSMFT/common/include/DataFormatsITSMFT/TimeDeadMap.h‎

Lines changed: 8 additions & 87 deletions
Original file line numberDiff line numberDiff line change
@@ -15,8 +15,6 @@
1515
#define ALICEO2_ITSMFT_TIMEDEADMAP_H
1616

1717
#include "Rtypes.h"
18-
#include "DetectorsCommonDataFormats/DetID.h"
19-
#include <iostream>
2018
#include <vector>
2119
#include <map>
2220

@@ -26,6 +24,8 @@ namespace o2
2624
namespace itsmft
2725
{
2826

27+
class NoiseMap;
28+
2929
class TimeDeadMap
3030
{
3131
public:
@@ -56,96 +56,17 @@ class TimeDeadMap
5656
mStaticDeadMap.clear();
5757
}
5858

59-
void decodeMap(o2::itsmft::NoiseMap& noisemap)
60-
{ // for static part only
61-
if (mMAP_VERSION == "3") {
62-
LOG(error) << "Trying to decode static part of deadmap version " << mMAP_VERSION << ". Not implemented, doing nothing.";
63-
return;
64-
}
65-
for (int iel = 0; iel < mStaticDeadMap.size(); iel++) {
66-
uint16_t w = mStaticDeadMap[iel];
67-
noisemap.maskFullChip(w & 0x7FFF);
68-
if (w & 0x8000) {
69-
for (int w2 = (w & 0x7FFF) + 1; w2 < mStaticDeadMap.at(iel + 1); w2++) {
70-
noisemap.maskFullChip(w2);
71-
}
72-
}
73-
}
74-
}
75-
76-
void decodeMap(unsigned long orbit, o2::itsmft::NoiseMap& noisemap, bool includeStaticMap = true, long orbitGapAllowed = 330000)
77-
{ // for time-dependent and (optionally) static part. Use orbitGapAllowed = -1 to ignore check on orbit difference
78-
79-
if (mMAP_VERSION != "3" && mMAP_VERSION != "4") {
80-
LOG(error) << "Trying to decode time-dependent deadmap version " << mMAP_VERSION << ". Not implemented, doing nothing.";
81-
return;
82-
}
83-
84-
if (mEvolvingDeadMap.empty()) {
85-
LOG(warning) << "Time-dependent dead map is empty. Doing nothing.";
86-
return;
87-
}
88-
89-
std::vector<uint16_t> closestVec;
90-
long dT = getMapAtOrbit(orbit, closestVec);
91-
92-
if (orbitGapAllowed >= 0 && std::abs(dT) > orbitGapAllowed) {
93-
LOG(warning) << "Requested orbit " << orbit << ", found " << orbit - dT << ". Orbit gap is too high, skipping time-dependent map.";
94-
closestVec.clear();
95-
}
96-
97-
// add static part if requested. something may be masked twice
98-
if (includeStaticMap && mMAP_VERSION != "3") {
99-
closestVec.insert(closestVec.end(), mStaticDeadMap.begin(), mStaticDeadMap.end());
100-
}
101-
102-
// vector encoding: if 1<<15 = 0x8000 is set, the word encodes the first element of a range, with mask (1<<15)-1 = 0x7FFF. The last element of the range is the next in the vector.
103-
104-
for (int iel = 0; iel < closestVec.size(); iel++) {
105-
uint16_t w = closestVec.at(iel);
106-
noisemap.maskFullChip(w & 0x7FFF);
107-
if (w & 0x8000) {
108-
for (int w2 = (w & 0x7FFF) + 1; w2 < closestVec.at(iel + 1); w2++) {
109-
noisemap.maskFullChip(w2);
110-
}
111-
}
112-
}
113-
};
114-
59+
void decodeMap(NoiseMap& noisemap) const;
60+
void decodeMap(unsigned long orbit, o2::itsmft::NoiseMap& noisemap, bool includeStaticMap = true, long orbitGapAllowed = 330000) const;
11561
std::string getMapVersion() const { return mMAP_VERSION; };
11662

11763
unsigned long getEvolvingMapSize() const { return mEvolvingDeadMap.size(); };
118-
119-
std::vector<unsigned long> getEvolvingMapKeys()
120-
{
121-
std::vector<unsigned long> keys;
122-
std::transform(mEvolvingDeadMap.begin(), mEvolvingDeadMap.end(), std::back_inserter(keys),
123-
[](const auto& O) { return O.first; });
124-
return keys;
125-
}
126-
127-
void getStaticMap(std::vector<uint16_t>& mmap) { mmap = mStaticDeadMap; };
128-
129-
long getMapAtOrbit(unsigned long orbit, std::vector<uint16_t>& mmap)
130-
{ // fills mmap and returns requested_orbit - found_orbit. Found orbit is the highest key lower or equal to the requested one
131-
if (mEvolvingDeadMap.empty()) {
132-
LOG(warning) << "Requested orbit " << orbit << "from an empty time-dependent map. Doing nothing";
133-
return (long)orbit;
134-
}
135-
auto closest = mEvolvingDeadMap.upper_bound(orbit);
136-
if (closest != mEvolvingDeadMap.begin()) {
137-
--closest;
138-
mmap = closest->second;
139-
return (long)orbit - closest->first;
140-
} else {
141-
mmap = mEvolvingDeadMap.begin()->second;
142-
return (long)(orbit)-mEvolvingDeadMap.begin()->first;
143-
}
144-
}
145-
64+
std::vector<unsigned long> getEvolvingMapKeys() const;
65+
void getStaticMap(std::vector<uint16_t>& mmap) const { mmap = mStaticDeadMap; };
66+
long getMapAtOrbit(unsigned long orbit, std::vector<uint16_t>& mmap) const;
14667
void setMapVersion(std::string version) { mMAP_VERSION = version; };
14768

148-
bool isDefault() { return mIsDefaultObject; };
69+
bool isDefault() const { return mIsDefaultObject; };
14970
void setAsDefault(bool isdef = true) { mIsDefaultObject = isdef; };
15071

15172
private:

‎DataFormats/Detectors/ITSMFT/common/src/TimeDeadMap.cxx‎

Lines changed: 100 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,100 @@
1+
// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2+
// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3+
// All rights not expressly granted are reserved.
4+
//
5+
// This software is distributed under the terms of the GNU General Public
6+
// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7+
//
8+
// In applying this license CERN does not waive the privileges and immunities
9+
// granted to it by virtue of its status as an Intergovernmental Organization
10+
// or submit itself to any jurisdiction.
11+
12+
/// \file TimeDeadMap.cxx
13+
/// \brief Implementation of the time-dependent map
14+
15+
#include "DataFormatsITSMFT/TimeDeadMap.h"
16+
#include "DataFormatsITSMFT/NoiseMap.h"
17+
#include "Framework/Logger.h"
18+
19+
using namespace o2::itsmft;
20+
21+
void TimeDeadMap::decodeMap(o2::itsmft::NoiseMap& noisemap) const
22+
{ // for static part only
23+
if (mMAP_VERSION == "3") {
24+
LOG(error) << "Trying to decode static part of deadmap version " << mMAP_VERSION << ". Not implemented, doing nothing.";
25+
return;
26+
}
27+
for (int iel = 0; iel < mStaticDeadMap.size(); iel++) {
28+
uint16_t w = mStaticDeadMap[iel];
29+
noisemap.maskFullChip(w & 0x7FFF);
30+
if (w & 0x8000) {
31+
for (int w2 = (w & 0x7FFF) + 1; w2 < mStaticDeadMap.at(iel + 1); w2++) {
32+
noisemap.maskFullChip(w2);
33+
}
34+
}
35+
}
36+
}
37+
38+
void TimeDeadMap::decodeMap(unsigned long orbit, o2::itsmft::NoiseMap& noisemap, bool includeStaticMap, long orbitGapAllowed) const
39+
{ // for time-dependent and (optionally) static part. Use orbitGapAllowed = -1 to ignore check on orbit difference
40+
41+
if (mMAP_VERSION != "3" && mMAP_VERSION != "4") {
42+
LOG(error) << "Trying to decode time-dependent deadmap version " << mMAP_VERSION << ". Not implemented, doing nothing.";
43+
return;
44+
}
45+
46+
if (mEvolvingDeadMap.empty()) {
47+
LOG(warning) << "Time-dependent dead map is empty. Doing nothing.";
48+
return;
49+
}
50+
51+
std::vector<uint16_t> closestVec;
52+
long dT = getMapAtOrbit(orbit, closestVec);
53+
54+
if (orbitGapAllowed >= 0 && std::abs(dT) > orbitGapAllowed) {
55+
LOG(warning) << "Requested orbit " << orbit << ", found " << orbit - dT << ". Orbit gap is too high, skipping time-dependent map.";
56+
closestVec.clear();
57+
}
58+
59+
// add static part if requested. something may be masked twice
60+
if (includeStaticMap && mMAP_VERSION != "3") {
61+
closestVec.insert(closestVec.end(), mStaticDeadMap.begin(), mStaticDeadMap.end());
62+
}
63+
64+
// vector encoding: if 1<<15 = 0x8000 is set, the word encodes the first element of a range, with mask (1<<15)-1 = 0x7FFF. The last element of the range is the next in the vector.
65+
66+
for (int iel = 0; iel < closestVec.size(); iel++) {
67+
uint16_t w = closestVec.at(iel);
68+
noisemap.maskFullChip(w & 0x7FFF);
69+
if (w & 0x8000) {
70+
for (int w2 = (w & 0x7FFF) + 1; w2 < closestVec.at(iel + 1); w2++) {
71+
noisemap.maskFullChip(w2);
72+
}
73+
}
74+
}
75+
}
76+
77+
std::vector<unsigned long> TimeDeadMap::getEvolvingMapKeys() const
78+
{
79+
std::vector<unsigned long> keys;
80+
std::transform(mEvolvingDeadMap.begin(), mEvolvingDeadMap.end(), std::back_inserter(keys),
81+
[](const auto& O) { return O.first; });
82+
return keys;
83+
}
84+
85+
long TimeDeadMap::getMapAtOrbit(unsigned long orbit, std::vector<uint16_t>& mmap) const
86+
{ // fills mmap and returns requested_orbit - found_orbit. Found orbit is the highest key lower or equal to the requested one
87+
if (mEvolvingDeadMap.empty()) {
88+
LOG(warning) << "Requested orbit " << orbit << "from an empty time-dependent map. Doing nothing";
89+
return (long)orbit;
90+
}
91+
auto closest = mEvolvingDeadMap.upper_bound(orbit);
92+
if (closest != mEvolvingDeadMap.begin()) {
93+
--closest;
94+
mmap = closest->second;
95+
return (long)orbit - closest->first;
96+
} else {
97+
mmap = mEvolvingDeadMap.begin()->second;
98+
return (long)(orbit)-mEvolvingDeadMap.begin()->first;
99+
}
100+
}

‎DataFormats/simulation/src/DigitizationContext.cxx‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -452,18 +452,18 @@ std::vector<std::tuple<int, int, int>> getTimeFrameBoundaries(std::vector<o2::In
452452
// in this range search the smallest index which precedes
453453
// timeframe ti by not more than "orbitsEarly" orbits
454454
// (could probably use binary search, in case optimization becomes necessary)
455-
int earlyOrbitIndex = prev_tf_range.second;
455+
int earlyOrbitIndex = -1; // init to start of this timeframe ... there may not be early orbits
456456

457457
// this is the orbit of the ti-th timeframe start
458458
auto orbit_timeframe_start = startOrbit + ti * orbitsPerTF;
459459

460-
auto orbit_timeframe_early_fractional = orbit_timeframe_start - orbitsEarly;
461-
auto orbit_timeframe_early_integral = (uint32_t)(orbit_timeframe_early_fractional);
460+
auto orbit_timeframe_early_fractional = 1. * orbit_timeframe_start - orbitsEarly;
461+
auto orbit_timeframe_early_integral = static_cast<long>(std::floor(orbit_timeframe_early_fractional));
462462

463463
auto bc_early = (uint32_t)((orbit_timeframe_early_fractional - orbit_timeframe_early_integral) * o2::constants::lhc::LHCMaxBunches);
464464

465465
// this is the interaction record of the ti-th timeframe start
466-
o2::InteractionRecord timeframe_start_record(0, orbit_timeframe_early_integral);
466+
o2::InteractionRecord timeframe_start_record(0, orbit_timeframe_start);
467467
// this is the interaction record in some previous timeframe after which interactions could still
468468
// influence the ti-th timeframe according to orbitsEarly
469469
o2::InteractionRecord timeframe_early_record(bc_early, orbit_timeframe_early_integral);

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

Lines changed: 19 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -130,7 +130,9 @@ class TimeFrameGPU : public TimeFrame<nLayers>
130130
// Host-specific getters
131131
gsl::span<int, nLayers - 1> getNTracklets() { return mNTracklets; }
132132
gsl::span<int, nLayers - 2> getNCells() { return mNCells; }
133-
std::array<int, nLayers - 2>& getArrayNCells() { return mNCells; }
133+
auto& getArrayNCells() { return mNCells; }
134+
gsl::span<int, nLayers - 3> getNNeighbours() { return mNNeighbours; }
135+
auto& getArrayNNeighbours() { return mNNeighbours; }
134136

135137
// Host-available device getters
136138
gsl::span<int*> getDeviceTrackletsLUTs() { return mTrackletsLUTDevice; }
@@ -139,7 +141,9 @@ class TimeFrameGPU : public TimeFrame<nLayers>
139141
gsl::span<CellSeed*> getDeviceCells() { return mCellsDevice; }
140142

141143
// Overridden getters
142-
int getNumberOfCells() const;
144+
int getNumberOfTracklets() const final;
145+
int getNumberOfCells() const final;
146+
int getNumberOfNeighbours() const final;
143147

144148
private:
145149
void allocMemAsync(void**, size_t, Stream*, bool); // Abstract owned and unowned memory allocations
@@ -149,6 +153,7 @@ class TimeFrameGPU : public TimeFrame<nLayers>
149153
// Host-available device buffer sizes
150154
std::array<int, nLayers - 1> mNTracklets;
151155
std::array<int, nLayers - 2> mNCells;
156+
std::array<int, nLayers - 3> mNNeighbours;
152157

153158
// Device pointers
154159
IndexTableUtils* mIndexTableUtilsDevice;
@@ -218,12 +223,24 @@ inline std::vector<unsigned int> TimeFrameGPU<nLayers>::getClusterSizes()
218223
return sizes;
219224
}
220225

226+
template <int nLayers>
227+
inline int TimeFrameGPU<nLayers>::getNumberOfTracklets() const
228+
{
229+
return std::accumulate(mNTracklets.begin(), mNTracklets.end(), 0);
230+
}
231+
221232
template <int nLayers>
222233
inline int TimeFrameGPU<nLayers>::getNumberOfCells() const
223234
{
224235
return std::accumulate(mNCells.begin(), mNCells.end(), 0);
225236
}
226237

238+
template <int nLayers>
239+
inline int TimeFrameGPU<nLayers>::getNumberOfNeighbours() const
240+
{
241+
return std::accumulate(mNNeighbours.begin(), mNNeighbours.end(), 0);
242+
}
243+
227244
} // namespace o2::its::gpu
228245

229246
#endif

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -342,6 +342,9 @@ void TimeFrameGPU<nLayers>::createNeighboursIndexTablesDevice()
342342
LOGP(debug, "gpu-transfer: loading neighbours LUT for {} elements on layer {}, for {} MB.", mNCells[iLayer], iLayer, mNCells[iLayer] * sizeof(CellSeed) / MB);
343343
allocMemAsync(reinterpret_cast<void**>(&mNeighboursIndexTablesDevice[iLayer]), (mNCells[iLayer] + 1) * sizeof(int), nullptr, this->getExtAllocator());
344344
GPUChkErrS(cudaMemsetAsync(mNeighboursIndexTablesDevice[iLayer], 0, (mNCells[iLayer] + 1) * sizeof(int), mGpuStreams[0]->get()));
345+
if (iLayer < nLayers - 3) {
346+
mNNeighbours[iLayer] = 0;
347+
}
345348
}
346349
STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get());
347350
}

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

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -169,7 +169,7 @@ template <int nLayers>
169169
void TrackerTraitsGPU<nLayers>::findCellsNeighbours(const int iteration)
170170
{
171171
mTimeFrameGPU->createNeighboursIndexTablesDevice();
172-
auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
172+
const auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance();
173173
for (int iLayer{0}; iLayer < this->mTrkParams[iteration].CellsPerRoad() - 1; ++iLayer) {
174174
const int nextLayerCellsNum{static_cast<int>(mTimeFrameGPU->getNCells()[iLayer + 1])};
175175

@@ -208,10 +208,11 @@ void TrackerTraitsGPU<nLayers>::findCellsNeighbours(const int iteration)
208208
conf.nBlocks,
209209
conf.nThreads);
210210

211-
filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
212-
mTimeFrameGPU->getDeviceNeighbours(iLayer),
213-
nNeigh,
214-
mTimeFrameGPU->getExternalAllocator());
211+
nNeigh = filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
212+
mTimeFrameGPU->getDeviceNeighbours(iLayer),
213+
nNeigh,
214+
mTimeFrameGPU->getExternalAllocator());
215+
mTimeFrameGPU->getArrayNNeighbours()[iLayer] = nNeigh;
215216
}
216217
mTimeFrameGPU->createNeighboursDeviceArray();
217218
mTimeFrameGPU->unregisterRest();

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1217,6 +1217,8 @@ void processNeighboursHandler(const int startLayer,
12171217
maxChi2ClusterAttachment,
12181218
propagator,
12191219
matCorrType);
1220+
GPUChkErrS(cudaPeekAtLastError());
1221+
GPUChkErrS(cudaDeviceSynchronize());
12201222
12211223
int level = startLevel;
12221224
thrust::device_vector<int, gpu::TypedAllocator<int>> lastCellId(allocInt);
@@ -1276,6 +1278,8 @@ void processNeighboursHandler(const int startLayer,
12761278
maxChi2ClusterAttachment,
12771279
propagator,
12781280
matCorrType);
1281+
GPUChkErrS(cudaPeekAtLastError());
1282+
GPUChkErrS(cudaDeviceSynchronize());
12791283
}
12801284
thrust::device_vector<CellSeed, gpu::TypedAllocator<CellSeed>> outSeeds(updatedCellSeed.size(), allocCellSeed);
12811285
auto end = thrust::copy_if(updatedCellSeed.begin(), updatedCellSeed.end(), outSeeds.begin(), gpu::seed_selector(1.e3, maxChi2NDF * ((startLevel + 2) * 2 - 5)));

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

Lines changed: 0 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -29,22 +29,6 @@ namespace o2::its
2929
class Cell final
3030
{
3131
public:
32-
GPUhdDefault() Cell() = default;
33-
GPUhd() Cell(const int firstClusterIndex, const int secondClusterIndex, const int thirdClusterIndex,
34-
const int firstTrackletIndex, const int secondTrackletIndex)
35-
: mFirstClusterIndex(firstClusterIndex),
36-
mSecondClusterIndex(secondClusterIndex),
37-
mThirdClusterIndex(thirdClusterIndex),
38-
mFirstTrackletIndex(firstTrackletIndex),
39-
mSecondTrackletIndex(secondTrackletIndex),
40-
mLevel(1) {}
41-
GPUhdDefault() Cell(const Cell&) = default;
42-
GPUhdDefault() Cell(Cell&&) = default;
43-
GPUhdDefault() ~Cell() = default;
44-
45-
GPUhdDefault() Cell& operator=(const Cell&) = default;
46-
GPUhdDefault() Cell& operator=(Cell&&) noexcept = default;
47-
4832
GPUhd() int getFirstClusterIndex() const { return mFirstClusterIndex; };
4933
GPUhd() int getSecondClusterIndex() const { return mSecondClusterIndex; };
5034
GPUhd() int getThirdClusterIndex() const { return mThirdClusterIndex; };

0 commit comments

Comments
 (0)