Skip to content

Commit cae4cfd

Browse files
antoniopetreantoniopetre
committed
[alpaka] change assert
1 parent f321e35 commit cae4cfd

33 files changed

Lines changed: 212 additions & 210 deletions

‎src/alpaka/AlpakaCore/HistoContainer.h‎

Lines changed: 21 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -24,10 +24,10 @@ namespace cms {
2424
const uint32_t nt = offsets[nh];
2525
cms::alpakatools::for_each_element_in_grid_strided(acc, nt, [&](uint32_t i) {
2626
auto off = alpaka_std::upper_bound(offsets, offsets + nh + 1, i);
27-
assert((*off) > 0);
27+
ALPAKA_ASSERT_OFFLOAD((*off) > 0);
2828
int32_t ih = off - offsets - 1;
29-
assert(ih >= 0);
30-
assert(ih < int(nh));
29+
ALPAKA_ASSERT_OFFLOAD(ih >= 0);
30+
ALPAKA_ASSERT_OFFLOAD(ih < int(nh));
3131
h->count(acc, v[i], ih);
3232
});
3333
}
@@ -44,10 +44,10 @@ namespace cms {
4444

4545
cms::alpakatools::for_each_element_in_grid_strided(acc, nt, [&](uint32_t i) {
4646
auto off = alpaka_std::upper_bound(offsets, offsets + nh + 1, i);
47-
assert((*off) > 0);
47+
ALPAKA_ASSERT_OFFLOAD((*off) > 0);
4848
int32_t ih = off - offsets - 1;
49-
assert(ih >= 0);
50-
assert(ih < int(nh));
49+
ALPAKA_ASSERT_OFFLOAD(ih >= 0);
50+
ALPAKA_ASSERT_OFFLOAD(ih < int(nh));
5151
h->fill(acc, v[i], i, ih);
5252
});
5353
}
@@ -125,7 +125,7 @@ namespace cms {
125125
int bs = Hist::bin(value);
126126
int be = std::min(int(Hist::nbins() - 1), bs + n);
127127
bs = std::max(0, bs - n);
128-
assert(be >= bs);
128+
ALPAKA_ASSERT_OFFLOAD(be >= bs);
129129
for (auto pj = hist.begin(bs); pj < hist.end(be); ++pj) {
130130
func(*pj);
131131
}
@@ -136,7 +136,7 @@ namespace cms {
136136
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void forEachInWindow(Hist const &hist, V wmin, V wmax, Func const &func) {
137137
auto bs = Hist::bin(wmin);
138138
auto be = Hist::bin(wmax);
139-
assert(be >= bs);
139+
ALPAKA_ASSERT_OFFLOAD(be >= bs);
140140
for (auto pj = hist.begin(bs); pj < hist.end(be); ++pj) {
141141
func(*pj);
142142
}
@@ -210,15 +210,15 @@ namespace cms {
210210

211211
template <typename T_Acc>
212212
ALPAKA_FN_ACC ALPAKA_FN_INLINE void countDirect(const T_Acc &acc, T b) {
213-
assert(b < nbins());
213+
ALPAKA_ASSERT_OFFLOAD(b < nbins());
214214
atomicIncrement(acc, off[b]);
215215
}
216216

217217
template <typename T_Acc>
218218
ALPAKA_FN_ACC ALPAKA_FN_INLINE void fillDirect(const T_Acc &acc, T b, index_type j) {
219-
assert(b < nbins());
219+
ALPAKA_ASSERT_OFFLOAD(b < nbins());
220220
auto w = atomicDecrement(acc, off[b]);
221-
assert(w > 0);
221+
ALPAKA_ASSERT_OFFLOAD(w > 0);
222222
bins[w - 1] = j;
223223
}
224224

@@ -255,44 +255,44 @@ namespace cms {
255255
template <typename T_Acc>
256256
ALPAKA_FN_ACC ALPAKA_FN_INLINE void count(const T_Acc &acc, T t) {
257257
uint32_t b = bin(t);
258-
assert(b < nbins());
258+
ALPAKA_ASSERT_OFFLOAD(b < nbins());
259259
atomicIncrement(acc, off[b]);
260260
}
261261

262262
template <typename T_Acc>
263263
ALPAKA_FN_ACC ALPAKA_FN_INLINE void fill(const T_Acc &acc, T t, index_type j) {
264264
uint32_t b = bin(t);
265-
assert(b < nbins());
265+
ALPAKA_ASSERT_OFFLOAD(b < nbins());
266266
auto w = atomicDecrement(acc, off[b]);
267-
assert(w > 0);
267+
ALPAKA_ASSERT_OFFLOAD(w > 0);
268268
bins[w - 1] = j;
269269
}
270270

271271
template <typename T_Acc>
272272
ALPAKA_FN_ACC ALPAKA_FN_INLINE void count(const T_Acc &acc, T t, uint32_t nh) {
273273
uint32_t b = bin(t);
274-
assert(b < nbins());
274+
ALPAKA_ASSERT_OFFLOAD(b < nbins());
275275
b += histOff(nh);
276-
assert(b < totbins());
276+
ALPAKA_ASSERT_OFFLOAD(b < totbins());
277277
atomicIncrement(acc, off[b]);
278278
}
279279

280280
template <typename T_Acc>
281281
ALPAKA_FN_ACC ALPAKA_FN_INLINE void fill(const T_Acc &acc, T t, index_type j, uint32_t nh) {
282282
uint32_t b = bin(t);
283-
assert(b < nbins());
283+
ALPAKA_ASSERT_OFFLOAD(b < nbins());
284284
b += histOff(nh);
285-
assert(b < totbins());
285+
ALPAKA_ASSERT_OFFLOAD(b < totbins());
286286
auto w = atomicDecrement(acc, off[b]);
287-
assert(w > 0);
287+
ALPAKA_ASSERT_OFFLOAD(w > 0);
288288
bins[w - 1] = j;
289289
}
290290

291291
template <typename T_Acc>
292292
ALPAKA_FN_ACC ALPAKA_FN_INLINE void finalize(const T_Acc &acc, Counter *ws = nullptr) {
293-
assert(off[totbins() - 1] == 0);
293+
ALPAKA_ASSERT_OFFLOAD(off[totbins() - 1] == 0);
294294
blockPrefixScan(acc, off, totbins(), ws);
295-
assert(off[totbins() - 1] == off[totbins() - 2]);
295+
ALPAKA_ASSERT_OFFLOAD(off[totbins() - 1] == off[totbins() - 2]);
296296
}
297297

298298
constexpr auto size() const { return uint32_t(off[totbins() - 1]); }

‎src/alpaka/AlpakaCore/prefixScan.h‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -53,16 +53,16 @@ namespace cms {
5353
uint32_t const gridBlockIdx(alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
5454
uint32_t const blockThreadIdx(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
5555
assert(ws);
56-
assert(size <= 1024);
57-
assert(0 == blockDimension % 32);
56+
ALPAKA_ASSERT_OFFLOAD(size <= 1024);
57+
ALPAKA_ASSERT_OFFLOAD(0 == blockDimension % 32);
5858
auto first = blockThreadIdx;
5959
auto mask = __ballot_sync(0xffffffff, first < size);
6060
auto laneId = blockThreadIdx & 0x1f;
6161

6262
for (auto i = first; i < size; i += blockDimension) {
6363
warpPrefixScan(laneId, ci, co, i, mask);
6464
auto warpId = i / 32;
65-
assert(warpId < 32);
65+
ALPAKA_ASSERT_OFFLOAD(warpId < 32);
6666
if (31 == laneId)
6767
ws[warpId] = co[i];
6868
mask = __ballot_sync(mask, i + blockDimension < size);
@@ -101,16 +101,16 @@ namespace cms {
101101
uint32_t const gridBlockIdx(alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
102102
uint32_t const blockThreadIdx(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
103103
assert(ws);
104-
assert(size <= 1024);
105-
assert(0 == blockDimension % 32);
104+
ALPAKA_ASSERT_OFFLOAD(size <= 1024);
105+
ALPAKA_ASSERT_OFFLOAD(0 == blockDimension % 32);
106106
auto first = blockThreadIdx;
107107
auto mask = __ballot_sync(0xffffffff, first < size);
108108
auto laneId = blockThreadIdx & 0x1f;
109109

110110
for (auto i = first; i < size; i += blockDimension) {
111111
warpPrefixScan(laneId, c, i, mask);
112112
auto warpId = i / 32;
113-
assert(warpId < 32);
113+
ALPAKA_ASSERT_OFFLOAD(warpId < 32);
114114
if (31 == laneId)
115115
ws[warpId] = c[i];
116116
mask = __ballot_sync(mask, i + blockDimension < size);
@@ -146,7 +146,7 @@ namespace cms {
146146
// first each block does a scan of size 1024; (better be enough blocks....)
147147
#ifndef NDEBUG
148148
uint32_t const gridDimension(alpaka::getWorkDiv<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
149-
assert(gridDimension / threadDimension <= 1024);
149+
ALPAKA_ASSERT_OFFLOAD(gridDimension / threadDimension <= 1024);
150150
#endif
151151
int off = blockDimension * blockIdx * threadDimension;
152152
auto& ws = alpaka::declareSharedVar<T[32], __COUNTER__>(acc);
@@ -165,11 +165,11 @@ namespace cms {
165165
if (!isLastBlockDone)
166166
return;
167167

168-
assert(int(gridDimension) == *pc);
168+
ALPAKA_ASSERT_OFFLOAD(int(gridDimension) == *pc);
169169

170170
auto& psum = alpaka::declareSharedVar<T[1024], __COUNTER__>(acc);
171171

172-
assert(static_cast<int32_t>(blockDimension * threadDimension) >= gridDimension);
172+
ALPAKA_ASSERT_OFFLOAD(static_cast<int32_t>(blockDimension * threadDimension) >= gridDimension);
173173

174174
for (int elemId = 0; elemId < static_cast<int>(threadDimension); ++elemId) {
175175
int index = +threadIdx * threadDimension + elemId;

‎src/alpaka/AlpakaCore/radixSort.h‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -85,10 +85,10 @@ namespace cms {
8585
auto& ibs = alpaka::declareSharedVar<int, __COUNTER__>(acc);
8686
auto& p = alpaka::declareSharedVar<int, __COUNTER__>(acc);
8787

88-
assert(size > 0);
88+
ALPAKA_ASSERT_OFFLOAD(size > 0);
8989

9090
const uint32_t blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Elems>(acc)[0u]);
91-
assert(blockDimension >= sb);
91+
ALPAKA_ASSERT_OFFLOAD(blockDimension >= sb);
9292

9393
// bool debug = false; // threadIdx.x==0 && blockIdx.x==5;
9494

@@ -190,7 +190,7 @@ namespace cms {
190190
*/
191191

192192
alpaka::syncBlockThreads(acc);
193-
assert(c[0] == 0);
193+
ALPAKA_ASSERT_OFFLOAD(c[0] == 0);
194194

195195
// swap (local, ok)
196196
auto t = j;
@@ -204,7 +204,7 @@ namespace cms {
204204
}
205205

206206
if ((w != 8) && (0 == (NS & 1)))
207-
assert(j == ind); // w/d is even so ind is correct
207+
ALPAKA_ASSERT_OFFLOAD(j == ind); // w/d is even so ind is correct
208208

209209
if (j != ind) // odd...
210210
cms::alpakatools::for_each_element_in_block_strided(acc, size, [&](uint32_t i) { ind[i] = ind2[i]; });

‎src/alpaka/AlpakaDataFormats/SiPixelDigiErrorsAlpaka.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
1818
formatterErrors_h{std::move(errors)} {
1919
auto perror_h = alpaka::getPtrNative(error_h);
2020
perror_h->construct(maxFedWords, alpaka::getPtrNative(data_d));
21-
assert(perror_h->empty());
22-
assert(perror_h->capacity() == static_cast<int>(maxFedWords));
21+
ALPAKA_ASSERT_OFFLOAD(perror_h->empty());
22+
ALPAKA_ASSERT_OFFLOAD(perror_h->capacity() == static_cast<int>(maxFedWords));
2323

2424
// TO DO: nothing really async in here for now... Pass the queue in constructor argument instead, and don't wait anymore!
2525
Queue queue(device);

‎src/alpaka/CondFormats/SiPixelGainForHLTonGPU.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -53,9 +53,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
5353
auto offset =
5454
range.first + col * lengthOfColumnData + lengthOfAveragedDataInEachColumn * numberOfDataBlocksToSkip;
5555

56-
assert(offset < range.second);
57-
assert(offset < 3088384);
58-
assert(0 == offset % 2);
56+
ALPAKA_ASSERT_OFFLOAD(offset < range.second);
57+
ALPAKA_ASSERT_OFFLOAD(offset < 3088384);
58+
ALPAKA_ASSERT_OFFLOAD(0 == offset % 2);
5959

6060
auto s = v_pedestals[offset / 2];
6161

‎src/alpaka/CondFormats/pixelCPEforGPU.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -198,8 +198,8 @@ namespace pixelCPEforGPU {
198198

199199
auto xsize = int(urxl) + 2 - int(llxl);
200200
auto ysize = int(uryl) + 2 - int(llyl);
201-
assert(xsize >= 0); // 0 if bixpix...
202-
assert(ysize >= 0);
201+
ALPAKA_ASSERT_OFFLOAD(xsize >= 0); // 0 if bixpix...
202+
ALPAKA_ASSERT_OFFLOAD(ysize >= 0);
203203

204204
if (phase1PixelTopology::isBigPixX(cp.minRow[ic]))
205205
++xsize;

‎src/alpaka/DataFormats/SiPixelDigisSoA.cc‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,13 @@
11
#include "DataFormats/SiPixelDigisSoA.h"
22

33
#include <cassert>
4+
#include <alpaka/alpaka.hpp>
45

56
SiPixelDigisSoA::SiPixelDigisSoA(
67
size_t nDigis, const uint32_t *pdigi, const uint32_t *rawIdArr, const uint16_t *adc, const int32_t *clus)
78
: pdigi_(pdigi, pdigi + nDigis),
89
rawIdArr_(rawIdArr, rawIdArr + nDigis),
910
adc_(adc, adc + nDigis),
1011
clus_(clus, clus + nDigis) {
11-
assert(pdigi_.size() == nDigis);
12+
ALPAKA_ASSERT_OFFLOAD(pdigi_.size() == nDigis);
1213
}

‎src/alpaka/Framework/ReusableObjectHolder.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -84,10 +84,10 @@ namespace edm {
8484
ReusableObjectHolder() : m_outstandingObjects(0) {}
8585
ReusableObjectHolder(ReusableObjectHolder&& iOther)
8686
: m_availableQueue(std::move(iOther.m_availableQueue)), m_outstandingObjects(0) {
87-
assert(0 == iOther.m_outstandingObjects);
87+
ALPAKA_ASSERT_OFFLOAD(0 == iOther.m_outstandingObjects);
8888
}
8989
~ReusableObjectHolder() {
90-
assert(0 == m_outstandingObjects);
90+
ALPAKA_ASSERT_OFFLOAD(0 == m_outstandingObjects);
9191
std::unique_ptr<T, Deleter> item;
9292
while (m_availableQueue.try_pop(item)) {
9393
item.reset();

‎src/alpaka/Framework/WaitingTaskList.cc‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616
// user include files
1717
#include "tbb/task.h"
1818
#include <cassert>
19+
#include <alpaka/alpaka.hpp>
1920

2021
#include "WaitingTaskList.h"
2122
#include "hardware_pause.h"
@@ -51,7 +52,7 @@ void WaitingTaskList::reset() {
5152
m_exceptionPtr = std::exception_ptr{};
5253
unsigned int nSeenTasks = m_lastAssignedCacheIndex;
5354
m_lastAssignedCacheIndex = 0;
54-
assert(m_head == nullptr);
55+
ALPAKA_ASSERT_OFFLOAD(m_head == nullptr);
5556
if (nSeenTasks > m_nodeCacheSize) {
5657
//need to expand so next time we don't have to do any
5758
// memory requests

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -345,7 +345,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
345345
-circle_results.q * (fast_fit(2) - sqrt(Rfit::sqr(fast_fit(2)) - 0.25 * (e - d).squaredNorm())),
346346
circle_results.q * (1. / fast_fit(2) + u(n));
347347

348-
assert(circle_results.q * circle_results.par(1) <= 0);
348+
ALPAKA_ASSERT_OFFLOAD(circle_results.q * circle_results.par(1) <= 0);
349349

350350
Rfit::Vector2d eMinusd = e - d;
351351
double tmp1 = eMinusd.squaredNorm();

0 commit comments

Comments
 (0)