gpuClustering Namespace Reference

Namespaces
	pixelStatus

Functions
template<typename TrackerTraits >
	__attribute__ ((always_inline)) void clusterChargeCut(SiPixelClusterThresholds clusterThresholds
	assert (nMaxModules< maxNumModules)
	for (auto module=firstModule;module< endModule;module+=gridDim.x)
constexpr uint32_t	maxHitsInIter ()

Variables
uint16_t *__restrict__ uint16_t const *__restrict__	adc
__shared__ int32_t	charge [maxNumClustersPerModules]
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__ uint32_t const *__restrict__ int32_t *__restrict__	clusterId
constexpr uint16_t	clusterThresholdLayerOne = 2000
constexpr uint16_t	clusterThresholdOtherLayers = 4000
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 uint32_t	maxNumDigis = 3 * 256 * 1024
constexpr uint16_t	maxNumModules = 4000
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
__shared__ int	msize
uint16_t *__restrict__ uint16_t const *__restrict__ uint32_t const *__restrict__ uint32_t *__restrict__	nClustersInModule
__shared__ uint16_t	newclusId [maxNumClustersPerModules]
constexpr int	nMaxModules = TrackerTraits::numberOfModules
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]
constexpr uint32_t	pixelSizeX = 160
constexpr uint32_t	pixelSizeY = 416
constexpr const uint32_t	pixelStatusSize = isPhase2 ? 1 : pixelStatus::size
constexpr int	startBPIX2 = TrackerTraits::layerStart[1]
__shared__ uint32_t	status [pixelStatusSize]
uint16_t *__restrict__ uint16_t const *__restrict__	x
uint16_t *__restrict__ uint16_t const *__restrict__ uint16_t const *__restrict__	y

Function Documentation

__attribute__()

template<typename TrackerTraits >

gpuClustering::__attribute__ ( (always_inline)  )

inline

assert()

gpuClustering::assert

(

)

Referenced by for().

for()

gpuClustering::for

(

)

Definition at line 42 of file gpuClusterChargeCut.h.

References cms::cudacompat::__syncthreads(), cms::cudacompat::__syncthreads_and(), adc, assert(), cms::cudacompat::atomicAdd(), cms::cudacompat::blockDim, cms::cudacompat::blockIdx, charge, DMR_cfg::chargeCut, clusterId, dqmdumpme::first, SiPixelClusterThresholds::getThresholdForLayerOnCondition(), caHitNtupletGeneratorKernels::good, mps_fire::i, invalidModuleId, phase1PixelTopology::maxNumClustersPerModules, callgraph::module, moduleId, moduleStart, nClustersInModule, newclusId, nMaxModules, numElements, hltrates_dqm_sourceclient-live_cfg::offset, ok, startBPIX2, cms::cudacompat::threadIdx, and gpuVertexFinder::ws.

                                                                             {
      auto firstPixel = moduleStart[1 + module];
      auto thisModuleId = id[firstPixel];
      while (thisModuleId == invalidModuleId and firstPixel < numElements) {
        // skip invalid or duplicate pixels
        ++firstPixel;
        thisModuleId = id[firstPixel];
      }
      if (firstPixel >= numElements) {
        // reached the end of the input while skipping the invalid pixels, nothing left to do
        break;
      }
      if (thisModuleId != moduleId[module]) {
        // reached the end of the module while skipping the invalid pixels, skip this module
        continue;
      }
      assert(thisModuleId < nMaxModules);

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

      bool good = true;
      for (auto i = threadIdx.x; i < nclus; i += blockDim.x) {
        newclusId[i] = ok[i] = charge[i] >= chargeCut ? 1 : 0;
        if (0 == ok[i])
          good = false;
      }

      // if all clusters above threshold do nothing
      if (__syncthreads_and(good))
        continue;

      // renumber
      __shared__ uint16_t ws[32];
      constexpr auto maxThreads = 1024;
      auto minClust = nclus > maxThreads ? maxThreads : nclus;

      cms::cuda::blockPrefixScan(newclusId, newclusId, minClust, ws);
      if constexpr (maxNumClustersPerModules > maxThreads)  //only if needed
      {
        for (uint32_t offset = maxThreads; offset < nclus; offset += maxThreads) {
          cms::cuda::blockPrefixScan(newclusId + offset, newclusId + offset, nclus - offset, ws);

          for (uint32_t i = threadIdx.x + offset; i < nclus; i += blockDim.x) {
            uint32_t prevBlockEnd = ((i / maxThreads) * maxThreads) - 1;
            newclusId[i] += newclusId[prevBlockEnd];
          }

          __syncthreads();
        }
      }
      assert(nclus > newclusId[nclus - 1]);

      nClustersInModule[thisModuleId] = newclusId[nclus - 1];

      // 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
        if (0 == ok[clusterId[i]])
          clusterId[i] = id[i] = invalidModuleId;
        else
          clusterId[i] = newclusId[clusterId[i]] - 1;
      }

      //done
      __syncthreads();
    }  // loop on modules

