Skip to content

Commit f08023c

Browse files
f3schf3sch
committed
ITS: allow to Individually set kernel parameters
1 parent 1aa2c14 commit f08023c

File tree

4 files changed

+71
-19
lines changed

4 files changed

+71
-19
lines changed

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

Lines changed: 16 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -80,8 +80,8 @@ void TrackerTraitsGPU<nLayers>::computeLayerTracklets(const int iteration, int i
8080
mTimeFrameGPU->getPositionResolutions(),
8181
this->mTrkParams[iteration].LayerRadii,
8282
mTimeFrameGPU->getMSangles(),
83-
conf.nBlocks,
84-
conf.nThreads,
83+
conf.nBlocksLayerTracklets[iteration],
84+
conf.nThreadsLayerTracklets[iteration],
8585
mTimeFrameGPU->getStreams());
8686
mTimeFrameGPU->createTrackletsBuffers();
8787
computeTrackletsInROFsHandler<nLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
@@ -113,8 +113,8 @@ void TrackerTraitsGPU<nLayers>::computeLayerTracklets(const int iteration, int i
113113
mTimeFrameGPU->getPositionResolutions(),
114114
this->mTrkParams[iteration].LayerRadii,
115115
mTimeFrameGPU->getMSangles(),
116-
conf.nBlocks,
117-
conf.nThreads,
116+
conf.nBlocksLayerTracklets[iteration],
117+
conf.nThreadsLayerTracklets[iteration],
118118
mTimeFrameGPU->getStreams());
119119
}
120120

@@ -144,8 +144,8 @@ void TrackerTraitsGPU<nLayers>::computeLayerCells(const int iteration)
144144
this->mTrkParams[iteration].MaxChi2ClusterAttachment,
145145
this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
146146
this->mTrkParams[iteration].NSigmaCut,
147-
conf.nBlocks,
148-
conf.nThreads);
147+
conf.nBlocksLayerCells[iteration],
148+
conf.nThreadsLayerCells[iteration]);
149149
mTimeFrameGPU->createCellsBuffers(iLayer);
150150
computeCellsHandler(mTimeFrameGPU->getDeviceArrayClusters(),
151151
mTimeFrameGPU->getDeviceArrayUnsortedClusters(),
@@ -161,8 +161,8 @@ void TrackerTraitsGPU<nLayers>::computeLayerCells(const int iteration)
161161
this->mTrkParams[iteration].MaxChi2ClusterAttachment,
162162
this->mTrkParams[iteration].CellDeltaTanLambdaSigma,
163163
this->mTrkParams[iteration].NSigmaCut,
164-
conf.nBlocks,
165-
conf.nThreads);
164+
conf.nBlocksLayerCells[iteration],
165+
conf.nThreadsLayerCells[iteration]);
166166
}
167167
}
168168

@@ -191,8 +191,8 @@ void TrackerTraitsGPU<nLayers>::findCellsNeighbours(const int iteration)
191191
currentLayerCellsNum,
192192
nextLayerCellsNum,
193193
1e2,
194-
conf.nBlocks,
195-
conf.nThreads);
194+
conf.nBlocksFindNeighbours[iteration],
195+
conf.nThreadsFindNeighbours[iteration]);
196196

197197
mTimeFrameGPU->createNeighboursDevice(iLayer, nNeigh);
198198

@@ -207,8 +207,8 @@ void TrackerTraitsGPU<nLayers>::findCellsNeighbours(const int iteration)
207207
currentLayerCellsNum,
208208
nextLayerCellsNum,
209209
1e2,
210-
conf.nBlocks,
211-
conf.nThreads);
210+
conf.nBlocksFindNeighbours[iteration],
211+
conf.nThreadsFindNeighbours[iteration]);
212212

213213
nNeigh = filterCellNeighboursHandler(mTimeFrameGPU->getDeviceNeighbourPairs(iLayer),
214214
mTimeFrameGPU->getDeviceNeighbours(iLayer),
@@ -247,8 +247,8 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
247247
this->mTrkParams[0].MaxChi2NDF,
248248
mTimeFrameGPU->getDevicePropagator(),
249249
this->mTrkParams[0].CorrType,
250-
conf.nBlocks,
251-
conf.nThreads);
250+
conf.nBlocksProcessNeighbours[iteration],
251+
conf.nThreadsProcessNeighbours[iteration]);
252252
}
253253
// fixme: I don't want to move tracks back and forth, but I need a way to use a thrust::allocator that is aware of our managed memory.
254254
if (trackSeeds.empty()) {
@@ -269,8 +269,8 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
269269
this->mTrkParams[0].MaxChi2NDF, // float maxChi2NDF
270270
mTimeFrameGPU->getDevicePropagator(), // const o2::base::Propagator* propagator
271271
this->mTrkParams[0].CorrType, // o2::base::PropagatorImpl<float>::MatCorrType
272-
conf.nBlocks,
273-
conf.nThreads);
272+
conf.nBlocksTracksSeeds[iteration],
273+
conf.nThreadsTracksSeeds[iteration]);
274274

