CMS 3D CMS Logo

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

#include <OneToManyAssoc.h>

Inheritance diagram for cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >:
cms::alpakatools::OneToManyAssocRandomAccess< I, ONES, SIZE > cms::alpakatools::OneToManyAssocSequential< I, ONES, SIZE > cms::alpakatools::OneToManyAssocRandomAccess< I, NHISTS *NBINS+1, SIZE > cms::alpakatools::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >

Classes

struct  View
 
struct  zeroAndInit
 

Public Types

using Counter = uint32_t
 
using CountersOnly = OneToManyAssocBase< I, ONES, 0 >
 
using index_type = I
 

Public Member Functions

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 ()
 

Static Public Member Functions

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

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::OneToManyAssocBase< I, ONES, SIZE >

Definition at line 22 of file OneToManyAssoc.h.

Member Typedef Documentation

◆ Counter

template<typename I, int32_t ONES, int32_t SIZE>
using cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::Counter = uint32_t

Definition at line 24 of file OneToManyAssoc.h.

◆ CountersOnly

template<typename I, int32_t ONES, int32_t SIZE>
using cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::CountersOnly = OneToManyAssocBase<I, ONES, 0>

Definition at line 26 of file OneToManyAssoc.h.

◆ index_type

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

Definition at line 28 of file OneToManyAssoc.h.

Member Function Documentation

◆ add()

template<typename I, int32_t ONES, int32_t SIZE>
template<typename TAcc >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::add ( const TAcc &  acc,
CountersOnly const &  co 
)
inline

Definition at line 65 of file OneToManyAssoc.h.

References cms::cudacompat::atomicAdd(), cms::cuda::co, cms::alpakatools::FlexiStorage< I, S >::data(), mps_fire::i, cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off, and cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::totOnes().

Referenced by counter.Counter::register().

65  {
66  for (uint32_t i = 0; static_cast<int>(i) < totOnes(); ++i) {
68  }
69  }
FlexiStorage< Counter, ONES > off
__host__ __device__ VT * co
Definition: prefixScan.h:47
std::vector< Block > Blocks
Definition: Block.h:99
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61

◆ atomicDecrement()

template<typename I, int32_t ONES, int32_t SIZE>
template<typename TAcc >
ALPAKA_FN_ACC static ALPAKA_FN_INLINE uint32_t cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::atomicDecrement ( const TAcc &  acc,
Counter x 
)
inlinestatic

Definition at line 77 of file OneToManyAssoc.h.

References cms::cudacompat::atomicSub(), and x.

Referenced by cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::fill().

77  {
78  return alpaka::atomicSub(acc, &x, 1u, alpaka::hierarchy::Blocks{});
79  }
T1 atomicSub(T1 *a, T2 b)
Definition: cudaCompat.h:73
std::vector< Block > Blocks
Definition: Block.h:99

◆ atomicIncrement()

template<typename I, int32_t ONES, int32_t SIZE>
template<typename TAcc >
ALPAKA_FN_ACC static ALPAKA_FN_INLINE uint32_t cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::atomicIncrement ( const TAcc &  acc,
Counter x 
)
inlinestatic

Definition at line 72 of file OneToManyAssoc.h.

References cms::cudacompat::atomicAdd(), and x.

Referenced by cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::count().

72  {
73  return alpaka::atomicAdd(acc, &x, 1u, alpaka::hierarchy::Blocks{});
74  }
std::vector< Block > Blocks
Definition: Block.h:99
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:61

◆ begin() [1/2]

template<typename I, int32_t ONES, int32_t SIZE>
constexpr index_type const* cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::begin ( void  ) const
inline

◆ begin() [2/2]

template<typename I, int32_t ONES, int32_t SIZE>
constexpr index_type const* cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::begin ( uint32_t  b) const
inline

◆ capacity()

template<typename I, int32_t ONES, int32_t SIZE>
constexpr auto cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::capacity ( ) const
inline

◆ count()

template<typename I, int32_t ONES, int32_t SIZE>
template<typename TAcc >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::count ( const TAcc &  acc,
b 
)
inline

◆ ctCapacity()

template<typename I, int32_t ONES, int32_t SIZE>
static constexpr int32_t cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::ctCapacity ( )
inlinestatic

◆ ctNOnes()

template<typename I, int32_t ONES, int32_t SIZE>
static constexpr int32_t cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::ctNOnes ( )
inlinestatic

◆ end() [1/2]

template<typename I, int32_t ONES, int32_t SIZE>
constexpr index_type const* cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::end ( void  ) const
inline

◆ end() [2/2]

template<typename I, int32_t ONES, int32_t SIZE>
constexpr index_type const* cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::end ( uint32_t  b) const
inline

◆ fill()

template<typename I, int32_t ONES, int32_t SIZE>
template<typename TAcc >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::fill ( const TAcc &  acc,
b,
index_type  j 
)
inline

◆ initStorage()

template<typename I, int32_t ONES, int32_t SIZE>
ALPAKA_FN_HOST_ACC void cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::initStorage ( View  view)
inline

