cms::cuda::OneToManyAssoc Class Reference

#include <OneToManyAssoc.h>

Public Types
using	Counter = uint32_t
using	CountersOnly = OneToManyAssoc< I, ONES, 0 >
using	index_type = I
using	View = OneToManyAssocView< OneToManyAssoc< I, ONES, SIZE >>

Public Member Functions
__host__ __device__	__attribute__ ((always_inline)) void add(CountersOnly const &co)
__host__ __device__	__attribute__ ((always_inline)) void count(int32_t b)
__host__ __device__	__attribute__ ((always_inline)) void fill(int32_t b
__host__ __device__	__attribute__ ((always_inline)) int32_t bulkFill(AtomicPairCounter &apc
__host__ __device__	__attribute__ ((always_inline)) void bulkFinalize(AtomicPairCounter const &apc)
__host__ __device__	__attribute__ ((always_inline)) void bulkFinalizeFill(AtomicPairCounter const &apc)
__host__ __device__	__attribute__ ((always_inline)) void final ize(Counter *ws
	assert (w > 0)
constexpr auto	capacity () const
	if (int(c.m) >=nOnes()) return-int32_t(c.m)
__host__ __device__ void	initStorage (View view)
constexpr auto	nOnes () const
constexpr auto	totOnes () const
__host__ __device__ void	zero ()

Static Public Member Functions
static __host__ __device__	__attribute__ ((always_inline)) uint32_t atomicIncrement(Counter &x)
static __host__ __device__	__attribute__ ((always_inline)) uint32_t atomicDecrement(Counter &x)
static constexpr int32_t	ctCapacity ()
static constexpr int32_t	ctNOnes ()
static constexpr uint32_t	ilog2 (uint32_t v)

Public Attributes
	content [w-1] = j
FlexiStorage< index_type, SIZE >	content
__host__ __device__ index_type	j
return c	m
__host__ __device__ index_type const uint32_t	n
	off [c.m] = c.n
__host__ __device__ ONES	off
int32_t	psws
__host__ __device__ index_type const *	v
auto	w = atomicDecrement(off[b])

Detailed Description

Definition at line 143 of file OneToManyAssoc.h.

Member Typedef Documentation

using cms::cuda::OneToManyAssoc::Counter = uint32_t

Definition at line 146 of file OneToManyAssoc.h.

using cms::cuda::OneToManyAssoc::CountersOnly = OneToManyAssoc<I, ONES, 0>

Definition at line 148 of file OneToManyAssoc.h.

using cms::cuda::OneToManyAssoc::index_type = I

Definition at line 150 of file OneToManyAssoc.h.

using cms::cuda::OneToManyAssoc::View = OneToManyAssocView<OneToManyAssoc<I, ONES, SIZE>>

Definition at line 145 of file OneToManyAssoc.h.

Member Function Documentation

__host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  ) const

inline

Definition at line 191 of file OneToManyAssoc.h.

                                                                           {
        for (int32_t i = 0; i < totOnes(); ++i) {
#ifdef __CUDA_ARCH__
          atomicAdd(off.data() + i, co.off[i]);
#else
          auto &a = (std::atomic<Counter> &)(off[i]);
          a += co.off[i];
#endif
        }
      }

static __host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  )

inlinestatic

Definition at line 202 of file OneToManyAssoc.h.

                                                                                      {
#ifdef __CUDA_ARCH__
        return atomicAdd(&x, 1);
#else
        auto &a = (std::atomic<Counter> &)(x);
        return a++;
#endif
      }

static __host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  )

inlinestatic

Definition at line 211 of file OneToManyAssoc.h.

                                                                                      {
#ifdef __CUDA_ARCH__
        return atomicSub(&x, 1);
#else
        auto &a = (std::atomic<Counter> &)(x);
        return a--;
#endif
      }

__host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  )

inline

Definition at line 220 of file OneToManyAssoc.h.

                                                                {
        assert(b < nOnes());
        atomicIncrement(off[b]);
      }

__host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  )

inline

__host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  )

inline

__host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  ) const

inline

Definition at line 242 of file OneToManyAssoc.h.

                                                                                          {
        off[apc.get().m] = apc.get().n;
      }

__host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  ) const

inline

Definition at line 246 of file OneToManyAssoc.h.

                                                                                              {
        int m = apc.get().m;
        auto n = apc.get().n;
        if (m >= nOnes()) {  // overflow!
          off[nOnes()] = uint32_t(off[nOnes() - 1]);
          return;
        }
        auto first = m + blockDim.x * blockIdx.x + threadIdx.x;
        for (int i = first; i < totOnes(); i += gridDim.x * blockDim.x) {
          off[i] = n;
        }
      }

__host__ __device__ cms::cuda::OneToManyAssoc::__attribute__ ( (always_inline)  )

inlinefinal

cms::cuda::OneToManyAssoc::assert	(	w	,
		0
	)

Referenced by cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__(), and cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::initStorage().

constexpr auto cms::cuda::OneToManyAssoc::capacity ( ) const

inline

Definition at line 169 of file OneToManyAssoc.h.