275275
mTimeFrameGPU->downloadTrackITSExtDevice(trackSeeds);
276276

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

Lines changed: 26 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -107,9 +107,32 @@ struct TrackerParamConfig : public o2::conf::ConfigurableParamHelper<TrackerPara
107107
};
108108

109109
struct ITSGpuTrackingParamConfig : public o2::conf::ConfigurableParamHelper<ITSGpuTrackingParamConfig> {
110-
// GPU-specific parameters
111-
int nBlocks = 20;
112-
int nThreads = 256;
110+
static constexpr int MaxIter = TrackerParamConfig::MaxIter;
111+
112+
/// Set nBlocks/nThreads to summarily override all kernel launch parameters in each iteration.
113+
/// Parameters must start with nBlocks/nThreads.
114+
static constexpr int OverrideValue{-1};
115+
static constexpr char const* BlocksName = "nBlocks";
116+
static constexpr char const* ThreadsName = "nThreads";
117+
int nBlocks = OverrideValue;
118+
int nThreads = OverrideValue;
119+
void maybeOverride() const;
120+
121+
/// Individual kernel launch parameter for each iteration
122+
int nBlocksLayerTracklets[MaxIter] = {30, 30, 30, 30};
123+
int nThreadsLayerTracklets[MaxIter] = {256, 256, 256, 256};
124+
125+
int nBlocksLayerCells[MaxIter] = {30, 30, 30, 30};
126+
int nThreadsLayerCells[MaxIter] = {256, 256, 256, 256};
127+
128+
int nBlocksFindNeighbours[MaxIter] = {30, 30, 30, 30};
129+
int nThreadsFindNeighbours[MaxIter] = {256, 256, 256, 256};
130+
131+
int nBlocksProcessNeighbours[MaxIter] = {30, 30, 30, 30};
132+
int nThreadsProcessNeighbours[MaxIter] = {256, 256, 256, 256};
133+
134+
int nBlocksTracksSeeds[MaxIter] = {30, 30, 30, 30};
135+
int nThreadsTracksSeeds[MaxIter] = {256, 256, 256, 256};
113136

114137
O2ParamDef(ITSGpuTrackingParamConfig, "ITSGpuTrackingParam");
115138
};

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -39,6 +39,7 @@ Tracker::Tracker(TrackerTraits7* traits) : mTraits(traits)
3939
/// Initialise standard configuration with 1 iteration
4040
mTrkParams.resize(1);
4141
if (traits->isGPU()) {
42+
ITSGpuTrackingParamConfig::Instance().maybeOverride();
4243
ITSGpuTrackingParamConfig::Instance().printKeyValues(true, true);
4344
}
4445
}

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

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9,8 +9,36 @@
99
// granted to it by virtue of its status as an Intergovernmental Organization
1010
// or submit itself to any jurisdiction.
1111

12+
#include <boost/property_tree/ptree.hpp>
13+
14+
#include "Framework/Logger.h"
1215
#include "ITStracking/TrackingConfigParam.h"
1316

1417
O2ParamImpl(o2::its::VertexerParamConfig);
1518
O2ParamImpl(o2::its::TrackerParamConfig);
1619
O2ParamImpl(o2::its::ITSGpuTrackingParamConfig);
20+
21+
namespace o2::its
22+
{
23+
24+
void ITSGpuTrackingParamConfig::maybeOverride() const
25+
{
26+
if (nBlocks == OverrideValue && nThreads == OverrideValue) {
27+
return;
28+
}
29+
const auto name = getName();
30+
auto members = getDataMembers();
31+
for (auto member : *members) {
32+
if (!member.name.ends_with(BlocksName) && !member.name.ends_with(ThreadsName)) {
33+
if (nBlocks != OverrideValue && member.name.starts_with(BlocksName) && (member.value != nBlocks)) {
34+
o2::conf::ConfigurableParam::setValue<int>(name, member.name, nBlocks);
35+
}
36+
if (nThreads != OverrideValue && member.name.starts_with(ThreadsName) && (member.value != nThreads)) {
37+
o2::conf::ConfigurableParam::setValue<int>(name, member.name, nThreads);
38+
}
39+
}
40+
}
41+
LOGP(info, "Overwriting gpu threading parameters");
42+
} // namespace o2::its
43+
44+
} // namespace o2::its

0 commit comments

Comments
 (0)