CMS 3D CMS Logo

List of all members | Classes | Public Types | Public Member Functions
cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE > Class Template Reference

#include <OneToManyAssoc.h>

Inheritance diagram for cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE >:
cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >

Classes

struct  finalizeBulk
 

Public Types

using index_type = typename OneToManyAssocBase< I, ONES, SIZE >::index_type
 
- Public Types inherited from cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >
using Counter = uint32_t
 
using CountersOnly = OneToManyAssocBase< I, ONES, 0 >
 
using index_type = I
 

Public Member Functions

template<typename TAcc >
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE int32_t bulkFill (const TAcc &acc, AtomicPairCounter &apc, index_type const *v, uint32_t n)
 
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void bulkFinalize (AtomicPairCounter const &apc)
 
template<typename TAcc >
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void bulkFinalizeFill (TAcc &acc, AtomicPairCounter const &apc)
 
- Public Member Functions inherited from cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >
template<typename TAcc >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void add (const TAcc &acc, CountersOnly const &co)
 
constexpr index_type const * begin () const
 
constexpr index_type const * begin (uint32_t b) const
 
constexpr auto capacity () const
 
template<typename TAcc >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void count (const TAcc &acc, I b)
 
constexpr index_type const * end () const
 
constexpr index_type const * end (uint32_t b) const
 
template<typename TAcc >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void fill (const TAcc &acc, I b, index_type j)
 
ALPAKA_FN_HOST_ACC void initStorage (View view)
 
constexpr auto nOnes () const
 
constexpr auto size () const
 
constexpr auto size (uint32_t b) const
 
constexpr auto totOnes () const
 
ALPAKA_FN_HOST_ACC void zero ()
 

Additional Inherited Members

- Static Public Member Functions inherited from cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >
template<typename TAcc >
ALPAKA_FN_ACC static ALPAKA_FN_INLINE uint32_t atomicDecrement (const TAcc &acc, Counter &x)
 
template<typename TAcc >
ALPAKA_FN_ACC static ALPAKA_FN_INLINE uint32_t atomicIncrement (const TAcc &acc, Counter &x)
 
static constexpr int32_t ctCapacity ()
 
static constexpr int32_t ctNOnes ()
 
template<typename TAcc , typename TQueue >
static ALPAKA_FN_INLINE void launchZero (OneToManyAssocBase *h, TQueue &queue)
 
template<typename TAcc , typename TQueue >
static ALPAKA_FN_INLINE void launchZero (View view, TQueue &queue)
 
- Public Attributes inherited from cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >
FlexiStorage< index_type, SIZE > content
 
FlexiStorage< Counter, ONES > off
 
int32_t psws
 

Detailed Description

template<typename I, int32_t ONES, int32_t SIZE>
class cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE >

Definition at line 161 of file OneToManyAssoc.h.

Member Typedef Documentation

◆ index_type

template<typename I , int32_t ONES, int32_t SIZE>
using cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE >::index_type = typename OneToManyAssocBase<I, ONES, SIZE>::index_type

Definition at line 163 of file OneToManyAssoc.h.

Member Function Documentation

◆ bulkFill()

template<typename I , int32_t ONES, int32_t SIZE>
template<typename TAcc >
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE int32_t cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE >::bulkFill ( const TAcc &  acc,
AtomicPairCounter apc,
index_type const *  v,
uint32_t  n 
)
inline

Definition at line 167 of file OneToManyAssoc.h.

References caHitNtupletGeneratorKernels::apc, HltBtagPostValidation_cff::c, cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::content, dqmiolumiharvest::j, dqmiodumpmetadata::n, and cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off.

167  {
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)
173  this->content[c.second + j] = v[j];
174  return c.first;
175  }
FlexiStorage< Counter, ONES > off
FlexiStorage< index_type, SIZE > content
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter * apc

◆ bulkFinalize()

template<typename I , int32_t ONES, int32_t SIZE>
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE >::bulkFinalize ( AtomicPairCounter const &  apc)
inline

Definition at line 177 of file OneToManyAssoc.h.

References caHitNtupletGeneratorKernels::apc, cms::alpakatools::AtomicPairCounter::Counters::first, cms::alpakatools::AtomicPairCounter::get(), cms::cuda::AtomicPairCounter::get(), and cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off.

177  {
178  this->off[apc.get().first] = apc.get().second;
179  }
FlexiStorage< Counter, ONES > off
__device__ __host__ Counters get() const
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter * apc

◆ bulkFinalizeFill()

template<typename I , int32_t ONES, int32_t SIZE>
template<typename TAcc >
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE >::bulkFinalizeFill ( TAcc &  acc,
AtomicPairCounter const &  apc 
)
inline

Definition at line 182 of file OneToManyAssoc.h.

References caHitNtupletGeneratorKernels::apc, f, dqmdumpme::first, cms::cuda::AtomicPairCounter::get(), mps_fire::i, cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::nOnes(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off, alignCSCRings::s, and cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::totOnes().

182  {
183  int f = apc.get().first;
184  auto s = apc.get().second;
185  if (f >= this->nOnes()) { // overflow!
186  this->off[this->nOnes()] = uint32_t(this->off[this->nOnes() - 1]);
187  return;
188  }
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]) {
191  this->off[i] = s;
192  }
193  }
FlexiStorage< Counter, ONES > off
__device__ __host__ Counters get() const
double f[11][100]
TupleMultiplicity< TrackerTraits > const HitToTuple< TrackerTraits > const cms::cuda::AtomicPairCounter * apc