gpuClustering Namespace Reference

Functions
	__attribute__ ((always_inline)) void clusterChargeCut(SiPixelClusterThresholds clusterThresholds
	for (auto module=firstModule;module< endModule;module+=gridDim.x)
constexpr uint32_t	maxHitsInIter ()
constexpr uint32_t	maxHitsInModule ()

Variables
uint16_t *__restrict__ uint16_t const *__restrict__	adc
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__ int32_t *__restrict__	clusterId
auto	endModule = moduleStart[0]
auto	firstModule = blockIdx.x
uint16_t *__restrict__	id
constexpr int	invalidClusterId = -9999
constexpr uint16_t	invalidModuleId = std::numeric_limits<uint16_t>::max() - 1
constexpr int32_t	maxNumClustersPerModules = maxHitsInModule()
constexpr uint16_t	maxNumModules = 2000
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__	moduleId
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__	moduleStart
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__	nClustersInModule
__shared__ uint16_t	newclusId [maxNumClustersPerModules]
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__ int32_t *__restrict__ uint32_t	numElements
__shared__ uint8_t	ok [maxNumClustersPerModules]
uint16_t const *__restrict__	x
uint16_t const *__restrict__ uint16_t const *__restrict__	y

Function Documentation

__attribute__()

gpuClustering::__attribute__ ( (always_inline)  )

inline

for()

gpuClustering::for

(

)

Definition at line 33 of file gpuClusterChargeCut.h.

                                                                             {
      auto firstPixel = moduleStart[1 + module];
      auto thisModuleId = id[firstPixel];
      assert(thisModuleId < maxNumModules);
      assert(thisModuleId == moduleId[module]);
 
      auto nclus = nClustersInModule[thisModuleId];
      if (nclus == 0)
        continue;
 
      if (threadIdx.x == 0 && nclus > maxNumClustersPerModules)
        printf("Warning too many clusters in module %d in block %d: %d > %d\n",
               thisModuleId,
               blockIdx.x,
               nclus,
               maxNumClustersPerModules);
 
      auto first = firstPixel + threadIdx.x;
 
      if (nclus > maxNumClustersPerModules) {
        // remove excess  FIXME find a way to cut charge first....
        for (auto i = first; i < numElements; i += blockDim.x) {
          if (id[i] == invalidModuleId)
            continue;  // not valid
          if (id[i] != thisModuleId)
            break;  // end of module
          if (clusterId[i] >= maxNumClustersPerModules) {
            id[i] = invalidModuleId;
            clusterId[i] = invalidModuleId;
          }
        }
        nclus = maxNumClustersPerModules;
      }
 
#ifdef GPU_DEBUG
      if (thisModuleId % 100 == 1)
        if (threadIdx.x == 0)
          printf("start cluster charge cut for module %d in block %d\n", thisModuleId, blockIdx.x);
#endif
 
      assert(nclus <= maxNumClustersPerModules);
      for (auto i = threadIdx.x; i < nclus; i += blockDim.x) {
        charge[i] = 0;
      }
      __syncthreads();
 
      for (auto i = first; i < numElements; i += blockDim.x) {
        if (id[i] == invalidModuleId)
          continue;  // not valid
        if (id[i] != thisModuleId)
          break;  // end of module
        atomicAdd(&charge[clusterId[i]], adc[i]);
      }
      __syncthreads();
 
      auto chargeCut =
          clusterThresholds.getThresholdForLayerOnCondition(thisModuleId < phase1PixelTopology::layerStart[1]);
      for (auto i = threadIdx.x; i < nclus; i += blockDim.x) {
        newclusId[i] = ok[i] = charge[i] > chargeCut ? 1 : 0;
      }
 
      __syncthreads();
 
      // renumber
      __shared__ uint16_t ws[32];
      cms::cuda::blockPrefixScan(newclusId, nclus, ws);
 
      assert(nclus >= newclusId[nclus - 1]);
 
      if (nclus == newclusId[nclus - 1])
        continue;
 
      nClustersInModule[thisModuleId] = newclusId[nclus - 1];
      __syncthreads();
 
      // mark bad cluster again
      for (auto i = threadIdx.x; i < nclus; i += blockDim.x) {
        if (0 == ok[i])
          newclusId[i] = invalidModuleId + 1;
      }
      __syncthreads();
 
      // reassign id
      for (auto i = first; i < numElements; i += blockDim.x) {
        if (id[i] == invalidModuleId)
          continue;  // not valid
        if (id[i] != thisModuleId)
          break;  // end of module
        clusterId[i] = newclusId[clusterId[i]] - 1;
        if (clusterId[i] == invalidModuleId)
          id[i] = invalidModuleId;
      }
 
      //done
    }  // loop on modules

References cms::cudacompat::__syncthreads(), adc, cms::cuda::assert(), cms::cudacompat::atomicAdd(), cms::cudacompat::blockDim, cms::cudacompat::blockIdx, ALCARECOTkAlJpsiMuMu_cff::charge, clusterId, first, SiPixelClusterThresholds::getThresholdForLayerOnCondition(), mps_fire::i, invalidModuleId, phase1PixelTopology::layerStart, maxNumClustersPerModules, maxNumModules, callgraph::module, moduleId, moduleStart, nClustersInModule, newclusId, numElements, ok, cms::cudacompat::threadIdx, and gpuVertexFinder::ws.

maxHitsInIter()

constexpr uint32_t gpuClustering::maxHitsInIter ( )

constexpr

Definition at line 14 of file gpuClusteringConstants.h.

{ return 160; }

maxHitsInModule()

constexpr uint32_t gpuClustering::maxHitsInModule ( )

constexpr

Definition at line 16 of file gpuClusteringConstants.h.

{ return 1024; }

Referenced by SiPixelRecHitFromCUDA::produce(), and SiPixelRecHitSoAFromLegacy::produce().

Variable Documentation

adc

uint16_t* __restrict__ uint16_t const* __restrict__ gpuClustering::adc

Definition at line 20 of file gpuClusterChargeCut.h.

Referenced by QIE10Task::_process(), QIE11Task::_process(), PedestalTask::_process(), LEDTask::_process(), DigiTask::_process(), LaserTask::_process(), CastorCoderDb::adc2fC_(), HcalCoderDb::adc2fC_(), CastorTPGCoder::adc2Linear(), HcalTPGCoder::adc2Linear(), HcaluLUTTPGCoder::adc2Linear(), hcal::adc_for_sample(), AutoCorrMat::add(), Chamber_AutoCorrMat::add(), SiPixelCluster::add(), SiPixelArrayBuffer::add_adc(), ThreeThresholdAlgorithm::addToCandidate(), ESElectronicsSimFast::analogToDigital(), CSCComparatorDigiValidation::analyze(), CMTRawAnalyzer::analyze(), ZDCQIE10Task::analyze(), HcalDigiStatistics::analyze(), CastorDigiStatistics::analyze(), ESTimingTask::analyze(), FCDTask::analyze(), ZDCTask::analyze(), SiStripBaselineValidator::analyze(), EcalPreshowerNoiseDistrib::analyze(), SiStripMonitorDigi::analyze(), CTPPSPixelDQMSource::analyze(), BtlDigiHitsValidation::analyze(), EcalSimRawData::analyze(), HGCalDigiValidation::analyze(), EcalLaserAnalyzerYousi::analyze(), HGCalTBAnalyzer::analyze(), RecAnalyzerMinbias::analyze(), EcnaAnalyzer::analyze(), HGCalTBAnalyzer::analyzeDigi(), SiStripClusterToDigiProducer::applyGain(), ThreeThresholdAlgorithm::applyGains(), CSCFindPeakTime::averageTime(), EcalClusterLazyToolsBase::BasicClusterTime(), FWSiPixelClusterProxyBuilder::build(), RPixDetClusterizer::buildClusters(), DCCTBDataMapper::buildXtalFields(), gpuCalibPixel::calibDigis(), RPixDetClusterizer::calibrate(), PixelThresholdClusterizer::calibrate(), SiStripAPVRestorer::cleaner_LocalMinimumAdder(), hcal::reconstruction::compute_coder_charge(), hcal::reconstruction::compute_diff_charge_gain(), DigiConverterFP420::convert(), SiTrivialDigitalConverter::convert(), RPixDummyROCSimulator::ConvertChargeToHits(), convertHB(), SiTrivialDigitalConverter::convertRaw(), PixelThresholdClusterizer::copy_to_buffer(), CastorSimpleRecAlgoImpl::corrSaturation(), HGCalDigiValidation::digiValidation(), PulseFitWithShape::doFit(), TPNFit::doFit(), PulseFitWithFunction::doFit(), ESElectronicsSim::encode(), EcalLiteDTUCoder::encode(), EcalCoder::encode(), ESRecHitAnalyticAlgo::EvalAmplitude(), ESRecHitFitAlgo::EvalAmplitude(), ESRecHitSimAlgo::evalAmplitude(), PedestalsTask::fill(), DaqScopeModeTask::fill(), PedsFullNoiseTask::fill(), SiPixelDigiModule::fill(), sistrip::FEDBufferPayloadCreator::fillClusterData(), sistrip::FEDBufferPayloadCreator::fillClusterDataPreMixMode(), CMTRawAnalyzer::fillDigiErrorsHFQIE10(), CMTRawAnalyzer::fillDigiErrorsQIE11(), reco::HcalNoiseInfoProducer::filldigis(), Phase2TrackerMonitorDigi::fillITPixelDigiHistos(), SiPixelCalibDigiProducer::fillPixel(), QcdLowPtDQM::fillPixels(), HLTHcalLaserMisfireFilter::filter(), CSCXonStrip_MatchGatti::findXOnStrip(), JetCoreClusterSplitter::fittingSplit(), CSCFindPeakTime::fivePoleFitCharge(), CSCFindPeakTime::fivePoleFitTime(), SiStripAPVRestorer::flatRegionsFinder(), for(), TSFit::fpol3dg(), generate(), EcalSimRawData::genFeData(), EcalSimRawData::getEbDigi(), HcalQIEManager::getHfQieTable(), HcalLutManager::getLinearizationLutXmlFromAsciiMasterEmap(), HcaluLUTTPGCoder::getMSB(), CSCMake2DRecHit::hitFromStripAndWire(), SiStripAPVRestorer::hybridFormatInspect(), CTPPSPixelDataFormatter::interpretRawData(), PixelDataFormatter::interpretRawData(), SiStripFedZeroSuppression::isAValidDigi(), CastorSimpleRecAlgoImpl::isSaturated(), HcalZSAlgoRealistic::keepMe(), CastorLedAnalysis::LedCastorHists(), HcalLedAnalysis::LedHBHEHists(), HcalLedAnalysis::LedHFHists(), HcalLedAnalysis::LedHOHists(), OptoScanTask::locateTicks(), HcaluLUTTPGCoder::lookupMSB(), SiPixelDigitizerAlgorithm::make_digis(), CSCHitFromStripOnly::makeStripData(), ShallowClustersProducer::NearDigis::NearDigis(), ESRecHitSimAlgo::oldEvalAmplitude(), pixelgpudetails::pack(), UHTRpacker::packQIE8sample(), CSCFindPeakTime::parabolaFitTime(), HcalTriggerPrimitiveAlgo::passTDC(), CSCFindPeakTime::peakTime(), PrintRecoObjects::print(), MatacqTBDataFormatter::printData(), MatacqDataFormatter::printData(), EcalFenixStripFgvbEE::process(), HBHEPhase1Reconstructor::processData(), SiPixelRecHitSoAFromLegacy::produce(), PreMixingSiPixelWorker::put(), PGlobalDigi::putBRL1Digis(), PGlobalDigi::putBRL2Digis(), PGlobalDigi::putBRL3Digis(), PGlobalDigi::putCSCstripDigis(), edm::DataMixingEMDigiWorker::putEM(), PGlobalDigi::putFWD1nDigis(), PGlobalDigi::putFWD1pDigis(), PGlobalDigi::putFWD2nDigis(), PGlobalDigi::putFWD2pDigis(), edm::DataMixingSiPixelWorker::putSiPixel(), PGlobalDigi::putTECW1Digis(), PGlobalDigi::putTECW2Digis(), PGlobalDigi::putTECW3Digis(), PGlobalDigi::putTECW4Digis(), PGlobalDigi::putTECW5Digis(), PGlobalDigi::putTECW6Digis(), PGlobalDigi::putTECW7Digis(), PGlobalDigi::putTECW8Digis(), PGlobalDigi::putTIBL1Digis(), PGlobalDigi::putTIBL2Digis(), PGlobalDigi::putTIBL3Digis(), PGlobalDigi::putTIBL4Digis(), PGlobalDigi::putTIDW1Digis(), PGlobalDigi::putTIDW2Digis(), PGlobalDigi::putTIDW3Digis(), PGlobalDigi::putTOBL1Digis(), PGlobalDigi::putTOBL2Digis(), PGlobalDigi::putTOBL3Digis(), PGlobalDigi::putTOBL4Digis(), TMatacq::rawPulseAnalysis(), EcalPedOffset::readDACs(), EcalPedHists::readEBdigis(), EcalPedHists::readEEdigis(), HcalDigisValidation::reco(), ecaldqm::PresampleTask::runOnDigis(), ecaldqm::LedTask::runOnDigis(), ecaldqm::LaserTask::runOnDigis(), HGCFEElectronics< DFr >::runShaperWithToT(), HGCFEElectronics< DFr >::runSimpleShaper(), BTLElectronicsSim::runTrivialShaper(), ETLElectronicsSim::runTrivialShaper(), HGCFEElectronics< DFr >::runTrivialShaper(), HcalSignalGenerator< HBHEDigitizerTraits >::samplesInPE(), HGCalConcentratorAutoEncoderImpl::select(), SiPixelArrayBuffer::set_adc(), DCUCapsuleTempRawDat::setCapsuleTempADC(), CastorQIEShape::setLowEdges(), HcalQIEShape::setLowEdges(), TPNPulse::setPulse(), TAPDPulse::setPulse(), QIE10DataFrame::setSample(), QIE11DataFrame::setSample(), SiPixelCluster::SiPixelCluster(), EBDataFrame::spikeEstimator(), ThreeThresholdAlgorithm::stripByStripAdd(), SiStripApvShotCleaner::subtractCM(), hcaldqm::utilities::sumQ_v10(), SiStripFedZeroSuppression::suppress(), HcalTDC::timing(), ZeroSuppressFP420::trkFEDclusterizer(), DigiConverterFP420::truncate(), SiTrivialDigitalConverter::truncate(), SiTrivialDigitalConverter::truncateRaw(), HcalFiberPattern::unpack(), HcaluLUTTPGCoder::update(), HcalSignalGenerator< HBHEDigitizerTraits >::validDigi(), EcalSignalGenerator< EEDigitizerTraits >::validDigi(), and ESUnpacker::word2digi().

clusterId

uint16_t const *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t *__restrict__ int32_t *__restrict__ gpuClustering::clusterId

Definition at line 20 of file gpuClusterChargeCut.h.

Referenced by for(), and RealisticHitToClusterAssociator::XYdistanceFromMaxHit().

endModule

auto gpuClustering::endModule = moduleStart[0]

Definition at line 32 of file gpuClusterChargeCut.h.

firstModule

auto gpuClustering::firstModule = blockIdx.x

Definition at line 31 of file gpuClusterChargeCut.h.

id

uint16_t* __restrict__ gpuClustering::id

Definition at line 20 of file gpuClusterChargeCut.h.

invalidClusterId

constexpr int gpuClustering::invalidClusterId = -9999

constexpr

Definition at line 21 of file gpuClusteringConstants.h.

invalidModuleId

constexpr uint16_t gpuClustering::invalidModuleId = std::numeric_limits<uint16_t>::max() - 1

constexpr

Definition at line 20 of file gpuClusteringConstants.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuCalibPixel::calibDigis(), for(), SiPixelRecHitSoAFromLegacy::produce(), and SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper().

maxNumClustersPerModules

constexpr int32_t gpuClustering::maxNumClustersPerModules = maxHitsInModule()

constexpr

Definition at line 19 of file gpuClusteringConstants.h.

Referenced by for(), and SiPixelDigisClustersFromSoA::produce().

maxNumModules

constexpr uint16_t gpuClustering::maxNumModules = 2000

constexpr

Definition at line 18 of file gpuClusteringConstants.h.

Referenced by TrackerAlignmentLevelBuilder::buildPXBAlignmentLevels(), TrackerAlignmentLevelBuilder::buildPXEAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTECAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTIBAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTIDAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTOBAlignmentLevels(), gpuCalibPixel::calibDigis(), for(), gpuPixelDoublets::for(), TrackingRecHit2DHeterogeneous< Traits >::hitsModuleStartToHostAsync(), SiPixelDigisClustersFromSoA::produce(), SiPixelRecHitSoAFromLegacy::produce(), SiPixelRecHitFromCUDA::produce(), and cms::SiPixelRecHitConverter::run().

moduleId

uint16_t const *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t *__restrict__ gpuClustering::moduleId

Definition at line 20 of file gpuClusterChargeCut.h.

Referenced by SummaryOutputProducer::GenericSummary::addEntry(), dqm::implementation::DQMStore::bookTransaction(), edm::SystemTimeKeeper::fillTriggerTimingReport(), for(), and SummaryOutputProducer::prepareSummaries().

moduleStart

uint16_t const *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ gpuClustering::moduleStart

Definition at line 20 of file gpuClusterChargeCut.h.

Referenced by gpuCalibPixel::calibDigis(), and for().

nClustersInModule

uint16_t const *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ gpuClustering::nClustersInModule

Definition at line 20 of file gpuClusterChargeCut.h.

Referenced by gpuCalibPixel::calibDigis(), and for().

newclusId

__shared__ uint16_t gpuClustering::newclusId[maxNumClustersPerModules]

Definition at line 29 of file gpuClusterChargeCut.h.

Referenced by for().

numElements

uint16_t const *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t *__restrict__ int32_t *__restrict__ int gpuClustering::numElements

Initial value:

{
    __shared__ int32_t charge[maxNumClustersPerModules]

Definition at line 26 of file gpuClusterChargeCut.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuCalibPixel::calibDigis(), edm::OwnArray< T, MAX_SIZE, P >::fillView(), for(), and TopologyWorker::setPartList().

ok

__shared__ uint8_t gpuClustering::ok[maxNumClustersPerModules]

Definition at line 28 of file gpuClusterChargeCut.h.

Referenced by for().

x

uint16_t const* __restrict__ gpuClustering::x

Definition at line 39 of file gpuClustering.h.

y

uint16_t const* __restrict__ uint16_t const* __restrict__ gpuClustering::y

Definition at line 39 of file gpuClustering.h.