[alpaka] change assert

cms-patatrack · Aug 3, 2021 · a8eb537 · a8eb537
1 parent e6da7bf
commit a8eb537
Show file tree

Hide file tree

Showing 33 changed files with 214 additions and 212 deletions.
diff --git a/src/alpaka/AlpakaCore/HistoContainer.h b/src/alpaka/AlpakaCore/HistoContainer.h
@@ -24,10 +24,10 @@ namespace cms {
         const uint32_t nt = offsets[nh];
         cms::alpakatools::for_each_element_in_grid_strided(acc, nt, [&](uint32_t i) {
           auto off = alpaka_std::upper_bound(offsets, offsets + nh + 1, i);
-          assert((*off) > 0);
+          ALPAKA_ASSERT_OFFLOAD((*off) > 0);
           int32_t ih = off - offsets - 1;
-          assert(ih >= 0);
-          assert(ih < int(nh));
+          ALPAKA_ASSERT_OFFLOAD(ih >= 0);
+          ALPAKA_ASSERT_OFFLOAD(ih < int(nh));
           h->count(acc, v[i], ih);
         });
       }
@@ -44,10 +44,10 @@ namespace cms {
 
         cms::alpakatools::for_each_element_in_grid_strided(acc, nt, [&](uint32_t i) {
           auto off = alpaka_std::upper_bound(offsets, offsets + nh + 1, i);
-          assert((*off) > 0);
+          ALPAKA_ASSERT_OFFLOAD((*off) > 0);
           int32_t ih = off - offsets - 1;
-          assert(ih >= 0);
-          assert(ih < int(nh));
+          ALPAKA_ASSERT_OFFLOAD(ih >= 0);
+          ALPAKA_ASSERT_OFFLOAD(ih < int(nh));
           h->fill(acc, v[i], i, ih);
         });
       }
