#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 count(T t)

__host__ __device__	__attribute__ ((always_inline)) void fill(T t

__host__ __device__	__attribute__ ((always_inline)) void fill(T t

	assert (b< nbins())

	assert (b< nbins())

	assert (b< nbins())

	assert (b< totbins())

	assert (b< totbins())

	assert (w > 0)

	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)) int32_t bulkFill(AtomicPairCounter &apc

__host__ __device__	__attribute__ ((always_inline)) void add(CountersOnly const &co)

__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 count(int32_t b)

__host__ __device__	__attribute__ ((always_inline)) void fill(int32_t b

__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 atomicDecrement(Counter &x)

static __host__ __device__	__attribute__ ((always_inline)) uint32_t atomicIncrement(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 >
FlexiStorage< index_type, SIZE >	content

	content [w - 1]

__host__ __device__ index_type	j

return c	m

__host__ __device__ const index_type uint32_t	n

__host__ __device__ ONES	off

	off [c.m]

int32_t	psws

__host__ __device__ const index_type *	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

◆ 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

Definition at line 138 of file HistoContainer.h.

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

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

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

◆ 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	(	w	,
		0
	)

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

inlinestaticconstexpr

Definition at line 132 of file HistoContainer.h.

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

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

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

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

inlinestaticconstexpr

Definition at line 130 of file HistoContainer.h.

130 { return NBINS * nh; }

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

Referenced by gpuPixelDoublets::for().

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

inlinestaticconstexpr

Definition at line 109 of file HistoContainer.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;
       }

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

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

inlinestaticconstexpr

Definition at line 123 of file HistoContainer.h.

123 { return NBINS; }

References NBINS.

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

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

inlinestaticconstexpr

Definition at line 126 of file HistoContainer.h.

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

References trklet::ilog2(), and NBINS.

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

inlinestaticconstexpr

Definition at line 124 of file HistoContainer.h.

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

inlinestaticconstexpr

Definition at line 122 of file HistoContainer.h.

122 { return S; }

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

inlinestaticconstexpr

Definition at line 125 of file HistoContainer.h.

125 { return NHISTS * NBINS + 1; }

References NBINS.

Referenced by gpuVertexFinder::__attribute__().

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.

Public Types

Public Member Functions

Static Public Member Functions

Public Attributes

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 >

Member Typedef Documentation

◆ Base

◆ Counter

◆ index_type

◆ UT

◆ View

Member Function Documentation

◆ __attribute__() [1/4]

◆ __attribute__() [2/4]

◆ __attribute__() [3/4]

◆ __attribute__() [4/4]

◆ assert() [1/7]

◆ assert() [2/7]

◆ assert() [3/7]

◆ assert() [4/7]

◆ assert() [5/7]

◆ assert() [6/7]

◆ assert() [7/7]

◆ Base::atomicIncrement()

◆ bin()

◆ histOff()

◆ ilog2()

◆ nbins()

◆ nbits()

◆ nhists()

◆ sizeT()

◆ totbins()

Member Data Documentation

◆ b

◆ content

◆ j

◆ nh [1/2]

◆ nh [2/2]

◆ w

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 >

◆ attribute() [1/4]

◆ attribute() [2/4]

◆ attribute() [3/4]

◆ attribute() [4/4]