31 const uint32_t
blockIdx(alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
33 auto endModule = clus_view[0].moduleStart();
37 auto&
charge = alpaka::declareSharedVar<int32_t[maxNumClustersPerModules], __COUNTER__>(acc);
38 auto&
ok = alpaka::declareSharedVar<uint8_t[maxNumClustersPerModules], __COUNTER__>(acc);
39 auto&
newclusId = alpaka::declareSharedVar<uint16_t[maxNumClustersPerModules], __COUNTER__>(acc);
41 const uint32_t
gridDimension(alpaka::getWorkDiv<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
44 auto firstPixel = clus_view[1 +
module].moduleStart();
45 auto thisModuleId = digi_view[firstPixel].moduleId();
50 uint32_t nclus = clus_view[thisModuleId].clusInModule();
55 printf(
"Warning too many clusters in module %d in block %d: %d > %d\n",
62 const uint32_t blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Elems>(acc)[0u]);
65 const auto& [firstElementIdxNoStride, endElementIdxNoStride] =
69 uint32_t firstElementIdx = firstElementIdxNoStride;
70 uint32_t endElementIdx = endElementIdxNoStride;
74 i, firstElementIdx, endElementIdx, blockDimension,
numElements))
78 if (digi_view[
i].
moduleId() != thisModuleId)
89 if (thisModuleId % 100 == 1)
91 printf(
"start cluster charge cut for module %d in block %d\n", thisModuleId,
module);
96 alpaka::syncBlockThreads(acc);
98 uint32_t firstElementIdx = firstElementIdxNoStride;
99 uint32_t endElementIdx = endElementIdxNoStride;
102 i, firstElementIdx, endElementIdx, blockDimension,
numElements))
106 if (digi_view[
i].
moduleId() != thisModuleId)
110 static_cast<int32_t>(digi_view[
i].
adc()),
111 alpaka::hierarchy::Threads{});
113 alpaka::syncBlockThreads(acc);
121 allGood = allGood &&
false;
127 alpaka::syncBlockThreads(acc);
134 auto&
ws = alpaka::declareSharedVar<uint16_t[32], __COUNTER__>(acc);
136 auto minClust =
std::min(nclus, maxThreads);
146 uint32_t prevBlockEnd = ((
i +
offset / maxThreads) * maxThreads) - 1;
149 alpaka::syncBlockThreads(acc);
158 clus_view[thisModuleId].clusInModule() =
newclusId[nclus - 1];
159 alpaka::syncBlockThreads(acc);
162 if (thisModuleId % 100 == 1)
164 printf(
"module %d -> chargeCut = %d; nclus (pre cut) = %d; nclus (after cut) = %d\n",
168 clus_view[thisModuleId].clusInModule());
176 alpaka::syncBlockThreads(acc);
179 firstElementIdx = firstElementIdxNoStride;
180 endElementIdx = endElementIdxNoStride;
183 i, firstElementIdx, endElementIdx, blockDimension,
numElements))
187 if (digi_view[
i].
moduleId() != thisModuleId)
189 if (0 ==
ok[digi_view[
i].clus()])
192 digi_view[
i].clus() =
newclusId[digi_view[
i].clus()] - 1;
198 alpaka::syncBlockThreads(acc);
const uint32_t gridDimension(alpaka::getWorkDiv< alpaka::Grid, alpaka::Blocks >(acc)[0u])
constexpr int nMaxModules
ALPAKA_ASSERT_OFFLOAD(offsets)
constexpr uint16_t numberOfModules
constexpr int32_t getThresholdForLayerOnCondition(bool isLayer1) const noexcept
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__ int32_t *__restrict__ uint32_t numElements
__shared__ uint16_t newclusId[maxNumClustersPerModules]
constexpr int32_t maxNumClustersPerModules
static constexpr uint32_t layerStart[numberOfLayers+1]
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__ moduleId
constexpr uint16_t maxNumModules
T1 atomicAdd(T1 *a, T2 b)
constexpr uint16_t invalidModuleId
uint16_t *__restrict__ uint16_t const *__restrict__ adc