CMS 3D CMS Logo
List of all members | Public Types | Public Member Functions | Static Public Member Functions | Public Attributes
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

 __attribute__ ((always_inline)) int32_t bulkFill(AtomicPairCounter &apc
 
 __attribute__ ((always_inline)) void add(CountersOnly const &co)
 
 __attribute__ ((always_inline)) void bulkFinalize(AtomicPairCounter const &apc)
 
 __attribute__ ((always_inline)) void bulkFinalizeFill(AtomicPairCounter const &apc)
 
 __attribute__ ((always_inline)) void count(T t
 
 __attribute__ ((always_inline)) void count(T t)
 
 __attribute__ ((always_inline)) void countDirect(T b)
 
 __attribute__ ((always_inline)) void fill(T t
 
 __attribute__ ((always_inline)) void fill(T t
 
 __attribute__ ((always_inline)) void fillDirect(T b
 
 __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)
 
void zero ()
 

Static Public Member Functions

static __attribute__ ((always_inline)) uint32_t atomicDecrement(Counter &x)
 
static __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
 
index_type j
 
return c m
 
const index_type uint32_t n
 
uint32_t nh
 
index_type uint32_t nh
 
 off [c.m] = c.n
 
int32_t psws
 
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>
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 cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  ) &
inlinestatic

Definition at line 214 of file HistoContainer.h.

214  {
215 #ifdef __CUDA_ARCH__
216  return atomicSub(&x, 1);
217 #else
218  auto &a = (std::atomic<Counter> &)(x);
219  return a--;
220 #endif
221  }

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

◆ __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 cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  ) &
inlinestatic

Definition at line 205 of file HistoContainer.h.

205  {
206 #ifdef __CUDA_ARCH__
207  return atomicAdd(&x, 1);
208 #else
209  auto &a = (std::atomic<Counter> &)(x);
210  return a++;
211 #endif
212  }

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

◆ __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>
cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 194 of file HistoContainer.h.

194  {
195  for (uint32_t i = 0; i < totbins(); ++i) {
196 #ifdef __CUDA_ARCH__
197  atomicAdd(off + i, co.off[i]);
198 #else
199  auto &a = (std::atomic<Counter> &)(off[i]);
200  a += co.off[i];
201 #endif
202  }
203  }

References a, cms::cudacompat::atomicAdd(), cms::cuda::co, and mps_fire::i.

◆ __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>
cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 245 of file HistoContainer.h.

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

References cms::cuda::AtomicPairCounter::get(), cms::cuda::AtomicPairCounter::Counters::m, and cms::cuda::AtomicPairCounter::Counters::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>
cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  ) const &
inline

Definition at line 249 of file HistoContainer.h.

249  {
250  auto m = apc.get().m;
251  auto n = apc.get().n;
252  if (m >= nbins()) { // overflow!
253  off[nbins()] = uint32_t(off[nbins() - 1]);
254  return;
255  }
256  auto first = m + blockDim.x * blockIdx.x + threadIdx.x;
257  for (auto i = first; i < totbins(); i += gridDim.x * blockDim.x) {
258  off[i] = n;
259  }
260  }

References cms::cudacompat::blockDim, cms::cudacompat::blockIdx, dqmdumpme::first, cms::cuda::AtomicPairCounter::get(), cms::cudacompat::gridDim, mps_fire::i, cms::cuda::AtomicPairCounter::Counters::m, visualization-live-secondInstance_cfg::m, cms::cuda::AtomicPairCounter::Counters::n, cms::cuda::n, LaserClient_cfi::nbins, cms::cudacompat::threadIdx, and cms::cudacompat::dim3::x.

◆ __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>
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>
cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  )
inline

Definition at line 262 of file HistoContainer.h.

262  {
263  uint32_t b = bin(t);
264  assert(b < nbins());
265  atomicIncrement(off[b]);
266  }

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

◆ __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>
cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::__attribute__ ( (always_inline)  )
inline

Definition at line 223 of file HistoContainer.h.

223  {
224  assert(b < nbins());
225  atomicIncrement(off[b]);
226  }

References cms::cuda::assert(), b, and LaserClient_cfi::nbins.

◆ __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>
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>
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>
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>
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.

183  {
184  constexpr uint32_t shift = sizeT() - nbits();
185  constexpr uint32_t mask = (1 << nbits()) - 1;
186  return (t >> shift) & mask;
187  }

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

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

