Migrate assert to to ALPAKA_ASSERT_ACC [14.0.x]

CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h

@@ -2,7 +2,9 @@
 #define CondFormats_SiPixelObjects_interface_alpaka_SiPixelGainCalibrationForHLTUtilities_h
 
 #include <cstdint>
+
 #include <alpaka/alpaka.hpp>
+
 #include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLTLayout.h"
 
 struct SiPixelGainUtilities {
@@ -22,9 +24,9 @@ struct SiPixelGainUtilities {
     unsigned int numberOfDataBlocksToSkip = row / view.numberOfRowsAveragedOver();
 
     auto offset = start + col * lengthOfColumnData + lengthOfAveragedDataInEachColumn * numberOfDataBlocksToSkip;
-    assert(offset < end);
-    assert(offset < 3088384);
-    assert(0 == offset % 2);
+    ALPAKA_ASSERT_ACC(offset < end);
+    ALPAKA_ASSERT_ACC(offset < 3088384);
+    ALPAKA_ASSERT_ACC(0 == offset % 2);
 
     auto lp = view.v_pedestals();
     auto s = lp[offset / 2];
@@ -38,4 +40,4 @@ struct SiPixelGainUtilities {
   };
 };
 
-#endif  //CondFormats_SiPixelObjects_interface_alpaka_SiPixelGainCalibrationForHLTUtilities_h
+#endif  // CondFormats_SiPixelObjects_interface_alpaka_SiPixelGainCalibrationForHLTUtilities_h

DataFormats/SiPixelClusterSoA/test/alpaka/Clusters_test.dev.cc

@@ -31,10 +31,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::testClusterSoA {
     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 : cms::alpakatools::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);
-        assert(clust_view[j].clusModuleStart() == j * 4);
+        ALPAKA_ASSERT_ACC(clust_view[j].moduleStart() == j);
+        ALPAKA_ASSERT_ACC(clust_view[j].clusInModule() == j * 2);
+        ALPAKA_ASSERT_ACC(clust_view[j].moduleId() == j * 3);
+        ALPAKA_ASSERT_ACC(clust_view[j].clusModuleStart() == j * 4);
       }
     }
   };

DataFormats/SiPixelDigiSoA/test/alpaka/DigiErrors_test.dev.cc

@@ -32,10 +32,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::testDigisSoA {
     template <typename TAcc, typename = std::enable_if_t<isAccelerator<TAcc>>>
     ALPAKA_FN_ACC void operator()(TAcc const& acc, SiPixelDigiErrorsSoAConstView digiErrors_view) const {
       for (uint32_t j : cms::alpakatools::uniform_elements(acc, digiErrors_view.metadata().size())) {
-        assert(digiErrors_view[j].pixelErrors().rawId == j);
-        assert(digiErrors_view[j].pixelErrors().word == j);
-        assert(digiErrors_view[j].pixelErrors().errorType == j % 256);
-        assert(digiErrors_view[j].pixelErrors().fedId == j % 256);
+        ALPAKA_ASSERT_ACC(digiErrors_view[j].pixelErrors().rawId == j);
+        ALPAKA_ASSERT_ACC(digiErrors_view[j].pixelErrors().word == j);
+        ALPAKA_ASSERT_ACC(digiErrors_view[j].pixelErrors().errorType == j % 256);
+        ALPAKA_ASSERT_ACC(digiErrors_view[j].pixelErrors().fedId == j % 256);
       }
     }
   };

DataFormats/SiPixelDigiSoA/test/alpaka/Digis_test.dev.cc

@@ -29,10 +29,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::testDigisSoA {
     template <typename TAcc, typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
     ALPAKA_FN_ACC void operator()(TAcc const& acc, SiPixelDigisSoAConstView digi_view) const {
       for (uint32_t j : cms::alpakatools::uniform_elements(acc, digi_view.metadata().size())) {
-        assert(digi_view[j].clus() == int(j));
-        assert(digi_view[j].rawIdArr() == j * 2);
-        assert(digi_view[j].xx() == j * 3);
-        assert(digi_view[j].moduleId() == j * 4);
+        ALPAKA_ASSERT_ACC(digi_view[j].clus() == int(j));
+        ALPAKA_ASSERT_ACC(digi_view[j].rawIdArr() == j * 2);
+        ALPAKA_ASSERT_ACC(digi_view[j].xx() == j * 3);
+        ALPAKA_ASSERT_ACC(digi_view[j].moduleId() == j * 4);
       }
     }
   };

DataFormats/VertexSoA/test/alpaka/ZVertexSoA_test.dev.cc

