Enumerations
enum	Status : uint32_t { kEmpty = 0x00, kFound = 0x01, kDuplicate = 0x03 }

Functions
static __device__ constexpr uint32_t	getIndex (uint16_t x, uint16_t y)

__device__ constexpr uint32_t	getShift (uint16_t x, uint16_t y)

__device__ constexpr Status	getStatus (uint32_t const *__restrict__ status, uint16_t x, uint16_t y)

__device__ constexpr bool	isDuplicate (uint32_t const *__restrict__ status, uint16_t x, uint16_t y)

__device__ constexpr void	promote (uint32_t *__restrict__ status, const uint16_t x, const uint16_t y)

Variables
constexpr uint32_t	bits = 2

constexpr uint32_t	mask = (0x01 << bits) - 1

constexpr uint32_t	size = pixelSizeX * pixelSizeY / valuesPerWord

constexpr uint32_t	valuesPerWord = sizeof(uint32_t) * 8 / bits

Enumeration Type Documentation

◆ Status

enum gpuClustering::pixelStatus::Status : uint32_t

Enumerator
kEmpty
kFound
kDuplicate

Definition at line 23 of file gpuClustering.h.

23 : uint32_t { kEmpty = 0x00, kFound = 0x01, kDuplicate = 0x03 };

gpuClustering::pixelStatus::kDuplicate

Definition: gpuClustering.h:23

gpuClustering::pixelStatus::kEmpty

Definition: gpuClustering.h:23

gpuClustering::pixelStatus::kFound

Definition: gpuClustering.h:23

Function Documentation

◆ getIndex()

static __device__ constexpr uint32_t gpuClustering::pixelStatus::getIndex	(	uint16_t	x,
		uint16_t	y
	)

inlinestatic

Definition at line 30 of file gpuClustering.h.

References gpuClustering::pixelSizeX, valuesPerWord, and x.

Referenced by l1t::L1TGlobalUtil::getAlgBitFromName(), getStatus(), promote(), l1t::L1TGlobalUtil::retrieveL1Event(), and l1t::L1TGlobalUtil::retrieveL1Setup().

                                                                                  {
       return (pixelSizeX * y + x) / valuesPerWord;
     }

◆ getShift()

__device__ constexpr uint32_t gpuClustering::pixelStatus::getShift	(	uint16_t	x,
		uint16_t	y
	)

inline

Definition at line 34 of file gpuClustering.h.

References valuesPerWord, and x.

Referenced by getStatus(), and promote().

34 { return (x % valuesPerWord) * 2; }

gpuClustering::pixelStatus::valuesPerWord

constexpr uint32_t valuesPerWord

Definition: gpuClustering.h:27

gpuClustering::x

uint16_t *__restrict__ uint16_t const *__restrict__ x

Definition: gpuClustering.h:97

◆ getStatus()

__device__ constexpr Status gpuClustering::pixelStatus::getStatus	(	uint32_t const *__restrict__	status,
		uint16_t	x,
		uint16_t	y
	)

inline

Definition at line 36 of file gpuClustering.h.

References getIndex(), getShift(), mask, edm::shift, gpuClustering::status, and x.

Referenced by DataCertificationJetMET::dqmEndJob(), and isDuplicate().

                                                                                                               {
       uint32_t index = getIndex(x, y);
       uint32_t shift = getShift(x, y);
       return Status{(status[index] >> shift) & mask};
     }

◆ isDuplicate()

__device__ constexpr bool gpuClustering::pixelStatus::isDuplicate	(	uint32_t const *__restrict__	status,
		uint16_t	x,
		uint16_t	y
	)

inline

Definition at line 42 of file gpuClustering.h.

References getStatus(), kDuplicate, gpuClustering::status, and x.

                                                                                                               {
       return getStatus(status, x, y) == kDuplicate;
     }

◆ promote()

__device__ constexpr void gpuClustering::pixelStatus::promote	(	uint32_t *__restrict__	status,
		const uint16_t	x,
		const uint16_t	y
	)