Referenced by gpuVertexFinder::__attribute__().

{ return content.capacity(); }

static constexpr int32_t cms::cuda::OneToManyAssoc::ctCapacity ( )

inlinestatic

Definition at line 168 of file OneToManyAssoc.h.

Referenced by cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::initStorage().

{ return SIZE; }

static constexpr int32_t cms::cuda::OneToManyAssoc::ctNOnes ( )

inlinestatic

Definition at line 165 of file OneToManyAssoc.h.

Referenced by cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::initStorage().

{ return ONES; }

cms::cuda::OneToManyAssoc::if

(

int(c.m) >=nOnes()

)

static constexpr uint32_t cms::cuda::OneToManyAssoc::ilog2 ( uint32_t  v )

inlinestatic

Definition at line 152 of file OneToManyAssoc.h.

                                                  {
        constexpr uint32_t b[] = {0x2, 0xC, 0xF0, 0xFF00, 0xFFFF0000};
        constexpr uint32_t s[] = {1, 2, 4, 8, 16};

        uint32_t r = 0;  // result of log2(v) will go here
        for (auto i = 4; i >= 0; i--)
          if (v & b[i]) {
            v >>= s[i];
            r |= s[i];
          }
        return r;
      }

__host__ __device__ void cms::cuda::OneToManyAssoc::initStorage ( View  view )

inline

Definition at line 171 of file OneToManyAssoc.h.

                                                      {
        assert(view.assoc == this);
        if constexpr (ctCapacity() < 0) {
          assert(view.contentStorage);
          assert(view.contentSize > 0);
          content.init(view.contentStorage, view.contentSize);
        }
        if constexpr (ctNOnes() < 0) {
          assert(view.offStorage);
          assert(view.offSize > 0);
          off.init(view.offStorage, view.offSize);
        }
      }

constexpr auto cms::cuda::OneToManyAssoc::nOnes ( ) const

inline

Definition at line 167 of file OneToManyAssoc.h.

Referenced by cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__().

{ return totOnes() - 1; }

constexpr auto cms::cuda::OneToManyAssoc::totOnes ( ) const

inline

Definition at line 166 of file OneToManyAssoc.h.

Referenced by cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::nOnes(), and cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::zero().

{ return off.capacity(); }

__host__ __device__ void cms::cuda::OneToManyAssoc::zero ( )

inline

Definition at line 185 of file OneToManyAssoc.h.

                                      {
        for (int32_t i = 0; i < totOnes(); ++i) {
          off[i] = 0;
        }
      }

Member Data Documentation

cms::cuda::OneToManyAssoc::content[c.n+j] = j

Definition at line 229 of file OneToManyAssoc.h.

Referenced by python.rootplot.utilities.Hist2D::_flat_content(), python.rootplot.root2matplotlib.Hist2D::box(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::capacity(), python.rootplot.root2matplotlib.Hist2D::col(), python.rootplot.root2matplotlib.Hist2D::contour(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::initStorage(), python.rootplot.utilities.Hist2D::TH2F(), and python.rootplot.root2matplotlib.Hist2D::TH2F().

FlexiStorage<index_type, SIZE> cms::cuda::OneToManyAssoc::content

Definition at line 276 of file OneToManyAssoc.h.

Referenced by python.rootplot.utilities.Hist2D::_flat_content(), python.rootplot.root2matplotlib.Hist2D::box(), python.rootplot.root2matplotlib.Hist2D::col(), python.rootplot.root2matplotlib.Hist2D::contour(), python.rootplot.utilities.Hist2D::TH2F(), and python.rootplot.root2matplotlib.Hist2D::TH2F().

__host__ __device__ index_type cms::cuda::OneToManyAssoc::j

Initial value:

{
        assert(b < nOnes())

Definition at line 225 of file OneToManyAssoc.h.

return c cms::cuda::OneToManyAssoc::m

Definition at line 239 of file OneToManyAssoc.h.

Referenced by cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__().

__host__ __device__ index_type const uint32_t cms::cuda::OneToManyAssoc::n

Initial value:

{
        auto c = apc.add(n)

Definition at line 232 of file OneToManyAssoc.h.

Referenced by cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__().

cms::cuda::OneToManyAssoc::off[c.m] = c.n

Definition at line 236 of file OneToManyAssoc.h.

Referenced by gpuVertexFinder::__attribute__(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::__attribute__(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::initStorage(), cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::totOnes(), and cms::cuda::OneToManyAssoc< hindex_type, S+1, 5 *S >::zero().

__host__ __device__ ONES cms::cuda::OneToManyAssoc::off

Definition at line 274 of file OneToManyAssoc.h.

int32_t cms::cuda::OneToManyAssoc::psws

Definition at line 275 of file OneToManyAssoc.h.

__host__ __device__ index_type const* cms::cuda::OneToManyAssoc::v

Definition at line 232 of file OneToManyAssoc.h.

auto cms::cuda::OneToManyAssoc::w = atomicDecrement(off[b])

Definition at line 227 of file OneToManyAssoc.h.