From 2e8e5fd76b0dbfb8abf86ba42de0615ef1e70918 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Wed, 30 Dec 2020 17:31:37 +0100 Subject: [PATCH] Clean up the pixel local reconstruction code (cms-patatrack#602) Address the pixel local reconstruction review comments: - remove obsolete comments; - consistently use named constants; - rename data members and methods to be more descriptive; - rename local variables according to the coding rules and for consistency with cms-sw#32591; - update transient dictionaries to match data types. --- CUDADataFormats/Common/src/classes_def.xml | 4 +- .../interface/TrackingRecHit2DHeterogeneous.h | 12 +++--- .../interface/TrackingRecHit2DSOAView.h | 8 ++-- .../SiPixelROCsStatusAndMappingWrapper.h | 1 - .../src/SiPixelROCsStatusAndMappingWrapper.cc | 9 ++-- .../interface/SiPixelROCsStatusAndMapping.h | 1 - .../SiPixelRecHits/interface/PixelCPEFast.h | 8 ++-- .../SiPixelRecHits/interface/pixelCPEforGPU.h | 42 +++++++++---------- .../SiPixelRecHits/plugins/gpuPixelRecHits.h | 16 +++---- .../SiPixelRecHits/src/PixelCPEFast.cc | 38 ++++++++--------- 10 files changed, 68 insertions(+), 71 deletions(-) diff --git a/CUDADataFormats/Common/src/classes_def.xml b/CUDADataFormats/Common/src/classes_def.xml index 024d927595914..d8514251c807a 100644 --- a/CUDADataFormats/Common/src/classes_def.xml +++ b/CUDADataFormats/Common/src/classes_def.xml @@ -1,4 +1,4 @@ - - + + diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h index 72a136ab5f5b6..967b5c6c8282f 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h @@ -10,7 +10,7 @@ class TrackingRecHit2DHeterogeneous { template using unique_ptr = typename Traits::template unique_ptr; - using Hist = TrackingRecHit2DSOAView::Hist; + using PhiBinner = TrackingRecHit2DSOAView::PhiBinner; TrackingRecHit2DHeterogeneous() = default; @@ -33,7 +33,7 @@ class TrackingRecHit2DHeterogeneous { auto hitsModuleStart() const { return m_hitsModuleStart; } auto hitsLayerStart() { return m_hitsLayerStart; } - auto phiBinner() { return m_hist; } + auto phiBinner() { return m_phiBinner; } auto iphi() { return m_iphi; } // only the local coord and detector index @@ -48,7 +48,7 @@ class TrackingRecHit2DHeterogeneous { unique_ptr m_store16; //! unique_ptr m_store32; //! - unique_ptr m_HistStore; //! + unique_ptr m_PhiBinnerStore; //! unique_ptr m_AverageGeometryStore; //! unique_ptr m_view; //! @@ -58,7 +58,7 @@ class TrackingRecHit2DHeterogeneous { uint32_t const* m_hitsModuleStart; // needed for legacy, this is on GPU! // needed as kernel params... - Hist* m_hist; + PhiBinner* m_phiBinner; uint32_t* m_hitsLayerStart; int16_t* m_iphi; }; @@ -98,13 +98,13 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous(uint32_t nH // so unless proven VERY inefficient we keep it ordered as generated m_store16 = Traits::template make_device_unique(nHits * n16, stream); m_store32 = Traits::template make_device_unique(nHits * n32 + 11, stream); - m_HistStore = Traits::template make_device_unique(stream); + m_PhiBinnerStore = Traits::template make_device_unique(stream); auto get16 = [&](int i) { return m_store16.get() + i * nHits; }; auto get32 = [&](int i) { return m_store32.get() + i * nHits; }; // copy all the pointers - m_hist = view->m_hist = m_HistStore.get(); + m_phiBinner = view->m_phiBinner = m_PhiBinnerStore.get(); view->m_xl = get32(0); view->m_yl = get32(1); diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h index 53297a78a428f..7f3c59cd70faf 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h @@ -20,8 +20,6 @@ class TrackingRecHit2DSOAView { using PhiBinner = cms::cuda::HistoContainer; - using Hist = PhiBinner; // FIXME - using AverageGeometry = phase1PixelTopology::AverageGeometry; template @@ -67,8 +65,8 @@ class TrackingRecHit2DSOAView { __device__ __forceinline__ uint32_t* hitsLayerStart() { return m_hitsLayerStart; } __device__ __forceinline__ uint32_t const* hitsLayerStart() const { return m_hitsLayerStart; } - __device__ __forceinline__ Hist& phiBinner() { return *m_hist; } - __device__ __forceinline__ Hist const& phiBinner() const { return *m_hist; } + __device__ __forceinline__ PhiBinner& phiBinner() { return *m_phiBinner; } + __device__ __forceinline__ PhiBinner const& phiBinner() const { return *m_phiBinner; } __device__ __forceinline__ AverageGeometry& averageGeometry() { return *m_averageGeometry; } __device__ __forceinline__ AverageGeometry const& averageGeometry() const { return *m_averageGeometry; } @@ -96,7 +94,7 @@ class TrackingRecHit2DSOAView { uint32_t* m_hitsLayerStart; - PhiBinner* m_hist; // FIXME use a more descriptive name consistently + PhiBinner* m_phiBinner; uint32_t m_nHits; }; diff --git a/CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h b/CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h index 5f875d7dff5a9..f7555a75d9bec 100644 --- a/CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h +++ b/CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h @@ -14,7 +14,6 @@ class SiPixelFedCablingMap; class TrackerGeometry; class SiPixelQuality; -// TODO: since this has more information than just cabling map, maybe we should invent a better name? class SiPixelROCsStatusAndMappingWrapper { public: SiPixelROCsStatusAndMappingWrapper(SiPixelFedCablingMap const &cablingMap, diff --git a/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc b/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc index 2437696656d25..665d31b97ead2 100644 --- a/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc +++ b/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc @@ -8,6 +8,7 @@ #include // CMSSW includes +#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h" #include "CondFormats/SiPixelObjects/interface/SiPixelFedCablingMap.h" #include "CondFormats/SiPixelObjects/interface/SiPixelFedCablingTree.h" @@ -51,8 +52,8 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFe else cablingMapHost->badRocs[index] = false; } else { // store some dummy number - cablingMapHost->rawId[index] = 9999; - cablingMapHost->rocInDet[index] = 9999; + cablingMapHost->rawId[index] = gpuClustering::invalidModuleId; + cablingMapHost->rocInDet[index] = gpuClustering::invalidModuleId; cablingMapHost->badRocs[index] = true; modToUnpDefault[index] = true; } @@ -70,8 +71,8 @@ SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFe // idinLnk varies between 1 to 8 for (int i = 1; i < index; i++) { - if (cablingMapHost->rawId[i] == 9999) { - cablingMapHost->moduleId[i] = 9999; + if (cablingMapHost->rawId[i] == gpuClustering::invalidModuleId) { + cablingMapHost->moduleId[i] = gpuClustering::invalidModuleId; } else { /* std::cout << cablingMapHost->rawId[i] << std::endl; diff --git a/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h b/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h index a0771aaefb366..f7cd8dedca941 100644 --- a/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h +++ b/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h @@ -11,7 +11,6 @@ namespace pixelgpudetails { constexpr unsigned int MAX_SIZE_BYTE_BOOL = MAX_SIZE * sizeof(unsigned char); } // namespace pixelgpudetails -// TODO: since this has more information than just cabling map, maybe we should invent a better name? struct SiPixelROCsStatusAndMapping { alignas(128) unsigned int fed[pixelgpudetails::MAX_SIZE]; alignas(128) unsigned int link[pixelgpudetails::MAX_SIZE]; diff --git a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h index f50db3af11868..bf85f6c74ebd9 100644 --- a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h +++ b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h @@ -60,10 +60,10 @@ class PixelCPEFast final : public PixelCPEBase { void errorFromTemplates(DetParam const &theDetParam, ClusterParamGeneric &theClusterParam, float qclus) const; static void collect_edge_charges(ClusterParam &theClusterParam, //!< input, the cluster - int &Q_f_X, //!< output, Q first in X - int &Q_l_X, //!< output, Q last in X - int &Q_f_Y, //!< output, Q first in Y - int &Q_l_Y, //!< output, Q last in Y + int &q_f_X, //!< output, Q first in X + int &q_l_X, //!< output, Q last in X + int &q_f_Y, //!< output, Q first in Y + int &q_l_Y, //!< output, Q last in Y bool truncate); const float edgeClusterErrorX_; diff --git a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h index f655329d02013..03e136d8d23ef 100644 --- a/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h +++ b/RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h @@ -81,10 +81,10 @@ namespace pixelCPEforGPU { uint32_t minCol[N]; uint32_t maxCol[N]; - int32_t Q_f_X[N]; - int32_t Q_l_X[N]; - int32_t Q_f_Y[N]; - int32_t Q_l_Y[N]; + int32_t q_f_X[N]; + int32_t q_l_X[N]; + int32_t q_f_Y[N]; + int32_t q_l_Y[N]; int32_t charge[N]; @@ -114,8 +114,8 @@ namespace pixelCPEforGPU { } constexpr inline float correction(int sizeM1, - int Q_f, //!< Charge in the first pixel. - int Q_l, //!< Charge in the last pixel. + int q_f, //!< Charge in the first pixel. + int q_l, //!< Charge in the last pixel. uint16_t upper_edge_first_pix, //!< As the name says. uint16_t lower_edge_last_pix, //!< As the name says. float lorentz_shift, //!< L-shift at half thickness @@ -134,16 +134,16 @@ namespace pixelCPEforGPU { //--- Width of the clusters minus the edge (first and last) pixels. //--- In the note, they are denoted x_F and x_L (and y_F and y_L) // assert(lower_edge_last_pix >= upper_edge_first_pix); - auto W_inner = pitch * float(lower_edge_last_pix - upper_edge_first_pix); // in cm + auto w_inner = pitch * float(lower_edge_last_pix - upper_edge_first_pix); // in cm //--- Predicted charge width from geometry - auto W_pred = theThickness * cot_angle // geometric correction (in cm) + auto w_pred = theThickness * cot_angle // geometric correction (in cm) - lorentz_shift; // (in cm) &&& check fpix! - w_eff = std::abs(W_pred) - W_inner; + w_eff = std::abs(w_pred) - w_inner; //--- If the observed charge width is inconsistent with the expectations - //--- based on the track, do *not* use W_pred-W_inner. Instead, replace + //--- based on the track, do *not* use w_pred-w_inner. Instead, replace //--- it with an *average* effective charge width, which is the average //--- length of the edge pixels. @@ -162,14 +162,14 @@ namespace pixelCPEforGPU { } //--- Finally, compute the position in this projection - float Qdiff = Q_l - Q_f; - float Qsum = Q_l + Q_f; + float qdiff = q_l - q_f; + float qsum = q_l + q_f; //--- Temporary fix for clusters with both first and last pixel with charge = 0 - if (Qsum == 0) - Qsum = 1.0f; + if (qsum == 0) + qsum = 1.0f; - return 0.5f * (Qdiff / Qsum) * w_eff; + return 0.5f * (qdiff / qsum) * w_eff; } constexpr inline void position(CommonParams const& __restrict__ comParams, @@ -206,8 +206,8 @@ namespace pixelCPEforGPU { if (phase1PixelTopology::isBigPixY(cp.maxCol[ic])) ++ysize; - int unbalanceX = 8. * std::abs(float(cp.Q_f_X[ic] - cp.Q_l_X[ic])) / float(cp.Q_f_X[ic] + cp.Q_l_X[ic]); - int unbalanceY = 8. * std::abs(float(cp.Q_f_Y[ic] - cp.Q_l_Y[ic])) / float(cp.Q_f_Y[ic] + cp.Q_l_Y[ic]); + int unbalanceX = 8. * std::abs(float(cp.q_f_X[ic] - cp.q_l_X[ic])) / float(cp.q_f_X[ic] + cp.q_l_X[ic]); + int unbalanceY = 8. * std::abs(float(cp.q_f_Y[ic] - cp.q_l_Y[ic])) / float(cp.q_f_Y[ic] + cp.q_l_Y[ic]); xsize = 8 * xsize - unbalanceX; ysize = 8 * ysize - unbalanceY; @@ -230,8 +230,8 @@ namespace pixelCPEforGPU { auto thickness = detParams.isBarrel ? comParams.theThicknessB : comParams.theThicknessE; auto xcorr = correction(cp.maxRow[ic] - cp.minRow[ic], - cp.Q_f_X[ic], - cp.Q_l_X[ic], + cp.q_f_X[ic], + cp.q_l_X[ic], llxl, urxl, detParams.chargeWidthX, // lorentz shift in cm @@ -242,8 +242,8 @@ namespace pixelCPEforGPU { phase1PixelTopology::isBigPixX(cp.maxRow[ic])); auto ycorr = correction(cp.maxCol[ic] - cp.minCol[ic], - cp.Q_f_Y[ic], - cp.Q_l_Y[ic], + cp.q_f_Y[ic], + cp.q_l_Y[ic], llyl, uryl, detParams.chargeWidthY, // lorentz shift in cm diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h index 89a40c8723ae3..2401fed6c5171 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h +++ b/RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h @@ -99,10 +99,10 @@ namespace gpuPixelRecHits { clusParams.minCol[ic] = std::numeric_limits::max(); clusParams.maxCol[ic] = 0; clusParams.charge[ic] = 0; - clusParams.Q_f_X[ic] = 0; - clusParams.Q_l_X[ic] = 0; - clusParams.Q_f_Y[ic] = 0; - clusParams.Q_l_Y[ic] = 0; + clusParams.q_f_X[ic] = 0; + clusParams.q_l_X[ic] = 0; + clusParams.q_f_Y[ic] = 0; + clusParams.q_l_Y[ic] = 0; } __syncthreads(); @@ -149,13 +149,13 @@ namespace gpuPixelRecHits { auto ch = std::min(digis.adc(i), pixmx); atomicAdd(&clusParams.charge[cl], ch); if (clusParams.minRow[cl] == x) - atomicAdd(&clusParams.Q_f_X[cl], ch); + atomicAdd(&clusParams.q_f_X[cl], ch); if (clusParams.maxRow[cl] == x) - atomicAdd(&clusParams.Q_l_X[cl], ch); + atomicAdd(&clusParams.q_l_X[cl], ch); if (clusParams.minCol[cl] == y) - atomicAdd(&clusParams.Q_f_Y[cl], ch); + atomicAdd(&clusParams.q_f_Y[cl], ch); if (clusParams.maxCol[cl] == y) - atomicAdd(&clusParams.Q_l_Y[cl], ch); + atomicAdd(&clusParams.q_l_Y[cl], ch); } __syncthreads(); diff --git a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc index 3a57ce120b545..0077c0748ca28 100644 --- a/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc +++ b/RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc @@ -354,11 +354,11 @@ LocalPoint PixelCPEFast::localPosition(DetParam const& theDetParam, ClusterParam theClusterParam.qBin_ = 0; } - int Q_f_X; //!< Q of the first pixel in X - int Q_l_X; //!< Q of the last pixel in X - int Q_f_Y; //!< Q of the first pixel in Y - int Q_l_Y; //!< Q of the last pixel in Y - collect_edge_charges(theClusterParam, Q_f_X, Q_l_X, Q_f_Y, Q_l_Y, useErrorsFromTemplates_ && truncatePixelCharge_); + int q_f_X; //!< Q of the first pixel in X + int q_l_X; //!< Q of the last pixel in X + int q_f_Y; //!< Q of the first pixel in Y + int q_l_Y; //!< Q of the last pixel in Y + collect_edge_charges(theClusterParam, q_f_X, q_l_X, q_f_Y, q_l_Y, useErrorsFromTemplates_ && truncatePixelCharge_); // do GPU like ... pixelCPEforGPU::ClusParams cp; @@ -368,10 +368,10 @@ LocalPoint PixelCPEFast::localPosition(DetParam const& theDetParam, ClusterParam cp.minCol[0] = theClusterParam.theCluster->minPixelCol(); cp.maxCol[0] = theClusterParam.theCluster->maxPixelCol(); - cp.Q_f_X[0] = Q_f_X; - cp.Q_l_X[0] = Q_l_X; - cp.Q_f_Y[0] = Q_f_Y; - cp.Q_l_Y[0] = Q_l_Y; + cp.q_f_X[0] = q_f_X; + cp.q_l_X[0] = q_l_X; + cp.q_f_Y[0] = q_f_Y; + cp.q_l_Y[0] = q_l_Y; auto ind = theDetParam.theDet->index(); pixelCPEforGPU::position(commonParamsGPU_, detParamsGPU_[ind], cp, 0); @@ -392,16 +392,16 @@ LocalPoint PixelCPEFast::localPosition(DetParam const& theDetParam, ClusterParam //! and the inner cluster charge, projected in x and y. //----------------------------------------------------------------------------- void PixelCPEFast::collect_edge_charges(ClusterParam& theClusterParamBase, //!< input, the cluster - int& Q_f_X, //!< output, Q first in X - int& Q_l_X, //!< output, Q last in X - int& Q_f_Y, //!< output, Q first in Y - int& Q_l_Y, //!< output, Q last in Y + int& q_f_X, //!< output, Q first in X + int& q_l_X, //!< output, Q last in X + int& q_f_Y, //!< output, Q first in Y + int& q_l_Y, //!< output, Q last in Y bool truncate) { ClusterParamGeneric& theClusterParam = static_cast(theClusterParamBase); // Initialize return variables. - Q_f_X = Q_l_X = 0; - Q_f_Y = Q_l_Y = 0; + q_f_X = q_l_X = 0; + q_f_Y = q_l_Y = 0; // Obtain boundaries in index units int xmin = theClusterParam.theCluster->minPixelRow(); @@ -421,15 +421,15 @@ void PixelCPEFast::collect_edge_charges(ClusterParam& theClusterParamBase, //!< // // X projection if (pixel.x == xmin) - Q_f_X += pix_adc; + q_f_X += pix_adc; if (pixel.x == xmax) - Q_l_X += pix_adc; + q_l_X += pix_adc; // // Y projection if (pixel.y == ymin) - Q_f_Y += pix_adc; + q_f_Y += pix_adc; if (pixel.y == ymax) - Q_l_Y += pix_adc; + q_l_Y += pix_adc; } }