1 #ifndef HeterogeneousCore_AlpakaInterface_interface_OneToManyAssoc_h 2 #define HeterogeneousCore_AlpakaInterface_interface_OneToManyAssoc_h 9 #include <alpaka/alpaka.hpp> 58 ALPAKA_FN_HOST_ACC
void zero() {
64 template <
typename TAcc>
66 for (uint32_t
i = 0;
static_cast<int>(
i) <
totOnes(); ++
i) {
71 template <
typename TAcc>
76 template <
typename TAcc>
81 template <
typename TAcc>
82 ALPAKA_FN_ACC ALPAKA_FN_INLINE
void count(
const TAcc &acc,
I b) {
87 template <
typename TAcc>
97 template <
typename TAcc>
99 ALPAKA_ASSERT_ACC((1 == alpaka::getWorkDiv<alpaka::Grid, alpaka::Blocks>(acc)[0]));
104 h->initStorage(
view);
106 alpaka::syncBlockThreads(acc);
113 template <
typename TAcc,
typename TQueue>
115 View view = {
h,
nullptr,
nullptr, -1, -1};
119 template <
typename TAcc,
typename TQueue>
129 if constexpr (!requires_single_thread_per_block_v<TAcc>) {
132 auto workDiv = cms::alpakatools::make_workdiv<TAcc>(nblocks,
nthreads);
137 h->initStorage(
view);
157 template <
typename I,
165 template <
typename TAcc>
166 ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE int32_t
168 auto c =
apc.inc_add(acc,
n);
169 if (
int(
c.first) >= this->nOnes())
170 return -int32_t(
c.first);
171 this->
off[
c.first] =
c.second;
172 for (uint32_t
j = 0;
j <
n; ++
j)
181 template <
typename TAcc>
189 auto first = f + alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc)[0];
190 for (
int i =
first;
i < this->
totOnes();
i += alpaka::getWorkDiv<alpaka::Grid, alpaka::Threads>(acc)[0]) {
196 template <
typename TAcc>
205 template <
typename I,
214 template <
typename TAcc>
221 ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE
void finalize() {
223 for (uint32_t
i = 1;
static_cast<int>(
i) < this->
totOnes(); ++
i)
227 template <
typename TAcc,
typename TQueue>
229 View view = {
h,
nullptr,
nullptr, -1, -1};
233 template <
typename TAcc,
typename TQueue>
238 if constexpr (!requires_single_thread_per_block_v<TAcc>) {
245 poff =
view.offStorage;
251 auto workDiv = cms::alpakatools::make_workdiv<TAcc>(nblocks,
nthreads);
252 alpaka::exec<TAcc>(
queue,
260 alpaka::getWarpSizes(alpaka::getDev(
queue))[0]);
269 #endif // HeterogeneousCore_CUDAUtilities_interface_HistoContainer_h
cudaStream_t int32_t ONES
T1 atomicSub(T1 *a, T2 b)
__host__ __device__ VT * co
__device__ __host__ Counters get() const
std::vector< Block > Blocks
const std::complex< double > I
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int nthreads
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter * apc
The Signals That Services Can Subscribe To This is based on ActivityRegistry h
Helper function to determine trigger accepts.
ALPAKA_ASSERT_ACC(offsets)
T1 atomicAdd(T1 *a, T2 b)