1 #ifndef HeterogeneousCore_AlpakaInterface_interface_HistoContainer_h 2 #define HeterogeneousCore_AlpakaInterface_interface_HistoContainer_h 9 #include <alpaka/alpaka.hpp> 21 template <
typename TAcc,
typename Histo,
typename T>
23 Histo *__restrict__
h,
25 T const *__restrict__
v,
26 uint32_t
const *__restrict__
offsets)
const {
34 h->count(acc,
v[
i], ih);
40 template <
typename TAcc,
typename Histo,
typename T>
42 Histo *__restrict__
h,
44 T const *__restrict__
v,
45 uint32_t
const *__restrict__
offsets)
const {
53 h->fill(acc,
v[
i],
i, ih);
58 template <
typename TAcc,
typename Histo,
typename T,
typename TQueue>
61 T const *__restrict__
v,
62 uint32_t
const *__restrict__
offsets,
68 const auto threadsPerBlockOrElementsPerThread =
nthreads;
70 const auto workDiv = make_workdiv<TAcc>(blocksPerGrid, threadsPerBlockOrElementsPerThread);
78 template <
typename TAcc,
typename Histo,
typename T,
typename TQueue>
82 T const *__restrict__
v,
83 uint32_t
const *__restrict__
offsets,
89 const auto threadsPerBlockOrElementsPerThread =
nthreads;
91 const auto workDiv = make_workdiv<TAcc>(blocksPerGrid, threadsPerBlockOrElementsPerThread);
100 template <
typename Hist,
typename V,
typename Func>
106 for (
auto pj =
hist.begin(
bs); pj <
hist.end(
be); ++pj) {
112 template <
typename Hist,
typename V,
typename Func>
117 for (
auto pj =
hist.begin(
bs); pj <
hist.end(
be); ++pj) {
122 template <
typename T,
125 uint32_t
S =
sizeof(
T) * 8,
126 typename I = uint32_t,
138 constexpr uint32_t
b[] = {0x2, 0xC, 0xF0, 0xFF00, 0xFFFF0000};
142 for (
auto i = 4;
i >= 0;
i--)
164 template <
typename TAcc>
165 ALPAKA_FN_ACC ALPAKA_FN_INLINE
void count(
const TAcc &acc,
T t) {
168 Base::atomicIncrement(acc, this->off[
b]);
171 template <
typename TAcc>
175 auto w = Base::atomicDecrement(acc, this->off[
b]);
180 template <
typename TAcc>
181 ALPAKA_FN_ACC ALPAKA_FN_INLINE
void count(
const TAcc &acc,
T t, uint32_t
nh) {
186 Base::atomicIncrement(acc, this->off[
b]);
189 template <
typename TAcc>
195 auto w = Base::atomicDecrement(acc, this->off[
b]);
201 #endif // HeterogeneousCore_AlpakaInterface_interface_HistoContainer_h cms::alpakatools::HistoContainer< uint8_t, 256, 16000, 8, uint16_t > Hist
ALPAKA_FN_HOST_ACC constexpr RandomIt upper_bound(RandomIt first, RandomIt last, const T &value, Compare comp={})
uint32_t T const *__restrict__ uint32_t const *__restrict__ offsets
std::function< unsigned int(align::ID)> Counter
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t V
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V int Func func
uint32_t T const *__restrict__ v
const std::complex< double > I
def template(fileName, svg, replaceme="REPLACEME")
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int nthreads
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t int Histo::index_type cudaStream_t Func __host__ __device__ V int n
static unsigned int const shift
__host__ __device__ V wmin
The Signals That Services Can Subscribe To This is based on ActivityRegistry h
Helper function to determine trigger accepts.
ALPAKA_ASSERT_ACC(offsets)
uint32_t T const *__restrict__ uint32_t const *__restrict__ int32_t totSize
__host__ __device__ V V wmax