#include <TrackingRecHit2DHeterogeneous.h>

Public Types
using	PhiBinner = TrackingRecHit2DSOAView::PhiBinner

template<typename T >
using	unique_ptr = typename Traits::template unique_ptr< T >

Public Member Functions
void	copyFromGPU (TrackingRecHit2DHeterogeneous< cms::cudacompat::GPUTraits > const *input, cudaStream_t stream)

auto	hitsLayerStart ()

auto	hitsModuleStart () const

cms::cuda::host::unique_ptr < uint32_t[]>	hitsModuleStartToHostAsync (cudaStream_t stream) const

auto	iphi ()

cms::cuda::host::unique_ptr < float[]>	localCoordToHostAsync (cudaStream_t stream) const

auto	nHits () const

auto	nMaxModules () const

auto	offsetBPIX2 () const

TrackingRecHit2DHeterogeneous &	operator= (const TrackingRecHit2DHeterogeneous &)=delete

TrackingRecHit2DHeterogeneous &	operator= (TrackingRecHit2DHeterogeneous &&)=default

auto	phiBinner ()

auto	phiBinnerStorage ()

	TrackingRecHit2DHeterogeneous ()=default

	TrackingRecHit2DHeterogeneous (uint32_t nHits, bool isPhase2, int32_t offsetBPIX2, pixelCPEforGPU::ParamsOnGPU const cpeParams, uint32_t const hitsModuleStart, cudaStream_t stream, TrackingRecHit2DHeterogeneous< cms::cudacompat::GPUTraits > const *input=nullptr)

	TrackingRecHit2DHeterogeneous (const TrackingRecHit2DHeterogeneous &)=delete

	TrackingRecHit2DHeterogeneous (TrackingRecHit2DHeterogeneous &&)=default

TrackingRecHit2DSOAView *	view ()

TrackingRecHit2DSOAView const *	view () const

	~TrackingRecHit2DHeterogeneous ()=default

Private Attributes
unique_ptr < TrackingRecHit2DSOAView::AverageGeometry >	m_AverageGeometryStore

uint32_t *	m_hitsLayerStart

uint32_t const *	m_hitsModuleStart

int16_t *	m_iphi

uint32_t	m_nHits

uint32_t	m_nMaxModules

int32_t	m_offsetBPIX2

PhiBinner *	m_phiBinner

PhiBinner::index_type *	m_phiBinnerStorage

unique_ptr < TrackingRecHit2DSOAView::PhiBinner >	m_PhiBinnerStore

unique_ptr< uint16_t[]>	m_store16

unique_ptr< float[]>	m_store32

unique_ptr < TrackingRecHit2DSOAView >	m_view

Static Private Attributes
static constexpr uint32_t	n16 = 4

static constexpr uint32_t	n32 = 10

Detailed Description

template<typename Traits>
class TrackingRecHit2DHeterogeneous< Traits >

Definition at line 9 of file TrackingRecHit2DHeterogeneous.h.

Member Typedef Documentation

template<typename Traits>

using TrackingRecHit2DHeterogeneous< Traits >::PhiBinner = TrackingRecHit2DSOAView::PhiBinner

Definition at line 14 of file TrackingRecHit2DHeterogeneous.h.

template<typename Traits>

template<typename T >

using TrackingRecHit2DHeterogeneous< Traits >::unique_ptr = typename Traits::template unique_ptr<T>

Definition at line 12 of file TrackingRecHit2DHeterogeneous.h.

Constructor & Destructor Documentation

template<typename Traits>

TrackingRecHit2DHeterogeneous< Traits >::TrackingRecHit2DHeterogeneous ( )

default

template<typename Traits >

TrackingRecHit2DHeterogeneous< Traits >::TrackingRecHit2DHeterogeneous	(	uint32_t	nHits,
		bool	isPhase2,
		int32_t	offsetBPIX2,
		pixelCPEforGPU::ParamsOnGPU const *	cpeParams,
		uint32_t const *	hitsModuleStart,
		cudaStream_t	stream,
		TrackingRecHit2DHeterogeneous< cms::cudacompat::GPUTraits > const *	input = `nullptr`
	)

explicit

