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