Skip to content

Commit d216d22

Browse files
antoniopetreantoniopetre
committed
[alpaka] Refactor CAHitNtupletGeneratorKernelsImpl using elements_with_stride class
1 parent 88482d4 commit d216d22

1 file changed

Lines changed: 55 additions & 46 deletions

File tree

‎src/alpaka/plugin-PixelTriplets/alpaka/CAHitNtupletGeneratorKernelsImpl.h‎

Lines changed: 55 additions & 46 deletions
Original file line numberDiff line numberDiff line change
@@ -72,13 +72,14 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
7272
}
7373

7474
const auto ntNbins = foundNtuplets->nbins();
75-
cms::alpakatools::for_each_element_in_grid_strided(acc, ntNbins, [&](uint32_t idx) {
75+
76+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, ntNbins)) {
7677
if (foundNtuplets->size(idx) > 5)
7778
printf("ERROR %d, %d\n", idx, foundNtuplets->size(idx));
7879
assert(foundNtuplets->size(idx) < 6);
7980
for (auto ih = foundNtuplets->begin(idx); ih != foundNtuplets->end(idx); ++ih)
8081
assert(*ih < nHits);
81-
});
82+
}
8283
#endif
8384

8485
if (0 == threadIdx) {
@@ -93,7 +94,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
9394
}
9495

9596
const auto ntNCells = (*nCells);
96-
cms::alpakatools::for_each_element_in_grid_strided(acc, ntNCells, [&](uint32_t idx) {
97+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, ntNCells)) {
9798
auto const &thisCell = cells[idx];
9899
if (thisCell.outerNeighbors().full()) //++tooManyNeighbors[thisCell.theLayerPairId];
99100
printf("OuterNeighbors overflow %d in %d\n", idx, thisCell.theLayerPairId);
@@ -105,12 +106,12 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
105106
alpaka::atomicAdd(acc, &c.nEmptyCells, 1ull, alpaka::hierarchy::Blocks{});
106107
if (thisCell.tracks().empty())
107108
alpaka::atomicAdd(acc, &c.nZeroTrackCells, 1ull, alpaka::hierarchy::Blocks{});
108-
});
109+
}
109110

110-
cms::alpakatools::for_each_element_in_grid_strided(acc, nHits, [&](uint32_t idx) {
111+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, nHits)) {
111112
if (isOuterHitOfCell[idx].full()) // ++tooManyOuterHitOfCell;
112113
printf("OuterHitOfCell overflow %d\n", idx);
113-
});
114+
}
114115
}
115116
};
116117

@@ -123,14 +124,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
123124
constexpr auto bad = trackQuality::bad;
124125

125126
const auto ntNCells = (*nCells);
126-
cms::alpakatools::for_each_element_in_grid_strided(acc, ntNCells, [&](uint32_t idx) {
127+
128+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, ntNCells)) {
127129
auto const &thisCell = cells[idx];
128130

129131
if (thisCell.theDoubletId < 0) {
130132
for (auto it : thisCell.tracks())
131133
quality[it] = bad;
132134
}
133-
});
135+
}
134136
}
135137
};
136138

@@ -147,7 +149,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
147149

148150
assert(nCells);
149151
const auto ntNCells = (*nCells);
150-
cms::alpakatools::for_each_element_in_grid_strided(acc, ntNCells, [&](uint32_t idx) {
152+
153+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, ntNCells)) {
151154
auto const &thisCell = cells[idx];
152155

153156
if (thisCell.tracks().size() >= 2) {
@@ -167,7 +170,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
167170
quality[it] = dup; //no race: simple assignment of the same constant
168171
}
169172
}
170-
});
173+
}
171174
}
172175
};
173176

@@ -184,7 +187,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
184187

185188
assert(nCells);
186189

187-
cms::alpakatools::for_each_element_in_grid_strided(acc, (*nCells), [&](uint32_t idx) {
190+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, *nCells)) {
188191
auto const &thisCell = cells[idx];
189192
if (thisCell.tracks().size() >= 2) {
190193
// if (thisCell.theDoubletId < 0) continue;
@@ -210,7 +213,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
210213
tracks->quality(it) = dup; //no race: simple assignment of the same constant
211214
}
212215
}
213-
});
216+
}
214217
}
215218
};
216219

@@ -232,6 +235,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
232235
float dcaCutOuterTriplet) const {
233236
auto const &hh = *hhp;
234237

238+
const Idx elementShift = 0;
235239
const uint32_t dimIndexY = 0u;
236240
const uint32_t dimIndexX = 1u;
237241
const uint32_t threadIdxY(alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[dimIndexY]);
@@ -242,11 +246,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
242246
(*apc2) = 0;
243247
} // ready for next kernel
244248