Definition at line 88 of file TrackingRecHit2DHeterogeneous.h.

     : m_nHits(nHits), m_offsetBPIX2(offsetBPIX2), m_hitsModuleStart(hitsModuleStart) {
   auto view = Traits::template make_host_unique<TrackingRecHit2DSOAView>(stream);
 
   m_nMaxModules = isPhase2 ? phase2PixelTopology::numberOfModules : phase1PixelTopology::numberOfModules;
 
   view->m_nHits = nHits;
   view->m_nMaxModules = m_nMaxModules;
   m_view = Traits::template make_unique<TrackingRecHit2DSOAView>(stream);  // leave it on host and pass it by value?
   m_AverageGeometryStore = Traits::template make_unique<TrackingRecHit2DSOAView::AverageGeometry>(stream);
   view->m_averageGeometry = m_AverageGeometryStore.get();
   view->m_cpeParams = cpeParams;
   view->m_hitsModuleStart = hitsModuleStart;
 
   // if empy do not bother
   if (0 == nHits) {
     if constexpr (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
       cms::cuda::copyAsync(m_view, view, stream);
     } else {
       m_view.reset(view.release());  // NOLINT: std::move() breaks CUDA version
     }
     return;
   }
 
   // the single arrays are not 128 bit alligned...
   // the hits are actually accessed in order only in building
   // if ordering is relevant they may have to be stored phi-ordered by layer or so
   // this will break 1to1 correspondence with cluster and module locality
   // so unless proven VERY inefficient we keep it ordered as generated
 
   // host copy is "reduced"  (to be reviewed at some point)
   if constexpr (std::is_same<Traits, cms::cudacompat::HostTraits>::value) {
     // it has to compile for ALL cases
     copyFromGPU(input, stream);
   } else {
     assert(input == nullptr);
 
     auto nL = isPhase2 ? phase2PixelTopology::numberOfLayers : phase1PixelTopology::numberOfLayers;
 
     m_store16 = Traits::template make_unique<uint16_t[]>(nHits * n16, stream);
     m_store32 = Traits::template make_unique<float[]>(nHits * n32 + nL + 1, stream);
     m_PhiBinnerStore = Traits::template make_unique<TrackingRecHit2DSOAView::PhiBinner>(stream);
   }
 
   static_assert(sizeof(TrackingRecHit2DSOAView::hindex_type) == sizeof(float));
   static_assert(sizeof(TrackingRecHit2DSOAView::hindex_type) == sizeof(TrackingRecHit2DSOAView::PhiBinner::index_type));
 
   auto get32 = [&](int i) { return m_store32.get() + i * nHits; };
 
   // copy all the pointers
   m_phiBinner = view->m_phiBinner = m_PhiBinnerStore.get();
   m_phiBinnerStorage = view->m_phiBinnerStorage =
       reinterpret_cast<TrackingRecHit2DSOAView::PhiBinner::index_type*>(get32(9));
 
   view->m_xl = get32(0);
   view->m_yl = get32(1);
   view->m_xerr = get32(2);
   view->m_yerr = get32(3);
   view->m_chargeAndStatus = reinterpret_cast<uint32_t*>(get32(4));
 
   if constexpr (!std::is_same<Traits, cms::cudacompat::HostTraits>::value) {
     assert(input == nullptr);
     view->m_xg = get32(5);
     view->m_yg = get32(6);
     view->m_zg = get32(7);
     view->m_rg = get32(8);
 
     auto get16 = [&](int i) { return m_store16.get() + i * nHits; };
     m_iphi = view->m_iphi = reinterpret_cast<int16_t*>(get16(1));
 
     view->m_xsize = reinterpret_cast<int16_t*>(get16(2));
     view->m_ysize = reinterpret_cast<int16_t*>(get16(3));
     view->m_detInd = get16(0);
 
     m_phiBinner = view->m_phiBinner = m_PhiBinnerStore.get();
     m_hitsLayerStart = view->m_hitsLayerStart = reinterpret_cast<uint32_t*>(get32(n32));
   }
 
   // transfer view
   if constexpr (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
     cms::cuda::copyAsync(m_view, view, stream);
   } else {
     m_view.reset(view.release());  // NOLINT: std::move() breaks CUDA version
   }
 }

