cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS > Class Template Reference

#include <HistoContainer.h>

Inheritance diagram for cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >:

Public Types
using	Base = OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >
using	Counter = typename Base::Counter
using	index_type = typename Base::index_type
using	UT = typename std::make_unsigned< T >::type
using	View = typename Base::View
Public Types inherited from cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >
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 count(T t)
__host__ __device__	__attribute__ ((always_inline)) void fill(T t
__host__ __device__	__attribute__ ((always_inline)) void count(T t
__host__ __device__	__attribute__ ((always_inline)) void fill(T t
	assert (b< nbins())
	assert (w > 0)
	assert (b< nbins())
	assert (b< totbins())
	assert (b< nbins())
	assert (b< totbins())
	assert (w > 0)
	Base::atomicIncrement (this->off[b])
Public Member Functions inherited from cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >
__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
	for (uint32_t j=0;j< n;++j) content[c.n+j]
	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 constexpr UT	bin (T t)
static constexpr auto	histOff (uint32_t nh)
static constexpr uint32_t	ilog2 (uint32_t v)
static constexpr uint32_t	nbins ()
static constexpr uint32_t	nbits ()
static constexpr uint32_t	nhists ()
static constexpr uint32_t	sizeT ()
static constexpr uint32_t	totbins ()
Static Public Member Functions inherited from cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >
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
	b = histOff(nh)
this	content [w - 1] = j
__host__ __device__ index_type	j
__host__ __device__ uint32_t	nh
__host__ __device__ index_type uint32_t	nh
auto	w = Base::atomicDecrement(this->off[b])
Public Attributes inherited from cms::cuda::OneToManyAssoc< I, NHISTS *NBINS+1, SIZE >
	content [w - 1]
FlexiStorage< index_type, SIZE >	content
__host__ __device__ index_type	j
return c	m
__host__ __device__ index_type const uint32_t	n
	off [c.m]
__host__ __device__ ONES	off
int32_t	psws
__host__ __device__ index_type const *	v
auto	w

Detailed Description

template<typename T, uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>
class cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >

Definition at line 101 of file HistoContainer.h.

Member Typedef Documentation

Base

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

using cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::Base = OneToManyAssoc<I, NHISTS * NBINS + 1, SIZE>

Definition at line 103 of file HistoContainer.h.

Counter

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

using cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::Counter = typename Base::Counter

Definition at line 105 of file HistoContainer.h.

index_type

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

using cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::index_type = typename Base::index_type

Definition at line 106 of file HistoContainer.h.

UT

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

using cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::UT = typename std::make_unsigned<T>::type

Definition at line 107 of file HistoContainer.h.

View

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

using cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::View = typename Base::View

Definition at line 104 of file HistoContainer.h.

Member Function Documentation

__attribute__() [1/4]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

__host__ __device__ cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  )

inline

Definition at line 138 of file HistoContainer.h.

References cms::cuda::assert(), b, newFWLiteAna::bin, LaserClient_cfi::nbins, and submitPVValidationJobs::t.

                                                          {
        uint32_t b = bin(t);
        assert(b < nbins());
        Base::atomicIncrement(this->off[b]);
      }

__attribute__() [2/4]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

__host__ __device__ cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  )

inline

__attribute__() [3/4]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

__host__ __device__ cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  )

inline

__attribute__() [4/4]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

__host__ __device__ cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  )

inline

assert() [1/7]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::assert

(

)

assert() [2/7]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::assert	(	w	,
		0
	)

assert() [3/7]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::assert

(

)

assert() [4/7]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::assert

(

)

assert() [5/7]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::assert

(

)

assert() [6/7]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::assert

(

)

assert() [7/7]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::assert	(	w	,
		0
	)

Base::atomicIncrement()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::Base::atomicIncrement

(

this->

off[b]

)

bin()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr UT cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::bin ( T  t )

inlinestatic

Definition at line 132 of file HistoContainer.h.

References edm::shift, and submitPVValidationJobs::t.

Referenced by gpuPixelDoublets::for(), and gpuVertexFinder::while().

                                   {
        constexpr uint32_t shift = sizeT() - nbits();
        constexpr uint32_t mask = (1 << nbits()) - 1;
        return (t >> shift) & mask;
      }

histOff()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr auto cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::histOff ( uint32_t  nh )

inlinestatic

Definition at line 130 of file HistoContainer.h.

References NBINS, and cms::cuda::nh.

Referenced by gpuPixelDoublets::for().

{ return NBINS * nh; }

ilog2()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::ilog2 ( uint32_t  v )

inlinestatic

Definition at line 109 of file HistoContainer.h.

References b, mps_fire::i, alignCSCRings::r, alignCSCRings::s, and cms::cuda::v.

                                                  {
        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;
      }

nbins()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbins ( )

inlinestatic

Definition at line 123 of file HistoContainer.h.

References NBINS.

Referenced by gpuVertexFinder::__attribute__(), gpuPixelDoublets::for(), and gpuVertexFinder::while().

{ return NBINS; }

nbits()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nbits ( )

inlinestatic

Definition at line 126 of file HistoContainer.h.

References trklet::ilog2(), and NBINS.

{ return ilog2(NBINS - 1) + 1; }

nhists()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nhists ( )

inlinestatic

Definition at line 124 of file HistoContainer.h.

{ return NHISTS; }

sizeT()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::sizeT ( )

inlinestatic

Definition at line 122 of file HistoContainer.h.

{ return S; }

totbins()

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

static constexpr uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::totbins ( )

inlinestatic

Definition at line 125 of file HistoContainer.h.

References NBINS.

Referenced by gpuVertexFinder::__attribute__().

{ return NHISTS * NBINS + 1; }

Member Data Documentation

b

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::b = histOff(nh)

Definition at line 155 of file HistoContainer.h.

content

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

this cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::content = j

Definition at line 149 of file HistoContainer.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().

j

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

__host__ __device__ index_type cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::j

Initial value:

{
        uint32_t b = bin(t)

Definition at line 144 of file HistoContainer.h.

nh [1/2]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

__host__ __device__ uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nh

Initial value:

{
        uint32_t b = bin(t)

Definition at line 152 of file HistoContainer.h.

nh [2/2]

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

__host__ __device__ index_type uint32_t cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::nh

Initial value:

{
        uint32_t b = bin(t)

Definition at line 160 of file HistoContainer.h.

w

template<typename T , uint32_t NBINS, int32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>

auto cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::w = Base::atomicDecrement(this->off[b])

Definition at line 147 of file HistoContainer.h.