@@ -37,13 +37,13 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::testZVertexSoAT {
       }
 
       for (int32_t j : cms::alpakatools::uniform_elements(acc, zvertex_view.nvFinal())) {
-        assert(zvertex_view[j].idv() == j);
-        assert(zvertex_view[j].zv() - (float)j < 0.0001);
-        assert(zvertex_view[j].wv() - (float)j < 0.0001);
-        assert(zvertex_view[j].chi2() - (float)j < 0.0001);
-        assert(zvertex_view[j].ptv2() - (float)j < 0.0001);
-        assert(zvertex_view[j].ndof() == j);
-        assert(zvertex_view[j].sortInd() == uint32_t(j));
+        ALPAKA_ASSERT_ACC(zvertex_view[j].idv() == j);
+        ALPAKA_ASSERT_ACC(zvertex_view[j].zv() - (float)j < 0.0001);
+        ALPAKA_ASSERT_ACC(zvertex_view[j].wv() - (float)j < 0.0001);
+        ALPAKA_ASSERT_ACC(zvertex_view[j].chi2() - (float)j < 0.0001);
+        ALPAKA_ASSERT_ACC(zvertex_view[j].ptv2() - (float)j < 0.0001);
+        ALPAKA_ASSERT_ACC(zvertex_view[j].ndof() == j);
+        ALPAKA_ASSERT_ACC(zvertex_view[j].sortInd() == uint32_t(j));
       }
     }
   };

HeterogeneousCore/AlpakaInterface/test/alpaka/testAtomicPairCounter.dev.cc

@@ -25,7 +25,7 @@ struct update {
       auto m = i % 11;
       m = m % 6 + 1;  // max 6, no 0
       auto c = dc->inc_add(acc, m);
-      assert(c.first < n);
+      ALPAKA_ASSERT_ACC(c.first < n);
       ind[c.first] = c.second;
       for (uint32_t j = c.second; j < c.second + m; ++j)
         cont[j] = i;
@@ -37,7 +37,7 @@ struct finalize {
   template <typename TAcc>
   ALPAKA_FN_ACC void operator()(
       const TAcc &acc, AtomicPairCounter const *dc, uint32_t *ind, uint32_t *cont, uint32_t n) const {
-    assert(dc->get().first == n);
+    ALPAKA_ASSERT_ACC(dc->get().first == n);
     ind[n] = dc->get().second;
   }
 };

HeterogeneousCore/AlpakaInterface/test/alpaka/testBuffer.dev.cc

@@ -35,7 +35,7 @@ namespace {
           div,
           [] ALPAKA_FN_ACC(Acc1D const& acc, int* data, size_t size) {
             for (auto index : cms::alpakatools::uniform_elements(acc, size)) {
-              assert(data[index] != 0);
+              ALPAKA_ASSERT_ACC(data[index] != 0);
             }
           },
           buf_d.data(),

HeterogeneousCore/AlpakaInterface/test/alpaka/testOneRadixSort.dev.cc

@@ -75,9 +75,9 @@ namespace {
       // radix sort works in a single block (and the assert macro does not like the comma of the template parameters).
       const auto blocksPerGrid = alpaka::getWorkDiv<alpaka::Grid, alpaka::Blocks>(acc)[0u];
       const auto blocksIdx = alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0u];
-      assert(1 == blocksPerGrid);
-      assert(0 == blocksIdx);
-      assert(elements <= 2048);
+      ALPAKA_ASSERT_ACC(1 == blocksPerGrid);
+      ALPAKA_ASSERT_ACC(0 == blocksIdx);
+      ALPAKA_ASSERT_ACC(elements <= 2048);
 
       auto& order = alpaka::declareSharedVar<uint16_t[2048], __COUNTER__>(acc);
       auto& sws = alpaka::declareSharedVar<uint16_t[2048], __COUNTER__>(acc);
@@ -100,7 +100,7 @@ namespace {
       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]));
+        ALPAKA_ASSERT_ACC(truncate<2>(z[ntrack]) <= truncate<2>(z[mtrack]));
       }
 
       alpaka::syncBlockThreads(acc);