179 { return SIZE; }

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

181 { return NBINS * nh; }

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

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

161  {
162  constexpr uint32_t b[] = {0x2, 0xC, 0xF0, 0xFF00, 0xFFFF0000};
163  constexpr uint32_t s[] = {1, 2, 4, 8, 16};
164 
165  uint32_t r = 0; // result of log2(v) will go here
166  for (auto i = 4; i >= 0; i--)
167  if (v & b[i]) {
168  v >>= s[i];
169  r |= s[i];
170  }
171  return r;
172  }

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

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

175 { return NBINS; }

References NBINS.

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

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

References NBINS.

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

176 { 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.

174 { return S; }

References 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.

177 { return NHISTS * NBINS + 1; }

References NBINS.

◆ zero()

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

Definition at line 189 of file HistoContainer.h.

189  {
190  for (auto &i : off)
191  i = 0;
192  }

References mps_fire::i.

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>
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.

◆ n

template<typename T , uint32_t NBINS, uint32_t SIZE, uint32_t S = sizeof(T) * 8, typename I = uint32_t, uint32_t NHISTS = 1>
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.

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

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

cms::cudacompat::atomicSub
T1 atomicSub(T1 *a, T2 b)
Definition: cudaCompat.h:53
cms::cuda::HistoContainer::off
off[c.m]
Definition: HistoContainer.h:239
cms::cuda::HistoContainer::nh
uint32_t nh
Definition: HistoContainer.h:276
cms::cuda::HistoContainer::atomicIncrement
atomicIncrement(off[b])
mps_fire.i
i
Definition: mps_fire.py:428
NBINS
const int NBINS
Definition: CaloCachedShapeIntegrator.cc:3
cms::cuda::HistoContainer::ilog2
static constexpr uint32_t ilog2(uint32_t v)
Definition: HistoContainer.h:161
cms::cuda::HistoContainer::m
return c m
Definition: HistoContainer.h:242
DDAxes::x
dqmdumpme.first
first
Definition: dqmdumpme.py:55
cms::cudacompat::gridDim
thread_local dim3 gridDim
Definition: cudaCompat.cc:6
cms::cuda::HistoContainer::assert
assert(b< nbins())
alignCSCRings.s
s
Definition: alignCSCRings.py:92
cms::cudacompat::atomicAdd
T1 atomicAdd(T1 *a, T2 b)
Definition: cudaCompat.h:46
cms::cuda::HistoContainer::nbits
static constexpr uint32_t nbits()
Definition: HistoContainer.h:178
S
double S(const TLorentzVector &, const TLorentzVector &)
Definition: Particle.cc:97
cms::cuda::HistoContainer::sizeT
static constexpr uint32_t sizeT()
Definition: HistoContainer.h:174
a
double a
Definition: hdecay.h:119
cms::cudacompat::blockDim
const dim3 blockDim
Definition: cudaCompat.h:25
cms::cudacompat::dim3::x
uint32_t x
Definition: cudaCompat.h:21
HltBtagPostValidation_cff.c
c
Definition: HltBtagPostValidation_cff.py:31
cms::cudacompat::threadIdx
const dim3 threadIdx
Definition: cudaCompat.h:24
cms::cuda::HistoContainer::n
const index_type uint32_t n
Definition: HistoContainer.h:235
alignCSCRings.r
r
Definition: alignCSCRings.py:93
StopReason::SIZE
cms::cuda::HistoContainer::nbins
static constexpr uint32_t nbins()
Definition: HistoContainer.h:175
edm::shift
static unsigned const int shift
Definition: LuminosityBlockID.cc:7
cms::cuda::HistoContainer::bin
static constexpr UT bin(T t)
Definition: HistoContainer.h:183
cms::cuda::HistoContainer::v
const index_type * v
Definition: HistoContainer.h:235
cms::cuda::HistoContainer::b
b
Definition: HistoContainer.h:279
cms::cuda::co
VT * co
Definition: prefixScan.h:46
cms::cuda::HistoContainer::totbins
static constexpr uint32_t totbins()
Definition: HistoContainer.h:177
submitPVValidationJobs.t
string t
Definition: submitPVValidationJobs.py:644
cms::cudacompat::blockIdx
thread_local dim3 blockIdx
Definition: cudaCompat.cc:5
Generated for CMSSW Reference Manual by doxygen 1.8.16