From b5c7e26b669a45089b0442d6d655b43feb3995d7 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 22 Apr 2024 12:08:17 +0200 Subject: [PATCH 1/5] Migrate assert() to ALPAKA_ASSERT_ACC() --- .../test/alpaka/testAtomicPairCounter.dev.cc | 4 ++-- .../AlpakaInterface/test/alpaka/testBuffer.dev.cc | 2 +- .../test/alpaka/testOneRadixSort.dev.cc | 10 +++++----- .../test/alpaka/testWorkDivision.dev.cc | 2 +- 4 files changed, 9 insertions(+), 9 deletions(-) diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testAtomicPairCounter.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testAtomicPairCounter.dev.cc index b48508c7153af..75488a1ce8b3a 100644 --- a/HeterogeneousCore/AlpakaInterface/test/alpaka/testAtomicPairCounter.dev.cc +++ b/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 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; } }; diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testBuffer.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testBuffer.dev.cc index 77ca80f639ca4..44d5171f77c88 100644 --- a/HeterogeneousCore/AlpakaInterface/test/alpaka/testBuffer.dev.cc +++ b/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(), diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneRadixSort.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneRadixSort.dev.cc index ef7bccb6f827d..218bc30ccca14 100644 --- a/HeterogeneousCore/AlpakaInterface/test/alpaka/testOneRadixSort.dev.cc +++ b/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(acc)[0u]; const auto blocksIdx = alpaka::getIdx(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(acc); auto& sws = alpaka::declareSharedVar(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) diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc index cd6cab3242d41..6b49d1692b361 100644 --- a/HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc +++ b/HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc @@ -84,7 +84,7 @@ struct testWordDivisionDefaultRange { if constexpr (loopScope == LoopScope::Block) { if (firstInLoopRange(acc)) { auto expected = expectedCount(acc, skip, size); - assert(counter == expected); + ALPAKA_ASSERT_ACC(counter == expected); } } } From 4fc762baba44d9b8b3cdaf34f7260580944cd84c Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 22 Apr 2024 12:08:32 +0200 Subject: [PATCH 2/5] Clean up comments --- .../AlpakaInterface/test/alpaka/testRadixSort.dev.cc | 5 ----- 1 file changed, 5 deletions(-) diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testRadixSort.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testRadixSort.dev.cc index 14efe123bf862..1b5c5875e0da7 100644 --- a/HeterogeneousCore/AlpakaInterface/test/alpaka/testRadixSort.dev.cc +++ b/HeterogeneousCore/AlpakaInterface/test/alpaka/testRadixSort.dev.cc @@ -87,7 +87,6 @@ void truncate(T& t) { template void go(Queue& queue, bool useShared) { std::mt19937 eng; - //std::mt19937 eng2; auto rgen = RS::ud(); std::chrono::high_resolution_clock::duration delta = 0ns; @@ -95,7 +94,6 @@ void go(Queue& queue, bool useShared) { constexpr int blockSize = 256 * 32; constexpr int N = blockSize * blocks; auto v_h = cms::alpakatools::make_host_buffer(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(blocks, ntXBl); if (useShared) // The original CUDA version used to call a kernel with __launch_bounds__(256, 4) specifier - // alpaka::enqueue(queue, alpaka::createTaskKernel(workdiv, radixSortMultiWrapper{}, @@ -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 From e08a9a843b2527032bd757573c5f270339fb1fc7 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 22 Apr 2024 12:08:51 +0200 Subject: [PATCH 3/5] Update test to use once_per_block --- .../test/alpaka/testWorkDivision.dev.cc | 15 +++------------ 1 file changed, 3 insertions(+), 12 deletions(-) diff --git a/HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc b/HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc index 6b49d1692b361..384e972ac469c 100644 --- a/HeterogeneousCore/AlpakaInterface/test/alpaka/testWorkDivision.dev.cc +++ b/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 -bool constexpr firstInLoopRange(TAcc const& acc) { - if constexpr (loopScope == LoopScope::Block) - return not alpaka::getIdx(acc)[0u]; - if constexpr (loopScope == LoopScope::Grid) - return not alpaka::getIdx(acc)[0u]; - assert(false); -} - template 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(acc)); // Init the counter for block range. Grid range does so my mean of memset. if constexpr (loopScope == LoopScope::Block) { - if (firstInLoopRange(acc)) + if (cms::alpakatools::once_per_block(acc)) { counter = 0; + } alpaka::syncBlockThreads(acc); } // The loop we are testing @@ -82,7 +73,7 @@ struct testWordDivisionDefaultRange { alpaka::syncBlockThreads(acc); // Check the result. Grid range will check by memcpy-ing the result. if constexpr (loopScope == LoopScope::Block) { - if (firstInLoopRange(acc)) { + if (cms::alpakatools::once_per_block(acc)) { auto expected = expectedCount(acc, skip, size); ALPAKA_ASSERT_ACC(counter == expected); } From 5f962bae270e070476c35cb363bc707bb2c639a8 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 22 Apr 2024 12:08:17 +0200 Subject: [PATCH 4/5] Migrate assert() to ALPAKA_ASSERT_ACC() --- .../plugins/alpaka/TimeComputationKernels.h | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h b/RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h index 05e01954215af..942a1500ead9a 100644 --- a/RecoLocalCalo/EcalRecProducers/plugins/alpaka/TimeComputationKernels.h +++ b/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(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(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; From 198c1267267cb5f0a8f4f9c582a23bba95fc6cc3 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Mon, 22 Apr 2024 12:08:17 +0200 Subject: [PATCH 5/5] Migrate assert() to ALPAKA_ASSERT_ACC() --- .../alpaka/SiPixelGainCalibrationForHLTUtilities.h | 10 ++++++---- .../test/alpaka/Clusters_test.dev.cc | 8 ++++---- .../test/alpaka/DigiErrors_test.dev.cc | 8 ++++---- .../SiPixelDigiSoA/test/alpaka/Digis_test.dev.cc | 8 ++++---- .../VertexSoA/test/alpaka/ZVertexSoA_test.dev.cc | 14 +++++++------- .../plugins/alpaka/PixelRecHitKernels.dev.cc | 2 +- .../SiPixelRecHits/plugins/alpaka/PixelRecHits.h | 9 ++++----- RecoTracker/PixelSeeding/plugins/alpaka/CACell.h | 8 ++++---- .../plugins/alpaka/CAHitNtupletGeneratorKernels.h | 7 ++++--- 9 files changed, 38 insertions(+), 36 deletions(-) diff --git a/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h b/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h index 1fbce15dbe231..3767ab655a2c5 100644 --- a/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h +++ b/CondFormats/SiPixelObjects/interface/alpaka/SiPixelGainCalibrationForHLTUtilities.h @@ -2,7 +2,9 @@ #define CondFormats_SiPixelObjects_interface_alpaka_SiPixelGainCalibrationForHLTUtilities_h #include + #include + #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 \ No newline at end of file +#endif // CondFormats_SiPixelObjects_interface_alpaka_SiPixelGainCalibrationForHLTUtilities_h diff --git a/DataFormats/SiPixelClusterSoA/test/alpaka/Clusters_test.dev.cc b/DataFormats/SiPixelClusterSoA/test/alpaka/Clusters_test.dev.cc index f98e7611c6229..77bbc73ad9f2e 100644 --- a/DataFormats/SiPixelClusterSoA/test/alpaka/Clusters_test.dev.cc +++ b/DataFormats/SiPixelClusterSoA/test/alpaka/Clusters_test.dev.cc @@ -31,10 +31,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::testClusterSoA { template >> 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); } } }; diff --git a/DataFormats/SiPixelDigiSoA/test/alpaka/DigiErrors_test.dev.cc b/DataFormats/SiPixelDigiSoA/test/alpaka/DigiErrors_test.dev.cc index 3829f7d755f6a..bf9b42c2c489d 100644 --- a/DataFormats/SiPixelDigiSoA/test/alpaka/DigiErrors_test.dev.cc +++ b/DataFormats/SiPixelDigiSoA/test/alpaka/DigiErrors_test.dev.cc @@ -32,10 +32,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::testDigisSoA { template >> 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); } } }; diff --git a/DataFormats/SiPixelDigiSoA/test/alpaka/Digis_test.dev.cc b/DataFormats/SiPixelDigiSoA/test/alpaka/Digis_test.dev.cc index 721b17912c65b..648b26ae3ddb7 100644 --- a/DataFormats/SiPixelDigiSoA/test/alpaka/Digis_test.dev.cc +++ b/DataFormats/SiPixelDigiSoA/test/alpaka/Digis_test.dev.cc @@ -29,10 +29,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::testDigisSoA { template >> 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); } } }; diff --git a/DataFormats/VertexSoA/test/alpaka/ZVertexSoA_test.dev.cc b/DataFormats/VertexSoA/test/alpaka/ZVertexSoA_test.dev.cc index 75a1ebf6b9269..1af92d5050943 100644 --- a/DataFormats/VertexSoA/test/alpaka/ZVertexSoA_test.dev.cc +++ b/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)); } } }; diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitKernels.dev.cc b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitKernels.dev.cc index fe3a703a8e713..1dbd77bb4ed09 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitKernels.dev.cc +++ b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHitKernels.dev.cc @@ -34,7 +34,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { uint32_t const* __restrict__ hitsModuleStart, pixelCPEforDevice::ParamsOnDeviceT 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]]; diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHits.h index aacdeb79a2749..7fd6234d0ba3b 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/alpaka/PixelRecHits.h +++ b/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)) { diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CACell.h b/RecoTracker/PixelSeeding/plugins/alpaka/CACell.h index d8af548109d29..44b7bea5075fb 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CACell.h +++ b/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 @@ -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)); } } diff --git a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.h b/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.h index 796bb3f36c586..6999169a46d73 100644 --- a/RecoTracker/PixelSeeding/plugins/alpaka/CAHitNtupletGeneratorKernels.h +++ b/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; } };