From a9e74eb710729de0b722e7fd8a0ed86fff10c2c5 Mon Sep 17 00:00:00 2001 From: Vincenzo Innocente Date: Thu, 21 Mar 2019 19:27:12 +0100 Subject: [PATCH] Make GPU-CPU cluster matching deterministic (#294) 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. --- .../SiPixelCluster/interface/SiPixelCluster.h | 8 ++- .../SiPixelCluster/src/classes_def.xml | 1 + .../plugins/SiPixelDigisClustersFromSoA.cc | 1 + .../plugins/SiPixelRawToClusterCUDA.cc | 7 +- .../plugins/SiPixelRawToClusterGPUKernel.cu | 67 +++---------------- .../plugins/SiPixelRawToClusterGPUKernel.h | 2 +- 6 files changed, 22 insertions(+), 64 deletions(-) diff --git a/DataFormats/SiPixelCluster/interface/SiPixelCluster.h b/DataFormats/SiPixelCluster/interface/SiPixelCluster.h index 22f9cb1020814..ba75447e945bb 100644 --- a/DataFormats/SiPixelCluster/interface/SiPixelCluster.h +++ b/DataFormats/SiPixelCluster/interface/SiPixelCluster.h @@ -21,6 +21,7 @@ #include #include #include +#include class PixelDigi; @@ -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 thePixelOffset; @@ -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::max(); float err_x=-99999.9f; float err_y=-99999.9f; diff --git a/DataFormats/SiPixelCluster/src/classes_def.xml b/DataFormats/SiPixelCluster/src/classes_def.xml index 55c9fd8538417..d43f062877eb0 100644 --- a/DataFormats/SiPixelCluster/src/classes_def.xml +++ b/DataFormats/SiPixelCluster/src/classes_def.xml @@ -4,6 +4,7 @@ + diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc index 4c405a8c85afd..2c7da14cf72af 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc @@ -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) diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc index 5dc04009f4832..b23faad9e78d3 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc @@ -67,7 +67,6 @@ class SiPixelRawToClusterCUDA: public edm::stream::EDProducer const bool includeErrors_; const bool useQuality_; const bool usePilotBlade_; - const bool convertADCtoElectrons_; }; SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfig): @@ -77,8 +76,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi cablingMapLabel_(iConfig.getParameter("CablingMapLabel")), includeErrors_(iConfig.getParameter("IncludeErrors")), useQuality_(iConfig.getParameter("UseQualityInfo")), - usePilotBlade_(iConfig.getParameter ("UsePilotBlade")), // Control the usage of pilot-blade data, FED=40 - convertADCtoElectrons_(iConfig.getParameter("ConvertADCtoElectrons")) + usePilotBlade_(iConfig.getParameter ("UsePilotBlade")) // Control the usage of pilot-blade data, FED=40 { if(includeErrors_) { digiErrorPutToken_ = produces>(); @@ -97,7 +95,6 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d desc.add("IncludeErrors",true); desc.add("UseQualityInfo",false); desc.add("UsePilotBlade",false)->setComment("## Use pilot blades"); - desc.add("ConvertADCtoElectrons", false)->setComment("## do the calibration ADC-> Electron and apply the threshold, requried for clustering"); desc.add("InputLabel",edm::InputTag("rawDataCollector")); { edm::ParameterSetDescription psd0; @@ -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()); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index fead8e59a0db3..3d4e377eb8221 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -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 * **/ @@ -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=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, @@ -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& stream) { diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h index 1ab8bc3fa5998..a0f89dc241c64 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h @@ -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);