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

#include <HistoContainer.h>

Public Types
using	Counter = uint32_t
using	CountersOnly = HistoContainer< T, NBINS, 0, S, I, NHISTS >
using	index_type = I
using	UT = typename std::make_unsigned< T >::type

Public Member Functions
__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(T t
__host__ __device__	__attribute__ ((always_inline)) void count(T t)
__host__ __device__	__attribute__ ((always_inline)) void countDirect(T b)
__host__ __device__	__attribute__ ((always_inline)) void fill(T t
__host__ __device__	__attribute__ ((always_inline)) void fill(T t
__host__ __device__	__attribute__ ((always_inline)) void fillDirect(T b
__host__ __device__	__attribute__ ((always_inline)) void final ize(Counter *ws
	assert (b< nbins())
	assert (b< nbins())
	assert (b< nbins())
	assert (b< totbins())
	assert (b< totbins())
	assert (w > 0)
	assert (w > 0)
	assert (w > 0)
	atomicIncrement (off[b])
	for (uint32_t j=0;j< n;++j) bins[c.n+j]
	if (c.m >=nbins()) return -int32_t(c.m)
__host__ __device__ void	zero ()

Static Public Member Functions
static __host__ __device__	__attribute__ ((always_inline)) uint32_t atomicDecrement(Counter &x)
static __host__ __device__	__attribute__ ((always_inline)) uint32_t atomicIncrement(Counter &x)
static constexpr UT	bin (T t)
static constexpr uint32_t	capacity ()
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 ()

Public Attributes
	b = histOff(nh)
index_type	bins [capacity()]
	bins [w - 1] = j
__host__ __device__ index_type	j
return c	m
__host__ __device__ const index_type uint32_t	n
__host__ __device__ uint32_t	nh
__host__ __device__ index_type uint32_t	nh
	off [c.m] = c.n
int32_t	psws
__host__ __device__ const index_type *	v
auto	w = atomicDecrement(off[b])

Detailed Description

template<typename T, uint32_t NBINS, uint32_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 152 of file HistoContainer.h.

Member Typedef Documentation

Counter

template<typename T, uint32_t NBINS, uint32_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 = uint32_t

Definition at line 154 of file HistoContainer.h.

CountersOnly

template<typename T, uint32_t NBINS, uint32_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 >::CountersOnly = HistoContainer<T, NBINS, 0, S, I, NHISTS>

Definition at line 156 of file HistoContainer.h.

index_type

template<typename T, uint32_t NBINS, uint32_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 = I

Definition at line 158 of file HistoContainer.h.

UT

template<typename T, uint32_t NBINS, uint32_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 159 of file HistoContainer.h.

Member Function Documentation

__attribute__() [1/13]

template<typename T, uint32_t NBINS, uint32_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/13]

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

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

inlinestatic

Definition at line 214 of file HistoContainer.h.

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

__attribute__() [3/13]

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

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

inlinestatic

Definition at line 205 of file HistoContainer.h.

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

__attribute__() [4/13]

template<typename T, uint32_t NBINS, uint32_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)  ) const &

inline

Definition at line 194 of file HistoContainer.h.

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

__attribute__() [5/13]

template<typename T, uint32_t NBINS, uint32_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)  ) const &

inline

Definition at line 245 of file HistoContainer.h.

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

__attribute__() [6/13]

template<typename T, uint32_t NBINS, uint32_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)  ) const &

inline

Definition at line 249 of file HistoContainer.h.

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

__attribute__() [7/13]

template<typename T, uint32_t NBINS, uint32_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__() [8/13]

template<typename T, uint32_t NBINS, uint32_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 262 of file HistoContainer.h.

                                                          {
        uint32_t b = bin(t);
        assert(b < nbins());
        atomicIncrement(off[b]);
      }

__attribute__() [9/13]

template<typename T, uint32_t NBINS, uint32_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 223 of file HistoContainer.h.

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

__attribute__() [10/13]

template<typename T, uint32_t NBINS, uint32_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__() [11/13]