template<typename Traits>

TrackingRecHit2DHeterogeneous< Traits >::~TrackingRecHit2DHeterogeneous ( )

default

template<typename Traits>

TrackingRecHit2DHeterogeneous< Traits >::TrackingRecHit2DHeterogeneous ( const TrackingRecHit2DHeterogeneous< Traits > & )

delete

template<typename Traits>

TrackingRecHit2DHeterogeneous< Traits >::TrackingRecHit2DHeterogeneous ( TrackingRecHit2DHeterogeneous< Traits > && )

default

Member Function Documentation

template<typename Traits>

void TrackingRecHit2DHeterogeneous< Traits >::copyFromGPU	(	TrackingRecHit2DHeterogeneous< cms::cudacompat::GPUTraits > const *	input,
		cudaStream_t	stream
	)

Referenced by TrackingRecHit2DHeterogeneous< Traits >::TrackingRecHit2DHeterogeneous().

template<typename Traits>

auto TrackingRecHit2DHeterogeneous< Traits >::hitsLayerStart ( )

inline

Definition at line 42 of file TrackingRecHit2DHeterogeneous.h.

References TrackingRecHit2DHeterogeneous< Traits >::m_hitsLayerStart.

42 { return m_hitsLayerStart; }

TrackingRecHit2DHeterogeneous::m_hitsLayerStart

uint32_t * m_hitsLayerStart

Definition: TrackingRecHit2DHeterogeneous.h:76

template<typename Traits>

auto TrackingRecHit2DHeterogeneous< Traits >::hitsModuleStart ( ) const

inline

Definition at line 41 of file TrackingRecHit2DHeterogeneous.h.

References TrackingRecHit2DHeterogeneous< Traits >::m_hitsModuleStart.

Referenced by TrackingRecHit2DHeterogeneous< Traits >::TrackingRecHit2DHeterogeneous().

41 { return m_hitsModuleStart; }

TrackingRecHit2DHeterogeneous::m_hitsModuleStart

uint32_t const * m_hitsModuleStart

Definition: TrackingRecHit2DHeterogeneous.h:70

cms::cuda::host::unique_ptr< uint32_t[]> TrackingRecHit2DGPU::hitsModuleStartToHostAsync ( cudaStream_t stream ) const

Definition at line 15 of file TrackingRecHit2DHeterogeneous.cc.

References cudaCheck, TrackingRecHit2DHeterogeneous< Traits >::m_hitsModuleStart, TrackingRecHit2DHeterogeneous< Traits >::nMaxModules(), runTheMatrix::ret, and cms::cuda::stream.

                                                                                                            {
   auto ret = cms::cuda::make_host_unique<uint32_t[]>(nMaxModules() + 1, stream);
   cudaCheck(
       cudaMemcpyAsync(ret.get(), m_hitsModuleStart, sizeof(uint32_t) * (nMaxModules() + 1), cudaMemcpyDefault, stream));
   return ret;
 }

template<typename Traits>

auto TrackingRecHit2DHeterogeneous< Traits >::iphi ( )

inline

Definition at line 45 of file TrackingRecHit2DHeterogeneous.h.

References TrackingRecHit2DHeterogeneous< Traits >::m_iphi.

45 { return m_iphi; }

TrackingRecHit2DHeterogeneous::m_iphi

int16_t * m_iphi

Definition: TrackingRecHit2DHeterogeneous.h:77

cms::cuda::host::unique_ptr< float[]> TrackingRecHit2DGPU::localCoordToHostAsync ( cudaStream_t stream ) const

Definition at line 8 of file TrackingRecHit2DHeterogeneous.cc.

References cms::cuda::copyAsync(), TrackingRecHit2DHeterogeneous< Traits >::m_store32, TrackingRecHit2DHeterogeneous< Traits >::nHits(), runTheMatrix::ret, and cms::cuda::stream.

                                                                                                    {
   auto ret = cms::cuda::make_host_unique<float[]>(5 * nHits(), stream);
   cms::cuda::copyAsync(ret, m_store32, 5 * nHits(), stream);
   return ret;
 }