inline

Definition at line 46 of file gpuClustering.h.

References cms::cudacompat::atomicCAS(), getIndex(), getShift(), kDuplicate, kEmpty, kFound, mask, edm::shift, gpuClustering::status, and x.

                                                                                                                 {
       uint32_t index = getIndex(x, y);
       uint32_t shift = getShift(x, y);
       uint32_t old_word = status[index];
       uint32_t expected = old_word;
       do {
         expected = old_word;
         Status old_status{(old_word >> shift) & mask};
         if (kDuplicate == old_status) {
           // nothing to do
           return;
         }
         Status new_status = (kEmpty == old_status) ? kFound : kDuplicate;
         uint32_t new_word = old_word | (static_cast<uint32_t>(new_status) << shift);
         old_word = atomicCAS(&status[index], expected, new_word);
       } while (expected != old_word);
     }

Variable Documentation

◆ bits

constexpr uint32_t gpuClustering::pixelStatus::bits = 2

Definition at line 25 of file gpuClustering.h.

◆ mask

constexpr uint32_t gpuClustering::pixelStatus::mask = (0x01 << bits) - 1

Definition at line 26 of file gpuClustering.h.

Referenced by pat::helper::MultiIsolator::addIsolator(), l1t::TriggerSystem::addMask(), cscdqm::Detector::AddressFromString(), DetIdSelector::addSelection(), TrackingNtuple::addStripMatchedHit(), AbsElectronicODERHS::allParametersSet(), SiPixelDynamicInefficiencyDB::analyze(), L1TGlobalSummary::analyze(), L1TdeGCT::analyze(), L1TDEMON::analyze(), SiPixelBadModuleReader::analyze(), L1GtRunSettingsViewer::analyze(), l1t::GtRecordDump::analyze(), FFTJetCorrectionProducer::applyCorrections(), MuonIdProducer::arbitrateMuons(), TotemT2DQMSource::areChannelsTriggered(), hcaldqm::DQClient::beginLuminosityBlock(), GEMRecHitProducer::beginRun(), cms::cuda::HistoContainer< T, NBINS, SIZE, S, I, NHISTS >::bin(), L1GlobalCaloTrigger::build(), PFMultiDepthClusterizer::buildClusters(), RPixDetClusterizer::buildClusters(), SiPixelClusterSource::buildStructure(), SiPixelDigiSource::buildStructure(), SiPixelRawDataErrorSource::buildStructure(), SiPixelRecHitSource::buildStructure(), PixelLumiDQM::calculateBunchMask(), VFATFrame::calculateCRC(), L1TMuonBarrelKalmanAlgo::calculateEta(), L1TMuonBarrelKalmanAlgo::chain(), L1GtConditionEvaluation::checkBit(), l1t::ConditionEvaluation::checkBit(), EcalRecHit::checkFlagMask(), GlobalTrackingRegion::checkTracks(), CosmicTrackingRegion::checkTracks(), RectangularEtaPhiTrackingRegion::checkTracks(), ECALPFSeedCleaner::clean(), FlagsCleanerECAL::clean(), RBXAndHPDCleaner::clean(), SpikeAndDoubleSpikeCleaner::clean(), MuonTrajectoryCleaner::clean(), Phase2L1GMT::TrackMuonMatchAlgorithm::clean(), Phase2L1GMT::TrackMuonMatchAlgorithm::cleanNeighbor(), PhiMemoryImage::clear_bit(), TrackingNtuple::clearVariables(), L1TMuonBarrelKalmanLUTs::coarseEta(), magfieldparam::poly2d_base::Collect(), L1TMuonBarrelParamsHelper::configFromDB(), L1TMuonBarrelParamsHelper::configFromPy(), l1t::MTF7Payload::count(), GEMVFAT::crc_cal(), SiPixelFEDChannelContainerFromQualityConverter::createFromSiPixelQuality(), FastSiPixelFEDChannelContainerFromQuality::createFromSiPixelQuality(), DEutils< T >::de_equal(), L1MuGMTCancelOutUnit::decide(), DEutils< T >::DEDigi(), FFTJetProducer::determineGriddedConstituents(), FFTJetProducer::determineVectorConstituents(), SiStripFineDelayHit::detId(), GEMClusterizer::doAction(), PulseChiSqSNNLS::DoFit(), hcaldqm::DQHarvester::dqmBeginLuminosityBlock(), EcalSeverityLevelAlgo::EcalSeverityLevelAlgo(), L1TGlobalSummary::endRun(), trklet::TrackletProcessor::execute(), PFMultiDepthClusterizer::expandCluster(), Phase2Tracker::extract64(), TFParams::f3deg(), FFTGenericScaleCalculator::FFTGenericScaleCalculator(), HGCROCChannelDataFrame< D >::fillPacket(), TrackingNtuple::fillPixelHits(), TrackingNtuple::fillStripRphiStereoHits(), l1t::Stage2Layer2ClusterAlgorithmFirmwareImp1::filtering(), LocalMaximumSeedFinder::findSeeds(), ticl::LinkingAlgoByDirectionGeometric::findTrackstersInWindow(), tmtt::Sector::forceBitWidth(), TSFit::fpol3dg(), CPPFMaskReClusterizer::get(), RPCMaskReClusterizer::get(), GEMMaskReClusterizer::get(), VFATFrame::getActiveChannels(), sistrip::fedchannelunpacker::detail::getADC_B1(), sistrip::fedchannelunpacker::detail::getADC_B2(), L1TMuonBarrelKalmanAlgo::getByCode(), vid::CutFlowResult::getCutFlowResultMasking(), DCCTBBlockPrototype::getDataWord(), L1GTTInputProducer::getEtaBits(), RPCAMCLinkEvents::getEventName(), CaloRecHitAuxSetter::getField(), TrackerTopology::getField(), RPCAMCLink::getMask(), RPCDCCLink::getMask(), ApvAnalysisFactory::getMask(), RPCLBLink::getMask(), l1t::L1TGlobalUtil::getMaskByBit(), l1t::L1TGlobalUtil::getMaskByName(), HGCTriggerDetId::getMaskedId(), FWGeometry::getMatchedIds(), L1MuGMTReadoutRecord::getMIPbit(), L1GTTInputProducer::getPtBits(), L1MuGMTReadoutRecord::getQuietbit(), getStatus(), reco::HitPattern::getTrackerLayerCase(), reco::HitPattern::getTrackerMonoStereo(), HGCSample::getWord(), hcaldqm::DQTask::globalBeginLuminosityBlock(), gluedId(), L1TMuonBarrelKalmanAlgo::hitPattern(), reco::ElectronSeed::hitsMask(), L1MuPseudoSignedPacking::idxFromPacked(), GeneralTracksImporter::importToBlock(), HcalObjRepresent::IntToBinary(), l1t::Stage2Layer2TauAlgorithmFirmwareImp1::is3x3Maximum(), DEutils< T >::is_empty(), mkfit::isFinite(), edm::isFinite(), l1t::TriggerSystem::isMasked(), CmsTrackerDetIdBuilder::iterate(), L1GctEmulator::L1GctEmulator(), L1TGlobalPrescalesVetosESProducer::L1TGlobalPrescalesVetosESProducer(), L1TGlobalPrescalesVetosFractESProducer::L1TGlobalPrescalesVetosFractESProducer(), HGVHistoProducerAlgo::layerClusters_to_SimClusters(), LzmaDec_DecodeReal(), HGCGraphT< TILES >::makeAndConnectDoublets(), EcalUncalibRecHitMultiFitAlgo::makeRecHit(), HGCalWaferMask::maskCell(), HGCalDDDConstants::maskCell(), GEMDQMEfficiencySourceBase::maskChamberWithError(), RPCHalfSorter::maskHSBInput(), pat::Flags::maskToString(), edm::EventSelector::maskTriggerResults(), FWGeometry::match_id(), L1GctChannelMaskOnlineProd::newObject(), L1MuGMTChannelMaskOnlineProducer::newObject(), L1RPCHsbConfigOnlineProd::newObject(), L1TGlobalPrescalesVetosOnlineProd::newObject(), DCCEventBlock::next_tower_search(), cscdqm::Detector::NextAddress(), cscdqm::Detector::NextAddressBox(), reco::parser::test_bit_f::operator()(), PhysicsTools::VarProcessor::ConfIterator::operator()(), pos::operator<<(), operator<<(), PhysicsTools::VarProcessor::ConfIterator::operator<<(), operator>>(), TotemDAQMappingESSourceXML::ParseTreeDiamond(), CTPPSPixelDAQMappingESSourceXML::ParseTreePixel(), TotemDAQMappingESSourceXML::ParseTreeRP(), TotemDAQMappingESSourceXML::ParseTreeTotemT2(), TotemDAQMappingESSourceXML::ParseTreeTotemTiming(), CTPPSPixelDAQMappingESSourceXML::ParseXML(), TotemDAQMappingESSourceXML::ParseXML(), EgammaCutBasedEleId::PassWP(), VFATFrame::Print(), EgammaCutBasedEleId::PrintDebug(), EcalFenixMaxof2::process(), EcalFenixStripFgvbEE::process(), EcalFenixEtTot::process(), EcalFenixTcpFgvbEE::process(), EcalFenixFgvbEB::process(), RecHitProcessor::process(), RecHitProcessor::processLook(), PFClusterProducer::produce(), RPCRecHitProducer::produce(), GEMRecHitProducer::produce(), TrackSelectorByRegion::produce(), CTPPSPixelClusterProducer::produce(), CTPPSPixelLocalTrackProducer::produce(), L1MuGlobalMuonTrigger::produce(), pat::PATTriggerProducer::produce(), TotemDAQMappingESSourceXML::produce(), L1GctConfigProducers::produceChanMask(), CTPPSPixelDAQMappingESSourceXML::produceCTPPSPixelAnalysisMask(), CTPPSPixelDAQMappingESSourceXML::produceCTPPSPixelDAQMapping(), SiStripFineDelayHit::produceNoTracking(), ecaldqm::LaserClient::producePlots(), ecaldqm::LedClient::producePlots(), ecaldqm::PedestalClient::producePlots(), ecaldqm::PNIntegrityClient::producePlots(), ecaldqm::TestPulseClient::producePlots(), ecaldqm::TrigPrimClient::producePlots(), ecaldqm::PresampleClient::producePlots(), ecaldqm::OccupancyClient::producePlots(), ecaldqm::RawDataClient::producePlots(), ecaldqm::TimingClient::producePlots(), ecaldqm::SummaryClient::producePlots(), ecaldqm::IntegrityClient::producePlots(), ecaldqm::MLClient::producePlots(), promote(), PFMultiDepthClusterizer::prune(), SiPixelDynamicInefficiency::putDetIdmask(), MatacqTBRawEvent::read32(), MatacqRawEvent::read32(), L1MuRegionalCand::readDataField(), L1MuGMTCand::readDataField(), HGCROCChannelDataFrame< D >::readPacket(), L1MuGMTPSB::receiveData(), RPCRecHitBaseAlgo::reconstruct(), GEMRecHitBaseAlgo::reconstruct(), MiniFloatConverter::reduceMantissaToNbits(), PhiMemoryImage::rotl(), PhiMemoryImage::rotr(), Matriplex::round_up_align64(), RPCLogCone::RPCLogCone(), CTPPSPixelClusterProducer::run(), ecaldqm::TimingTask::runOnRecHits(), ecaldqm::OccupancyTask::runOnRecHits(), rpcrawtodigi::EventRecords::samePartition(), CSCDCCExaminer::scanbuf(), PhiMemoryImage::set_bit(), FixedFlatOccupancy1d::set_ExclusionMask(), sistrip::TrackerSpecialHeader::setAPVAddressErrorForFEUnit(), sistrip::FEDAPVErrorHeader::setAPVStatusBit(), L1GctRegion::setBit(), sistrip::FEDStatusRegister::setBit(), sistrip::FEDBackendStatusRegister::setBit(), sistrip::FEDFullDebugHeader::setBit(), L1MuGMTDebugBlock::SetCancelBits(), DTStatusFlag::setCellFEMask(), DTStatusFlag::setCellTDCMask(), DTStatusFlag::setCellTrigMask(), l1t::CaloParamsHelper::setCentralityRegionMask(), l1t::CaloParamsHelperO2O::setCentralityRegionMask(), pos::PixelFEDCard::setChannel(), L1GlobalCaloTrigger::setChannelMask(), BTLSample::setDataWord(), EcalSampleMask::setEcalSampleMaskRecordEB(), EcalSampleMask::setEcalSampleMaskRecordEE(), l1t::CaloParamsHelperO2O::setEgEtaCut(), l1t::CaloParamsHelper::setEgEtaCut(), sistrip::TrackerSpecialHeader::setFEEnableForFEUnit(), sistrip::TrackerSpecialHeader::setFEOverflowForFEUnit(), CaloRecHitAuxSetter::setField(), BTLSample::setFlagWord(), L1RPCHsbConfig::setHsbMask(), l1t::CaloParamsHelperO2O::setJetRegionMask(), l1t::CaloParamsHelper::setJetRegionMask(), FFTJetProducer::setJetStatusBit(), CmsShowMain::setLiveMode(), mathSSE::Vec4< double >::setMask(), CSCDCCExaminer::setMask(), EcalRecHit::setMasked(), HGCTriggerDetId::setMaskedId(), L1MuGMTReadoutRecord::setMIPbit(), DTConfigBti::setPTMSflag(), L1MuGMTReadoutRecord::setQuietbit(), FFTJetPFPileupCleaner::setRemovalBit(), DataAdderTableManager::setSelection(), l1t::CaloParamsHelperO2O::setTauRegionMask(), l1t::CaloParamsHelper::setTauRegionMask(), pos::PixelROCTrimBits::setTrim(), DTConfigTraco::setUsedBti(), DTConfigBti::setWENflag(), FTLSample::setWord(), ETLSample::setWord(), HGCSample::setWord(), MkFitProducer::stripClusterChargeCut(), pat::Flags::test(), PhiMemoryImage::test_bit(), EgammaCutBasedEleId::TestWP(), TrackingNtuple::TrackingNtuple(), Phase2L1GMT::twos_complement(), L1MuGMTLUT::u2vec(), sistrip::fedchannelunpacker::detail::unpackRawB(), sistrip::fedchannelunpacker::detail::unpackZSB(), L1TMuonBarrelKalmanAlgo::update(), L1TMuonBarrelKalmanAlgo::updateLUT(), HGCalTopology::valid(), HcalTriggerPrimitiveAlgo::validChannel(), cscdqm::Summary::WriteChamberState(), L1MuRegionalCand::writeDataField(), L1MuGMTCand::writeDataField(), and HcalHTRData::zsBunchMask().

◆ size

constexpr uint32_t gpuClustering::pixelStatus::size = pixelSizeX * pixelSizeY / valuesPerWord

Definition at line 28 of file gpuClustering.h.

◆ valuesPerWord

constexpr uint32_t gpuClustering::pixelStatus::valuesPerWord = sizeof(uint32_t) * 8 / bits

Definition at line 27 of file gpuClustering.h.

Referenced by getIndex(), and getShift().

Enumerations

Functions

Variables

Enumeration Type Documentation

◆ Status

Function Documentation

◆ getIndex()

◆ getShift()

◆ getStatus()

◆ isDuplicate()

◆ promote()

Variable Documentation

◆ bits

◆ mask

◆ size

◆ valuesPerWord