cms::alpakatools Namespace Reference

Namespaces
	config
	detail
	impl
	traits

Classes
class	AtomicPairCounter
class	CachingAllocator
struct	CopyToDevice
struct	CopyToDevice< cms::alpakatest::AlpakaESTestDataB< alpaka_common::DevHost > >
struct	CopyToDevice< cms::alpakatest::AlpakaESTestDataEHost >
struct	CopyToDevice< PixelCPEFastParamsHost< TrackerTraits > >
struct	CopyToDevice< PortableHostCollection< TLayout > >
struct	CopyToDevice< PortableHostMultiCollection< TDev, T0, Args... > >
struct	CopyToDevice< PortableHostObject< TProduct > >
struct	CopyToHost
struct	CopyToHost< PortableDeviceCollection< TLayout, TDevice > >
struct	CopyToHost< PortableDeviceMultiCollection< TDev, T0, Args... > >
struct	CopyToHost< PortableDeviceObject< TProduct, TDevice > >
struct	CopyToHost< SiPixelClustersDevice< TDevice > >
struct	CopyToHost< SiPixelDigiErrorsDevice< TDevice > >
struct	CopyToHost< SiPixelDigisDevice< TDevice > >
struct	CopyToHost< TrackingRecHitDevice< TrackerTraits, TDevice > >
struct	CopyToHost< TracksDevice< TrackerTraits, TDevice > >
struct	CopyToHost< ZVertexDevice< TDevice > >
struct	countFromVector
struct	ElementIndex
class	EventCache
struct	fillFromVector
class	FlexiStorage
class	FlexiStorage< I, -1 >
class	HistoContainer
class	independent_group_elements_along
class	independent_groups_along
struct	multiBlockPrefixScan
class	OneToManyAssocBase
class	OneToManyAssocRandomAccess
class	OneToManyAssocSequential
class	QueueCache
struct	radixSortMultiWrapper
struct	radixSortMultiWrapper2
struct	requires_single_thread_per_block
class	ScopedContextAcquire
class	ScopedContextAnalyze
class	ScopedContextProduce
class	ScopedContextTask
struct	SimpleVector
class	uniform_elements_along
class	uniform_elements_nd
class	uniform_group_elements_along
class	uniform_groups_along
class	VecArray

Typedefs
template<typename TDev , typename T , typename = std::enable_if_t<alpaka::isDevice<TDev>>>
using	const_device_buffer = alpaka::ViewConst< device_buffer< TDev, T > >
template<typename T >
using	const_host_buffer = alpaka::ViewConst< host_buffer< T > >
template<typename TDev , typename T , typename = std::enable_if_t<alpaka::isDevice<TDev>>>
using	device_buffer = typename detail::buffer_type< TDev, T >::type
template<typename TDev , typename T , typename = std::enable_if_t<alpaka::isDevice<TDev>>>
using	device_view = typename detail::view_type< TDev, T >::type
template<typename T >
using	host_buffer = typename detail::buffer_type< DevHost, T >::type
template<typename T >
using	host_view = typename detail::view_type< DevHost, T >::type

Enumerations
enum	AllocatorPolicy { AllocatorPolicy::Synchronous = 0, AllocatorPolicy::Asynchronous = 1, AllocatorPolicy::Caching = 2 }
enum	Backend : unsigned short {   Backend::SerialSync = 0, Backend::CudaAsync = 1, Backend::ROCmAsync = 2, Backend::TbbAsync = 3,   Backend::size }

