Skip to content

Commit

Permalink
Make GPU-CPU cluster matching deterministic (#294)
Browse files Browse the repository at this point in the history 
Makes the GPU-CPU cluster matching deterministic by intrusively marking CPU clusters with the cluster index.
Reuse existing padding space to store the extra transient field, so that the size of SiPixelCluster does not increase.
There is still a warning in case of mismatch of the content of the cluster (based on charge comparison), that can eventually be downgraded to a debug message.

Properly rewrite the loop in the RawToDigi_kernel .

Remove obsolete code (comments and configuration parameters) in SiPixelRawToClusterCUDA and SiPixelRawToClusterGPUKernel.
  • Loading branch information
@VinInn @fwyzard
VinInn authored and fwyzard committed Oct 19, 2020
1 parent 3e697d5 commit a9e74eb
Show file tree
Hide file tree
Showing 6 changed files with 22 additions and 64 deletions.
8 changes: 7 additions & 1 deletion DataFormats/SiPixelCluster/interface/SiPixelCluster.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@
#include <vector>
#include <cstdint>
#include <cassert>
#include <limits>

class PixelDigi;

Expand Down Expand Up @@ -196,7 +197,10 @@ class SiPixelCluster {
float getSplitClusterErrorX() const { return err_x; }
float getSplitClusterErrorY() const { return err_y; }


// the original id (they get sorted)
auto originalId() const { return theOriginalClusterId;}
void setOriginalId(uint16_t id) { theOriginalClusterId=id;}

private:

std::vector<uint8_t> thePixelOffset;
Expand All @@ -207,6 +211,8 @@ class SiPixelCluster {
uint16_t theMinPixelCol=MAXPOS; // Minimum pixel index in the y direction (left edge).
uint8_t thePixelRowSpan=0; // Span pixel index in the x direction (low edge).
uint8_t thePixelColSpan=0; // Span pixel index in the y direction (left edge).

uint16_t theOriginalClusterId=std::numeric_limits<uint16_t>::max();

float err_x=-99999.9f;
float err_y=-99999.9f;
Expand Down
1 change: 1 addition & 0 deletions DataFormats/SiPixelCluster/src/classes_def.xml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
<version ClassVersion="11" checksum="1473312403"/>
<version ClassVersion="12" checksum="2042538185"/>
<version ClassVersion="13" checksum="2314992115"/>
<field name="theOriginalClusterId" transient="true"/>
<field name="err_x" transient="true"/>
<field name="err_y" transient="true"/>
<ioread sourceClass="SiPixelCluster" version="[11-12]" targetClass="SiPixelCluster" source="uint16_t thePixelCol" target="theMinPixelCol">
Expand Down
1 change: 1 addition & 0 deletions RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -109,6 +109,7 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con
auto const & acluster = aclusters[ic];
if ( acluster.charge < clusterThreshold) continue;
SiPixelCluster cluster(acluster.isize,acluster.adc, acluster.x,acluster.y, acluster.xmin,acluster.ymin);
cluster.setOriginalId(ic);
++totCluseFilled;
// std::cout << "putting in this cluster " << ic << " " << cluster.charge() << " " << cluster.pixelADC().size() << endl;
// sort by row (x)
Expand Down
7 changes: 2 additions & 5 deletions RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -67,7 +67,6 @@ class SiPixelRawToClusterCUDA: public edm::stream::EDProducer<edm::ExternalWork>
const bool includeErrors_;
const bool useQuality_;
const bool usePilotBlade_;
const bool convertADCtoElectrons_;
};

SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfig):
Expand All @@ -77,8 +76,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi
cablingMapLabel_(iConfig.getParameter<std::string>("CablingMapLabel")),
includeErrors_(iConfig.getParameter<bool>("IncludeErrors")),
useQuality_(iConfig.getParameter<bool>("UseQualityInfo")),
usePilotBlade_(iConfig.getParameter<bool> ("UsePilotBlade")), // Control the usage of pilot-blade data, FED=40
convertADCtoElectrons_(iConfig.getParameter<bool>("ConvertADCtoElectrons"))
usePilotBlade_(iConfig.getParameter<bool> ("UsePilotBlade")) // Control the usage of pilot-blade data, FED=40
{
if(includeErrors_) {
digiErrorPutToken_ = produces<CUDAProduct<SiPixelDigiErrorsCUDA>>();
Expand All @@ -97,7 +95,6 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d
desc.add<bool>("IncludeErrors",true);
desc.add<bool>("UseQualityInfo",false);
desc.add<bool>("UsePilotBlade",false)->setComment("## Use pilot blades");
desc.add<bool>("ConvertADCtoElectrons", false)->setComment("## do the calibration ADC-> Electron and apply the threshold, requried for clustering");
desc.add<edm::InputTag>("InputLabel",edm::InputTag("rawDataCollector"));
{
edm::ParameterSetDescription psd0;
Expand Down Expand Up @@ -220,7 +217,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent, const edm::Event
gpuAlgo_.makeClustersAsync(gpuMap, gpuModulesToUnpack, gpuGains,
wordFedAppender,
std::move(errors_),
wordCounterGPU, fedCounter, convertADCtoElectrons_,
wordCounterGPU, fedCounter,
useQuality_, includeErrors_,
edm::MessageDrop::instance()->debugEnabled,
ctx.stream());
Expand Down
67 changes: 10 additions & 57 deletions RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2,8 +2,6 @@
*
* File Name: RawToClusterGPU.cu
* Description: It converts Raw data into Digi Format on GPU
* then it converts adc -> electron and
* applies the adc threshold to needed for clustering
* Finaly the Output of RawToDigi data is given to pixelClusterizer
*
**/
Expand Down Expand Up @@ -341,51 +339,6 @@ namespace pixelgpudetails {
return rID;
}

/*----------
* Name: applyADCthreshold_kernel()
* Desc: converts adc count to electrons and then applies the
* threshold on each channel.
* make pixel to 0 if it is below the threshold
* Input: xx_d[], yy_d[], layer_d[], wordCounter, adc[], ADCThreshold
*-----------
* Output: xx_adc[], yy_adc[] with pixel threshold applied
*/
// kernel to apply adc threshold on the channels


// Felice: gains and pedestals are not the same for each pixel. This code should be rewritten to take
// in account local gains/pedestals
// __global__ void applyADCthreshold_kernel(const uint32_t *xx_d, const uint32_t *yy_d, const uint32_t *layer_d, uint32_t *adc, const uint32_t wordCounter,
// const ADCThreshold adcThreshold, uint32_t *xx_adc, uint32_t *yy_adc ) {
// int tid = threadIdx.x;
// int gIndex = blockDim.x*blockIdx.x+tid;
// if (gIndex<wordCounter) {
// uint32_t adcOld = adc[gIndex];
// const float gain = adcThreshold.theElectronPerADCGain_; // default: 1 adc = 135 electrons
// const float pedestal = 0; //
// int adcNew = int(adcOld*gain+pedestal);
// // rare chance of entering into the if ()
// if (layer_d[gIndex]>=adcThreshold.theFirstStack_) {
// if (adcThreshold.theStackADC_==1 && adcOld==1) {
// adcNew = int(255*135); // Arbitrarily use overflow value.
// }
// if (adcThreshold.theStackADC_ >1 && adcThreshold.theStackADC_!=255 && adcOld>=1){
// adcNew = int((adcOld-1) * gain * 255/float(adcThreshold.theStackADC_-1));
// }
// }
//
// if (adcNew >adcThreshold.thePixelThreshold ) {
// xx_adc[gIndex]=xx_d[gIndex];
// yy_adc[gIndex]=yy_d[gIndex];
// }
// else {
// xx_adc[gIndex]=0; // 0: dead pixel
// yy_adc[gIndex]=0;
// }
// adc[gIndex] = adcNew;
// }
// }


// Kernel to perform Raw to Digi conversion
__global__ void RawToDigi_kernel(const SiPixelFedCablingMapGPU *cablingMap, const unsigned char *modToUnp,
Expand All @@ -397,14 +350,15 @@ namespace pixelgpudetails {
{
//if (threadIdx.x==0) printf("Event: %u blockIdx.x: %u start: %u end: %u\n", eventno, blockIdx.x, begin, end);

auto gIndex = threadIdx.x + blockIdx.x * blockDim.x;
xx[gIndex] = 0;
yy[gIndex] = 0;
adc[gIndex] = 0;
bool skipROC = false;
int32_t first = threadIdx.x + blockIdx.x*blockDim.x;
for (int32_t iloop=first, nend=wordCounter; iloop<nend; iloop+=blockDim.x*gridDim.x) {

auto gIndex = iloop;
xx[gIndex] = 0;
yy[gIndex] = 0;
adc[gIndex] = 0;
bool skipROC = false;

do { // too many coninue below.... (to be fixed)
if (gIndex < wordCounter) {
uint8_t fedId = fedIds[gIndex/2]; // +1200;

// initialize (too many coninue below)
Expand Down Expand Up @@ -499,8 +453,8 @@ namespace pixelgpudetails {
pdigi[gIndex] = pixelgpudetails::pack(globalPix.row, globalPix.col, adc[gIndex]);
moduleId[gIndex] = detId.moduleId;
rawIdArr[gIndex] = rawId;
} // end of if (gIndex < end)
} while (false); // end fake loop
} // end of loop (gIndex < end)

} // end of Raw to Digi kernel

// Interface to outside
Expand All @@ -511,7 +465,6 @@ namespace pixelgpudetails {
const WordFedAppender& wordFed,
PixelFormatterErrors&& errors,
const uint32_t wordCounter, const uint32_t fedCounter,
bool convertADCtoElectrons,
bool useQualityInfo, bool includeErrors, bool debug,
cuda::stream_t<>& stream)
{
Expand Down
2 changes: 1 addition & 1 deletion RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -185,7 +185,7 @@ namespace pixelgpudetails {
const SiPixelGainForHLTonGPU *gains,
const WordFedAppender& wordFed,
PixelFormatterErrors&& errors,
const uint32_t wordCounter, const uint32_t fedCounter, bool convertADCtoElectrons,
const uint32_t wordCounter, const uint32_t fedCounter,
bool useQualityInfo, bool includeErrors, bool debug,
cuda::stream_t<>& stream);

Expand Down

0 comments on commit a9e74eb

Please sign in to comment.