maxHitsInIter()

constexpr uint32_t gpuClustering::maxHitsInIter

(

)

Definition at line 14 of file gpuClusteringConstants.h.

{ return 160; }  //TODO better tuning for PU 140-200

Variable Documentation

adc

uint16_t* __restrict__ uint16_t const* __restrict__ gpuClustering::adc

Definition at line 21 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(), SiPixelCluster::add(), SiPixelArrayBuffer::add_adc(), ThreeThresholdAlgorithm::addToCandidate(), ESElectronicsSimFast::analogToDigital(), CSCComparatorDigiValidation::analyze(), ZDCQIE10Task::analyze(), ESTimingTask::analyze(), FCDTask::analyze(), SiStripBaselineValidator::analyze(), ZDCTask::analyze(), EcalPreshowerNoiseDistrib::analyze(), CastorDigiStatistics::analyze(), SiStripMonitorDigi::analyze(), HcalDigiStatistics::analyze(), CTPPSPixelDQMSource::analyze(), BtlDigiHitsValidation::analyze(), HGCalDigiValidation::analyze(), EcalSimRawData::analyze(), EcalLaserAnalyzerYousi::analyze(), HGCalTBAnalyzer::analyze(), RecAnalyzerMinbias::analyze(), EcnaAnalyzer::analyze(), CMTRawAnalyzer::analyze(), HGCalTBAnalyzer::analyzeDigi(), SiStripClusterToDigiProducer::applyGain(), ThreeThresholdAlgorithm::applyGains(), CSCFindPeakTime::averageTime(), EcalClusterLazyToolsBase::BasicClusterTime(), FWSiPixelClusterProxyBuilder::build(), RPixDetClusterizer::buildClusters(), DCCTBDataMapper::buildXtalFields(), gpuCalibPixel::calibDigis(), gpuCalibPixel::calibDigisPhase2(), RPixDetClusterizer::calibrate(), PixelThresholdClusterizer::calibrate(), SiStripAPVRestorer::cleaner_LocalMinimumAdder(), hcal::reconstruction::compute_coder_charge(), hcal::reconstruction::compute_diff_charge_gain(), 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(), SiPixelDigitizerAlgorithm::lateSignalReweight(), 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(), EcalFenixStrip::process_part1(), HBHEPhase1Reconstructor::processData(), 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(), SiTrivialDigitalConverter::truncate(), SiTrivialDigitalConverter::truncateRaw(), HcalFiberPattern::unpack(), HcaluLUTTPGCoder::update(), HcalSignalGenerator< HBHEDigitizerTraits >::validDigi(), EcalSignalGenerator< EEDigitizerTraits >::validDigi(), and ESUnpacker::word2digi().

charge

__shared__ int32_t gpuClustering::charge[maxNumClustersPerModules]

Definition at line 30 of file gpuClusterChargeCut.h.

Referenced by for().

clusterId

uint16_t *__restrict__ 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 21 of file gpuClusterChargeCut.h.

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

clusterThresholdLayerOne

constexpr uint16_t gpuClustering::clusterThresholdLayerOne = 2000

Definition at line 17 of file gpuClusteringConstants.h.

Referenced by SiPixelDigisClustersFromSoAT< TrackerTraits >::fillDescriptions(), and SiPixelRawToClusterCUDAT< TrackerTraits >::fillDescriptions().

clusterThresholdOtherLayers

constexpr uint16_t gpuClustering::clusterThresholdOtherLayers = 4000

Definition at line 18 of file gpuClusteringConstants.h.

Referenced by SiPixelDigisClustersFromSoAT< TrackerTraits >::fillDescriptions(), and SiPixelRawToClusterCUDAT< TrackerTraits >::fillDescriptions().

endModule

auto gpuClustering::endModule = moduleStart[0]