@@ -128,7 +128,7 @@ namespace {
       for (auto itrack : uniform_elements(acc, elements - 1)) {
         auto ntrack = order[itrack];
         auto mtrack = order[itrack + 1];
-        assert(iz[ntrack] <= iz[mtrack]);
+        ALPAKA_ASSERT_ACC(iz[ntrack] <= iz[mtrack]);
       }
 
       if (doPrint)

HeterogeneousCore/AlpakaInterface/test/alpaka/testRadixSort.dev.cc

@@ -87,15 +87,13 @@ void truncate(T& t) {
 template <typename T, int NS = sizeof(T), typename U = T, typename LL = long long>
 void go(Queue& queue, bool useShared) {
   std::mt19937 eng;
-  //std::mt19937 eng2;
   auto rgen = RS<T>::ud();
 
   std::chrono::high_resolution_clock::duration delta = 0ns;
   constexpr int blocks = 10;
   constexpr int blockSize = 256 * 32;
   constexpr int N = blockSize * blocks;
   auto v_h = cms::alpakatools::make_host_buffer<T[]>(queue, N);
-  //uint16_t ind_h[N];
 
   constexpr bool sgn = T(-1) < T(0);
   std::cout << "Will sort " << N << (sgn ? " signed" : " unsigned")
@@ -161,7 +159,6 @@ void go(Queue& queue, bool useShared) {
     auto workdiv = make_workdiv<Acc1D>(blocks, ntXBl);
     if (useShared)
       // The original CUDA version used to call a kernel with __launch_bounds__(256, 4) specifier
-      //
       alpaka::enqueue(queue,
                       alpaka::createTaskKernel<Acc1D>(workdiv,
                                                       radixSortMultiWrapper<U, NS>{},
@@ -225,7 +222,6 @@ void go(Queue& queue, bool useShared) {
                     << offsets_h[ib + 1] - 1 << "] j=" << j << " ind[j]=" << ind_h[j]
                     << " (k1 < k2) : a1=" << (int64_t)a[ind_h[j]] << " k1=" << (int64_t)k1
                     << " a2= " << (int64_t)a[ind_h[j - 1]] << " k2=" << (int64_t)k2 << std::endl;
-          //sleep(2);
           assert(false);
         }
       }
@@ -236,7 +232,6 @@ void go(Queue& queue, bool useShared) {
       if (inds.size() != (offsets_h[ib + 1] - offsets_h[ib]))
         std::cout << "error " << i << ' ' << ib << ' ' << inds.size() << "!=" << (offsets_h[ib + 1] - offsets_h[ib])
                   << std::endl;
-      //
       assert(inds.size() == (offsets_h[ib + 1] - offsets_h[ib]));
     }
   }  // 50 times

HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc

@@ -21,16 +21,6 @@ enum class RangeType { Default, ExtentLimited, ExtentLimitedWithShift };
 // The concurrency scope between threads
 enum class LoopScope { Block, Grid };
 
-// Utility for one time initializations
-template <LoopScope loopScope, typename TAcc>
-bool constexpr firstInLoopRange(TAcc const& acc) {
-  if constexpr (loopScope == LoopScope::Block)
-    return not alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u];
-  if constexpr (loopScope == LoopScope::Grid)
-    return not alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0u];
-  assert(false);
-}
-
 template <RangeType rangeType, LoopScope loopScope, typename TAcc>
 size_t constexpr expectedCount(TAcc const& acc, size_t skip, size_t size) {
   if constexpr (rangeType == RangeType::ExtentLimitedWithShift)
@@ -65,8 +55,9 @@ struct testWordDivisionDefaultRange {
         (loopScope == LoopScope::Grid ? *globalCounter : alpaka::declareSharedVar<size_t, __COUNTER__>(acc));
     // Init the counter for block range. Grid range does so my mean of memset.
     if constexpr (loopScope == LoopScope::Block) {
-      if (firstInLoopRange<loopScope>(acc))
+      if (cms::alpakatools::once_per_block(acc)) {
         counter = 0;
+      }
       alpaka::syncBlockThreads(acc);
     }
     // The loop we are testing
@@ -82,9 +73,9 @@ struct testWordDivisionDefaultRange {
     alpaka::syncBlockThreads(acc);
     // Check the result. Grid range will check by memcpy-ing the result.
     if constexpr (loopScope == LoopScope::Block) {
-      if (firstInLoopRange<loopScope>(acc)) {
+      if (cms::alpakatools::once_per_block(acc)) {
         auto expected = expectedCount<rangeType, loopScope>(acc, skip, size);
-        assert(counter == expected);
+        ALPAKA_ASSERT_ACC(counter == expected);
       }
     }
   }

RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h

@@ -149,7 +149,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit {
       auto const totalElements = nthreads_per_channel * nchannels;
 
       auto const elemsPerBlock = alpaka::getWorkDiv<alpaka::Block, alpaka::Elems>(acc)[0u];
-      assert(nthreads_per_channel * nchannels_per_block == elemsPerBlock);
+      ALPAKA_ASSERT_ACC(nthreads_per_channel * nchannels_per_block == elemsPerBlock);
 
       auto* shr_chi2s = alpaka::getDynSharedMem<ScalarType>(acc);
       auto* shr_time_wgt = shr_chi2s + elemsPerBlock;
@@ -213,8 +213,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit {
           } else if (ltx <= 44) {
             sample_i = 8;
             sample_j = 9;
-          } else
-            assert(false);
+          } else {
+            // FIXME this needs a more portable solution, that wraps abort() / __trap() / throw depending on the back-end
+            ALPAKA_ASSERT_ACC(false);
+          }
 
           auto const tx_i = ch_start + sample_i;
           auto const tx_j = ch_start + sample_j;

RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitKernels.dev.cc

@@ -34,7 +34,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
                                   uint32_t const* __restrict__ hitsModuleStart,
                                   pixelCPEforDevice::ParamsOnDeviceT<TrackerTraits> const* __restrict__ cpeParams,
                                   uint32_t* __restrict__ hitsLayerStart) const {
-      assert(0 == hitsModuleStart[0]);
+      ALPAKA_ASSERT_ACC(0 == hitsModuleStart[0]);
 
       for (int32_t i : cms::alpakatools::uniform_elements(acc, TrackerTraits::numberOfLayers + 1)) {
         hitsLayerStart[i] = hitsModuleStart[cpeParams->layerGeometry().layerStart[i]];

RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHits.h

@@ -97,11 +97,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
 
             int nClusInIter = alpaka::math::min(acc, maxHitsInIter, endClus - startClus);
             int lastClus = startClus + nClusInIter;
-            assert(nClusInIter <= nclus);
-            assert(nClusInIter > 0);
-            assert(lastClus <= nclus);
-
-            assert(nclus > maxHitsInIter || (0 == startClus && nClusInIter == nclus && lastClus == nclus));
+            ALPAKA_ASSERT_ACC(nClusInIter <= nclus);
+            ALPAKA_ASSERT_ACC(nClusInIter > 0);
+            ALPAKA_ASSERT_ACC(lastClus <= nclus);
+            ALPAKA_ASSERT_ACC(nclus > maxHitsInIter || (0 == startClus && nClusInIter == nclus && lastClus == nclus));
 
             // init
             for (uint32_t ic : cms::alpakatools::independent_group_elements(acc, nClusInIter)) {

RecoTracker/PixelSeeding/plugins/alpaka/CACell.h

@@ -67,8 +67,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
       // link to default empty
       theOuterNeighbors = &cellNeighbors[0];
       theTracks = &cellTracks[0];
-      assert(outerNeighbors().empty());
-      assert(tracks().empty());
+      ALPAKA_ASSERT_ACC(outerNeighbors().empty());
+      ALPAKA_ASSERT_ACC(tracks().empty());
     }
 
     template <typename TAcc>
@@ -334,7 +334,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
                   hits[nh++] = cells[c].theFishboneId;  // Fishbone hit is always outer than inner hit
                 }
               }
-              assert(nh < TrackerTraits::maxHitsOnTrack);
+              ALPAKA_ASSERT_ACC(nh < TrackerTraits::maxHitsOnTrack);
               hits[nh] = theOuterHitId;
               auto it = foundNtuplets.bulkFill(acc, apc, hits, nh + 1);
               if (it >= 0) {  // if negative is overflow....
@@ -346,7 +346,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
           }
         }
         tmpNtuplet.pop_back();
-        assert(tmpNtuplet.size() < int(TrackerTraits::maxHitsOnTrack - 1));
+        ALPAKA_ASSERT_ACC(tmpNtuplet.size() < int(TrackerTraits::maxHitsOnTrack - 1));
       }
     }
 

RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.h

@@ -65,8 +65,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
 
       /// Is this a pair with inner == 0?
       ALPAKA_FN_ACC ALPAKA_FN_INLINE bool startAt0(int16_t pid) const {
-        assert((pixelTopology::Phase1::layerPairs[pid * 2] == 0) ==
-               (pid < 3 || pid == 13 || pid == 15 || pid == 16));  // to be 100% sure it's working, may be removed
+        ALPAKA_ASSERT_ACC(
+            (pixelTopology::Phase1::layerPairs[pid * 2] == 0) ==
+            (pid < 3 || pid == 13 || pid == 15 || pid == 16));  // to be 100% sure it's working, may be removed
         return pixelTopology::Phase1::layerPairs[pid * 2] == 0;
       }
     };
@@ -81,7 +82,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
 
       /// Is this a pair with inner == 0
       ALPAKA_FN_ACC ALPAKA_FN_INLINE bool startAt0(int16_t pid) const {
-        assert((pixelTopology::Phase2::layerPairs[pid * 2] == 0) == ((pid < 3) | (pid >= 23 && pid < 28)));
+        ALPAKA_ASSERT_ACC((pixelTopology::Phase2::layerPairs[pid * 2] == 0) == ((pid < 3) | (pid >= 23 && pid < 28)));
         return pixelTopology::Phase2::layerPairs[pid * 2] == 0;
       }
     };