Functions
template<typename TElem , typename TIdx , typename TExtent , typename TQueue , typename TDev , typename = std::enable_if_t<alpaka::isDevice<TDev> and alpaka::isQueue<TQueue>>>
ALPAKA_FN_HOST auto	allocCachedBuf (TDev const &dev, TQueue queue, TExtent const &extent=TExtent())
template<typename TAcc , typename T >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	blockPrefixScan (const TAcc &acc, T const *ci, T *co, int32_t size, T *ws=nullptr)
template<typename TAcc , typename T >
ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void	blockPrefixScan (const TAcc &acc, T *__restrict__ c, int32_t size, T *__restrict__ ws=nullptr)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	blocks_with_stride (TAcc const &acc, TArgs... args)
template<typename TPlatform , typename = std::enable_if_t<alpaka::isPlatform<TPlatform>>>
alpaka::Dev< TPlatform > const &	chooseDevice (edm::StreamID id)
template<typename TPlatform , typename = std::enable_if_t<alpaka::isPlatform<TPlatform>>>
std::vector< alpaka::Dev< TPlatform > > const &	devices ()
constexpr Idx	divide_up_by (Idx value, Idx divisor)
template<typename TAcc , typename T >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	dummyReorder (const TAcc &acc, T const *a, uint16_t *ind, uint16_t *ind2, uint32_t size)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	elements_in_block (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	elements_with_stride (TAcc const &acc, TArgs... args)
template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>
ALPAKA_FN_ACC auto	elements_with_stride_nd (TAcc const &acc)
template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>
ALPAKA_FN_ACC auto	elements_with_stride_nd (TAcc const &acc, alpaka::Vec< alpaka::Dim< TAcc >, Idx > extent)
template<typename TAcc , typename Histo , typename T , typename TQueue >
ALPAKA_FN_INLINE void	fillManyFromVector (Histo *__restrict__ h, uint32_t nh, T const *__restrict__ v, uint32_t const *__restrict__ offsets, uint32_t totSize, uint32_t nthreads, TQueue &queue)
template<typename TAcc , typename Histo , typename T , typename TQueue >
ALPAKA_FN_INLINE void	fillManyFromVector (Histo *__restrict__ h, typename Histo::View hv, uint32_t nh, T const *__restrict__ v, uint32_t const *__restrict__ offsets, uint32_t totSize, uint32_t nthreads, TQueue &queue)
template<typename Hist , typename V , typename Func >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	forEachInBins (Hist const &hist, V value, int n, Func func)
template<typename Hist , typename V , typename Func >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	forEachInWindow (Hist const &hist, V wmin, V wmax, Func const &func)
template<typename TDev , typename TQueue , typename = std::enable_if_t<alpaka::isDevice<TDev> and alpaka::isQueue<TQueue>>>
CachingAllocator< TDev, TQueue > &	getDeviceCachingAllocator (TDev const &device)
template<typename Event >
EventCache< Event > &	getEventCache ()
template<typename TQueue , typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
CachingAllocator< alpaka_common::DevHost, TQueue > &	getHostCachingAllocator ()
template<typename Queue >
QueueCache< Queue > &	getQueueCache ()
alpaka::DevCpu const &	host ()
alpaka::PlatformCpu const &	host_platform ()
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	independent_group_elements (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>
ALPAKA_FN_ACC auto	independent_group_elements_x (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>
ALPAKA_FN_ACC auto	independent_group_elements_y (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>
ALPAKA_FN_ACC auto	independent_group_elements_z (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	independent_groups (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>
ALPAKA_FN_ACC auto	independent_groups_x (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>
ALPAKA_FN_ACC auto	independent_groups_y (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>
ALPAKA_FN_ACC auto	independent_groups_z (TAcc const &acc, TArgs... args)
template<typename T , typename = std::enable_if_t<std::is_integral_v<T>>>
constexpr bool	isPowerOf2 (T v)
template<typename T , typename TDev >
std::enable_if_t< alpaka::isDevice< TDev > and not std::is_array_v< T >, device_buffer< TDev, T > >	make_device_buffer (TDev const &device)
template<typename T , typename TDev >
std::enable_if_t< alpaka::isDevice< TDev > and cms::is_unbounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, device_buffer< TDev, T > >	make_device_buffer (TDev const &device, Extent extent)
template<typename T , typename TDev >
std::enable_if_t< alpaka::isDevice< TDev > and cms::is_bounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, device_buffer< TDev, T > >	make_device_buffer (TDev const &device)
template<typename T , typename TQueue >
std::enable_if_t< alpaka::isQueue< TQueue > and not std::is_array_v< T >, device_buffer< alpaka::Dev< TQueue >, T > >	make_device_buffer (TQueue const &queue)
template<typename T , typename TQueue >
std::enable_if_t< alpaka::isQueue< TQueue > and cms::is_unbounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, device_buffer< alpaka::Dev< TQueue >, T > >	make_device_buffer (TQueue const &queue, Extent extent)
template<typename T , typename TQueue >
std::enable_if_t< alpaka::isQueue< TQueue > and cms::is_bounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, device_buffer< alpaka::Dev< TQueue >, T > >	make_device_buffer (TQueue const &queue)
template<typename T , typename TDev >
std::enable_if_t< not std::is_array_v< T >, device_view< TDev, T > >	make_device_view (TDev const &device, T &data)
template<typename T , typename TDev >
device_view< TDev, T[]>	make_device_view (TDev const &device, T *data, Extent extent)
template<typename T , typename TDev >
std::enable_if_t< cms::is_unbounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, device_view< TDev, T > >	make_device_view (TDev const &device, T &data, Extent extent)
template<typename T , typename TDev >
std::enable_if_t< cms::is_bounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, device_view< TDev, T > >	make_device_view (TDev const &device, T &data)
template<typename T >
std::enable_if_t< not std::is_array_v< T >, host_buffer< T > >	make_host_buffer ()
template<typename T >
std::enable_if_t< cms::is_unbounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_buffer< T > >	make_host_buffer (Extent extent)
template<typename T >
std::enable_if_t< cms::is_bounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_buffer< T > >	make_host_buffer ()
template<typename T , typename TPlatform >
std::enable_if_t< not std::is_array_v< T >, host_buffer< T > >	make_host_buffer ()
template<typename T , typename TPlatform >
std::enable_if_t< cms::is_unbounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_buffer< T > >	make_host_buffer (Extent extent)
template<typename T , typename TPlatform >
std::enable_if_t< cms::is_bounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_buffer< T > >	make_host_buffer ()
template<typename T , typename TQueue >
std::enable_if_t< alpaka::isQueue< TQueue > and not std::is_array_v< T >, host_buffer< T > >	make_host_buffer (TQueue const &queue)
template<typename T , typename TQueue >
std::enable_if_t< alpaka::isQueue< TQueue > and cms::is_unbounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_buffer< T > >	make_host_buffer (TQueue const &queue, Extent extent)
template<typename T , typename TQueue >
std::enable_if_t< alpaka::isQueue< TQueue > and cms::is_bounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_buffer< T > >	make_host_buffer (TQueue const &queue)
template<typename T >
std::enable_if_t< not std::is_array_v< T >, host_view< T > >	make_host_view (T &data)
template<typename T >
host_view< T[]>	make_host_view (T *data, Extent extent)
template<typename T >
std::enable_if_t< cms::is_unbounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_view< T > >	make_host_view (T &data, Extent extent)
template<typename T >
std::enable_if_t< cms::is_bounded_array_v< T > and not std::is_array_v< std::remove_extent_t< T > >, host_view< T > >	make_host_view (T &data)
template<class T >
SimpleVector< T >	make_SimpleVector (int capacity, T *data)
template<class T >
SimpleVector< T > *	make_SimpleVector (SimpleVector< T > *mem, int capacity, T *data)
template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
WorkDiv< Dim1D >	make_workdiv (Idx blocks, Idx elements)
template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
WorkDiv< alpaka::Dim< TAcc > >	make_workdiv (const Vec< alpaka::Dim< TAcc >> &blocks, const Vec< alpaka::Dim< TAcc >> &elements)
void	module_backend_config (edm::ConfigurationDescriptions &iDesc)
template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
ALPAKA_FN_ACC constexpr bool	once_per_block (TAcc const &acc)
template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
ALPAKA_FN_ACC constexpr bool	once_per_grid (TAcc const &acc)
template<typename TPlatform , typename = std::enable_if_t<alpaka::isPlatform<TPlatform>>>
TPlatform const &	platform ()
template<typename TAcc , typename T , int NS = sizeof(T), typename std::enable_if< std::is_unsigned< T >::value &&!requires_single_thread_per_block_v< TAcc >, T >::type * = nullptr>
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	radixSort (const TAcc &acc, T const *a, uint16_t *ind, uint16_t *ind2, uint32_t size)
template<typename TAcc , typename T , int NS, typename RF >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	radixSortImpl (const TAcc &acc, T const *__restrict__ a, uint16_t *ind, uint16_t *ind2, uint32_t size, RF reorder)
template<typename TAcc , typename T , int NS = sizeof(T)>
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	radixSortMulti (const TAcc &acc, T const *v, uint16_t *index, uint32_t const *offsets, uint16_t *workspace)
template<typename TAcc , typename T >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	reorderFloat (const TAcc &acc, T const *a, uint16_t *ind, uint16_t *ind2, uint32_t size)
template<typename TAcc , typename T >
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	reorderSigned (const TAcc &acc, T const *a, uint16_t *ind, uint16_t *ind2, uint32_t size)
constexpr Idx	round_up_by (Idx value, Idx divisor)
Backend	toBackend (std::string_view name)
std::string_view	toString (Backend backend)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	uniform_elements (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>
ALPAKA_FN_ACC auto	uniform_elements_x (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>
ALPAKA_FN_ACC auto	uniform_elements_y (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>
ALPAKA_FN_ACC auto	uniform_elements_z (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	uniform_group_elements (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>
ALPAKA_FN_ACC auto	uniform_group_elements_x (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>
ALPAKA_FN_ACC auto	uniform_group_elements_y (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>
ALPAKA_FN_ACC auto	uniform_group_elements_z (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>
ALPAKA_FN_ACC auto	uniform_groups (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>
ALPAKA_FN_ACC auto	uniform_groups_x (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>
ALPAKA_FN_ACC auto	uniform_groups_y (TAcc const &acc, TArgs... args)
template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>
ALPAKA_FN_ACC auto	uniform_groups_z (TAcc const &acc, TArgs... args)
template<typename TAcc , typename T , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	warpPrefixScan (const TAcc &acc, int32_t laneId, T const *ci, T *co, uint32_t i, bool active=true)
template<typename TAcc , typename T , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
ALPAKA_FN_ACC ALPAKA_FN_INLINE void	warpPrefixScan (const TAcc &acc, int32_t laneId, T *c, uint32_t i, bool active=true)

Variables
template<typename TDev , typename = std::enable_if_t<alpaka::isDevice<TDev>>>
constexpr AllocatorPolicy	allocator_policy = AllocatorPolicy::Synchronous
template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>
constexpr bool	requires_single_thread_per_block_v = requires_single_thread_per_block<TAcc>::value

Typedef Documentation

const_device_buffer

template<typename TDev , typename T , typename = std::enable_if_t<alpaka::isDevice<TDev>>>

using cms::alpakatools::const_device_buffer = typedef alpaka::ViewConst<device_buffer<TDev, T> >

Definition at line 180 of file memory.h.

const_host_buffer

template<typename T >

using cms::alpakatools::const_host_buffer = typedef alpaka::ViewConst<host_buffer<T> >

Definition at line 60 of file memory.h.

device_buffer

template<typename TDev , typename T , typename = std::enable_if_t<alpaka::isDevice<TDev>>>

using cms::alpakatools::device_buffer = typedef typename detail::buffer_type<TDev, T>::type

Definition at line 177 of file memory.h.

device_view

template<typename TDev , typename T , typename = std::enable_if_t<alpaka::isDevice<TDev>>>

using cms::alpakatools::device_view = typedef typename detail::view_type<TDev, T>::type

Definition at line 257 of file memory.h.

host_buffer

template<typename T >

using cms::alpakatools::host_buffer = typedef typename detail::buffer_type<DevHost, T>::type

Definition at line 57 of file memory.h.

host_view

template<typename T >

using cms::alpakatools::host_view = typedef typename detail::view_type<DevHost, T>::type

Definition at line 150 of file memory.h.

Enumeration Type Documentation

AllocatorPolicy

enum cms::alpakatools::AllocatorPolicy

strong

Enumerator
Synchronous
Asynchronous
Caching

Definition at line 14 of file AllocatorPolicy.h.

{ Synchronous = 0, Asynchronous = 1, Caching = 2 };

Backend

enum cms::alpakatools::Backend : unsigned short

strong

Enumerator
SerialSync
CudaAsync
ROCmAsync
TbbAsync
size

Definition at line 8 of file Backend.h.

: unsigned short { SerialSync = 0, CudaAsync = 1, ROCmAsync = 2, TbbAsync = 3, size };

Function Documentation

allocCachedBuf()

template<typename TElem , typename TIdx , typename TExtent , typename TQueue , typename TDev , typename = std::enable_if_t<alpaka::isDevice<TDev> and alpaka::isQueue<TQueue>>>

ALPAKA_FN_HOST auto cms::alpakatools::allocCachedBuf	(	TDev const &	dev,
		TQueue	queue,
		TExtent const &	extent = TExtent()
	)

Definition at line 197 of file CachedBufAlloc.h.

References createBeamHaloJobs::queue.

                                                                                                       {
    return traits::CachedBufAlloc<TElem, alpaka::Dim<TExtent>, TIdx, TDev, TQueue>::allocCachedBuf(dev, queue, extent);
  }

blockPrefixScan() [1/2]

template<typename TAcc , typename T >

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::blockPrefixScan	(	const TAcc &	acc,
		T const *	ci,
		T *	co,
		int32_t	size,
		T *	ws = nullptr
	)

Definition at line 47 of file prefixScan.h.

References ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), cms::cuda::co, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), dqmdumpme::first, mps_fire::i, isPowerOf2(), warpPrefixScan(), and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ws.

Referenced by cms::alpakatools::OneToManyAssocRandomAccess< I, NHISTS *NBINS+1, SIZE >::finalize(), gpuClustering::for(), pixelClustering::ClusterChargeCut< TrackerTraits >::operator()(), cms::alpakatools::multiBlockPrefixScan< T >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::FillHitsModuleStart< TrackerTraits >::operator()().

                                                                          {
    if constexpr (!requires_single_thread_per_block_v<TAcc>) {
      const auto warpSize = alpaka::warp::getSize(acc);
      int32_t const blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Threads>(acc)[0u]);
      int32_t const blockThreadIdx(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
      ALPAKA_ASSERT_ACC(ws);
      ALPAKA_ASSERT_ACC(size <= warpSize * warpSize);
      ALPAKA_ASSERT_ACC(0 == blockDimension % warpSize);
      auto first = blockThreadIdx;
      ALPAKA_ASSERT_ACC(isPowerOf2(warpSize));
      auto laneId = blockThreadIdx & (warpSize - 1);
      auto warpUpRoundedSize = (size + warpSize - 1) / warpSize * warpSize;

      for (auto i = first; i < warpUpRoundedSize; i += blockDimension) {
        // When padding the warp, warpPrefixScan is a noop
        warpPrefixScan(acc, laneId, ci, co, i, i < size);
        if (i < size) {
          // Skipped in warp padding threads.
          auto warpId = i / warpSize;
          ALPAKA_ASSERT_ACC(warpId < warpSize);
          if ((warpSize - 1) == laneId)
            ws[warpId] = co[i];
        }
      }
      alpaka::syncBlockThreads(acc);
      if (size <= warpSize)
        return;
      if (blockThreadIdx < warpSize) {
        warpPrefixScan(acc, laneId, ws, blockThreadIdx);
      }
      alpaka::syncBlockThreads(acc);
      for (auto i = first + warpSize; i < size; i += blockDimension) {
        int32_t warpId = i / warpSize;
        co[i] += ws[warpId - 1];
      }
      alpaka::syncBlockThreads(acc);
    } else {
      co[0] = ci[0];
      for (int32_t i = 1; i < size; ++i)
        co[i] = ci[i] + co[i - 1];
    }
  }

blockPrefixScan() [2/2]

template<typename TAcc , typename T >

ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE void cms::alpakatools::blockPrefixScan	(	const TAcc &	acc,
		T *__restrict__	c,
		int32_t	size,
		T *__restrict__	ws = nullptr
	)

Definition at line 92 of file prefixScan.h.

References ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), HltBtagPostValidation_cff::c, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), dqmdumpme::first, mps_fire::i, warpPrefixScan(), and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ws.

                                                                                         {
    if constexpr (!requires_single_thread_per_block_v<TAcc>) {
      const auto warpSize = alpaka::warp::getSize(acc);
      int32_t const blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Threads>(acc)[0u]);
      int32_t const blockThreadIdx(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
      ALPAKA_ASSERT_ACC(ws);
      ALPAKA_ASSERT_ACC(size <= warpSize * warpSize);
      ALPAKA_ASSERT_ACC(0 == blockDimension % warpSize);
      auto first = blockThreadIdx;
      auto laneId = blockThreadIdx & (warpSize - 1);
      auto warpUpRoundedSize = (size + warpSize - 1) / warpSize * warpSize;

      for (auto i = first; i < warpUpRoundedSize; i += blockDimension) {
        // When padding the warp, warpPrefixScan is a noop
        warpPrefixScan(acc, laneId, c, i, i < size);
        if (i < size) {
          // Skipped in warp padding threads.
          auto warpId = i / warpSize;
          ALPAKA_ASSERT_ACC(warpId < warpSize);
          if ((warpSize - 1) == laneId)
            ws[warpId] = c[i];
        }
      }
      alpaka::syncBlockThreads(acc);
      if (size <= warpSize)
        return;
      if (blockThreadIdx < warpSize) {
        warpPrefixScan(acc, laneId, ws, blockThreadIdx);
      }
      alpaka::syncBlockThreads(acc);
      for (auto i = first + warpSize; i < size; i += blockDimension) {
        auto warpId = i / warpSize;
        c[i] += ws[warpId - 1];
      }
      alpaka::syncBlockThreads(acc);
    } else {
      for (int32_t i = 1; i < size; ++i)
        c[i] += c[i - 1];
    }
  }

blocks_with_stride()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::blocks_with_stride ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 784 of file workdivision.h.

References writedatasetfile::args.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_prep_1d_and_initialize::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_time_compute_makeratio::operator()().

                                                                               {
    return uniform_groups_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

chooseDevice()

template<typename TPlatform , typename = std::enable_if_t<alpaka::isPlatform<TPlatform>>>

alpaka::Dev<TPlatform> const& cms::alpakatools::chooseDevice

(

edm::StreamID

id

)

Definition at line 16 of file chooseDevice.h.

References cms::Exception::addContext(), and devices().

                                                           {
    edm::Service<ALPAKA_TYPE_ALIAS(AlpakaService)> service;
    if (not service->enabled()) {
      cms::Exception ex("RuntimeError");
      ex << "Unable to choose current device because " << ALPAKA_TYPE_ALIAS_NAME(AlpakaService) << " is disabled.\n"
         << "If " << ALPAKA_TYPE_ALIAS_NAME(AlpakaService) << " was not explicitly disabled in the configuration,\n"
         << "the probable cause is that there is no accelerator or there is some problem\n"
         << "with the accelerator runtime or drivers.";
      ex.addContext("Calling cms::alpakatools::chooseDevice()");
      throw ex;
    }

    // For startes we "statically" assign the device based on
    // edm::Stream number. This is suboptimal if the number of
    // edm::Streams is not a multiple of the number of devices
    // (and even then there is no load balancing).

    // TODO: improve the "assignment" logic
    auto const& devices = cms::alpakatools::devices<TPlatform>();
    return devices[id % devices.size()];
  }

devices()

template<typename TPlatform , typename = std::enable_if_t<alpaka::isPlatform<TPlatform>>>

std::vector<alpaka::Dev<TPlatform> > const& cms::alpakatools::devices ( )

inline

Definition at line 22 of file devices.h.

Referenced by cms::alpakatools::detail::allocate_device_allocators(), ALPAKA_ACCELERATOR_NAMESPACE::AlpakaService::AlpakaService(), SiStripFedCablingBuilderFromDb::buildFecCabling(), ALPAKA_ACCELERATOR_NAMESPACE::detail::chooseDevice(), chooseDevice(), FastFedCablingHistosUsingDb::connections(), getDeviceCachingAllocator(), SiStripConfigDb::printDeviceDescriptions(), ALPAKA_ACCELERATOR_NAMESPACE::ESProducer::setWhatProducedDevice(), cms::cudatest::testDevices(), cms::rocmtest::testDevices(), ApvTimingHistosUsingDb::update(), OptoScanHistosUsingDb::update(), VpspScanHistosUsingDb::update(), LatencyHistosUsingDb::update(), CalibrationHistosUsingDb::update(), FineDelayHistosUsingDb::update(), VpspScanHistosUsingDb::uploadConfigurations(), OptoScanHistosUsingDb::uploadConfigurations(), ApvTimingHistosUsingDb::uploadConfigurations(), LatencyHistosUsingDb::uploadConfigurations(), CalibrationHistosUsingDb::uploadConfigurations(), and FineDelayHistosUsingDb::uploadConfigurations().

                                                          {
    // enumerate all devices the first time that this function is called
    static const auto devices = alpaka::getDevs(platform<TPlatform>());
    return devices;
  }

divide_up_by()

constexpr Idx cms::alpakatools::divide_up_by ( Idx  value, Idx  divisor  )

inline

Definition at line 19 of file workdivision.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::buildDoublets(), ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::classifyTuples(), ALPAKA_ACCELERATOR_NAMESPACE::PFClusterProducerKernel::execute(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgo::fill(), fillManyFromVector(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgo::fillMulti2(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgo::fillMulti3(), ALPAKA_ACCELERATOR_NAMESPACE::HelixFit< TrackerTraits >::launchBrokenLineKernels(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::launchKernels(), ALPAKA_ACCELERATOR_NAMESPACE::CAHitNtupletGeneratorKernels< TTTraits >::launchKernels(), ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::SiPixelRawToClusterKernel< pixelTopology::Phase2 >::makePhase1ClustersAsync(), ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::SiPixelRawToClusterKernel< pixelTopology::Phase2 >::makePhase2ClustersAsync(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::minimization_procedure(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitProducerKernel< CAL >::PFRecHitProducerKernel(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgo::update(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgo::updateMulti2(), and ALPAKA_ACCELERATOR_NAMESPACE::TestAlgo::updateMulti3().

{ return (value + divisor - 1) / divisor; }

dummyReorder()

template<typename TAcc , typename T >

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::dummyReorder	(	const TAcc &	acc,
		T const *	a,
		uint16_t *	ind,
		uint16_t *	ind2,
		uint32_t	size
	)

Definition at line 17 of file radixSort.h.

                                                                                 {}

elements_in_block()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::elements_in_block ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 999 of file workdivision.h.

References writedatasetfile::args.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_prep_1d_and_initialize::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_time_compute_makeratio::operator()().

                                                                              {
    return uniform_group_elements_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

elements_with_stride()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::elements_with_stride ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 346 of file workdivision.h.

References writedatasetfile::args.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitProducerKernelConstruct< CAL >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernel::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_time_compute_nullhypot::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoMultiKernel2::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_minimize::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoMultiKernel3::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ECLCCInit::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ECLCCCompute1::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitProducerKernelTopology< CAL >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ECLCCFlatten::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernelUpdate::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernelUpdateMulti2::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernelUpdateMulti3::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_time_compute_findamplchi2_and_finish::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_time_compute_fixMGPAslew::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_time_computation_init::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_time_correction_and_finalize::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::SeedingTopoThresh::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::PrepareTopoInputs::operator()().

                                                                                 {
    return uniform_elements_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

elements_with_stride_nd() [1/2]

template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>

ALPAKA_FN_ACC auto cms::alpakatools::elements_with_stride_nd ( TAcc const &  acc )

inline

Definition at line 565 of file workdivision.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::ecal::multifit::Kernel_prep_2d::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::FillRhfIndex::operator()().

                                                                     {
    return uniform_elements_nd<TAcc>(acc);
  }

elements_with_stride_nd() [2/2]

template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>

ALPAKA_FN_ACC auto cms::alpakatools::elements_with_stride_nd ( TAcc const &  acc, alpaka::Vec< alpaka::Dim< TAcc >, Idx >  extent  )

inline

Definition at line 570 of file workdivision.h.

                                                                                                             {
    return uniform_elements_nd<TAcc>(acc, extent);
  }

fillManyFromVector() [1/2]

template<typename TAcc , typename Histo , typename T , typename TQueue >

ALPAKA_FN_INLINE void cms::alpakatools::fillManyFromVector	(	Histo *__restrict__	h,
		uint32_t	nh,
		T const *__restrict__	v,
		uint32_t const *__restrict__	offsets,
		uint32_t	totSize,
		uint32_t	nthreads,
		TQueue &	queue
	)

Definition at line 59 of file HistoContainer.h.

References divide_up_by(), h, cms::cuda::nh, cms::cuda::nthreads, cms::cuda::offsets, createBeamHaloJobs::queue, svgfig::template(), cms::cuda::totSize, and cms::cuda::v.

Referenced by SiPixelRecHitSoAFromLegacyT< TrackerTraits >::produce().

                                                          {
    Histo::template launchZero<TAcc>(h, queue);

    const auto threadsPerBlockOrElementsPerThread = nthreads;
    const auto blocksPerGrid = divide_up_by(totSize, nthreads);
    const auto workDiv = make_workdiv<TAcc>(blocksPerGrid, threadsPerBlockOrElementsPerThread);

    alpaka::exec<TAcc>(queue, workDiv, countFromVector(), h, nh, v, offsets);
    Histo::template launchFinalize<TAcc>(h, queue);

    alpaka::exec<TAcc>(queue, workDiv, fillFromVector(), h, nh, v, offsets);
  }

fillManyFromVector() [2/2]

template<typename TAcc , typename Histo , typename T , typename TQueue >

ALPAKA_FN_INLINE void cms::alpakatools::fillManyFromVector	(	Histo *__restrict__	h,
		typename Histo::View	hv,
		uint32_t	nh,
		T const *__restrict__	v,
		uint32_t const *__restrict__	offsets,
		uint32_t	totSize,
		uint32_t	nthreads,
		TQueue &	queue
	)

Definition at line 79 of file HistoContainer.h.

References divide_up_by(), h, cms::cuda::nh, cms::cuda::nthreads, cms::cuda::offsets, createBeamHaloJobs::queue, svgfig::template(), cms::cuda::totSize, and cms::cuda::v.

                                                          {
    Histo::template launchZero<TAcc>(hv, queue);

    const auto threadsPerBlockOrElementsPerThread = nthreads;
    const auto blocksPerGrid = divide_up_by(totSize, nthreads);
    const auto workDiv = make_workdiv<TAcc>(blocksPerGrid, threadsPerBlockOrElementsPerThread);

    alpaka::exec<TAcc>(queue, workDiv, countFromVector(), h, nh, v, offsets);
    Histo::template launchFinalize<TAcc>(h, queue);

    alpaka::exec<TAcc>(queue, workDiv, fillFromVector(), h, nh, v, offsets);
  }

forEachInBins()

template<typename Hist , typename V , typename Func >

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::forEachInBins	(	Hist const &	hist,
		V	value,
		int	n,
		Func	func
	)

Definition at line 101 of file HistoContainer.h.

References ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), cms::cuda::be, newFWLiteAna::bin, cms::cuda::bs, cms::cuda::func, compareTotals::hist, SiStripPI::max, SiStripPI::min, cms::cuda::n, and LaserClient_cfi::nbins.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksIterative::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksDBSCAN::operator()().

                                                                                                 {
    int bs = Hist::bin(value);
    int be = std::min(int(Hist::nbins() - 1), bs + n);
    bs = std::max(0, bs - n);
    ALPAKA_ASSERT_ACC(be >= bs);
    for (auto pj = hist.begin(bs); pj < hist.end(be); ++pj) {
      func(*pj);
    }
  }

forEachInWindow()

template<typename Hist , typename V , typename Func >

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::forEachInWindow	(	Hist const &	hist,
		V	wmin,
		V	wmax,
		Func const &	func
	)

Definition at line 113 of file HistoContainer.h.

References ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), cms::cuda::be, newFWLiteAna::bin, cms::cuda::bs, cms::cuda::func, compareTotals::hist, cms::cuda::wmax, and cms::cuda::wmin.

                                                                                                          {
    auto bs = Hist::bin(wmin);
    auto be = Hist::bin(wmax);
    ALPAKA_ASSERT_ACC(be >= bs);
    for (auto pj = hist.begin(bs); pj < hist.end(be); ++pj) {
      func(*pj);
    }
  }

getDeviceCachingAllocator()

template<typename TDev , typename TQueue , typename = std::enable_if_t<alpaka::isDevice<TDev> and alpaka::isQueue<TQueue>>>

CachingAllocator<TDev, TQueue>& cms::alpakatools::getDeviceCachingAllocator ( TDev const &  device )

inline

Definition at line 78 of file getDeviceCachingAllocator.h.

References cms::cuda::assert(), CMS_THREAD_SAFE, and devices().

                                                                                       {
    // initialise all allocators, one per device
    CMS_THREAD_SAFE static auto allocators = detail::allocate_device_allocators<TDev, TQueue>();

    size_t const index = alpaka::getNativeHandle(device);
    assert(index < cms::alpakatools::devices<alpaka::Platform<TDev>>().size());

    // the public interface is thread safe
    return allocators[index];
  }

getEventCache()

template<typename Event >

EventCache<Event>& cms::alpakatools::getEventCache

(

)

Definition at line 92 of file EventCache.h.

References utilities::cache(), and CMS_THREAD_SAFE.

                                     {
    // the public interface is thread safe
    CMS_THREAD_SAFE static EventCache<Event> cache;
    return cache;
  }

getHostCachingAllocator()

template<typename TQueue , typename = std::enable_if_t<alpaka::isQueue<TQueue>>>

CachingAllocator<alpaka_common::DevHost, TQueue>& cms::alpakatools::getHostCachingAllocator ( )

inline

Definition at line 16 of file getHostCachingAllocator.h.

References cms::alpakatools::config::binGrowth, CMS_THREAD_SAFE, host(), cms::alpakatools::config::maxBin, cms::alpakatools::config::maxCachedBytes, cms::alpakatools::config::maxCachedFraction, and cms::alpakatools::config::minBin.

                                                                                   {
    // thread safe initialisation of the host allocator
    CMS_THREAD_SAFE static CachingAllocator<alpaka_common::DevHost, TQueue> allocator(
        host(),
        config::binGrowth,
        config::minBin,
        config::maxBin,
        config::maxCachedBytes,
        config::maxCachedFraction,
        false,   // reuseSameQueueAllocations
        false);  // debug

    // the public interface is thread safe
    return allocator;
  }

getQueueCache()

template<typename Queue >

QueueCache<Queue>& cms::alpakatools::getQueueCache

(

)

Definition at line 65 of file QueueCache.h.

References utilities::cache(), and CMS_THREAD_SAFE.

                                     {
    // the public interface is thread safe
    CMS_THREAD_SAFE static QueueCache<Queue> cache;
    return cache;
  }

host()

alpaka::DevCpu const& cms::alpakatools::host ( )

inline

Definition at line 14 of file host.h.

Referenced by getHostCachingAllocator(), make_host_buffer(), make_host_view(), ALPAKA_ACCELERATOR_NAMESPACE::PFClusterParamsESProducer::PFClusterParamsESProducer(), ALPAKA_ACCELERATOR_NAMESPACE::EcalMultifitParametersHostESProducer::produce(), ALPAKA_ACCELERATOR_NAMESPACE::EcalElectronicsMappingHostESProducer::produce(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitECALParamsESProducer::produce(), ALPAKA_ACCELERATOR_NAMESPACE::SiPixelGainCalibrationForHLTSoAESProducer::produce(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerC::produce(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitHCALParamsESProducer::produce(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerA::produce(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlpakaESProducerE::produce(), ALPAKA_ACCELERATOR_NAMESPACE::SiPixelCablingSoAESProducer::produce(), ALPAKA_ACCELERATOR_NAMESPACE::PFRecHitTopologyESProducer< CAL >::produce(), ALPAKA_ACCELERATOR_NAMESPACE::EcalMultifitConditionsHostESProducer::produce(), PortableHostObject< T >::ROOTReadStreamer(), PortableHostCollection< TestSoA >::ROOTReadStreamer(), and PortableHostMultiCollection< TestSoA, TestSoA2 >::ROOTReadStreamer().

{ return devices<alpaka::PlatformCpu>()[0]; }

host_platform()

alpaka::PlatformCpu const& cms::alpakatools::host_platform ( )

inline

Definition at line 11 of file host.h.

{ return platform<alpaka::PlatformCpu>(); }

independent_group_elements()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::independent_group_elements ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1281 of file workdivision.h.

References writedatasetfile::args.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::for(), pixelClustering::ClusterChargeCut< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelRecHits::GetHits< TrackerTraits >::operator()(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::zeroAndInit::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::FindClus< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::FillHitsModuleStart< TrackerTraits >::operator()(), radixSortImpl(), reorderFloat(), and reorderSigned().

                                                                                       {
    return independent_group_elements_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

independent_group_elements_x()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>

ALPAKA_FN_ACC auto cms::alpakatools::independent_group_elements_x ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1294 of file workdivision.h.

References writedatasetfile::args.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::for(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::CAFishbone< TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_connect< TrackerTraits >::operator()().

                                                                                         {
    return independent_group_elements_along<TAcc, alpaka::Dim<TAcc>::value - 1>(acc, static_cast<Idx>(args)...);
  }

independent_group_elements_y()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>

ALPAKA_FN_ACC auto cms::alpakatools::independent_group_elements_y ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1301 of file workdivision.h.

References writedatasetfile::args.

                                                                                         {
    return independent_group_elements_along<TAcc, alpaka::Dim<TAcc>::value - 2>(acc, static_cast<Idx>(args)...);
  }

independent_group_elements_z()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>

ALPAKA_FN_ACC auto cms::alpakatools::independent_group_elements_z ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1308 of file workdivision.h.

References writedatasetfile::args.

                                                                                         {
    return independent_group_elements_along<TAcc, alpaka::Dim<TAcc>::value - 3>(acc, static_cast<Idx>(args)...);
  }

independent_groups()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::independent_groups ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1143 of file workdivision.h.

References writedatasetfile::args.

Referenced by pixelClustering::ClusterChargeCut< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelRecHits::GetHits< TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::FindClus< TrackerTraits >::operator()().

                                                                               {
    return independent_groups_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

independent_groups_x()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>

ALPAKA_FN_ACC auto cms::alpakatools::independent_groups_x ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1156 of file workdivision.h.

References writedatasetfile::args.

                                                                                 {
    return independent_groups_along<TAcc, alpaka::Dim<TAcc>::value - 1>(acc, static_cast<Idx>(args)...);
  }

independent_groups_y()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>

ALPAKA_FN_ACC auto cms::alpakatools::independent_groups_y ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1163 of file workdivision.h.

References writedatasetfile::args.

                                                                                 {
    return independent_groups_along<TAcc, alpaka::Dim<TAcc>::value - 2>(acc, static_cast<Idx>(args)...);
  }

independent_groups_z()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>

ALPAKA_FN_ACC auto cms::alpakatools::independent_groups_z ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 1170 of file workdivision.h.

References writedatasetfile::args.

                                                                                 {
    return independent_groups_along<TAcc, alpaka::Dim<TAcc>::value - 3>(acc, static_cast<Idx>(args)...);
  }

isPowerOf2()

template<typename T , typename = std::enable_if_t<std::is_integral_v<T>>>

constexpr bool cms::alpakatools::isPowerOf2

(

T

v

)

Definition at line 11 of file prefixScan.h.

References findQualityFiles::v.

Referenced by blockPrefixScan().

                                 {
    // returns true iif v has only one bit set.
    while (v) {
      if (v & 1)
        return !(v >> 1);
      else
        v >>= 1;
    }
    return false;
  }

make_device_buffer() [1/6]

template<typename T , typename TDev >

std::enable_if_t<alpaka::isDevice<TDev> and not std::is_array_v<T>, device_buffer<TDev, T> > cms::alpakatools::make_device_buffer

(

TDev const &

device

)

Definition at line 185 of file memory.h.

                          {
    return alpaka::allocBuf<T, Idx>(device, Scalar{});
  }

make_device_buffer() [2/6]

template<typename T , typename TDev >

std::enable_if_t<alpaka::isDevice<TDev> and cms::is_unbounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, device_buffer<TDev, T> > cms::alpakatools::make_device_buffer	(	TDev const &	device,
		Extent	extent
	)

Definition at line 194 of file memory.h.

                                                        {
    return alpaka::allocBuf<std::remove_extent_t<T>, Idx>(device, Vec1D{extent});
  }

make_device_buffer() [3/6]

template<typename T , typename TDev >

std::enable_if_t<alpaka::isDevice<TDev> and cms::is_bounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, device_buffer<TDev, T> > cms::alpakatools::make_device_buffer

(

TDev const &

device

)

Definition at line 202 of file memory.h.

                                         {
    return alpaka::allocBuf<std::remove_extent_t<T>, Idx>(device, Vec1D{std::extent_v<T>});
  }

make_device_buffer() [4/6]

template<typename T , typename TQueue >

std::enable_if_t<alpaka::isQueue<TQueue> and not std::is_array_v<T>, device_buffer<alpaka::Dev<TQueue>, T> > cms::alpakatools::make_device_buffer

(

TQueue const &

queue

)

Definition at line 210 of file memory.h.

References allocator_policy, Asynchronous, Caching, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), createBeamHaloJobs::queue, and Synchronous.

                                          {
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Caching) {
      return allocCachedBuf<T, Idx>(alpaka::getDev(queue), queue, Scalar{});
    }
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Asynchronous) {
      return alpaka::allocAsyncBuf<T, Idx>(queue, Scalar{});
    }
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Synchronous) {
      return alpaka::allocBuf<T, Idx>(alpaka::getDev(queue), Scalar{});
    }
  }

make_device_buffer() [5/6]

template<typename T , typename TQueue >

std::enable_if_t<alpaka::isQueue<TQueue> and cms::is_unbounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, device_buffer<alpaka::Dev<TQueue>, T> > cms::alpakatools::make_device_buffer	(	TQueue const &	queue,
		Extent	extent
	)

Definition at line 226 of file memory.h.

References allocator_policy, Asynchronous, Caching, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), createBeamHaloJobs::queue, and Synchronous.

                                                         {
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Caching) {
      return allocCachedBuf<std::remove_extent_t<T>, Idx>(alpaka::getDev(queue), queue, Vec1D{extent});
    }
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Asynchronous) {
      return alpaka::allocAsyncBuf<std::remove_extent_t<T>, Idx>(queue, Vec1D{extent});
    }
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Synchronous) {
      return alpaka::allocBuf<std::remove_extent_t<T>, Idx>(alpaka::getDev(queue), Vec1D{extent});
    }
  }

make_device_buffer() [6/6]

template<typename T , typename TQueue >

std::enable_if_t<alpaka::isQueue<TQueue> and cms::is_bounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, device_buffer<alpaka::Dev<TQueue>, T> > cms::alpakatools::make_device_buffer

(

TQueue const &

queue

)

Definition at line 242 of file memory.h.

References allocator_policy, Asynchronous, Caching, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), createBeamHaloJobs::queue, and Synchronous.

                                          {
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Caching) {
      return allocCachedBuf<std::remove_extent_t<T>, Idx>(alpaka::getDev(queue), queue, Vec1D{std::extent_v<T>});
    }
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Asynchronous) {
      return alpaka::allocAsyncBuf<std::remove_extent_t<T>, Idx>(queue, Vec1D{std::extent_v<T>});
    }
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Synchronous) {
      return alpaka::allocBuf<std::remove_extent_t<T>, Idx>(alpaka::getDev(queue), Vec1D{std::extent_v<T>});
    }
  }

make_device_view() [1/4]

template<typename T , typename TDev >

std::enable_if_t<not std::is_array_v<T>, device_view<TDev, T> > cms::alpakatools::make_device_view	(	TDev const &	device,
		T &	data
	)

Definition at line 260 of file memory.h.

References data.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::SiPixelRawToClusterKernel< pixelTopology::Phase2 >::makePhase1ClustersAsync(), ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::SiPixelRawToClusterKernel< pixelTopology::Phase2 >::makePhase2ClustersAsync(), and TrackingRecHitDevice< TrackerTraits, TDev >::updateFromDevice().

                                                                                                           {
    return alpaka::ViewPlainPtr<TDev, T, Dim0D, Idx>(&data, device, Scalar{});
  }

make_device_view() [2/4]

template<typename T , typename TDev >

device_view<TDev, T[]> cms::alpakatools::make_device_view	(	TDev const &	device,
		T *	data,
		Extent	extent
	)

Definition at line 265 of file memory.h.

References data.

                                                                                      {
    return alpaka::ViewPlainPtr<TDev, T, Dim1D, Idx>(data, device, Vec1D{extent});
  }

make_device_view() [3/4]

template<typename T , typename TDev >

std::enable_if_t<cms::is_unbounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, device_view<TDev, T> > cms::alpakatools::make_device_view	(	TDev const &	device,
		T &	data,
		Extent	extent
	)

Definition at line 271 of file memory.h.

References data.

                                                               {
    return alpaka::ViewPlainPtr<TDev, std::remove_extent_t<T>, Dim1D, Idx>(data, device, Vec1D{extent});
  }

make_device_view() [4/4]

template<typename T , typename TDev >

std::enable_if_t<cms::is_bounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, device_view<TDev, T> > cms::alpakatools::make_device_view	(	TDev const &	device,
		T &	data
	)

Definition at line 277 of file memory.h.

References data.

                                                {
    return alpaka::ViewPlainPtr<TDev, std::remove_extent_t<T>, Dim1D, Idx>(data, device, Vec1D{std::extent_v<T>});
  }

make_host_buffer() [1/9]

template<typename T >

std::enable_if_t<not std::is_array_v<T>, host_buffer<T> > cms::alpakatools::make_host_buffer

(

)

Definition at line 65 of file memory.h.

References host().

                                                                          {
    return alpaka::allocBuf<T, Idx>(host(), Scalar{});
  }

make_host_buffer() [2/9]

template<typename T >

std::enable_if_t<cms::is_unbounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_buffer<T> > cms::alpakatools::make_host_buffer

(

Extent

extent

)

Definition at line 71 of file memory.h.

References host().

                                  {
    return alpaka::allocBuf<std::remove_extent_t<T>, Idx>(host(), Vec1D{extent});
  }

make_host_buffer() [3/9]

template<typename T >

std::enable_if_t<cms::is_bounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_buffer<T> > cms::alpakatools::make_host_buffer

(

)

Definition at line 77 of file memory.h.

References host().

                     {
    return alpaka::allocBuf<std::remove_extent_t<T>, Idx>(host(), Vec1D{std::extent_v<T>});
  }

make_host_buffer() [4/9]

template<typename T , typename TPlatform >

std::enable_if_t<not std::is_array_v<T>, host_buffer<T> > cms::alpakatools::make_host_buffer

(

)

Definition at line 85 of file memory.h.

References host().

                                                                          {
    using Platform = TPlatform;
    return alpaka::allocMappedBuf<Platform, T, Idx>(host(), platform<Platform>(), Scalar{});
  }

make_host_buffer() [5/9]

template<typename T , typename TPlatform >

std::enable_if_t<cms::is_unbounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_buffer<T> > cms::alpakatools::make_host_buffer

(

Extent

extent

)

Definition at line 92 of file memory.h.

References host().

                                  {
    using Platform = TPlatform;
    return alpaka::allocMappedBuf<Platform, std::remove_extent_t<T>, Idx>(host(), platform<Platform>(), Vec1D{extent});
  }

make_host_buffer() [6/9]

template<typename T , typename TPlatform >

std::enable_if_t<cms::is_bounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_buffer<T> > cms::alpakatools::make_host_buffer

(

)

Definition at line 99 of file memory.h.

References host().

                     {
    using Platform = TPlatform;
    return alpaka::allocMappedBuf<Platform, std::remove_extent_t<T>, Idx>(
        host(), platform<Platform>(), Vec1D{std::extent_v<T>});
  }

make_host_buffer() [7/9]

template<typename T , typename TQueue >

std::enable_if_t<alpaka::isQueue<TQueue> and not std::is_array_v<T>, host_buffer<T> > cms::alpakatools::make_host_buffer

(

TQueue const &

queue

)

Definition at line 109 of file memory.h.

References allocator_policy, Caching, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), host(), and createBeamHaloJobs::queue.

                           {
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Caching) {
      return allocCachedBuf<T, Idx>(host(), queue, Scalar{});
    } else {
      using Platform = alpaka::Platform<alpaka::Dev<TQueue>>;
      return alpaka::allocMappedBuf<Platform, T, Idx>(host(), platform<Platform>(), Scalar{});
    }
  }

make_host_buffer() [8/9]

template<typename T , typename TQueue >

std::enable_if_t<alpaka::isQueue<TQueue> and cms::is_unbounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_buffer<T> > cms::alpakatools::make_host_buffer	(	TQueue const &	queue,
		Extent	extent
	)

Definition at line 123 of file memory.h.

References allocator_policy, Caching, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), host(), and createBeamHaloJobs::queue.

                                                       {
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Caching) {
      return allocCachedBuf<std::remove_extent_t<T>, Idx>(host(), queue, Vec1D{extent});
    } else {
      using Platform = alpaka::Platform<alpaka::Dev<TQueue>>;
      return alpaka::allocMappedBuf<Platform, std::remove_extent_t<T>, Idx>(
          host(), platform<Platform>(), Vec1D{extent});
    }
  }

make_host_buffer() [9/9]

template<typename T , typename TQueue >

std::enable_if_t<alpaka::isQueue<TQueue> and cms::is_bounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_buffer<T> > cms::alpakatools::make_host_buffer

(

TQueue const &

queue

)

Definition at line 137 of file memory.h.

References allocator_policy, Caching, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), host(), and createBeamHaloJobs::queue.

                                        {
    if constexpr (allocator_policy<alpaka::Dev<TQueue>> == AllocatorPolicy::Caching) {
      return allocCachedBuf<std::remove_extent_t<T>, Idx>(host(), queue, Vec1D{std::extent_v<T>});
    } else {
      using Platform = alpaka::Platform<alpaka::Dev<TQueue>>;
      return alpaka::allocMappedBuf<Platform, std::remove_extent_t<T>, Idx>(
          host(), platform<Platform>(), Vec1D{std::extent_v<T>});
    }
  }

make_host_view() [1/4]

template<typename T >

std::enable_if_t<not std::is_array_v<T>, host_view<T> > cms::alpakatools::make_host_view

(

T &

data

)

Definition at line 153 of file memory.h.

References data, and host().

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::SiPixelRawToClusterKernel< pixelTopology::Phase2 >::makePhase1ClustersAsync(), ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::SiPixelRawToClusterKernel< pixelTopology::Phase2 >::makePhase2ClustersAsync(), and TrackingRecHitDevice< TrackerTraits, TDev >::updateFromDevice().

                                                                             {
    return alpaka::ViewPlainPtr<DevHost, T, Dim0D, Idx>(&data, host(), Scalar{});
  }

make_host_view() [2/4]

template<typename T >

host_view<T[]> cms::alpakatools::make_host_view	(	T *	data,
		Extent	extent
	)

Definition at line 158 of file memory.h.

References data, and host().

                                                        {
    return alpaka::ViewPlainPtr<DevHost, T, Dim1D, Idx>(data, host(), Vec1D{extent});
  }

make_host_view() [3/4]

template<typename T >

std::enable_if_t<cms::is_unbounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_view<T> > cms::alpakatools::make_host_view	(	T &	data,
		Extent	extent
	)

Definition at line 164 of file memory.h.

References data, and host().

                                         {
    return alpaka::ViewPlainPtr<DevHost, std::remove_extent_t<T>, Dim1D, Idx>(data, host(), Vec1D{extent});
  }

make_host_view() [4/4]

template<typename T >

std::enable_if_t<cms::is_bounded_array_v<T> and not std::is_array_v<std::remove_extent_t<T> >, host_view<T> > cms::alpakatools::make_host_view

(

T &

data

)

Definition at line 170 of file memory.h.

References data, and host().

                          {
    return alpaka::ViewPlainPtr<DevHost, std::remove_extent_t<T>, Dim1D, Idx>(data, host(), Vec1D{std::extent_v<T>});
  }

make_SimpleVector() [1/2]

template<class T >

SimpleVector<T> cms::alpakatools::make_SimpleVector	(	int	capacity,
		T *	data
	)

Definition at line 126 of file SimpleVector.h.

References gpuVertexFinder::capacity(), data, and runTheMatrix::ret.

                                                           {
    SimpleVector<T> ret;
    ret.construct(capacity, data);
    return ret;
  }

make_SimpleVector() [2/2]

template<class T >

SimpleVector<T>* cms::alpakatools::make_SimpleVector	(	SimpleVector< T > *	mem,
		int	capacity,
		T *	data
	)

Definition at line 134 of file SimpleVector.h.

References gpuVertexFinder::capacity(), data, mem, and runTheMatrix::ret.

                                                                                  {
    auto ret = new (mem) SimpleVector<T>();
    ret->construct(capacity, data);
    return ret;
  }

make_workdiv() [1/2]

template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

WorkDiv<Dim1D> cms::alpakatools::make_workdiv ( Idx  blocks, Idx  elements  )

inline

Definition at line 46 of file workdivision.h.

References gather_cfg::blocks, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), and bookConverter::elements.

                                                               {
    if constexpr (not requires_single_thread_per_block_v<TAcc>) {
      // On GPU backends, each thread is looking at a single element:
      //   - the number of threads per block is "elements";
      //   - the number of elements per thread is always 1.
      return WorkDiv<Dim1D>(blocks, elements, Idx{1});
    } else {
      // On CPU backends, run serially with a single thread per block:
      //   - the number of threads per block is always 1;
      //   - the number of elements per thread is "elements".
      return WorkDiv<Dim1D>(blocks, Idx{1}, elements);
    }
  }

make_workdiv() [2/2]

template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>

WorkDiv<alpaka::Dim<TAcc> > cms::alpakatools::make_workdiv ( const Vec< alpaka::Dim< TAcc >> &  blocks, const Vec< alpaka::Dim< TAcc >> &  elements  )

inline

Definition at line 62 of file workdivision.h.

References gather_cfg::blocks, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), and bookConverter::elements.

                                                                                       {
    using Dim = alpaka::Dim<TAcc>;
    if constexpr (not requires_single_thread_per_block_v<TAcc>) {
      // On GPU backends, each thread is looking at a single element:
      //   - the number of threads per block is "elements";
      //   - the number of elements per thread is always 1.
      return WorkDiv<Dim>(blocks, elements, Vec<Dim>::ones());
    } else {
      // On CPU backends, run serially with a single thread per block:
      //   - the number of threads per block is always 1;
      //   - the number of elements per thread is "elements".
      return WorkDiv<Dim>(blocks, Vec<Dim>::ones(), elements);
    }
  }

module_backend_config()

void cms::alpakatools::module_backend_config

(

edm::ConfigurationDescriptions &

iDesc

)

Definition at line 13 of file module_backend_config.cc.

References edm::ParameterSetDescription::addUntracked(), edm::ConfigurationDescriptions::defaultDescription(), edm::ParameterSetDescription::isLabelUnused(), kComment, AlCaHLTBitMon_QueryRunRegistry::string, and findQualityFiles::v.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::ESProducer::prevalidate(), and ALPAKA_ACCELERATOR_NAMESPACE::ProducerBase< BaseT, Args >::prevalidate().

                                                                  {
    // the code below leads to 'alpaka = untracked.PSet(backend = untracked.string)' to be added to the generated cfi files
    // TODO: I don't know if this is a desired behavior for HLT
    edm::ParameterSetDescription descAlpaka;
    descAlpaka.addUntracked<std::string>("backend", "")
        ->setComment(
            "Alpaka backend for this module. Can be empty string (for the global default), 'serial_sync', or "
            " - depending on the architecture and available hardware - 'cuda_async', 'rocm_async'");

    if (iDesc.defaultDescription()) {
      if (iDesc.defaultDescription()->isLabelUnused(kPSetName)) {
        iDesc.defaultDescription()
            ->addUntracked<edm::ParameterSetDescription>(kPSetName, descAlpaka)
            ->setComment(kComment);
      }
    }
    for (auto& v : iDesc) {
      if (v.second.isLabelUnused(kPSetName)) {
        v.second.addUntracked<edm::ParameterSetDescription>(kPSetName, descAlpaka)->setComment(kComment);
      }
    }
  }

once_per_block()

template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>

ALPAKA_FN_ACC constexpr bool cms::alpakatools::once_per_block ( TAcc const &  acc )

inline

Definition at line 1332 of file workdivision.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::constexpr(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_exotic(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_multiSeedIterative(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_multiSeedParallel(), ALPAKA_ACCELERATOR_NAMESPACE::hcalFastCluster_singleSeed(), ALPAKA_ACCELERATOR_NAMESPACE::ecal::raw::Kernel_unpack::operator()(), pixelClustering::ClusterChargeCut< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksIterative::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksDBSCAN::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelRecHits::GetHits< TrackerTraits >::operator()(), cms::alpakatools::OneToManyAssocBase< I, ONES, SIZE >::zeroAndInit::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::FindClus< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TopoClusterContraction::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::FastCluster::operator()().

                                                                      {
    return alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc) == Vec<alpaka::Dim<TAcc>>::zeros();
  }

once_per_grid()

template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>

ALPAKA_FN_ACC constexpr bool cms::alpakatools::once_per_grid ( TAcc const &  acc )

inline

Definition at line 1320 of file workdivision.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernel::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::InitDoublets< TrackerTraits >::operator()(), calibPixel::CalibDigis< debug >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_BLFastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_FastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_checkOverflows< TrackerTraits >::operator()(), calibPixel::CalibDigisPhase2::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::CountModules< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoStructKernel::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernelUpdate::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernelUpdateMulti2::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::TestAlgoKernelUpdateMulti3::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::RawToDigi_kernel< debug >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_connect< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_find_ntuplets< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillNLayers< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::SeedingTopoThresh::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::PrepareTopoInputs::operator()().

                                                                     {
    return alpaka::getIdx<alpaka::Grid, alpaka::Threads>(acc) == Vec<alpaka::Dim<TAcc>>::zeros();
  }

platform()

template<typename TPlatform , typename = std::enable_if_t<alpaka::isPlatform<TPlatform>>>

TPlatform const& cms::alpakatools::platform ( )

inline

Definition at line 14 of file devices.h.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::AlpakaService::AlpakaService().

                                     {
    // initialise the platform the first time that this function is called
    static const auto platform = TPlatform{};
    return platform;
  }

radixSort()

template<typename TAcc , typename T , int NS = sizeof(T), typename std::enable_if< std::is_unsigned< T >::value &&!requires_single_thread_per_block_v< TAcc >, T >::type * = nullptr>

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::radixSort	(	const TAcc &	acc,
		T const *	a,
		uint16_t *	ind,
		uint16_t *	ind2,
		uint32_t	size
	)

Definition at line 324 of file radixSort.h.

References a.

                                                                                 {
    radixSortImpl<TAcc, T, NS>(acc, a, ind, ind2, size, dummyReorder<TAcc, T>);
  }

radixSortImpl()

template<typename TAcc , typename T , int NS, typename RF >

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::radixSortImpl	(	const TAcc &	acc,
		T const *__restrict__	a,
		uint16_t *	ind,
		uint16_t *	ind2,
		uint32_t	size,
		RF	reorder
	)

Definition at line 97 of file radixSort.h.

References a, ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::ALPAKA_ASSERT_ACC(), cms::cuda::assert(), cms::cudacompat::atomicAdd(), cms::cudacompat::atomicMax(), newFWLiteAna::bin, HltBtagPostValidation_cff::c, ALPAKA_ACCELERATOR_NAMESPACE::brokenline::constexpr(), mps_fire::i, heavyIonCSV_trainingSettings::idx, independent_group_elements(), createfilelist::int, dqmiolumiharvest::j, dqmdumpme::k, hltrates_dqm_sourceclient-live_cfg::offset, trackingPlots::reorder, contentValuesCheck::ss, submitPVValidationJobs::t, ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::threadIdxLocal(), dqmMemoryStats::total, and x.

                                                                                                          {
    if constexpr (!requires_single_thread_per_block_v<TAcc>) {
      const auto warpSize = alpaka::warp::getSize(acc);
      const uint32_t threadIdxLocal(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
      [[maybe_unused]] const uint32_t blockDimension(alpaka::getWorkDiv<alpaka::Block, alpaka::Elems>(acc)[0u]);
      // we expect a power of 2 here
      assert(warpSize && (0 == (warpSize & (warpSize - 1))));
      const std::size_t warpMask = warpSize - 1;

      // Define the bin size (d=8 => 1 byte bin).
      constexpr int binBits = 8, dataBits = 8 * sizeof(T), totalSortingPassses = dataBits / binBits;
      // Make sure the slices are data aligned
      static_assert(0 == dataBits % binBits);
      // Make sure the NS parameter makes sense
      static_assert(NS > 0 && NS <= sizeof(T));
      constexpr int binsNumber = 1 << binBits;
      constexpr int binsMask = binsNumber - 1;
      // Prefix scan iterations. NS is counted in full bytes and not slices.
      constexpr int initialSortingPass = int(sizeof(T)) - NS;

      // Count/index for the prefix scan
      // TODO: rename
      auto& c = alpaka::declareSharedVar<int32_t[binsNumber], __COUNTER__>(acc);
      // Temporary storage for prefix scan. Only really needed for first-of-warp keeping
      // Then used for thread to bin mapping TODO: change type to byte and remap to
      auto& ct = alpaka::declareSharedVar<int32_t[binsNumber], __COUNTER__>(acc);
      // Bin to thread index mapping (used to store the highest thread index within a bin number
      // batch of threads.
      // TODO: currently initialized to an invalid value, but could also be initialized to the
      // lowest possible value (change to bytes?)
      auto& cu = alpaka::declareSharedVar<int32_t[binsNumber], __COUNTER__>(acc);
      // TODO we could also have an explicit caching of the current index for each thread.

      // TODO: do those have to be shared?
      auto& ibs = alpaka::declareSharedVar<int, __COUNTER__>(acc);
      auto& currentSortingPass = alpaka::declareSharedVar<int, __COUNTER__>(acc);

      ALPAKA_ASSERT_ACC(size > 0);
      // TODO: is this a hard requirement?
      ALPAKA_ASSERT_ACC(blockDimension >= binsNumber);

      currentSortingPass = initialSortingPass;

      auto j = ind;
      auto k = ind2;

      // Initializer index order to trivial increment.
      for (auto idx : independent_group_elements(acc, size)) {
        j[idx] = idx;
      }
      alpaka::syncBlockThreads(acc);

      // Iterate on the slices of the data.
      while (alpaka::syncBlockThreadsPredicate<alpaka::BlockAnd>(acc, (currentSortingPass < totalSortingPassses))) {
        for (auto idx : independent_group_elements(acc, binsNumber)) {
          c[idx] = 0;
        }
        alpaka::syncBlockThreads(acc);
        const auto sortingPassShift = binBits * currentSortingPass;

        // fill bins (count elements in each bin)
        for (auto idx : independent_group_elements(acc, size)) {
          auto bin = (a[j[idx]] >> sortingPassShift) & binsMask;
          alpaka::atomicAdd(acc, &c[bin], 1, alpaka::hierarchy::Threads{});
        }
        alpaka::syncBlockThreads(acc);

        if (!threadIdxLocal && 1 == alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0]) {
          //          printf("Pass=%d, Block=%d, ", currentSortingPass - 1, alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0]);
          size_t total = 0;
          for (int i = 0; i < (int)binsNumber; i++) {
            //            printf("count[%d]=%d ", i, c[i] );
            total += c[i];
          }
          //          printf("total=%zu\n", total);
          assert(total == size);
        }
        // prefix scan "optimized"???...
        // TODO: we might be able to reuse the warpPrefixScan function
        // Warp level prefix scan
        for (auto idx : independent_group_elements(acc, binsNumber)) {
          auto x = c[idx];
          auto laneId = idx & warpMask;

          for (int offset = 1; offset < warpSize; offset <<= 1) {
            auto y = alpaka::warp::shfl(acc, x, laneId - offset);
            if (laneId >= (uint32_t)offset)
              x += y;
          }
          ct[idx] = x;
        }
        alpaka::syncBlockThreads(acc);

        // Block level completion of prefix scan (add last sum of each preceding warp)
        for (auto idx : independent_group_elements(acc, binsNumber)) {
          auto ss = (idx / warpSize) * warpSize - 1;
          c[idx] = ct[idx];
          for (int i = ss; i > 0; i -= warpSize)
            c[idx] += ct[i];
        }
        // Post prefix scan, c[bin] contains the offsets in index counts to the last index +1 for each bin

        /*
        //prefix scan for the nulls  (for documentation)
        if (threadIdxLocal==0)
          for (int i = 1; i < sb; ++i) c[i] += c[i-1];
        */

        // broadcast: we will fill the new index array downward, from offset c[bin], with one thread per
        // bin, working on one set of bin size elements at a time.
        // This will reorder the indices by the currently considered slice, otherwise preserving the previous order.
        ibs = size - 1;
        alpaka::syncBlockThreads(acc);

        // Iterate on bin-sized slices to (size - 1) / binSize + 1 iterations
        while (alpaka::syncBlockThreadsPredicate<alpaka::BlockAnd>(acc, ibs >= 0)) {
          // Init
          for (auto idx : independent_group_elements(acc, binsNumber)) {
            cu[idx] = -1;
            ct[idx] = -1;
          }
          alpaka::syncBlockThreads(acc);

          // Find the highest index for all the threads dealing with a given bin (in cu[])
          // Also record the bin for each thread (in ct[])
          for (auto idx : independent_group_elements(acc, binsNumber)) {
            int i = ibs - idx;
            int32_t bin = -1;
            if (i >= 0) {
              bin = (a[j[i]] >> sortingPassShift) & binsMask;
              ct[idx] = bin;
              alpaka::atomicMax(acc, &cu[bin], int(i), alpaka::hierarchy::Threads{});
            }
          }
          alpaka::syncBlockThreads(acc);

          // FIXME: we can slash a memory access.
          for (auto idx : independent_group_elements(acc, binsNumber)) {
            int i = ibs - idx;
            // Are we still in inside the data?
            if (i >= 0) {
              int32_t bin = ct[idx];
              // Are we the thread with the highest index (from previous pass)?
              if (cu[bin] == i) {
                // With the highest index, we are actually the lowest thread number. We will
                // work "on behalf of" the higher thread numbers (including ourselves)
                // No way around scanning and testing for bin in ct[otherThread] number to find the other threads
                for (int peerThreadIdx = idx; peerThreadIdx < binsNumber; peerThreadIdx++) {
                  if (ct[peerThreadIdx] == bin) {
                    k[--c[bin]] = j[ibs - peerThreadIdx];
                  }
                }
              }
            }
            /*
            int32_t bin = (i >= 0 ? ((a[j[i]] >> sortingPassShift) & binsMask) : -1);
            if (i >= 0 && i == cu[bin])  // ensure to keep them in order: only one thread per bin is active, rest is idle.
              // 
              for (int ii = idx; ii < sb; ++ii)
                if (ct[ii] == bin) {
                  auto oi = ii - idx;
                  // assert(i>=oi);if(i>=oi)
                  k[--c[bin]] = j[i - oi]; // i = ibs - idx, oi = ii - idx => i - oi = ibs - ii;
                }
            */
          }
          alpaka::syncBlockThreads(acc);

          if (threadIdxLocal == 0) {
            ibs -= binsNumber;
            // https://github.com/cms-patatrack/pixeltrack-standalone/pull/210
            // TODO: is this really needed?
            alpaka::mem_fence(acc, alpaka::memory_scope::Grid{});
          }
          alpaka::syncBlockThreads(acc);
        }

        /*
        // broadcast for the nulls  (for documentation)
        if (threadIdxLocal==0)
        for (int i=size-first-1; i>=0; i--) { // =blockDim.x) {
          auto bin = (a[j[i]] >> d*p)&(sb-1);
          auto ik = atomicSub(&c[bin],1);
          k[ik-1] = j[i];
        }
        */

        alpaka::syncBlockThreads(acc);
        ALPAKA_ASSERT_ACC(c[0] == 0);

        // swap (local, ok)
        auto t = j;
        j = k;
        k = t;

        const uint32_t threadIdxLocal(alpaka::getIdx<alpaka::Block, alpaka::Threads>(acc)[0u]);
        if (threadIdxLocal == 0)
          ++currentSortingPass;
        alpaka::syncBlockThreads(acc);
      }

      if ((dataBits != 8) && (0 == (NS & 1)))
        ALPAKA_ASSERT_ACC(j ==
                          ind);  // dataBits/binBits is even so ind is correct (the result is in the right location)

      // TODO this copy is (doubly?) redundant with the reorder
      if (j != ind)  // odd number of sorting passes, we need to move the result to the right array (ind[])
        for (auto idx : independent_group_elements(acc, size)) {
          ind[idx] = ind2[idx];
        };

      alpaka::syncBlockThreads(acc);

      // now move negative first... (if signed)
      // TODO: the ind2 => ind copy should have beed deferred. We should pass (j != ind) as an extra parameter
      reorder(acc, a, ind, ind2, size);
    } else {
      //static_assert(false);
    }
  }

radixSortMulti()

template<typename TAcc , typename T , int NS = sizeof(T)>

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::radixSortMulti	(	const TAcc &	acc,
		T const *	v,
		uint16_t *	index,
		uint32_t const *	offsets,
		uint16_t *	workspace
	)

Definition at line 382 of file radixSort.h.

References a, cms::cuda::assert(), cms::cudacompat::blockIdx, unpackBuffers-CaloStage1::offsets, findQualityFiles::v, and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ws.

                                                                                                  {
    // TODO: check
    // Sort multiple blocks of data in v[] separated by in chunks located at offsets[]
    // extern __shared__ uint16_t ws[];
    uint16_t* ws = alpaka::getDynSharedMem<uint16_t>(acc);

    const uint32_t blockIdx(alpaka::getIdx<alpaka::Grid, alpaka::Blocks>(acc)[0u]);
    auto a = v + offsets[blockIdx];
    auto ind = index + offsets[blockIdx];
    auto ind2 = nullptr == workspace ? ws : workspace + offsets[blockIdx];
    auto size = offsets[blockIdx + 1] - offsets[blockIdx];
    assert(offsets[blockIdx + 1] >= offsets[blockIdx]);
    if (size > 0)
      radixSort<TAcc, T, NS>(acc, a, ind, ind2, size);
  }

reorderFloat()

template<typename TAcc , typename T >

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::reorderFloat	(	const TAcc &	acc,
		T const *	a,
		uint16_t *	ind,
		uint16_t *	ind2,
		uint32_t	size
	)

Definition at line 54 of file radixSort.h.

References a, heavyIonCSV_trainingSettings::idx, and independent_group_elements().

                                                                                 {
    //move negative first...

    auto& firstNeg = alpaka::declareSharedVar<uint32_t, __COUNTER__>(acc);
    firstNeg = a[ind[0]] < 0 ? 0 : size;
    alpaka::syncBlockThreads(acc);

    // find first negative
    for (auto idx : independent_group_elements(acc, size - 1)) {
      if ((a[ind[idx]] ^ a[ind[idx + 1]]) < 0)
        firstNeg = idx + 1;
    }
    alpaka::syncBlockThreads(acc);

    for (auto idx : independent_group_elements(acc, firstNeg, size)) {
      ind2[size - idx - 1] = ind[idx];
    }
    alpaka::syncBlockThreads(acc);

    for (auto idx : independent_group_elements(acc, firstNeg)) {
      ind2[idx + size - firstNeg] = ind[idx];
    }
    alpaka::syncBlockThreads(acc);

    for (auto idx : independent_group_elements(acc, size)) {
      ind[idx] = ind2[idx];
    }
  }

reorderSigned()

template<typename TAcc , typename T >

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::reorderSigned	(	const TAcc &	acc,
		T const *	a,
		uint16_t *	ind,
		uint16_t *	ind2,
		uint32_t	size
	)

Definition at line 21 of file radixSort.h.

References a, heavyIonCSV_trainingSettings::idx, and independent_group_elements().

                                                                                 {
    //move negative first...

    auto& firstNeg = alpaka::declareSharedVar<uint32_t, __COUNTER__>(acc);
    firstNeg = a[ind[0]] < 0 ? 0 : size;
    alpaka::syncBlockThreads(acc);

    // find first negative
    for (auto idx : independent_group_elements(acc, size - 1)) {
      if ((a[ind[idx]] ^ a[ind[idx + 1]]) < 0) {
        firstNeg = idx + 1;
      }
    }

    alpaka::syncBlockThreads(acc);

    for (auto idx : independent_group_elements(acc, firstNeg, size)) {
      ind2[idx - firstNeg] = ind[idx];
    }
    alpaka::syncBlockThreads(acc);

    for (auto idx : independent_group_elements(acc, firstNeg)) {
      ind2[idx + size - firstNeg] = ind[idx];
    }
    alpaka::syncBlockThreads(acc);

    for (auto idx : independent_group_elements(acc, size)) {
      ind[idx] = ind2[idx];
    }
  }

round_up_by()

constexpr Idx cms::alpakatools::round_up_by ( Idx  value, Idx  divisor  )

inline

Definition at line 16 of file workdivision.h.

{ return (value + divisor - 1) / divisor * divisor; }

toBackend()

Backend cms::alpakatools::toBackend

(

std::string_view

name

)

Definition at line 13 of file Backend.cc.

References cms::Exception::addContext(), HLT_2024v10_cff::distance, spr::find(), newFWLiteAna::found, and Skims_PA_cff::name.

                                         {
    auto found = std::find(backendNames.begin(), backendNames.end(), name);
    if (found == backendNames.end()) {
      cms::Exception ex("EnumNotFound");
      ex << "Invalid backend name '" << name << "'";
      ex.addContext("Calling cms::alpakatools::toBackend()");
      throw ex;
    }
    return static_cast<Backend>(std::distance(backendNames.begin(), found));
  }

toString()

std::string_view cms::alpakatools::toString

(

Backend

backend

)

Definition at line 24 of file Backend.cc.

References cms::Exception::addContext(), HLT_2024v10_cff::backend, size, and heppy_batch::val.

Referenced by TestAlpakaObjectAnalyzer::analyze(), and TestAlpakaAnalyzer::analyze().

                                           {
    auto val = static_cast<unsigned short>(backend);
    if (val >= static_cast<unsigned short>(Backend::size)) {
      cms::Exception ex("InvalidEnumValue");
      ex << "Invalid backend enum value " << val;
      ex.addContext("Calling cms::alpakatools::toString()");
      throw ex;
    }
    return backendNames[val];
  }

uniform_elements()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_elements ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 303 of file workdivision.h.

References writedatasetfile::args.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::for(), cms::alpakatools::countFromVector::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::InitDoublets< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksIterative::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::ClusterTracksDBSCAN::operator()(), calibPixel::CalibDigis< debug >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_BLFastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_FastFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::LoadTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::setHitsLayerStart< TrackerTraits >::operator()(), cms::alpakatools::fillFromVector::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_checkOverflows< TrackerTraits >::operator()(), calibPixel::CalibDigisPhase2::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelClustering::CountModules< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_CircleFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_LineFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fishboneCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::Kernel_BLFit< N, TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_earlyDuplicateRemover< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fastDuplicateRemover< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::pixelDetails::RawToDigi_kernel< debug >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_find_ntuplets< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_mark_used< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countMultiplicity< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillMultiplicity< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_classifyTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillHitDetIndices< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_fillNLayers< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_doStatsForHitInTracks< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_countSharedHit< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_markSharedHit< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_rejectDuplicate< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_sharedHitCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_tripletCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_simpleTripletCleaner< TrackerTraits >::operator()(), ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_print_found_ntuplets< TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::vertexFinder::sortByPt2().

                                                                             {
    return uniform_elements_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

uniform_elements_x()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_elements_x ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 315 of file workdivision.h.

References writedatasetfile::args.

                                                                               {
    return uniform_elements_along<TAcc, alpaka::Dim<TAcc>::value - 1>(acc, static_cast<Idx>(args)...);
  }

uniform_elements_y()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_elements_y ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 322 of file workdivision.h.

References writedatasetfile::args.

Referenced by ALPAKA_ACCELERATOR_NAMESPACE::caPixelDoublets::CAFishbone< TrackerTraits >::operator()(), and ALPAKA_ACCELERATOR_NAMESPACE::caHitNtupletGeneratorKernels::Kernel_connect< TrackerTraits >::operator()().

                                                                               {
    return uniform_elements_along<TAcc, alpaka::Dim<TAcc>::value - 2>(acc, static_cast<Idx>(args)...);
  }

uniform_elements_z()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_elements_z ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 329 of file workdivision.h.

References writedatasetfile::args.

                                                                               {
    return uniform_elements_along<TAcc, alpaka::Dim<TAcc>::value - 3>(acc, static_cast<Idx>(args)...);
  }

uniform_group_elements()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_group_elements ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 956 of file workdivision.h.

References writedatasetfile::args.

                                                                                   {
    return uniform_group_elements_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

uniform_group_elements_x()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_group_elements_x ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 969 of file workdivision.h.

References writedatasetfile::args.

                                                                                     {
    return uniform_group_elements_along<TAcc, alpaka::Dim<TAcc>::value - 1>(acc, static_cast<Idx>(args)...);
  }

uniform_group_elements_y()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_group_elements_y ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 976 of file workdivision.h.

References writedatasetfile::args.

                                                                                     {
    return uniform_group_elements_along<TAcc, alpaka::Dim<TAcc>::value - 2>(acc, static_cast<Idx>(args)...);
  }

uniform_group_elements_z()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_group_elements_z ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 983 of file workdivision.h.

References writedatasetfile::args.

                                                                                     {
    return uniform_group_elements_along<TAcc, alpaka::Dim<TAcc>::value - 3>(acc, static_cast<Idx>(args)...);
  }

uniform_groups()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and alpaka::Dim<TAcc>::value == 1>>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_groups ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 742 of file workdivision.h.

References writedatasetfile::args.

                                                                           {
    return uniform_groups_along<TAcc, 0>(acc, static_cast<Idx>(args)...);
  }

uniform_groups_x()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 0>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_groups_x ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 754 of file workdivision.h.

References writedatasetfile::args.

                                                                             {
    return uniform_groups_along<TAcc, alpaka::Dim<TAcc>::value - 1>(acc, static_cast<Idx>(args)...);
  }

uniform_groups_y()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 1>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_groups_y ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 761 of file workdivision.h.

References writedatasetfile::args.

                                                                             {
    return uniform_groups_along<TAcc, alpaka::Dim<TAcc>::value - 2>(acc, static_cast<Idx>(args)...);
  }

uniform_groups_z()

template<typename TAcc , typename... TArgs, typename = std::enable_if_t<alpaka::isAccelerator<TAcc> and (alpaka::Dim<TAcc>::value > 2>

ALPAKA_FN_ACC auto cms::alpakatools::uniform_groups_z ( TAcc const &  acc, TArgs...  args  )

inline

Definition at line 768 of file workdivision.h.

References writedatasetfile::args.

                                                                             {
    return uniform_groups_along<TAcc, alpaka::Dim<TAcc>::value - 3>(acc, static_cast<Idx>(args)...);
  }

warpPrefixScan() [1/2]

template<typename TAcc , typename T , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::warpPrefixScan	(	const TAcc &	acc,
		int32_t	laneId,
		T const *	ci,
		T *	co,
		uint32_t	i,
		bool	active = true
	)

Definition at line 23 of file prefixScan.h.

References CMS_UNROLL_LOOP, cms::cuda::co, DTskim_cfg::dataType, mps_fire::i, hltrates_dqm_sourceclient-live_cfg::offset, and x.

Referenced by blockPrefixScan(), and warpPrefixScan().

                                                                                           {
    // ci and co may be the same
    T x = active ? ci[i] : 0;
    CMS_UNROLL_LOOP
    for (int32_t offset = 1; offset < alpaka::warp::getSize(acc); offset <<= 1) {
      // Force the exact type for integer types otherwise the compiler will find the template resolution ambiguous.
      using dataType = std::conditional_t<std::is_floating_point_v<T>, T, std::int32_t>;
      T y = alpaka::warp::shfl(acc, static_cast<dataType>(x), laneId - offset);
      if (laneId >= offset)
        x += y;
    }
    if (active)
      co[i] = x;
  }

warpPrefixScan() [2/2]

template<typename TAcc , typename T , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>

ALPAKA_FN_ACC ALPAKA_FN_INLINE void cms::alpakatools::warpPrefixScan	(	const TAcc &	acc,
		int32_t	laneId,
		T *	c,
		uint32_t	i,
		bool	active = true
	)

Definition at line 40 of file prefixScan.h.

References HltBtagPostValidation_cff::c, mps_fire::i, and warpPrefixScan().

                                                                             {
    warpPrefixScan(acc, laneId, c, c, i, active);
  }

Variable Documentation

allocator_policy

template<typename TDev , typename = std::enable_if_t<alpaka::isDevice<TDev>>>

constexpr AllocatorPolicy cms::alpakatools::allocator_policy = AllocatorPolicy::Synchronous

inline

Definition at line 17 of file AllocatorPolicy.h.

Referenced by make_device_buffer(), and make_host_buffer().

requires_single_thread_per_block_v

template<typename TAcc , typename = std::enable_if_t<alpaka::isAccelerator<TAcc>>>

constexpr bool cms::alpakatools::requires_single_thread_per_block_v = requires_single_thread_per_block<TAcc>::value

inline

Definition at line 42 of file workdivision.h.