Skip to content

Commit

Permalink
GPU2CPU for clusters and RecHIts (#18)
Browse files Browse the repository at this point in the history
  • Loading branch information
@VinInn @fwyzard
VinInn authored and fwyzard committed Nov 6, 2020
1 parent d46d528 commit e7e4027
Show file tree
Hide file tree
Showing 6 changed files with 116 additions and 153 deletions.
17 changes: 0 additions & 17 deletions RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -49,23 +49,6 @@ class PixelCPEFast final : public PixelCPEBase
LocalPoint localPosition (DetParam const & theDetParam, ClusterParam & theClusterParam) const override;
LocalError localError (DetParam const & theDetParam, ClusterParam & theClusterParam) const override;

//--------------------------------------------------------------------
// Methods.
//------------------------------------------------------------------
static float
generic_position_formula( int size, //!< Size of this projection.
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-width
float theThickness, //detector thickness
float cot_angle, //!< cot of alpha_ or beta_
float pitch, //!< thePitchX or thePitchY
bool first_is_big, //!< true if the first is big
bool last_is_big //!< true if the last is big
);

static void
collect_edge_charges(ClusterParam & theClusterParam, //!< input, the cluster
int & Q_f_X, //!< output, Q first in X
Expand Down
52 changes: 52 additions & 0 deletions RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
#include "DataFormats/GeometrySurface/interface/SOARotation.h"
#include <cstdint>
#include <cmath>
#include <iterator>

#include<cassert>

Expand Down Expand Up @@ -68,6 +69,9 @@ namespace pixelCPEforGPU {

float xpos[N];
float ypos[N];

float xerr[N];
float yerr[N];
};


Expand Down Expand Up @@ -203,4 +207,52 @@ namespace pixelCPEforGPU {

}

// FIXME these are errors form Run1
constexpr inline
void error(CommonParams const & comParams, DetParams const & detParams, ClusParams & cp, uint32_t ic) {
// Edge cluster errors
cp.xerr[ic]= 0.0050;
cp.yerr[ic]= 0.0085;


constexpr float xerr_barrel_l1[] = {0.00115, 0.00120, 0.00088};
constexpr float xerr_barrel_l1_def = 0.01030;
constexpr float yerr_barrel_l1[] = {0.00375,0.00230,0.00250,0.00250,0.00230,0.00230,0.00210,0.00210,0.00240};
constexpr float yerr_barrel_l1_def=0.00210;
constexpr float xerr_barrel_ln[]= {0.00115, 0.00120, 0.00088};
constexpr float xerr_barrel_ln_def=0.01030;
constexpr float yerr_barrel_ln[]= {0.00375,0.00230,0.00250,0.00250,0.00230,0.00230,0.00210,0.00210,0.00240};
constexpr float yerr_barrel_ln_def=0.00210;
constexpr float xerr_endcap[]= {0.0020, 0.0020};
constexpr float xerr_endcap_def=0.0020;
constexpr float yerr_endcap[]= {0.00210};
constexpr float yerr_endcap_def=0.00210;

// is edgy?
bool isEdgeX = cp.minRow[ic]==0 || cp.maxRow[ic]==phase1PixelTopology::lastRowInModule;
bool isEdgeY = cp.minCol[ic]==0 || cp.maxCol[ic]==phase1PixelTopology::lastColInModule;

if (!isEdgeX) {
auto sx = cp.maxRow[ic]-cp.minRow[ic];
if (!detParams.isBarrel ) {
cp.xerr[ic] = sx <std::size(xerr_endcap) ? xerr_endcap[sx] : xerr_endcap_def;
} else if (detParams.layer==1) {
cp.xerr[ic] = sx <std::size(xerr_barrel_l1) ? xerr_barrel_l1[sx]: xerr_barrel_l1_def;
} else {
cp.xerr[ic] = sx <std::size(xerr_barrel_ln) ? xerr_barrel_ln[sx]: xerr_barrel_ln_def;
}
}

if (!isEdgeY) {
auto sy =cp.maxCol[ic]-cp.minCol[ic];;
if (!detParams.isBarrel ) {
cp.yerr[ic] = sy <std::size(yerr_endcap) ? yerr_endcap[sy] : yerr_endcap_def;
} else if (detParams.layer==1) {
cp.yerr[ic] = sy <std::size(yerr_barrel_l1) ? yerr_barrel_l1[sy]: yerr_barrel_l1_def;
} else {
cp.yerr[ic] = sy <std::size(yerr_barrel_ln) ? yerr_barrel_ln[sy]: yerr_barrel_ln_def;
}
}
}

}
31 changes: 24 additions & 7 deletions RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -19,13 +19,17 @@ HitsOnGPU allocHitsOnGPU() {
cudaCheck(cudaMalloc((void**) & hh.xg_d,(gpuClustering::MaxNumModules*256)*sizeof(float)));
cudaCheck(cudaMalloc((void**) & hh.yg_d,(gpuClustering::MaxNumModules*256)*sizeof(float)));
cudaCheck(cudaMalloc((void**) & hh.zg_d,(gpuClustering::MaxNumModules*256)*sizeof(float)));
cudaCheck(cudaMalloc((void**) & hh.xerr_d,(gpuClustering::MaxNumModules*256)*sizeof(float)));
cudaCheck(cudaMalloc((void**) & hh.yerr_d,(gpuClustering::MaxNumModules*256)*sizeof(float)));
cudaCheck(cudaMalloc((void**) & hh.mr_d,(gpuClustering::MaxNumModules*256)*sizeof(uint16_t)));
cudaDeviceSynchronize();

return hh;
}


void pixelRecHits_wrapper(
HitsOnCPU
pixelRecHits_wrapper(
context const & c,
pixelCPEforGPU::ParamsOnGPU const * cpeParams,
uint32_t ndigis,
Expand All @@ -42,7 +46,7 @@ void pixelRecHits_wrapper(
std::partial_sum(std::begin(hitsModuleStart),std::end(hitsModuleStart),std::begin(hitsModuleStart));

auto nhits = hitsModuleStart[gpuClustering::MaxNumModules];
std::cout << " total number of clusters " << nhits << std::endl;
// std::cout << " total number of clusters " << nhits << std::endl;

cudaCheck(cudaMemcpyAsync(hh.hitsModuleStart_d, &hitsModuleStart, (gpuClustering::MaxNumModules+1)*sizeof(uint32_t), cudaMemcpyHostToDevice, c.stream));

Expand All @@ -60,15 +64,28 @@ void pixelRecHits_wrapper(
hh.hitsModuleStart_d,
hh.charge_d,
hh.xg_d,hh.yg_d,hh.zg_d,
false
hh.xerr_d,hh.yerr_d, hh.mr_d,
true // for the time being stay local...
);

int32_t charge[nhits];
cudaCheck(cudaMemcpyAsync(charge, hh.charge_d, nhits*sizeof(uint32_t), cudaMemcpyDeviceToHost, c.stream));

// all this needed only if hits on CPU are required....
HitsOnCPU hoc(nhits);
memcpy(hoc.hitsModuleStart,hitsModuleStart,2001*sizeof(uint32_t));
cudaCheck(cudaMemcpyAsync(hoc.charge.data(), hh.charge_d, nhits*sizeof(uint32_t), cudaMemcpyDeviceToHost, c.stream));

/*
int ngood=0;
auto l1 = hitsModuleStart[96];
for (auto i=0U; i<nhits; ++i) if( charge[i]>4000 || (i<l1 &&charge[i]>2000) ) ++ngood;
for (auto i=0U; i<nhits; ++i) if( hoc.charge[i]>4000 || (i<l1 &&hoc.charge[i]>2000) ) ++ngood;
std::cout << " total number of good clusters " << ngood << std::endl;
*/

cudaCheck(cudaMemcpyAsync(hoc.xl.data(), hh.xg_d, nhits*sizeof(uint32_t), cudaMemcpyDeviceToHost, c.stream));
cudaCheck(cudaMemcpyAsync(hoc.yl.data(), hh.yg_d, nhits*sizeof(uint32_t), cudaMemcpyDeviceToHost, c.stream));
cudaCheck(cudaMemcpyAsync(hoc.xe.data(), hh.xerr_d, nhits*sizeof(uint32_t), cudaMemcpyDeviceToHost, c.stream));
cudaCheck(cudaMemcpyAsync(hoc.ye.data(), hh.yerr_d, nhits*sizeof(uint32_t), cudaMemcpyDeviceToHost, c.stream));
cudaCheck(cudaMemcpyAsync(hoc.mr.data(), hh.mr_d, nhits*sizeof(uint16_t), cudaMemcpyDeviceToHost, c.stream));

return hoc;
}
17 changes: 15 additions & 2 deletions RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.h
View file
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
#pragma once

#include<cstdint>
#include<vector>

namespace pixelCPEforGPU {
struct ParamsOnGPU;
Expand All @@ -9,15 +10,27 @@ namespace pixelCPEforGPU {
struct context;

struct HitsOnGPU{

uint32_t * hitsModuleStart_d;
int32_t * charge_d;
float *xg_d, *yg_d, *zg_d;
float *xerr_d, *yerr_d;
uint16_t * mr_d;
};

struct HitsOnCPU {
explicit HitsOnCPU(uint32_t nhits) :
charge(nhits),xl(nhits),yl(nhits),xe(nhits),ye(nhits), mr(nhits){}
uint32_t hitsModuleStart[2001];
std::vector<int32_t> charge;
std::vector<float> xl, yl;
std::vector<float> xe, ye;
std::vector<uint16_t> mr;
};


HitsOnGPU allocHitsOnGPU();

void pixelRecHits_wrapper(
HitsOnCPU pixelRecHits_wrapper(
context const & c,
pixelCPEforGPU::ParamsOnGPU const * cpeParams,
uint32_t ndigis,
Expand Down
8 changes: 6 additions & 2 deletions RecoLocalTracker/SiPixelRecHits/plugins/gpuPixelRecHits.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@ namespace gpuPixelRecHits {
uint32_t const * hitsModuleStart,
int32_t * chargeh,
float * xh, float * yh, float * zh,
float * xe, float * ye, uint16_t * mr,
bool local // if true fill just x & y in local coord...
){

Expand Down Expand Up @@ -112,7 +113,8 @@ namespace gpuPixelRecHits {
assert(h<2000*256);

pixelCPEforGPU::position(cpeParams->commonParams(), cpeParams->detParams(me), clusParams,ic);

pixelCPEforGPU::error(cpeParams->commonParams(), cpeParams->detParams(me), clusParams,ic);

chargeh[h] = clusParams.charge[ic];

if (local) {
Expand All @@ -123,7 +125,9 @@ namespace gpuPixelRecHits {
xh[h],yh[h],zh[h]
);
}

xe[h]= clusParams.xerr[ic];
ye[h]= clusParams.yerr[ic];
mr[h]= clusParams.minRow[ic];
}

}
Expand Down
144 changes: 19 additions & 125 deletions RecoLocalTracker/SiPixelRecHits/src/PixelCPEFast.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -166,11 +166,6 @@ PixelCPEFast::localPosition(DetParam const & theDetParam, ClusterParam & theClus

assert(!theClusterParam.with_track_angle);

float chargeWidthX = (theDetParam.lorentzShiftInCmX * theDetParam.widthLAFractionX);
float chargeWidthY = (theDetParam.lorentzShiftInCmY * theDetParam.widthLAFractionY);
float shiftX = 0.5f*theDetParam.lorentzShiftInCmX;
float shiftY = 0.5f*theDetParam.lorentzShiftInCmY;

if ( UseErrorsFromTemplates_ ) {

float qclus = theClusterParam.theCluster->charge();
Expand Down Expand Up @@ -223,135 +218,34 @@ PixelCPEFast::localPosition(DetParam const & theDetParam, ClusterParam & theClus
UseErrorsFromTemplates_ && TruncatePixelCharge_
);

//--- Find the inner widths along X and Y in one shot. We
//--- compute the upper right corner of the inner pixels
//--- (== lower left corner of upper right pixel) and
//--- the lower left corner of the inner pixels
//--- (== upper right corner of lower left pixel), and then
//--- subtract these two points in the formula.

//--- Upper Right corner of Lower Left pixel -- in measurement frame
uint16_t llx = theClusterParam.theCluster->minPixelRow()+1;
uint16_t lly = theClusterParam.theCluster->minPixelCol()+1;

//--- Lower Left corner of Upper Right pixel -- in measurement frame
uint16_t urx = theClusterParam.theCluster->maxPixelRow();
uint16_t ury = theClusterParam.theCluster->maxPixelCol();

auto llxl = phase1PixelTopology::localX(llx);
auto llyl = phase1PixelTopology::localY(lly);
auto urxl = phase1PixelTopology::localX(urx);
auto uryl = phase1PixelTopology::localY(ury);


float xPos =
generic_position_formula( theClusterParam.theCluster->sizeX(),
Q_f_X, Q_l_X,
llxl, urxl,
chargeWidthX, // lorentz shift in cm
theDetParam.theThickness,
theClusterParam.cotalpha,
theDetParam.thePitchX,
phase1PixelTopology::isBigPixX( theClusterParam.theCluster->minPixelRow() ),
phase1PixelTopology::isBigPixX( theClusterParam.theCluster->maxPixelRow() )
);

// apply the lorentz offset correction
xPos = xPos + shiftX + theDetParam.thePitchX*float(phase1PixelTopology::xOffset);

float yPos =
generic_position_formula( theClusterParam.theCluster->sizeY(),
Q_f_Y, Q_l_Y,
llyl, uryl,
chargeWidthY, // lorentz shift in cm
theDetParam.theThickness,
theClusterParam.cotbeta,
theDetParam.thePitchY,
phase1PixelTopology::isBigPixY( theClusterParam.theCluster->minPixelCol() ),
phase1PixelTopology::isBigPixY( theClusterParam.theCluster->maxPixelCol() )
);
// apply the lorentz offset correction
yPos = yPos + shiftY + theDetParam.thePitchY*float(phase1PixelTopology::yOffset);

//--- Now put the two together
LocalPoint pos_in_local( xPos, yPos );
return pos_in_local;
}
// do GPU like ...

pixelCPEforGPU::ClusParams cp;


cp.minRow[0] = theClusterParam.theCluster->minPixelRow();
cp.maxRow[0] = theClusterParam.theCluster->maxPixelRow();
cp.minCol[0] = theClusterParam.theCluster->minPixelCol();
cp.maxCol[0] = theClusterParam.theCluster->maxPixelCol();

//-----------------------------------------------------------------------------
//! A generic version of the position formula. Since it works for both
//! X and Y, in the interest of the simplicity of the code, all parameters
//! are passed by the caller. The only class variable used by this method
//! is the theThickness, since that's common for both X and Y.
//-----------------------------------------------------------------------------
float
PixelCPEFast::
generic_position_formula( int size, //!< Size of this projection.
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
float theThickness, //detector thickness
float cot_angle, //!< cot of alpha_ or beta_
float pitch, //!< thePitchX or thePitchY
bool first_is_big, //!< true if the first is big
bool last_is_big //!< true if the last is big
)
{

float geom_center = 0.5f * pitch*float( upper_edge_first_pix + lower_edge_last_pix );

//--- The case of only one pixel in this projection is separate. Note that
//--- here first_pix == last_pix, so the average of the two is still the
//--- center of the pixel.
if ( size == 1 ) {return geom_center;}

float W_eff; // the compiler detects the logic below (and warns if buggy!!!!0
bool simple=true;
if (size==2) {
//--- 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);
float W_inner = pitch * float(lower_edge_last_pix-upper_edge_first_pix); // in cm

//--- Predicted charge width from geometry
float W_pred = theThickness * cot_angle // geometric correction (in cm)
- lorentz_shift; // (in cm) &&& check fpix!

W_eff = std::abs( W_pred ) - W_inner;
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;

//--- If the observed charge width is inconsistent with the expectations
//--- based on the track, do *not* use W_pred-W_innner. Instead, replace
//--- it with an *average* effective charge width, which is the average
//--- length of the edge pixels.
//
simple = ( W_eff < 0.0f ) | ( W_eff > pitch ); // this produces "large" regressions for very small numeric differences...
auto ind = theDetParam.theDet->index();
pixelCPEforGPU::position(m_commonParamsGPU, m_detParamsGPU[ind],cp,0);
auto xPos = cp.xpos[0];
auto yPos = cp.ypos[0];

}
if (simple) {
//--- Total length of the two edge pixels (first+last)
float sum_of_edge = 2.0f;
if (first_is_big) sum_of_edge += 1.0f;
if (last_is_big) sum_of_edge += 1.0f;
W_eff = pitch * 0.5f * sum_of_edge; // ave. length of edge pixels (first+last) (cm)
}


//--- Finally, compute the position in this projection
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;
float hit_pos = geom_center + 0.5f*(Qdiff/Qsum) * W_eff;

return hit_pos;
//--- Now put the two together
LocalPoint pos_in_local( xPos, yPos );
return pos_in_local;
}



//-----------------------------------------------------------------------------
//! Collect the edge charges in x and y, in a single pass over the pixel vector.
//! Calculate charge in the first and last pixel projected in x and y
Expand Down

0 comments on commit e7e4027

Please sign in to comment.