Skip to content

Commit bc89e07

Browse files
f3schf3sch
committed
fix4
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent c7643cf commit bc89e07

3 files changed

Lines changed: 7 additions & 51 deletions

File tree

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

Lines changed: 1 addition & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -18,10 +18,8 @@
1818
#include "ITStracking/BoundedAllocator.h"
1919
#include "ITStracking/ROFLookupTables.h"
2020
#include "ITStracking/TrackingTopology.h"
21-
#include "ITStracking/Definitions.h"
2221
#include "ITStrackingGPU/Utils.h"
2322
#include "DetectorsBase/Propagator.h"
24-
#include "GPUCommonDef.h"
2523

2624
namespace o2::its
2725
{
@@ -147,11 +145,9 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
147145
int** cellsLUTs,
148146
const int sourceCellTopologyId,
149147
const int targetCellTopologyId,
150-
const typename TrackingTopology<NLayers>::View topology,
151148
const float maxChi2ClusterAttachment,
152149
const float bz,
153150
const unsigned int nCells,
154-
const int maxCellNeighbours,
155151
gpu::Stream& stream);
156152

157153
void scanCellNeighboursHandler(int* neighboursCursor,
@@ -167,11 +163,9 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice,
167163
CellNeighbour* cellNeighbours,
168164
const int sourceCellTopologyId,
169165
const int targetCellTopologyId,
170-
const typename TrackingTopology<NLayers>::View topology,
171166
const float maxChi2ClusterAttachment,
172167
const float bz,
173168
const unsigned int nCells,
174-
const int maxCellNeighbours,
175169
gpu::Stream& stream);
176170