245-
cms::alpakatools::for_each_element_in_grid_strided(
246-
acc,
247-
(*nCells),
248-
0u,
249-
[&](uint32_t idx) {
249+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, *nCells, elementShift, dimIndexY)) {
250250
auto cellIndex = idx;
251251
auto &thisCell = cells[idx];
252252
//if (thisCell.theDoubletId < 0 || thisCell.theUsed>1)
@@ -297,8 +297,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
297297
}
298298
},
299299
dimIndexX); // loop on inner cells
300-
},
301-
dimIndexY); // loop on outer cells
300+
} // loop on outer cells
302301
}
303302
};
304303

@@ -318,7 +317,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
318317

319318
//auto first = threadIdx.x + blockIdx.x * blockDim.x;
320319
//for (int idx = first, nt = (*nCells); idx < nt; idx += gridDim.x * blockDim.x) {
321-
cms::alpakatools::for_each_element_in_grid_strided(acc, (*nCells), [&](uint32_t idx) {
320+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, *nCells)) {
322321
auto const &thisCell = cells[idx];
323322
if (thisCell.theDoubletId >= 0) { // cut by earlyFishbone
324323

@@ -333,7 +332,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
333332
// printf("in %d found quadruplets: %d\n", cellIndex, apc->get());
334333
}
335334
}
336-
});
335+
}
337336
}
338337
};
339338

@@ -344,11 +343,12 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
344343
GPUCACell *__restrict__ cells,
345344
uint32_t const *nCells) const {
346345
// auto const &hh = *hhp;
347-
cms::alpakatools::for_each_element_in_grid_strided(acc, (*nCells), [&](uint32_t idx) {
346+
347+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, *nCells)) {
348348
auto &thisCell = cells[idx];
349349
if (!thisCell.tracks().empty())
350350
thisCell.theUsed |= 2;
351-
});
351+
}
352352
}
353353
};
354354

@@ -358,7 +358,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
358358
HitContainer const *__restrict__ foundNtuplets,
359359
Quality const *__restrict__ quality,
360360
CAConstants::TupleMultiplicity *tupleMultiplicity) const {
361-
cms::alpakatools::for_each_element_in_grid_strided(acc, foundNtuplets->nbins(), [&](uint32_t it) {
361+
362+
for (uint32_t it : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, foundNtuplets->nbins())) {
362363
auto nhits = foundNtuplets->size(it);
363364
if (nhits >= 3 && quality[it] != trackQuality::dup) {
364365
assert(quality[it] == trackQuality::bad);
@@ -367,7 +368,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
367368
assert(nhits < 8);
368369
tupleMultiplicity->countDirect(acc, nhits);
369370
}
370-
});
371+
}
371372
}
372373
};
373374

@@ -377,7 +378,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
377378
HitContainer const *__restrict__ foundNtuplets,
378379
Quality const *__restrict__ quality,
379380
CAConstants::TupleMultiplicity *tupleMultiplicity) const {
380-
cms::alpakatools::for_each_element_in_grid_strided(acc, foundNtuplets->nbins(), [&](uint32_t it) {
381+
382+
for (uint32_t it : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, foundNtuplets->nbins())) {
381383
auto nhits = foundNtuplets->size(it);
382384
if (nhits >= 3 && quality[it] != trackQuality::dup) {
383385
assert(quality[it] == trackQuality::bad);
@@ -386,7 +388,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
386388
assert(nhits < 8);
387389
tupleMultiplicity->fillDirect(acc, nhits, it);
388390
}
389-
});
391+
}
390392
}
391393
};
392394

@@ -397,7 +399,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
397399
TkSoA const *__restrict__ tracks,
398400
CAHitNtupletGeneratorKernels::QualityCuts cuts,
399401
Quality *__restrict__ quality) const {
400-
cms::alpakatools::for_each_element_in_grid_strided(acc, tuples->nbins(), [&](uint32_t it) {
402+
403+
for (uint32_t it : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, tuples->nbins())) {
401404
auto nhits = tuples->size(it);
402405
if (nhits == 0)
403406
return; // guard
@@ -453,7 +456,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
453456
} // chi2Cut
454457
} // !isNaN
455458
} // trackQuality and nhits
456-
});
459+
}
457460
}
458461
};
459462

@@ -463,13 +466,14 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
463466
HitContainer const *__restrict__ tuples,
464467
Quality const *__restrict__ quality,
465468
CAHitNtupletGeneratorKernels::Counters *counters) const {
466-
cms::alpakatools::for_each_element_in_grid_strided(acc, tuples->nbins(), [&](uint32_t idx) {
469+
470+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, tuples->nbins())) {
467471
if (tuples->size(idx) == 0)
468472
return; //guard
469473
if (quality[idx] == trackQuality::loose) {
470474
alpaka::atomicAdd(acc, &(counters->nGoodTracks), 1ull, alpaka::hierarchy::Blocks{});
471475
}
472-
});
476+
}
473477
}
474478
};
475479

