Skip to content

Clean up the pixel local reconstruction code #602

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
fwyzard merged 4 commits into from
Dec 30, 2020
Merged

Clean up the pixel local reconstruction code #602

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
5d6c35e
Clean up the pixel local reconstruction code
fwyzard Dec 30, 2020
a0d4763
Clean up the pixel local reconstruction code
fwyzard Dec 30, 2020
210db0f
Clean up the pixel local reconstruction code
fwyzard Dec 30, 2020
e503ab5
Clean up the pixel local reconstruction code
fwyzard Dec 30, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions CUDADataFormats/Common/src/classes_def.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Open in desktop
Original file line number Diff line number Diff line change
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
Loading