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;
}
}