template<typename T, uint32_t NBINS, uint32_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__() [12/13]

template<typename T, uint32_t NBINS, uint32_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__() [13/13]

template<typename T, uint32_t NBINS, uint32_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)  )

inlinefinal

assert() [1/8]

template<typename T, uint32_t NBINS, uint32_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/8]

template<typename T, uint32_t NBINS, uint32_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/8]

template<typename T, uint32_t NBINS, uint32_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/8]

template<typename T, uint32_t NBINS, uint32_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/8]

template<typename T, uint32_t NBINS, uint32_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/8]

template<typename T, uint32_t NBINS, uint32_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/8]

template<typename T, uint32_t NBINS, uint32_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() [8/8]

template<typename T, uint32_t NBINS, uint32_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
	)

atomicIncrement()

template<typename T, uint32_t NBINS, uint32_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 >::atomicIncrement

(

off

[b]

)

bin()

template<typename T, uint32_t NBINS, uint32_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 183 of file HistoContainer.h.

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

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

capacity()

template<typename T, uint32_t NBINS, uint32_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 >::capacity ( )

inlinestaticconstexpr

Definition at line 179 of file HistoContainer.h.

{ return SIZE; }

Referenced by gpuVertexFinder::__attribute__().

for()

template<typename T, uint32_t NBINS, uint32_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 >::for

(

)

histOff()

template<typename T, uint32_t NBINS, uint32_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 181 of file HistoContainer.h.

{ return NBINS * nh; }

Referenced by gpuPixelDoublets::for().

if()

template<typename T, uint32_t NBINS, uint32_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 >::if

(

c.m >=

nbins()

)

ilog2()

template<typename T, uint32_t NBINS, uint32_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 161 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;
      }

nbins()

template<typename T, uint32_t NBINS, uint32_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 175 of file HistoContainer.h.

{ return NBINS; }

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

nbits()

template<typename T, uint32_t NBINS, uint32_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 178 of file HistoContainer.h.

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

nhists()

template<typename T, uint32_t NBINS, uint32_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 176 of file HistoContainer.h.

{ return NHISTS; }

sizeT()

template<typename T, uint32_t NBINS, uint32_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 174 of file HistoContainer.h.

{ return S; }

totbins()

template<typename T, uint32_t NBINS, uint32_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 177 of file HistoContainer.h.

{ return NHISTS * NBINS + 1; }

Referenced by gpuVertexFinder::__attribute__().

zero()

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

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

inline

Definition at line 189 of file HistoContainer.h.

                                      {
        for (auto &i : off)
          i = 0;
      }

Member Data Documentation

b

template<typename T, uint32_t NBINS, uint32_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 279 of file HistoContainer.h.

bins [1/2]

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

index_type cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::bins[capacity()]

Definition at line 311 of file HistoContainer.h.

bins [2/2]

template<typename T, uint32_t NBINS, uint32_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 >::bins[w - 1] = j

Definition at line 232 of file HistoContainer.h.

j

template<typename T, uint32_t NBINS, uint32_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:

{
        assert(b < nbins())

Definition at line 228 of file HistoContainer.h.

m

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

return c cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::m

Definition at line 242 of file HistoContainer.h.

Referenced by cms::cuda::HistoContainer< hindex_type, S, 5 *S >::__attribute__(), and if().

n

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

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

Initial value:

{
        auto c = apc.add(n)

Definition at line 235 of file HistoContainer.h.

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

nh [1/2]

template<typename T, uint32_t NBINS, uint32_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 276 of file HistoContainer.h.

nh [2/2]

template<typename T, uint32_t NBINS, uint32_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 284 of file HistoContainer.h.

off

template<typename T, uint32_t NBINS, uint32_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 >::off[c.m] = c.n

Definition at line 239 of file HistoContainer.h.

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

psws

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

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

Definition at line 310 of file HistoContainer.h.

v

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

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

Definition at line 235 of file HistoContainer.h.

w

template<typename T, uint32_t NBINS, uint32_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 = atomicDecrement(off[b])

Definition at line 230 of file HistoContainer.h.