@@ -125,7 +125,7 @@ namespace cms {
       int bs = Hist::bin(value);
       int be = std::min(int(Hist::nbins() - 1), bs + n);
       bs = std::max(0, bs - n);
-      assert(be >= bs);
+      ALPAKA_ASSERT_OFFLOAD(be >= bs);
       for (auto pj = hist.begin(bs); pj < hist.end(be); ++pj) {
         func(*pj);
       }
@@ -136,7 +136,7 @@ namespace cms {
     ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void forEachInWindow(Hist const &hist, V wmin, V wmax, Func const &func) {
       auto bs = Hist::bin(wmin);
       auto be = Hist::bin(wmax);
-      assert(be >= bs);
+      ALPAKA_ASSERT_OFFLOAD(be >= bs);
       for (auto pj = hist.begin(bs); pj < hist.end(be); ++pj) {
         func(*pj);
       }
@@ -210,15 +210,15 @@ namespace cms {
 
       template <typename T_Acc>
       ALPAKA_FN_ACC ALPAKA_FN_INLINE void countDirect(const T_Acc &acc, T b) {
-        assert(b < nbins());
+        ALPAKA_ASSERT_OFFLOAD(b < nbins());
         atomicIncrement(acc, off[b]);
       }
 
       template <typename T_Acc>
       ALPAKA_FN_ACC ALPAKA_FN_INLINE void fillDirect(const T_Acc &acc, T b, index_type j) {
-        assert(b < nbins());
+        ALPAKA_ASSERT_OFFLOAD(b < nbins());
         auto w = atomicDecrement(acc, off[b]);
-        assert(w > 0);
+        ALPAKA_ASSERT_OFFLOAD(w > 0);
         bins[w - 1] = j;
       }
 
@@ -255,44 +255,44 @@ namespace cms {
       template <typename T_Acc>
       ALPAKA_FN_ACC ALPAKA_FN_INLINE void count(const T_Acc &acc, T t) {
         uint32_t b = bin(t);
-        assert(b < nbins());
+        ALPAKA_ASSERT_OFFLOAD(b < nbins());
         atomicIncrement(acc, off[b]);
       }
 
       template <typename T_Acc>
       ALPAKA_FN_ACC ALPAKA_FN_INLINE void fill(const T_Acc &acc, T t, index_type j) {
         uint32_t b = bin(t);
-        assert(b < nbins());
+       ALPAKA_ASSERT_OFFLOAD(b < nbins());
         auto w = atomicDecrement(acc, off[b]);
-        assert(w > 0);
+        ALPAKA_ASSERT_OFFLOAD(w > 0);
         bins[w - 1] = j;
       }
 
       template <typename T_Acc>
       ALPAKA_FN_ACC ALPAKA_FN_INLINE void count(const T_Acc &acc, T t, uint32_t nh) {
         uint32_t b = bin(t);
-        assert(b < nbins());
+        ALPAKA_ASSERT_OFFLOAD(b < nbins());
         b += histOff(nh);
-        assert(b < totbins());
+        ALPAKA_ASSERT_OFFLOAD(b < totbins());
         atomicIncrement(acc, off[b]);
       }
 
       template <typename T_Acc>
       ALPAKA_FN_ACC ALPAKA_FN_INLINE void fill(const T_Acc &acc, T t, index_type j, uint32_t nh) {
         uint32_t b = bin(t);
-        assert(b < nbins());
+        ALPAKA_ASSERT_OFFLOAD(b < nbins());
         b += histOff(nh);
-        assert(b < totbins());
+        ALPAKA_ASSERT_OFFLOAD(b < totbins());
         auto w = atomicDecrement(acc, off[b]);
-        assert(w > 0);
+        ALPAKA_ASSERT_OFFLOAD(w > 0);
         bins[w - 1] = j;
       }
 
       template <typename T_Acc>
       ALPAKA_FN_ACC ALPAKA_FN_INLINE void finalize(const T_Acc &acc, Counter *ws = nullptr) {
-        assert(off[totbins() - 1] == 0);
+        ALPAKA_ASSERT_OFFLOAD(off[totbins() - 1] == 0);
         blockPrefixScan(acc, off, totbins(), ws);
-        assert(off[totbins() - 1] == off[totbins() - 2]);
+        ALPAKA_ASSERT_OFFLOAD(off[totbins() - 1] == off[totbins() - 2]);
       }
 
       constexpr auto size() const { return uint32_t(off[totbins() - 1]); }

diff --git a/src/alpaka/AlpakaCore/prefixScan.h b/src/alpaka/AlpakaCore/prefixScan.h
@@ -52,17 +52,17 @@ namespace cms {
       uint32_t const blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Threads>(acc)[0u]);
       uint32_t const gridBlockIdx(alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
       uint32_t const blockThreadIdx(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
-      assert(ws);
-      assert(size <= 1024);
-      assert(0 == blockDimension % 32);
+      ALPAKA_ASSERT_OFFLOAD(ws);
+      ALPAKA_ASSERT_OFFLOAD(size <= 1024);
+      ALPAKA_ASSERT_OFFLOAD(0 == blockDimension % 32);
       auto first = blockThreadIdx;
       auto mask = __ballot_sync(0xffffffff, first < size);
       auto laneId = blockThreadIdx & 0x1f;
 
       for (auto i = first; i < size; i += blockDimension) {
         warpPrefixScan(laneId, ci, co, i, mask);
         auto warpId = i / 32;
-        assert(warpId < 32);
+        ALPAKA_ASSERT_OFFLOAD(warpId < 32);
         if (31 == laneId)
           ws[warpId] = co[i];
         mask = __ballot_sync(mask, i + blockDimension < size);
@@ -100,17 +100,17 @@ namespace cms {
       uint32_t const blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Threads>(acc)[0u]);
       uint32_t const gridBlockIdx(alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
       uint32_t const blockThreadIdx(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
-      assert(ws);
-      assert(size <= 1024);
-      assert(0 == blockDimension % 32);
+      ALPAKA_ASSERT_OFFLOAD(ws);
+      ALPAKA_ASSERT_OFFLOAD(size <= 1024);
+      ALPAKA_ASSERT_OFFLOAD(0 == blockDimension % 32);
       auto first = blockThreadIdx;
       auto mask = __ballot_sync(0xffffffff, first < size);
       auto laneId = blockThreadIdx & 0x1f;
 
       for (auto i = first; i < size; i += blockDimension) {
         warpPrefixScan(laneId, c, i, mask);
         auto warpId = i / 32;
-        assert(warpId < 32);
+        ALPAKA_ASSERT_OFFLOAD(warpId < 32);
         if (31 == laneId)
           ws[warpId] = c[i];
         mask = __ballot_sync(mask, i + blockDimension < size);
@@ -146,7 +146,7 @@ namespace cms {
         // first each block does a scan of size 1024; (better be enough blocks....)
 #ifndef NDEBUG
         uint32_t const gridDimension(alpaka::getWorkDiv<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
-        assert(gridDimension / threadDimension <= 1024);
+        ALPAKA_ASSERT_OFFLOAD(gridDimension / threadDimension <= 1024);
 #endif
         int off = blockDimension * blockIdx * threadDimension;
         auto& ws = alpaka::declareSharedVar<T[32], __COUNTER__>(acc);
@@ -165,11 +165,11 @@ namespace cms {
         if (!isLastBlockDone)
           return;
 
-        assert(int(gridDimension) == *pc);
+        ALPAKA_ASSERT_OFFLOAD(int(gridDimension) == *pc);
 
         auto& psum = alpaka::declareSharedVar<T[1024], __COUNTER__>(acc);
 
-        assert(static_cast<int32_t>(blockDimension * threadDimension) >= gridDimension);
+        ALPAKA_ASSERT_OFFLOAD(static_cast<int32_t>(blockDimension * threadDimension) >= gridDimension);
 
         for (int elemId = 0; elemId < static_cast<int>(threadDimension); ++elemId) {
           int index = +threadIdx * threadDimension + elemId;

diff --git a/src/alpaka/AlpakaCore/radixSort.h b/src/alpaka/AlpakaCore/radixSort.h
@@ -85,10 +85,10 @@ namespace cms {
       auto& ibs = alpaka::declareSharedVar<int, __COUNTER__>(acc);
       auto& p = alpaka::declareSharedVar<int, __COUNTER__>(acc);
 
-      assert(size > 0);
+      ALPAKA_ASSERT_OFFLOAD(size > 0);
 
       const uint32_t blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Elems>(acc)[0u]);
-      assert(blockDimension >= sb);
+      ALPAKA_ASSERT_OFFLOAD(blockDimension >= sb);
 
       // bool debug = false; // threadIdx.x==0 && blockIdx.x==5;
 
@@ -190,7 +190,7 @@ namespace cms {
     */
 
         alpaka::syncBlockThreads(acc);
-        assert(c[0] == 0);
+        ALPAKA_ASSERT_OFFLOAD(c[0] == 0);
 
         // swap (local, ok)
         auto t = j;
@@ -204,7 +204,7 @@ namespace cms {
       }
 
       if ((w != 8) && (0 == (NS & 1)))
-        assert(j == ind);  // w/d is even so ind is correct
+        ALPAKA_ASSERT_OFFLOAD(j == ind);  // w/d is even so ind is correct
 
       if (j != ind)  // odd...
         cms::alpakatools::for_each_element_in_block_strided(acc, size, [&](uint32_t i) { ind[i] = ind2[i]; });

diff --git a/src/alpaka/AlpakaDataFormats/SiPixelDigiErrorsAlpaka.h b/src/alpaka/AlpakaDataFormats/SiPixelDigiErrorsAlpaka.h
@@ -18,8 +18,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
           formatterErrors_h{std::move(errors)} {
       auto perror_h = alpaka::getPtrNative(error_h);
       perror_h->construct(maxFedWords, alpaka::getPtrNative(data_d));
-      assert(perror_h->empty());
-      assert(perror_h->capacity() == static_cast<int>(maxFedWords));
+      ALPAKA_ASSERT_OFFLOAD(perror_h->empty());
+      ALPAKA_ASSERT_OFFLOAD(perror_h->capacity() == static_cast<int>(maxFedWords));
 
       // TO DO: nothing really async in here for now... Pass the queue in constructor argument instead, and don't wait anymore!
       Queue queue(device);

diff --git a/src/alpaka/CondFormats/SiPixelGainForHLTonGPU.h b/src/alpaka/CondFormats/SiPixelGainForHLTonGPU.h
@@ -53,9 +53,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
       auto offset =
           range.first + col * lengthOfColumnData + lengthOfAveragedDataInEachColumn * numberOfDataBlocksToSkip;
 
-      assert(offset < range.second);
-      assert(offset < 3088384);
-      assert(0 == offset % 2);
+      ALPAKA_ASSERT_OFFLOAD(offset < range.second);
+      ALPAKA_ASSERT_OFFLOAD(offset < 3088384);
+      ALPAKA_ASSERT_OFFLOAD(0 == offset % 2);
 
       auto s = v_pedestals[offset / 2];
 

diff --git a/src/alpaka/CondFormats/pixelCPEforGPU.h b/src/alpaka/CondFormats/pixelCPEforGPU.h
@@ -198,8 +198,8 @@ namespace pixelCPEforGPU {
 
     auto xsize = int(urxl) + 2 - int(llxl);
     auto ysize = int(uryl) + 2 - int(llyl);
-    assert(xsize >= 0);  // 0 if bixpix...
-    assert(ysize >= 0);
+    ALPAKA_ASSERT_OFFLOAD(xsize >= 0);  // 0 if bixpix...
+    ALPAKA_ASSERT_OFFLOAD(ysize >= 0);
 
     if (phase1PixelTopology::isBigPixX(cp.minRow[ic]))
       ++xsize;

diff --git a/src/alpaka/DataFormats/SiPixelDigisSoA.cc b/src/alpaka/DataFormats/SiPixelDigisSoA.cc
@@ -1,12 +1,13 @@
 #include "DataFormats/SiPixelDigisSoA.h"
 
 #include <cassert>
+#include <alpaka/alpaka.hpp>
 
 SiPixelDigisSoA::SiPixelDigisSoA(
     size_t nDigis, const uint32_t *pdigi, const uint32_t *rawIdArr, const uint16_t *adc, const int32_t *clus)
     : pdigi_(pdigi, pdigi + nDigis),
       rawIdArr_(rawIdArr, rawIdArr + nDigis),
       adc_(adc, adc + nDigis),
       clus_(clus, clus + nDigis) {
-  assert(pdigi_.size() == nDigis);
+  ALPAKA_ASSERT_OFFLOAD(pdigi_.size() == nDigis);
 }
diff --git a/src/alpaka/Framework/ReusableObjectHolder.h b/src/alpaka/Framework/ReusableObjectHolder.h
@@ -84,10 +84,10 @@ namespace edm {
     ReusableObjectHolder() : m_outstandingObjects(0) {}
     ReusableObjectHolder(ReusableObjectHolder&& iOther)
         : m_availableQueue(std::move(iOther.m_availableQueue)), m_outstandingObjects(0) {
-      assert(0 == iOther.m_outstandingObjects);
+      ALPAKA_ASSERT_OFFLOAD(0 == iOther.m_outstandingObjects);
     }
     ~ReusableObjectHolder() {
-      assert(0 == m_outstandingObjects);
+      ALPAKA_ASSERT_OFFLOAD(0 == m_outstandingObjects);
       std::unique_ptr<T, Deleter> item;
       while (m_availableQueue.try_pop(item)) {
         item.reset();

diff --git a/src/alpaka/Framework/WaitingTaskList.cc b/src/alpaka/Framework/WaitingTaskList.cc
@@ -16,6 +16,7 @@
 // user include files
 #include "tbb/task.h"
 #include <cassert>
+#include <alpaka/alpaka.hpp>
 
 #include "WaitingTaskList.h"
 #include "hardware_pause.h"
@@ -51,7 +52,7 @@ void WaitingTaskList::reset() {
   m_exceptionPtr = std::exception_ptr{};
   unsigned int nSeenTasks = m_lastAssignedCacheIndex;
   m_lastAssignedCacheIndex = 0;
-  assert(m_head == nullptr);
+  ALPAKA_ASSERT_OFFLOAD(m_head == nullptr);
   if (nSeenTasks > m_nodeCacheSize) {
     //need to expand so next time we don't have to do any
     // memory requests

diff --git a/src/alpaka/plugin-PixelTriplets/alpaka/BrokenLine.h b/src/alpaka/plugin-PixelTriplets/alpaka/BrokenLine.h
@@ -345,7 +345,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
           -circle_results.q * (fast_fit(2) - sqrt(Rfit::sqr(fast_fit(2)) - 0.25 * (e - d).squaredNorm())),
           circle_results.q * (1. / fast_fit(2) + u(n));
 
-      assert(circle_results.q * circle_results.par(1) <= 0);
+      ALPAKA_ASSERT_OFFLOAD(circle_results.q * circle_results.par(1) <= 0);
 
       Rfit::Vector2d eMinusd = e - d;
       double tmp1 = eMinusd.squaredNorm();

diff --git a/src/alpaka/plugin-PixelTriplets/alpaka/BrokenLineFitOnGPU.cc b/src/alpaka/plugin-PixelTriplets/alpaka/BrokenLineFitOnGPU.cc
@@ -8,7 +8,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
                                               uint32_t hitsInFit,
                                               uint32_t maxNumberOfTuples,
                                               Queue& queue) {
-    assert(tuples_d);
+    ALPAKA_ASSERT_OFFLOAD(tuples_d);
 
     const auto blockSize = 64;
     const auto numberOfBlocks = (maxNumberOfConcurrentFits_ + blockSize - 1) / blockSize;

diff --git a/src/alpaka/plugin-PixelTriplets/alpaka/BrokenLineFitOnGPU.h b/src/alpaka/plugin-PixelTriplets/alpaka/BrokenLineFitOnGPU.h
@@ -36,12 +36,12 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
                                   uint32_t offset) const {
       constexpr uint32_t hitsInFit = N;
 
-      assert(hitsInFit <= nHits);
+      ALPAKA_ASSERT_OFFLOAD(hitsInFit <= nHits);
 
-      assert(hhp);
-      assert(pfast_fit);
-      assert(foundNtuplets);
-      assert(tupleMultiplicity);
+      ALPAKA_ASSERT_OFFLOAD(hhp);
+      ALPAKA_ASSERT_OFFLOAD(pfast_fit);
+      ALPAKA_ASSERT_OFFLOAD(foundNtuplets);
+      ALPAKA_ASSERT_OFFLOAD(tupleMultiplicity);
 
       // look in bin for this hit multiplicity
 #ifdef BROKENLINE_DEBUG
@@ -60,9 +60,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
 
         // get it from the ntuple container (one to one to helix)
         auto tkid = *(tupleMultiplicity->begin(nHits) + tuple_idx);
-        assert(tkid < foundNtuplets->nbins());
+        ALPAKA_ASSERT_OFFLOAD(tkid < foundNtuplets->nbins());
 
-        assert(foundNtuplets->size(tkid) == nHits);
+        ALPAKA_ASSERT_OFFLOAD(foundNtuplets->size(tkid) == nHits);
 
         Rfit::Map3xNd<N> hits(phits + local_idx);
         Rfit::Map4d fast_fit(pfast_fit + local_idx);
@@ -111,10 +111,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
         BrokenLine::BL_Fast_fit(hits, fast_fit);
 
         // no NaN here....
-        assert(fast_fit(0) == fast_fit(0));
-        assert(fast_fit(1) == fast_fit(1));
-        assert(fast_fit(2) == fast_fit(2));
-        assert(fast_fit(3) == fast_fit(3));
+        ALPAKA_ASSERT_OFFLOAD(fast_fit(0) == fast_fit(0));
+        ALPAKA_ASSERT_OFFLOAD(fast_fit(1) == fast_fit(1));
+        ALPAKA_ASSERT_OFFLOAD(fast_fit(2) == fast_fit(2));
+        ALPAKA_ASSERT_OFFLOAD(fast_fit(3) == fast_fit(3));
       });
 
     }  // kernel operator()
@@ -132,10 +132,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
                                   double *__restrict__ pfast_fit,
                                   uint32_t nHits,
                                   uint32_t offset) const {
-      assert(N <= nHits);
+      ALPAKA_ASSERT_OFFLOAD(N <= nHits);
 
-      assert(results);
-      assert(pfast_fit);
+      ALPAKA_ASSERT_OFFLOAD(results);
+      ALPAKA_ASSERT_OFFLOAD(pfast_fit);
 
       // same as above...