Rename elements_with_stride to uniform_elements in user code

bsunanda · Feb 12, 2024 · 008ca51 · 008ca51
1 parent 1bd5d4a
commit 008ca51
Show file tree

Hide file tree

Showing 25 changed files with 117 additions and 117 deletions.
diff --git a/DataFormats/SiPixelClusterSoA/test/alpaka/Clusters_test.dev.cc b/DataFormats/SiPixelClusterSoA/test/alpaka/Clusters_test.dev.cc
@@ -15,7 +15,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
     public:
       template <typename TAcc, typename = std::enable_if_t<isAccelerator<TAcc>>>
       ALPAKA_FN_ACC void operator()(TAcc const& acc, SiPixelClustersSoAView clust_view) const {
-        for (int32_t j : elements_with_stride(acc, clust_view.metadata().size())) {
+        for (int32_t j : uniform_elements(acc, clust_view.metadata().size())) {
           clust_view[j].moduleStart() = j;
           clust_view[j].clusInModule() = j * 2;
           clust_view[j].moduleId() = j * 3;
@@ -28,7 +28,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
     public:
       template <typename TAcc, typename = std::enable_if_t<isAccelerator<TAcc>>>
       ALPAKA_FN_ACC void operator()(TAcc const& acc, SiPixelClustersSoAConstView clust_view) const {
-        for (uint32_t j : elements_with_stride(acc, clust_view.metadata().size())) {
+        for (uint32_t j : uniform_elements(acc, clust_view.metadata().size())) {
           assert(clust_view[j].moduleStart() == j);
           assert(clust_view[j].clusInModule() == j * 2);
           assert(clust_view[j].moduleId() == j * 3);

diff --git a/DataFormats/TrackSoA/test/alpaka/TrackSoAHeterogeneous_test.dev.cc b/DataFormats/TrackSoA/test/alpaka/TrackSoAHeterogeneous_test.dev.cc
@@ -27,7 +27,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
           tracks_view.nTracks() = nTracks;
         }
 
-        for (int32_t j : elements_with_stride(acc, nTracks)) {
+        for (int32_t j : uniform_elements(acc, nTracks)) {
           tracks_view[j].pt() = (float)j;
           tracks_view[j].eta() = (float)j;
           tracks_view[j].chi2() = (float)j;
@@ -50,7 +50,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
         if (cms::alpakatools::once_per_grid(acc)) {
           ALPAKA_ASSERT(tracks_view.nTracks() == nTracks);
         }
-        for (int32_t j : elements_with_stride(acc, tracks_view.nTracks())) {
+        for (int32_t j : uniform_elements(acc, tracks_view.nTracks())) {
           ALPAKA_ASSERT(abs(tracks_view[j].pt() - (float)j) < .0001);
           ALPAKA_ASSERT(abs(tracks_view[j].eta() - (float)j) < .0001);
           ALPAKA_ASSERT(abs(tracks_view[j].chi2() - (float)j) < .0001);

diff --git a/DataFormats/TrackingRecHitSoA/test/alpaka/Hits_test.dev.cc b/DataFormats/TrackingRecHitSoA/test/alpaka/Hits_test.dev.cc
@@ -45,7 +45,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
         }
 
         // can be increased to soa.nHits() for debugging
-        for (uint32_t i : cms::alpakatools::elements_with_stride(acc, 10)) {
+        for (uint32_t i : cms::alpakatools::uniform_elements(acc, 10)) {
           printf("iPhi %d -> %d\n", i, soa[i].iphi());
         }
       }

diff --git a/HeterogeneousCore/AlpakaInterface/interface/HistoContainer.h b/HeterogeneousCore/AlpakaInterface/interface/HistoContainer.h
@@ -25,7 +25,7 @@ namespace cms::alpakatools {
                                   T const *__restrict__ v,
                                   uint32_t const *__restrict__ offsets) const {
       const uint32_t nt = offsets[nh];
-      for (uint32_t i : elements_with_stride(acc, nt)) {
+      for (uint32_t i : uniform_elements(acc, nt)) {
         auto off = alpaka_std::upper_bound(offsets, offsets + nh + 1, i);
         ALPAKA_ASSERT_OFFLOAD((*off) > 0);
         int32_t ih = off - offsets - 1;
@@ -44,7 +44,7 @@ namespace cms::alpakatools {
                                   T const *__restrict__ v,
                                   uint32_t const *__restrict__ offsets) const {
       const uint32_t nt = offsets[nh];
-      for (uint32_t i : elements_with_stride(acc, nt)) {
+      for (uint32_t i : uniform_elements(acc, nt)) {
         auto off = alpaka_std::upper_bound(offsets, offsets + nh + 1, i);
         ALPAKA_ASSERT_OFFLOAD((*off) > 0);
         int32_t ih = off - offsets - 1;

diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testAtomicPairCounter.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testAtomicPairCounter.dev.cc
@@ -18,7 +18,7 @@ struct update {
   template <typename TAcc>
   ALPAKA_FN_ACC void operator()(
       const TAcc &acc, AtomicPairCounter *dc, uint32_t *ind, uint32_t *cont, uint32_t n) const {
-    for (auto i : elements_with_stride(acc, n)) {
+    for (auto i : uniform_elements(acc, n)) {
       auto m = i % 11;
       m = m % 6 + 1;  // max 6, no 0
       auto c = dc->inc_add(acc, m);

diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneHistoContainer.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneHistoContainer.dev.cc
@@ -29,19 +29,19 @@ struct mykernel {
     auto& ws = alpaka::declareSharedVar<typename Hist::Counter[32], __COUNTER__>(acc);
 
     // set off zero
-    for (auto j : elements_with_stride(acc, Hist::totbins())) {
+    for (auto j : uniform_elements(acc, Hist::totbins())) {
       hist.off[j] = 0;
     }
     alpaka::syncBlockThreads(acc);
 
     // set bins zero
-    for (auto j : elements_with_stride(acc, Hist::totbins())) {
+    for (auto j : uniform_elements(acc, Hist::totbins())) {
       hist.content[j] = 0;
     }
     alpaka::syncBlockThreads(acc);
 
     // count
-    for (auto j : elements_with_stride(acc, N)) {
+    for (auto j : uniform_elements(acc, N)) {
       hist.count(acc, v[j]);
     }
     alpaka::syncBlockThreads(acc);
@@ -56,18 +56,18 @@ struct mykernel {
     ALPAKA_ASSERT_OFFLOAD(N == hist.size());
 
     // verify
-    for ([[maybe_unused]] auto j : elements_with_stride(acc, Hist::nbins())) {
+    for ([[maybe_unused]] auto j : uniform_elements(acc, Hist::nbins())) {
       ALPAKA_ASSERT_OFFLOAD(hist.off[j] <= hist.off[j + 1]);
     }
     alpaka::syncBlockThreads(acc);
 
-    for (auto j : elements_with_stride(acc, 32)) {
+    for (auto j : uniform_elements(acc, 32)) {
       ws[j] = 0;  // used by prefix scan...
     }
     alpaka::syncBlockThreads(acc);
 
     // fill
-    for (auto j : elements_with_stride(acc, N)) {
+    for (auto j : uniform_elements(acc, N)) {
       hist.fill(acc, v[j], j);
     }
     alpaka::syncBlockThreads(acc);
@@ -77,7 +77,7 @@ struct mykernel {
 
     // bin
 #ifndef NDEBUG
-    for (auto j : elements_with_stride(acc, hist.size() - 1)) {
+    for (auto j : uniform_elements(acc, hist.size() - 1)) {
       auto p = hist.begin() + j;
       ALPAKA_ASSERT_OFFLOAD((*p) < N);
       [[maybe_unused]] auto k1 = Hist::bin(v[*p]);
@@ -87,7 +87,7 @@ struct mykernel {
 #endif
 
     // forEachInWindow
-    for (auto i : elements_with_stride(acc, hist.size())) {
+    for (auto i : uniform_elements(acc, hist.size())) {
       auto p = hist.begin() + i;
       auto j = *p;
 #ifndef NDEBUG

diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneRadixSort.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneRadixSort.dev.cc
@@ -85,7 +85,7 @@ namespace {
       //      __shared__ uint16_t sws[2048];
       //      __shared__ float z[2048];
       //      __shared__ int iz[2048];
-      for (auto itrack : elements_with_stride(acc, elements)) {
+      for (auto itrack : uniform_elements(acc, elements)) {
         z[itrack] = gpu_input[itrack];
         iz[itrack] = 10000 * gpu_input[itrack];
         // order[itrack] = itrack;
@@ -95,7 +95,7 @@ namespace {
       alpaka::syncBlockThreads(acc);
 
       //verify
-      for (auto itrack : elements_with_stride(acc, elements - 1)) {
+      for (auto itrack : uniform_elements(acc, elements - 1)) {
         auto ntrack = order[itrack];
         auto mtrack = order[itrack + 1];
         assert(truncate<2>(z[ntrack]) <= truncate<2>(z[mtrack]));
@@ -123,7 +123,7 @@ namespace {
       radixSort<TAcc, int, 4>(acc, iz, order, sws, elements);
       alpaka::syncBlockThreads(acc);
 
-      for (auto itrack : elements_with_stride(acc, elements - 1)) {
+      for (auto itrack : uniform_elements(acc, elements - 1)) {
         auto ntrack = order[itrack];
         auto mtrack = order[itrack + 1];
         assert(iz[ntrack] <= iz[mtrack]);

diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneToManyAssoc.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneToManyAssoc.dev.cc
@@ -36,7 +36,7 @@ struct countMultiLocal {
                                 TK const* __restrict__ tk,
                                 Multiplicity* __restrict__ assoc,
                                 uint32_t n) const {
-    for (auto i : elements_with_stride(acc, n)) {
+    for (auto i : uniform_elements(acc, n)) {
       auto& local = alpaka::declareSharedVar<Multiplicity::CountersOnly, __COUNTER__>(acc);
       const uint32_t threadIdxLocal(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
       const bool oncePerSharedMemoryAccess = (threadIdxLocal == 0);
@@ -59,7 +59,7 @@ struct countMulti {
                                 TK const* __restrict__ tk,
                                 Multiplicity* __restrict__ assoc,
                                 uint32_t n) const {
-    for (auto i : elements_with_stride(acc, n)) {
+    for (auto i : uniform_elements(acc, n)) {
       assoc->count(acc, 2 + i % 4);
     }
   }
@@ -68,7 +68,7 @@ struct countMulti {
 struct verifyMulti {
   template <typename TAcc>
   ALPAKA_FN_ACC void operator()(const TAcc& acc, Multiplicity* __restrict__ m1, Multiplicity* __restrict__ m2) const {
-    for ([[maybe_unused]] auto i : elements_with_stride(acc, Multiplicity{}.totOnes())) {
+    for ([[maybe_unused]] auto i : uniform_elements(acc, Multiplicity{}.totOnes())) {
       ALPAKA_ASSERT_OFFLOAD(m1->off[i] == m2->off[i]);
     }
   }
@@ -80,7 +80,7 @@ struct count {
                                 TK const* __restrict__ tk,
                                 AssocRandomAccess* __restrict__ assoc,
                                 uint32_t n) const {
-    for (auto i : elements_with_stride(acc, 4 * n)) {
+    for (auto i : uniform_elements(acc, 4 * n)) {
       auto k = i / 4;
       auto j = i - 4 * k;
       ALPAKA_ASSERT_OFFLOAD(j < 4);
@@ -100,7 +100,7 @@ struct fill {
                                 TK const* __restrict__ tk,
                                 AssocRandomAccess* __restrict__ assoc,
                                 uint32_t n) const {
-    for (auto i : elements_with_stride(acc, 4 * n)) {
+    for (auto i : uniform_elements(acc, 4 * n)) {
       auto k = i / 4;
       auto j = i - 4 * k;
       ALPAKA_ASSERT_OFFLOAD(j < 4);
@@ -125,7 +125,7 @@ struct fillBulk {
   template <typename TAcc, typename Assoc>
   ALPAKA_FN_ACC void operator()(
       const TAcc& acc, AtomicPairCounter* apc, TK const* __restrict__ tk, Assoc* __restrict__ assoc, uint32_t n) const {
-    for (auto k : elements_with_stride(acc, n)) {
+    for (auto k : uniform_elements(acc, n)) {
       auto m = tk[k][3] < MaxElem ? 4 : 3;
       assoc->bulkFill(acc, *apc, &tk[k][0], m);
     }

diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testPrefixScan.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testPrefixScan.dev.cc
@@ -34,7 +34,7 @@ struct testPrefixScan {
     auto& c = alpaka::declareSharedVar<T[1024], __COUNTER__>(acc);
     auto& co = alpaka::declareSharedVar<T[1024], __COUNTER__>(acc);
 
-    for (auto i : elements_with_stride(acc, size)) {
+    for (auto i : uniform_elements(acc, size)) {
       c[i] = 1;
     };
 
@@ -49,7 +49,7 @@ struct testPrefixScan {
     // TODO: not needed? Not in multi kernel version, not in CUDA version
     alpaka::syncBlockThreads(acc);
 
-    for (auto i : elements_with_stride(acc, size)) {
+    for (auto i : uniform_elements(acc, size)) {
       if (0 == i)
         continue;
       if constexpr (!std::is_floating_point_v<T>) {
@@ -109,7 +109,7 @@ struct testWarpPrefixScan {
 struct init {
   template <typename TAcc>
   ALPAKA_FN_ACC void operator()(const TAcc& acc, uint32_t* v, uint32_t val, uint32_t n) const {
-    for (auto index : elements_with_stride(acc, n)) {
+    for (auto index : uniform_elements(acc, n)) {
       v[index] = val;
 
       if (index == 0)
@@ -121,7 +121,7 @@ struct init {
 struct verify {
   template <typename TAcc>
   ALPAKA_FN_ACC void operator()(const TAcc& acc, uint32_t const* v, uint32_t n) const {
-    for (auto index : elements_with_stride(acc, n)) {
+    for (auto index : uniform_elements(acc, n)) {
       ALPAKA_ASSERT_OFFLOAD(v[index] == index + 1);
 
       if (index == 0)

diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testSimpleVector.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testSimpleVector.dev.cc
@@ -15,7 +15,7 @@ using namespace ALPAKA_ACCELERATOR_NAMESPACE;
 struct vector_pushback {
   template <typename TAcc>
   ALPAKA_FN_ACC void operator()(const TAcc& acc, SimpleVector<int>* foo) const {
-    for (auto index : elements_with_stride(acc))
+    for (auto index : uniform_elements(acc))
       foo->push_back(acc, index);
   }
 };
@@ -30,7 +30,7 @@ struct vector_reset {
 struct vector_emplace_back {
   template <typename TAcc>
   ALPAKA_FN_ACC void operator()(const TAcc& acc, SimpleVector<int>* foo) const {
-    for (auto index : elements_with_stride(acc))
+    for (auto index : uniform_elements(acc))
       foo->emplace_back(acc, index);
   }
 };

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/CalibPixel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/CalibPixel.h
@@ -41,11 +41,11 @@ namespace calibPixel {
         clus_view[0].clusModuleStart() = 0;
         clus_view[0].moduleStart() = 0;
       }
-      for (auto i : cms::alpakatools::elements_with_stride(acc, phase1PixelTopology::numberOfModules)) {
+      for (auto i : cms::alpakatools::uniform_elements(acc, phase1PixelTopology::numberOfModules)) {
         clus_view[i].clusInModule() = 0;
       }
 
-      for (auto i : cms::alpakatools::elements_with_stride(acc, numElements)) {
+      for (auto i : cms::alpakatools::uniform_elements(acc, numElements)) {
         auto dvgi = view[i];
         if (dvgi.moduleId() == ::pixelClustering::invalidModuleId)
           continue;
@@ -106,11 +106,11 @@ namespace calibPixel {
         clus_view[0].clusModuleStart() = clus_view[0].moduleStart() = 0;
       }
 
-      for (uint32_t i : cms::alpakatools::elements_with_stride(acc, phase2PixelTopology::numberOfModules)) {
+      for (uint32_t i : cms::alpakatools::uniform_elements(acc, phase2PixelTopology::numberOfModules)) {
         clus_view[i].clusInModule() = 0;
       }
 
-      for (uint32_t i : cms::alpakatools::elements_with_stride(acc, numElements)) {
+      for (uint32_t i : cms::alpakatools::uniform_elements(acc, numElements)) {
         auto dvgi = view[i];
         if (pixelClustering::invalidModuleId != dvgi.moduleId()) {
           const int mode = (Phase2ReadoutMode < -1 ? -1 : Phase2ReadoutMode);

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/PixelClustering.h
@@ -104,7 +104,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering {
         printf("Starting to count modules to set module starts:");
       }
 #endif
-      for (int32_t i : cms::alpakatools::elements_with_stride(acc, numElements)) {
+      for (int32_t i : cms::alpakatools::uniform_elements(acc, numElements)) {
         digi_view[i].clus() = i;
         if (::pixelClustering::invalidModuleId == digi_view[i].moduleId())
           continue;

diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/alpaka/SiPixelRawToClusterKernel.dev.cc
@@ -304,7 +304,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
         if (cms::alpakatools::once_per_grid(acc))
           err.size() = 0;
 
-        for (auto gIndex : cms::alpakatools::elements_with_stride(acc, wordCounter)) {
+        for (auto gIndex : cms::alpakatools::uniform_elements(acc, wordCounter)) {
           auto dvgi = digisView[gIndex];
           dvgi.xx() = 0;
           dvgi.yy() = 0;

diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitKernels.dev.cc b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitKernels.dev.cc
@@ -36,7 +36,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
                                   uint32_t* __restrict__ hitsLayerStart) const {
       assert(0 == hitsModuleStart[0]);
 
-      for (int32_t i : cms::alpakatools::elements_with_stride(acc, TrackerTraits::numberOfLayers + 1)) {
+      for (int32_t i : cms::alpakatools::uniform_elements(acc, TrackerTraits::numberOfLayers + 1)) {
         hitsLayerStart[i] = hitsModuleStart[cpeParams->layerGeometry().layerStart[i]];
 #ifdef GPU_DEBUG
         int old = i == 0 ? 0 : hitsModuleStart[cpeParams->layerGeometry().layerStart[i - 1]];

diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc b/RecoTracker/PixelSeeding/plugins/alpaka/BrokenLineFit.dev.cc
@@ -62,7 +62,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
       }
 #endif
       const auto nt = riemannFit::maxNumberOfConcurrentFits;
-      for (auto local_idx : cms::alpakatools::elements_with_stride(acc, nt)) {
+      for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
         auto tuple_idx = local_idx + offset;
         if ((int)tuple_idx >= totTK) {
           ptkids[local_idx] = invalidTkId;
@@ -190,7 +190,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
       // same as above...
       // look in bin for this hit multiplicity
       const auto nt = riemannFit::maxNumberOfConcurrentFits;
-      for (auto local_idx : cms::alpakatools::elements_with_stride(acc, nt)) {
+      for (auto local_idx : cms::alpakatools::uniform_elements(acc, nt)) {
         if (invalidTkId == ptkids[local_idx])
           break;
         auto tkid = ptkids[local_idx];