177171
int filterCellNeighboursHandler(gpuPair<int, int>*,
@@ -181,10 +175,8 @@ int filterCellNeighboursHandler(gpuPair<int, int>*,
181175
o2::its::ExternalAllocator* = nullptr);
182176

183177
template <int NLayers>
184-
void processNeighboursHandler(const int startLayer,
185-
const int startLevel,
178+
void processNeighboursHandler(const int startLevel,
186179
const int defaultCellTopologyId,
187-
const typename TrackingTopology<NLayers>::View topology,
188180
CellSeed** allCellSeeds,
189181
CellSeed* currentCellSeeds,
190182
const int* currentCellTopologyIds,

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

Lines changed: 1 addition & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -64,7 +64,6 @@ void TrackerTraitsGPU<NLayers>::adoptTimeFrame(TimeFrame<NLayers>* tf)
6464
template <int NLayers>
6565
void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int iVertex)
6666
{
67-
const auto topology = mTimeFrameGPU->getDeviceTrackingTopologyView();
6867
const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
6968
for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) {
7069
if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) {
@@ -154,7 +153,6 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
154153
template <int NLayers>
155154
void TrackerTraitsGPU<NLayers>::computeLayerCells(const int iteration)
156155
{
157-
const auto topology = mTimeFrameGPU->getDeviceTrackingTopologyView();
158156
const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
159157
for (int iLayer{0}; iLayer < this->mTrkParams[iteration].NLayers; ++iLayer) {
160158
if (this->mTrkParams[iteration].PassFlags[IterationStep::FirstPass]) {
@@ -257,11 +255,9 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
257255
mTimeFrameGPU->getDeviceArrayCellsLUT(),
258256
sourceCellTopologyId,
259257
targetCellTopologyId,
260-
topology,
261258
this->mTrkParams[iteration].MaxChi2ClusterAttachment,
262259
this->mBz,
263260
sourceCellsNum,
264-
1e2,
265261
mTimeFrameGPU->getStream(targetCellTopologyId));
266262
}
267263

@@ -289,11 +285,9 @@ void TrackerTraitsGPU<NLayers>::findCellsNeighbours(const int iteration)
289285
mTimeFrameGPU->getDeviceNeighbours(targetCellTopologyId),
290286
sourceCellTopologyId,
291287
targetCellTopologyId,
292-
topology,
293288
this->mTrkParams[iteration].MaxChi2ClusterAttachment,
294289
this->mBz,
295290
sourceCellsNum,
296-
1e2,
297291
mTimeFrameGPU->getStream(targetCellTopologyId));
298292
}
299293
mTimeFrameGPU->recordEvent(targetCellTopologyId);
@@ -309,7 +303,6 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
309303
bounded_vector<bounded_vector<int>> sharedFirstClusters(this->mTrkParams[iteration].NLayers, bounded_vector<int>(this->getMemoryPool().get()), this->getMemoryPool().get());
310304
firstClusters.resize(this->mTrkParams[iteration].NLayers);
311305
sharedFirstClusters.resize(this->mTrkParams[iteration].NLayers);
312-
const auto topology = mTimeFrameGPU->getDeviceTrackingTopologyView();
313306
const auto hostTopology = mTimeFrameGPU->getTrackingTopologyView();
314307
for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
315308
bounded_vector<TrackSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
@@ -319,10 +312,8 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
319312
mTimeFrameGPU->getNCells()[startCellTopologyId] == 0) {
320313
continue;
321314
}
322-
processNeighboursHandler<NLayers>(startLayer,
323-
startLevel,
315+
processNeighboursHandler<NLayers>(startLevel,
324316
startCellTopologyId,
325-
topology,
326317
mTimeFrameGPU->getDeviceArrayCells(),
327318
mTimeFrameGPU->getDeviceCells()[startCellTopologyId],
328319
nullptr,

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

Lines changed: 5 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -167,14 +167,10 @@ GPUg() void __launch_bounds__(256, 1) computeLayerCellNeighboursKernel(
167167
CellNeighbour* cellNeighbours,
168168
const int sourceCellTopologyId,
169169
const int targetCellTopologyId,
170-
const typename TrackingTopology<NLayers>::View topology,
171170
const float maxChi2ClusterAttachment,
172171
const float bz,
173-
const unsigned int nCells,
174-
const int maxCellNeighbours = 1e2)
172+
const unsigned int nCells)
175173
{
176-
(void)topology;
177-
(void)maxCellNeighbours;
178174
for (int iCurrentCellIndex = blockIdx.x * blockDim.x + threadIdx.x; iCurrentCellIndex < nCells; iCurrentCellIndex += blockDim.x * gridDim.x) {
179175
const auto& currentCellSeed{cellSeedArray[sourceCellTopologyId][iCurrentCellIndex]};
180176
const int nextLayerTrackletIndex{currentCellSeed.getSecondTrackletIndex()};
@@ -470,7 +466,6 @@ template <bool dryRun, int NLayers, typename CurrentSeed>
470466
GPUg() void __launch_bounds__(256, 1) processNeighboursKernel(
471467
const int defaultCellTopologyId,
472468
const int level,
473-
const typename TrackingTopology<NLayers>::View topology,
474469
CellSeed** allCellSeeds,
475470
CurrentSeed* currentCellSeeds,
476471
const int* currentCellIds,
@@ -490,7 +485,6 @@ GPUg() void __launch_bounds__(256, 1) processNeighboursKernel(
490485
const o2::base::Propagator* propagator,
491486
const o2::base::PropagatorF::MatCorrType matCorrType)
492487
{
493-
(void)topology;
494488
for (unsigned int iCurrentCell = blockIdx.x * blockDim.x + threadIdx.x; iCurrentCell < nCurrentCells; iCurrentCell += blockDim.x * gridDim.x) {
495489
if constexpr (!dryRun) {
496490
if (foundSeedsTable[iCurrentCell] == foundSeedsTable[iCurrentCell + 1]) {
@@ -808,11 +802,9 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
808802
int** cellsLUTs,
809803
const int sourceCellTopologyId,
810804
const int targetCellTopologyId,
811-
const typename TrackingTopology<NLayers>::View topology,
812805
const float maxChi2ClusterAttachment,
813806
const float bz,
814807
const unsigned int nCells,
815-
const int maxCellNeighbours,
816808
gpu::Stream& stream)
817809
{
818810
gpu::computeLayerCellNeighboursKernel<true, NLayers><<<60, 256, 0, stream.get()>>>(
@@ -822,11 +814,9 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
822814
nullptr,
823815
sourceCellTopologyId,
824816
targetCellTopologyId,
825-
topology,
826817
maxChi2ClusterAttachment,
827818
bz,
828-
nCells,
829-
maxCellNeighbours);
819+
nCells);
830820
}
831821

832822
void scanCellNeighboursHandler(int* neighboursCursor,
@@ -847,11 +837,9 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice,
847837
CellNeighbour* cellNeighbours,
848838
const int sourceCellTopologyId,
849839
const int targetCellTopologyId,
850-
const typename TrackingTopology<NLayers>::View topology,
851840
const float maxChi2ClusterAttachment,
852841
const float bz,
853842
const unsigned int nCells,
854-
const int maxCellNeighbours,
855843
gpu::Stream& stream)
856844
{
857845
gpu::computeLayerCellNeighboursKernel<false, NLayers><<<60, 256, 0, stream.get()>>>(
@@ -861,11 +849,9 @@ void computeCellNeighboursHandler(CellSeed** cellsLayersDevice,
861849
cellNeighbours,
862850
sourceCellTopologyId,
863851
targetCellTopologyId,
864-
topology,
865852
maxChi2ClusterAttachment,
866853
bz,
867-
nCells,
868-
maxCellNeighbours);
854+
nCells);
869855
}
870856

871857
int filterCellNeighboursHandler(gpuPair<int, int>* cellNeighbourPairs,
@@ -885,10 +871,8 @@ int filterCellNeighboursHandler(gpuPair<int, int>* cellNeighbourPairs,
885871
}
886872

887873
template <int NLayers>
888-
void processNeighboursHandler(const int startLayer,
889-
const int startLevel,
874+
void processNeighboursHandler(const int startLevel,
890875
const int defaultCellTopologyId,
891-
const typename TrackingTopology<NLayers>::View topology,
892876
CellSeed** allCellSeeds,
893877
CellSeed* currentCellSeeds,
894878
const int* currentCellTopologyIds,
@@ -909,7 +893,6 @@ void processNeighboursHandler(const int startLayer,
909893
const o2::base::PropagatorF::MatCorrType matCorrType,
910894
o2::its::ExternalAllocator* alloc)
911895
{
912-
(void)startLayer;
913896
constexpr uint64_t Tag = qStr2Tag("ITS_PNH1");
914897
alloc->pushTagOnStack(Tag);
915898
auto allocInt = gpu::TypedAllocator<int>(alloc);
@@ -921,7 +904,6 @@ void processNeighboursHandler(const int startLayer,
921904
gpu::processNeighboursKernel<true, NLayers, CellSeed><<<60, 256>>>(
922905
defaultCellTopologyId,
923906
startLevel,
924-
topology,
925907
allCellSeeds,
926908
currentCellSeeds,
927909
nullptr,
@@ -948,7 +930,6 @@ void processNeighboursHandler(const int startLayer,
948930
gpu::processNeighboursKernel<false, NLayers, CellSeed><<<60, 256>>>(
949931
defaultCellTopologyId,
950932
startLevel,
951-
topology,
952933
allCellSeeds,
953934
currentCellSeeds,
954935
nullptr,
@@ -988,7 +969,6 @@ void processNeighboursHandler(const int startLayer,
988969
gpu::processNeighboursKernel<true, NLayers, TrackSeed<NLayers>><<<60, 256>>>(
989970
constants::UnusedIndex,
990971
level,
991-
topology,
992972
allCellSeeds,
993973
thrust::raw_pointer_cast(&lastCellSeed[0]),
994974
thrust::raw_pointer_cast(&lastCellId[0]),
@@ -1020,7 +1000,6 @@ void processNeighboursHandler(const int startLayer,
10201000
gpu::processNeighboursKernel<false, NLayers, TrackSeed<NLayers>><<<60, 256>>>(
10211001
constants::UnusedIndex,
10221002
level,
1023-
topology,
10241003
allCellSeeds,
10251004
thrust::raw_pointer_cast(&lastCellSeed[0]),
10261005
thrust::raw_pointer_cast(&lastCellId[0]),
@@ -1253,11 +1232,9 @@ template void countCellNeighboursHandler<7>(CellSeed** cellsLayersDevice,
12531232
int** cellsLUTs,
12541233
const int sourceCellTopologyId,
12551234
const int targetCellTopologyId,
1256-
const TrackingTopology<7>::View topology,
12571235
const float maxChi2ClusterAttachment,
12581236
const float bz,
12591237
const unsigned int nCells,
1260-
const int maxCellNeighbours,
12611238
gpu::Stream& stream);
12621239

12631240
template void computeCellNeighboursHandler<7>(CellSeed** cellsLayersDevice,
@@ -1266,17 +1243,13 @@ template void computeCellNeighboursHandler<7>(CellSeed** cellsLayersDevice,
12661243
CellNeighbour* cellNeighbours,
12671244
const int sourceCellTopologyId,
12681245
const int targetCellTopologyId,
1269-
const TrackingTopology<7>::View topology,
12701246
const float maxChi2ClusterAttachment,
12711247
const float bz,
12721248
const unsigned int nCells,
1273-
const int maxCellNeighbours,
12741249
gpu::Stream& stream);
12751250

1276-
template void processNeighboursHandler<7>(const int startLayer,
1277-
const int startLevel,
1251+
template void processNeighboursHandler<7>(const int startLevel,
12781252
const int defaultCellTopologyId,
1279-
const TrackingTopology<7>::View topology,
12801253
CellSeed** allCellSeeds,
12811254
CellSeed* currentCellSeeds,
12821255
const int* currentCellTopologyIds,

0 commit comments

Comments
 (0)