Definition at line 44 of file OneToManyAssoc.h.

References ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::content, cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::ctCapacity(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::ctNOnes(), and cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off.

44  {
45  ALPAKA_ASSERT_ACC(view.assoc == this);
46  if constexpr (ctCapacity() < 0) {
47  ALPAKA_ASSERT_ACC(view.contentStorage);
48  ALPAKA_ASSERT_ACC(view.contentSize > 0);
49  content.init(view.contentStorage, view.contentSize);
50  }
51  if constexpr (ctNOnes() < 0) {
52  ALPAKA_ASSERT_ACC(view.offStorage);
53  ALPAKA_ASSERT_ACC(view.offSize > 0);
54  off.init(view.offStorage, view.offSize);
55  }
56  }
FlexiStorage< Counter, ONES > off
static constexpr int32_t ctNOnes()
static constexpr int32_t ctCapacity()
FlexiStorage< index_type, SIZE > content

◆ launchZero() [1/2]

template<typename I, int32_t ONES, int32_t SIZE>
template<typename TAcc , typename TQueue >
static ALPAKA_FN_INLINE void cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::launchZero ( OneToManyAssocBase< I, ONES, SIZE > *  h,
TQueue &  queue 
)
inlinestatic

Definition at line 114 of file OneToManyAssoc.h.

References h, and createBeamHaloJobs::queue.

114  {
115  View view = {h, nullptr, nullptr, -1, -1};
116  launchZero<TAcc>(view, queue);
117  }
The Signals That Services Can Subscribe To This is based on ActivityRegistry h
Helper function to determine trigger accepts.
Definition: Activities.doc:4

◆ launchZero() [2/2]

template<typename I, int32_t ONES, int32_t SIZE>
template<typename TAcc , typename TQueue >
static ALPAKA_FN_INLINE void cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::launchZero ( View  view,
TQueue &  queue 
)
inlinestatic

Definition at line 120 of file OneToManyAssoc.h.

References ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::ctCapacity(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::ctNOnes(), h, cms::cuda::nthreads, and createBeamHaloJobs::queue.

120  {
121  if constexpr (ctCapacity() < 0) {
122  ALPAKA_ASSERT_ACC(view.contentStorage);
123  ALPAKA_ASSERT_ACC(view.contentSize > 0);
124  }
125  if constexpr (ctNOnes() < 0) {
126  ALPAKA_ASSERT_ACC(view.offStorage);
127  ALPAKA_ASSERT_ACC(view.offSize > 0);
128  }
129  if constexpr (!requires_single_thread_per_block_v<TAcc>) {
130  auto nthreads = 1024;
131  auto nblocks = 1; // MUST BE ONE as memory is initialize in thread 0 (alternative is two kernels);
132  auto workDiv = cms::alpakatools::make_workdiv<TAcc>(nblocks, nthreads);
133  alpaka::exec<TAcc>(queue, workDiv, zeroAndInit{}, view);
134  } else {
135  auto h = view.assoc;
137  h->initStorage(view);
138  h->zero();
139  h->psws = 0;
140  }
141  }
static constexpr int32_t ctNOnes()
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int nthreads
static constexpr int32_t ctCapacity()
The Signals That Services Can Subscribe To This is based on ActivityRegistry h
Helper function to determine trigger accepts.
Definition: Activities.doc:4

◆ nOnes()

template<typename I, int32_t ONES, int32_t SIZE>
constexpr auto cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::nOnes ( ) const
inline

◆ size() [1/2]

template<typename I, int32_t ONES, int32_t SIZE>
constexpr auto cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::size ( void  ) const
inline

◆ size() [2/2]

template<typename I, int32_t ONES, int32_t SIZE>
constexpr auto cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::size ( uint32_t  b) const
inline

Definition at line 144 of file OneToManyAssoc.h.

References b, and cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off.

Referenced by ntupleDataFormat._Collection::__iter__(), and ntupleDataFormat._Collection::__len__().

144 { return off[b + 1] - off[b]; }
FlexiStorage< Counter, ONES > off
double b
Definition: hdecay.h:120

◆ totOnes()

template<typename I, int32_t ONES, int32_t SIZE>
constexpr auto cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::totOnes ( ) const
inline

◆ zero()

template<typename I, int32_t ONES, int32_t SIZE>
ALPAKA_FN_HOST_ACC void cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::zero ( )
inline

Definition at line 58 of file OneToManyAssoc.h.

References mps_fire::i, cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off, and cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::totOnes().

58  {
59  for (int32_t i = 0; i < totOnes(); ++i) {
60  off[i] = 0;
61  }
62  }
FlexiStorage< Counter, ONES > off

Member Data Documentation

◆ content

template<typename I, int32_t ONES, int32_t SIZE>
FlexiStorage<index_type, SIZE> cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::content

◆ off

template<typename I, int32_t ONES, int32_t SIZE>
FlexiStorage<Counter, ONES> cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::off

◆ psws

template<typename I, int32_t ONES, int32_t SIZE>
int32_t cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::psws