Definition at line 41 of file gpuClusterChargeCut.h.

firstModule

auto gpuClustering::firstModule = blockIdx.x

Definition at line 40 of file gpuClusterChargeCut.h.

id

uint16_t *__restrict__ gpuClustering::id

Definition at line 21 of file gpuClusterChargeCut.h.

invalidClusterId

constexpr int gpuClustering::invalidClusterId = -9999

Definition at line 24 of file gpuClusteringConstants.h.

invalidModuleId

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

Definition at line 23 of file gpuClusteringConstants.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuCalibPixel::calibDigis(), gpuCalibPixel::calibDigisPhase2(), for(), SiPixelRecHitSoAFromLegacyT< TrackerTraits >::produce(), and SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper().

maxNumDigis

constexpr uint32_t gpuClustering::maxNumDigis = 3 * 256 * 1024

Definition at line 20 of file gpuClusteringConstants.h.

Referenced by SiPixelPhase2DigiToClusterCUDA::acquire().

maxNumModules

constexpr uint16_t gpuClustering::maxNumModules = 4000

Definition at line 21 of file gpuClusteringConstants.h.

Referenced by TrackerAlignmentLevelBuilder::buildPXBAlignmentLevels(), TrackerAlignmentLevelBuilder::buildPXEAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTECAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTIBAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTIDAlignmentLevels(), TrackerAlignmentLevelBuilder::buildTOBAlignmentLevels(), gpuPixelDoublets::for(), and cms::SiPixelRecHitConverter::run().

moduleId

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

Definition at line 21 of file gpuClusterChargeCut.h.

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

moduleStart

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

Definition at line 21 of file gpuClusterChargeCut.h.

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

msize

__shared__ int gpuClustering::msize

Definition at line 110 of file gpuClustering.h.

Referenced by SiPixelDigiModule::fill(), FWCompactVerticalLayout::GetDefaultSize(), and FWCompactVerticalLayout::Layout().

nClustersInModule

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

Definition at line 21 of file gpuClusterChargeCut.h.

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

newclusId

__shared__ uint16_t gpuClustering::newclusId[maxNumClustersPerModules]

Definition at line 32 of file gpuClusterChargeCut.h.

Referenced by for().

nMaxModules

constexpr int gpuClustering::nMaxModules = TrackerTraits::numberOfModules

Definition at line 35 of file gpuClusterChargeCut.h.

Referenced by for(), and SiPixelRecHitFromCUDAT< TrackerTraits >::produce().

numElements

uint16_t *__restrict__ 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:

{
    constexpr int32_t maxNumClustersPerModules = TrackerTraits::maxNumClustersPerModules

Definition at line 27 of file gpuClusterChargeCut.h.

Referenced by gpuPixelRecHits::__attribute__(), gpuCalibPixel::calibDigis(), gpuCalibPixel::calibDigisPhase2(), for(), and edmtest::ViewAnalyzer::testProduct().

ok

__shared__ uint8_t gpuClustering::ok[maxNumClustersPerModules]

Definition at line 31 of file gpuClusterChargeCut.h.

Referenced by for().

pixelSizeX

constexpr uint32_t gpuClustering::pixelSizeX = 160

Definition at line 18 of file gpuClustering.h.

Referenced by gpuClustering::pixelStatus::getIndex().

pixelSizeY

constexpr uint32_t gpuClustering::pixelSizeY = 416

Definition at line 19 of file gpuClustering.h.

pixelStatusSize

constexpr const uint32_t gpuClustering::pixelStatusSize = isPhase2 ? 1 : pixelStatus::size

Definition at line 108 of file gpuClustering.h.

startBPIX2

constexpr int gpuClustering::startBPIX2 = TrackerTraits::layerStart[1]

Definition at line 34 of file gpuClusterChargeCut.h.

Referenced by for(), and SiPixelRecHitSoAFromLegacyT< TrackerTraits >::produce().

status

__shared__ uint32_t gpuClustering::status[pixelStatusSize]

Definition at line 109 of file gpuClustering.h.

Referenced by gpuClustering::pixelStatus::getStatus(), gpuClustering::pixelStatus::isDuplicate(), and gpuClustering::pixelStatus::promote().

x

uint16_t* __restrict__ uint16_t const* __restrict__ gpuClustering::x

Definition at line 97 of file gpuClustering.h.

y

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

Definition at line 97 of file gpuClustering.h.