Skip to content

Commit 2b154a9

Browse files
mpucciompuccio
committed
Adapt GPU code to the new mult mask
1 parent dd608d6 commit 2b154a9

4 files changed

Lines changed: 13 additions & 10 deletions

File tree

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -143,7 +143,6 @@ class TimeFrameGPU final : public TimeFrame<NLayers>
143143
o2::track::TrackParCovF** getDeviceArrayTrackSeeds() { return mCellSeedsDeviceArray; }
144144
float** getDeviceArrayTrackSeedsChi2() { return mCellSeedsChi2DeviceArray; }
145145
int* getDeviceNeighboursIndexTables(const int layer) { return mNeighboursIndexTablesDevice[layer]; }
146-
uint8_t* getDeviceMultCutMask() { return mMultMaskDevice; }
147146

148147
void setDevicePropagator(const o2::base::PropagatorImpl<float>* p) final { this->mPropagatorDevice = p; }
149148

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -36,7 +36,7 @@ class ExternalAllocator;
3636

3737
template <int NLayers = 7>
3838
void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
39-
const uint8_t* multMask,
39+
const typename ROFMaskTable<NLayers>::View& multMask,
4040
const int layer,
4141
const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
4242
const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,
@@ -66,7 +66,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
6666

6767
template <int NLayers = 7>
6868
void computeTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
69-
const uint8_t* multMask,
69+
const typename ROFMaskTable<NLayers>::View& multMask,
7070
const int layer,
7171
const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
7272
const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -85,7 +85,7 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
8585
mTimeFrameGPU->createTrackletsLUTDevice(iteration, iLayer);
8686
mTimeFrameGPU->waitEvent(iLayer, iLayer + 1); // wait stream until all data is available
8787
countTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
88-
mTimeFrameGPU->getDeviceMultCutMask(),
88+
mTimeFrameGPU->getDeviceROFMaskTableView(),
8989
iLayer,
9090
mTimeFrameGPU->getDeviceROFOverlapTableView(),
9191
mTimeFrameGPU->getDeviceROFVertexLookupTableView(),
@@ -117,7 +117,7 @@ void TrackerTraitsGPU<NLayers>::computeLayerTracklets(const int iteration, int i
117117
continue;
118118
}
119119
computeTrackletsInROFsHandler<NLayers>(mTimeFrameGPU->getDeviceIndexTableUtils(),
120-
mTimeFrameGPU->getDeviceMultCutMask(),
120+
mTimeFrameGPU->getDeviceROFMaskTableView(),
121121
iLayer,
122122
mTimeFrameGPU->getDeviceROFOverlapTableView(),
123123
mTimeFrameGPU->getDeviceROFVertexLookupTableView(),

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

Lines changed: 9 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -536,7 +536,7 @@ GPUg() void __launch_bounds__(256, 1) computeLayerCellsKernel(
536536
template <bool initRun, int NLayers>
537537
GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel(
538538
const IndexTableUtils<NLayers>* utils,
539-
const uint8_t* multMask,
539+
const typename ROFMaskTable<NLayers>::View multMask,
540540
const int layerIndex,
541541
const typename ROFOverlapTable<NLayers>::View rofOverlaps,
542542
const typename ROFVertexLookupTable<NLayers>::View vertexLUT,
@@ -565,6 +565,10 @@ GPUg() void __launch_bounds__(256, 1) computeLayerTrackletsMultiROFKernel(
565565
const int totalROFs0 = rofOverlaps.getLayer(layerIndex).mNROFsTF;
566566
const int totalROFs1 = rofOverlaps.getLayer(layerIndex + 1).mNROFsTF;
567567
for (unsigned int pivotROF{blockIdx.x}; pivotROF < totalROFs0; pivotROF += gridDim.x) {
568+
if (!multMask.isROFEnabled(layerIndex, pivotROF)) {
569+
continue;
570+
}
571+
568572
const auto& pvs = vertexLUT.getVertices(layerIndex, pivotROF);
569573
auto primaryVertices = gpuSpan<const Vertex>(&vertices[pvs.getFirstEntry()], pvs.getEntries());
570574
if (primaryVertices.empty()) {
@@ -782,7 +786,7 @@ GPUg() void __launch_bounds__(256, 1) processNeighboursKernel(
782786

783787
template <int NLayers>
784788
void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
785-
const uint8_t* multMask,
789+
const typename ROFMaskTable<NLayers>::View& multMask,
786790
const int layer,
787791
const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
788792
const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,
@@ -840,7 +844,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
840844

841845
template <int NLayers>
842846
void computeTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
843-
const uint8_t* multMask,
847+
const typename ROFMaskTable<NLayers>::View& multMask,
844848
const int layer,
845849
const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
846850
const typename ROFVertexLookupTable<NLayers>::View& vertexLUT,
@@ -1300,7 +1304,7 @@ void computeTrackSeedHandler(CellSeed<NLayers>* trackSeeds,
13001304

13011305
/// Explicit instantiation of ITS2 handlers
13021306
template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
1303-
const uint8_t* multMask,
1307+
const ROFMaskTable<7>::View& multMask,
13041308
const int layer,
13051309
const ROFOverlapTable<7>::View& rofOverlaps,
13061310
const ROFVertexLookupTable<7>::View& vertexLUT,
@@ -1329,7 +1333,7 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
13291333
gpu::Streams& streams);
13301334

13311335
template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils,
1332-
const uint8_t* multMask,
1336+
const ROFMaskTable<7>::View& multMask,
13331337
const int layer,
13341338
const ROFOverlapTable<7>::View& rofOverlaps,
13351339
const ROFVertexLookupTable<7>::View& vertexLUT,

0 commit comments

Comments
 (0)