Skip to content

Commit

Permalink
Clean up the pixel local reconstruction code (#602)
Browse files Browse the repository at this point in the history 
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.
  • Loading branch information
@fwyzard
fwyzard committed Dec 30, 2020
1 parent 7717f1b commit 2e8e5fd
Show file tree
Hide file tree
Showing 10 changed files with 68 additions and 71 deletions.
4 changes: 2 additions & 2 deletions CUDADataFormats/Common/src/classes_def.xml
View file
@@ -1,4 +1,4 @@
<lcgdict>
<class name="HostProduct<unsigned int[]>" persistent="false"/>
<class name="edm::Wrapper<HostProduct<unsigned int[]>>" persistent="false"/>
<class name="HostProduct<uint32_t[]>" persistent="false"/>
<class name="edm::Wrapper<HostProduct<uint32_t[]>>" persistent="false"/>
</lcgdict>
12 changes: 6 additions & 6 deletions CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h
View file
Expand Up @@ -10,7 +10,7 @@ class TrackingRecHit2DHeterogeneous {
template <typename T>
using unique_ptr = typename Traits::template unique_ptr<T>;

using Hist = TrackingRecHit2DSOAView::Hist;
using PhiBinner = TrackingRecHit2DSOAView::PhiBinner;

TrackingRecHit2DHeterogeneous() = default;

Expand All @@ -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
Expand All @@ -48,7 +48,7 @@ class TrackingRecHit2DHeterogeneous {
unique_ptr<uint16_t[]> m_store16; //!
unique_ptr<float[]> m_store32; //!

unique_ptr<TrackingRecHit2DSOAView::Hist> m_HistStore; //!
unique_ptr<TrackingRecHit2DSOAView::PhiBinner> m_PhiBinnerStore; //!
unique_ptr<TrackingRecHit2DSOAView::AverageGeometry> m_AverageGeometryStore; //!

unique_ptr<TrackingRecHit2DSOAView> m_view; //!
Expand All @@ -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;
};
Expand Down Expand Up @@ -98,13 +98,13 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(uint32_t nH
// so unless proven VERY inefficient we keep it ordered as generated
m_store16 = Traits::template make_device_unique<uint16_t[]>(nHits * n16, stream);
m_store32 = Traits::template make_device_unique<float[]>(nHits * n32 + 11, stream);
m_HistStore = Traits::template make_device_unique<TrackingRecHit2DSOAView::Hist>(stream);
m_PhiBinnerStore = Traits::template make_device_unique<TrackingRecHit2DSOAView::PhiBinner>(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);
Expand Down
8 changes: 3 additions & 5 deletions CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h
View file
Expand Up @@ -20,8 +20,6 @@ class TrackingRecHit2DSOAView {
using PhiBinner =
cms::cuda::HistoContainer<int16_t, 128, gpuClustering::maxNumClusters, 8 * sizeof(int16_t), hindex_type, 10>;

using Hist = PhiBinner; // FIXME

using AverageGeometry = phase1PixelTopology::AverageGeometry;

template <typename>
Expand Down Expand Up @@ -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; }
Expand Down Expand Up @@ -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;
};
Expand Down
1 change: 0 additions & 1 deletion CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h
View file
Expand Up @@ -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,
Expand Down
9 changes: 5 additions & 4 deletions CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc
View file
Expand Up @@ -8,6 +8,7 @@
#include <cuda_runtime.h>

// 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"
Expand Down Expand Up @@ -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;
}
Expand All @@ -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;
Expand Down
1 change: 0 additions & 1 deletion CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h
View file
Expand Up @@ -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];
Expand Down
8 changes: 4 additions & 4 deletions RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
View file
Expand Up @@ -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_;
Expand Down
42 changes: 21 additions & 21 deletions RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h
View file
Expand Up @@ -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];

Expand Down Expand Up @@ -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
Expand All @@ -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.

Expand All @@ -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,
Expand Down Expand Up @@ -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;

Expand All @@ -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
Expand All @@ -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
Expand Down
16 changes: 8 additions & 8 deletions RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h
View file
Expand Up @@ -99,10 +99,10 @@ namespace gpuPixelRecHits {
clusParams.minCol[ic] = std::numeric_limits<uint32_t>::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();
Expand Down Expand Up @@ -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();
Expand Down
38 changes: 19 additions & 19 deletions RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc
View file
Expand Up @@ -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;
Expand All @@ -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);
Expand All @@ -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<ClusterParamGeneric&>(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();
Expand All @@ -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;
}
}

Expand Down

0 comments on commit 2e8e5fd

Please sign in to comment.