@@ -479,14 +483,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
479483
HitContainer const *__restrict__ tuples,
480484
Quality const *__restrict__ quality,
481485
CAHitNtupletGeneratorKernels::HitToTuple *hitToTuple) const {
482-
cms::alpakatools::for_each_element_in_grid_strided(acc, tuples->nbins(), [&](uint32_t idx) {
486+
487+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, tuples->nbins())) {
483488
if (tuples->size(idx) == 0)
484489
return; // guard
485490
if (quality[idx] == trackQuality::loose) {
486491
for (auto h = tuples->begin(idx); h != tuples->end(idx); ++h)
487492
hitToTuple->countDirect(acc, *h);
488493
}
489-
});
494+
}
490495
}
491496
};
492497

@@ -496,14 +501,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
496501
HitContainer const *__restrict__ tuples,
497502
Quality const *__restrict__ quality,
498503
CAHitNtupletGeneratorKernels::HitToTuple *hitToTuple) const {
499-
cms::alpakatools::for_each_element_in_grid_strided(acc, tuples->nbins(), [&](uint32_t idx) {
504+
505+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, tuples->nbins())) {
500506
if (tuples->size(idx) == 0)
501507
return; // guard
502508
if (quality[idx] == trackQuality::loose) {
503509
for (auto h = tuples->begin(idx); h != tuples->end(idx); ++h)
504510
hitToTuple->fillDirect(acc, *h, idx);
505511
}
506-
});
512+
}
507513
}
508514
};
509515

@@ -514,17 +520,18 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
514520
TrackingRecHit2DSOAView const *__restrict__ hhp,
515521
HitContainer *__restrict__ hitDetIndices) const {
516522
// copy offsets
517-
cms::alpakatools::for_each_element_in_grid_strided(
518-
acc, tuples->totbins(), [&](uint32_t idx) { hitDetIndices->off[idx] = tuples->off[idx]; });
523+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, tuples->totbins())) {
524+
hitDetIndices->off[idx] = tuples->off[idx];
525+
}
519526
// fill hit indices
520527
auto const &hh = *hhp;
521528
#ifndef NDEBUG
522529
auto nhits = hh.nHits();
523530
#endif
524-
cms::alpakatools::for_each_element_in_grid_strided(acc, tuples->size(), [&](uint32_t idx) {
531+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, tuples->size())) {
525532
assert(tuples->bins[idx] < nhits);
526533
hitDetIndices->bins[idx] = hh.detectorIndex(tuples->bins[idx]);
527-
});
534+
}
528535
}
529536
};
530537

@@ -534,13 +541,14 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
534541
CAHitNtupletGeneratorKernels::HitToTuple const *__restrict__ hitToTuple,
535542
CAHitNtupletGeneratorKernels::Counters *counters) const {
536543
auto &c = *counters;
537-
cms::alpakatools::for_each_element_in_grid_strided(acc, hitToTuple->nbins(), [&](uint32_t idx) {
544+
545+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, hitToTuple->nbins())) {
538546
if (hitToTuple->size(idx) != 0) { // SHALL NOT BE break
539547
alpaka::atomicAdd(acc, &c.nUsedHits, 1ull, alpaka::hierarchy::Blocks{});
540548
if (hitToTuple->size(idx) > 1)
541549
alpaka::atomicAdd(acc, &c.nDupHits, 1ull, alpaka::hierarchy::Blocks{});
542550
}
543-
});
551+
}
544552
}
545553
};
546554

@@ -563,7 +571,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
563571
// auto const & hh = *hhp;
564572
// auto l1end = hh.hitsLayerStart_d[1];
565573

566-
cms::alpakatools::for_each_element_in_grid_strided(acc, phitToTuple->nbins(), [&](uint32_t idx) {
574+
for (uint32_t idx : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, phitToTuple->nbins())) {
567575
if (hitToTuple.size(idx) >= 2) {
568576
float mc = 10000.f;
569577
uint16_t im = 60000;
@@ -600,7 +608,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
600608

601609
} // maxNh
602610
} // hitToTuple.size
603-
}); // loop over hits
611+
} // loop over hits
604612
}
605613
};
606614

@@ -617,7 +625,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
617625
auto const &foundNtuplets = *ptuples;
618626
auto const &tracks = *ptracks;
619627
const auto np = std::min(maxPrint, foundNtuplets.nbins());
620-
cms::alpakatools::for_each_element_in_grid_strided(acc, np, [&](uint32_t i) {
628+
629+
for (uint32_t i : cms::alpakatools::elements_with_stride<uint32_t, T_Acc>(acc, np)) {
621630
auto nh = foundNtuplets.size(i);
622631
if (nh >= 3) {
623632
printf("TK: %d %d %d %f %f %f %f %f %f %f %d %d %d %d %d\n",
@@ -638,7 +647,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
638647
nh > 3 ? int(*(foundNtuplets.begin(i) + 3)) : -1,
639648
nh > 4 ? int(*(foundNtuplets.begin(i) + 4)) : -1);
640649
} // nh
641-
});
650+
}
642651
}
643652
};
644653

@@ -677,4 +686,4 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
677686
}
678687
};
679688

680-
} // namespace ALPAKA_ACCELERATOR_NAMESPACE
689+
} // namespace ALPAKA_ACCELERATOR_NAMESPACE

0 commit comments

Comments
 (0)