#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

TrackingRecHit2DHeterogeneous &	operator= (const TrackingRecHit2DHeterogeneous &)=delete

TrackingRecHit2DHeterogeneous &	operator= (TrackingRecHit2DHeterogeneous &&)=default

auto	phiBinner ()

auto	phiBinnerStorage ()

	TrackingRecHit2DHeterogeneous ()=default

	TrackingRecHit2DHeterogeneous (const TrackingRecHit2DHeterogeneous &)=delete

	TrackingRecHit2DHeterogeneous (TrackingRecHit2DHeterogeneous &&)=default

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

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

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 8 of file TrackingRecHit2DHeterogeneous.h.

Member Typedef Documentation

◆ PhiBinner

template<typename Traits>

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

Definition at line 13 of file TrackingRecHit2DHeterogeneous.h.

◆ unique_ptr

template<typename Traits>

template<typename T >

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

Definition at line 11 of file TrackingRecHit2DHeterogeneous.h.

Constructor & Destructor Documentation

◆ TrackingRecHit2DHeterogeneous() [1/4]

template<typename Traits>

TrackingRecHit2DHeterogeneous< Traits >::TrackingRecHit2DHeterogeneous ( )

default

◆ TrackingRecHit2DHeterogeneous() [2/4]

template<typename Traits >

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

explicit

Definition at line 81 of file TrackingRecHit2DHeterogeneous.h.

     : m_nHits(nHits), m_hitsModuleStart(hitsModuleStart) {
   auto view = Traits::template make_host_unique<TrackingRecHit2DSOAView>(stream);
  
   view->m_nHits = nHits;
   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);
     m_store16 = Traits::template make_unique<uint16_t[]>(nHits * n16, stream);
     m_store32 = Traits::template make_unique<float[]>(nHits * n32 + phase1PixelTopology::numberOfLayers + 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
   }
 }

◆ ~TrackingRecHit2DHeterogeneous()

template<typename Traits>

TrackingRecHit2DHeterogeneous< Traits >::~TrackingRecHit2DHeterogeneous ( )

default

◆ TrackingRecHit2DHeterogeneous() [3/4]

template<typename Traits>

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

delete

◆ TrackingRecHit2DHeterogeneous() [4/4]

template<typename Traits>

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

default

Member Function Documentation

◆ copyFromGPU()

template<typename Traits>

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

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

◆ hitsLayerStart()

template<typename Traits>

auto TrackingRecHit2DHeterogeneous< Traits >::hitsLayerStart ( )

inline

Definition at line 37 of file TrackingRecHit2DHeterogeneous.h.

37 { return m_hitsLayerStart; }

References TrackingRecHit2DHeterogeneous< Traits >::m_hitsLayerStart.

◆ hitsModuleStart()

template<typename Traits>

auto TrackingRecHit2DHeterogeneous< Traits >::hitsModuleStart ( ) const

inline

Definition at line 36 of file TrackingRecHit2DHeterogeneous.h.

36 { return m_hitsModuleStart; }

References TrackingRecHit2DHeterogeneous< Traits >::m_hitsModuleStart.

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

◆ hitsModuleStartToHostAsync()

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

Definition at line 15 of file TrackingRecHit2DHeterogeneous.cc.

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

References cudaCheck, TrackingRecHit2DHeterogeneous< Traits >::m_hitsModuleStart, gpuClustering::maxNumModules, runTheMatrix::ret, and cms::cuda::stream.

◆ iphi()

template<typename Traits>

auto TrackingRecHit2DHeterogeneous< Traits >::iphi ( )

inline

Definition at line 40 of file TrackingRecHit2DHeterogeneous.h.

40 { return m_iphi; }

References TrackingRecHit2DHeterogeneous< Traits >::m_iphi.

◆ localCoordToHostAsync()

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

Definition at line 8 of file TrackingRecHit2DHeterogeneous.cc.

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