diff --git a/Dumpers/BinaryDumpers/CMakeLists.txt b/Dumpers/BinaryDumpers/CMakeLists.txt index f6ea56e1bacabdcdded748a25c3ca2881fff59d0..37508a25ac4f2d07fe479a665a21988d889d569c 100644 --- a/Dumpers/BinaryDumpers/CMakeLists.txt +++ b/Dumpers/BinaryDumpers/CMakeLists.txt @@ -24,7 +24,8 @@ gaudi_add_library(BinaryDumpers LHCb::LHCbDetLib LHCb::MagnetLib LHCb::MuonDAQLib - LHCb::RecEvent) + LHCb::RecEvent + LHCb::RichFutureDAQLib) gaudi_add_module(BinaryDumpersModule SOURCES @@ -45,6 +46,8 @@ gaudi_add_module(BinaryDumpersModule src/DumpUTHits.cpp src/DumpUTLookupTables.cpp src/DumpVPGeometry.cpp + src/DumpRichCableMapping.cpp + src/DumpRichPDMDBMapping.cpp src/PVDumper.cpp src/ProvideConstants.cpp src/TestMuonTable.cpp diff --git a/Dumpers/BinaryDumpers/include/Dumpers/Identifiers.h b/Dumpers/BinaryDumpers/include/Dumpers/Identifiers.h index a688507b0f40256eaafe48bd01acfb5080f6b748..e07143f5c30ecfc83e93457990105fa3f9aa57ee 100644 --- a/Dumpers/BinaryDumpers/include/Dumpers/Identifiers.h +++ b/Dumpers/BinaryDumpers/include/Dumpers/Identifiers.h @@ -87,5 +87,19 @@ namespace Allen { inline static std::string const id = "EcalGeometry"; }; + /** @class RichPDMDBMapping + * Identifier for the RICH PDMDB decode mapping for Allen + */ + struct RichPDMDBMapping : Identifier { + inline static std::string const id = "RichPDMDBMapping"; + }; + + /** @class RichCableMapping + * Identifier for the RICH cable mapping for Allen + */ + struct RichCableMapping : Identifier { + inline static std::string const id = "RichCableMapping"; + }; + } // namespace NonEventData } // namespace Allen diff --git a/Dumpers/BinaryDumpers/src/AllenUpdater.h b/Dumpers/BinaryDumpers/src/AllenUpdater.h index 7b2a5d21843d6228310d5e0ddbba904f8b117f73..7e27da48d63cdde189742b34ae32831085bc766b 100644 --- a/Dumpers/BinaryDumpers/src/AllenUpdater.h +++ b/Dumpers/BinaryDumpers/src/AllenUpdater.h @@ -83,7 +83,9 @@ public: private: Gaudi::Property m_triggerEventLoop {this, "TriggerEventLoop", false}; - Gaudi::Property> m_bankTypes {this, "BankTypes", {"VP", "UT", "FTCluster", "ECal", "Muon"}}; + Gaudi::Property> m_bankTypes {this, + "BankTypes", + {"VP", "UT", "FTCluster", "ECal", "Muon", "Rich1", "Rich2"}}; std::map< std::string, std::tuple>>> diff --git a/Dumpers/BinaryDumpers/src/DumpRichCableMapping.cpp b/Dumpers/BinaryDumpers/src/DumpRichCableMapping.cpp new file mode 100644 index 0000000000000000000000000000000000000000..98b05126be9810809a7331ae911b6afae33a4e9d --- /dev/null +++ b/Dumpers/BinaryDumpers/src/DumpRichCableMapping.cpp @@ -0,0 +1,84 @@ +/*****************************************************************************\ +* (c) Copyright 2000-2019 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ + +// Gaudi Array properties ( must be first ...) +#include "GaudiKernel/ParsersFactory.h" +#include "GaudiKernel/StdArrayAsProperty.h" + +// Rich Kernel +#include "RichFutureKernel/RichAlgBase.h" + +// Gaudi Functional +#include "LHCbAlgs/Transformer.h" + +// Rich Utils +#include "RichFutureUtils/RichDecodedData.h" +#include "RichUtils/RichException.h" +#include "RichUtils/RichHashMap.h" +#include "RichUtils/RichMap.h" +#include "RichUtils/RichSmartIDSorter.h" + +// RICH DAQ +#include "RichFutureDAQ/RichPackedFrameSizes.h" +#include "RichFutureDAQ/RichTel40CableMapping.h" + +// Dumper +#include "Dumper.h" +#include + +namespace { + struct RichCableMapping { + RichCableMapping() = default; + RichCableMapping(std::vector& data, const Rich::Future::DAQ::Tel40CableMapping& tel40Maps) + { + Rich::Future::DAQ::Allen::Tel40CableMapping allenTel40Maps {tel40Maps}; + DumpUtils::Writer output {}; + output.write(allenTel40Maps); + data = output.buffer(); + } + }; +} // namespace + +/** + * @brief Dump cable mapping for the RICH detector. + */ +class DumpRichCableMapping final + : public Allen::Dumpers::Dumper> { +public: + DumpRichCableMapping(const std::string& name, ISvcLocator* svcLoc); + + void operator()(const RichCableMapping&) const override; + + StatusCode initialize() override; + +private: + std::vector m_data; +}; + +DECLARE_COMPONENT(DumpRichCableMapping) + +DumpRichCableMapping::DumpRichCableMapping(const std::string& name, ISvcLocator* svcLoc) : + Dumper(name, svcLoc, {KeyValue {"RichCableMappingLocation", location(name, "cablemapping")}}) +{} + +StatusCode DumpRichCableMapping::initialize() +{ + return Dumper::initialize().andThen([&] { + register_producer(Allen::NonEventData::RichCableMapping::id, "rich_tel40maps", m_data); + + Rich::Future::DAQ::Tel40CableMapping::addConditionDerivation( + this, Rich::Future::DAQ::Tel40CableMapping::DefaultConditionKey); + + addConditionDerivation( + {Rich::Future::DAQ::Tel40CableMapping::DefaultConditionKey}, + inputLocation(), + [&](const Rich::Future::DAQ::Tel40CableMapping& det) { + auto cableMapping = RichCableMapping {m_data, det}; + dump(); + return cableMapping; + }); + }); +} + +void DumpRichCableMapping::operator()(const RichCableMapping&) const {} diff --git a/Dumpers/BinaryDumpers/src/DumpRichPDMDBMapping.cpp b/Dumpers/BinaryDumpers/src/DumpRichPDMDBMapping.cpp new file mode 100644 index 0000000000000000000000000000000000000000..516cd8c973b30b08b2b428eab407c3a01949d9d8 --- /dev/null +++ b/Dumpers/BinaryDumpers/src/DumpRichPDMDBMapping.cpp @@ -0,0 +1,83 @@ +/*****************************************************************************\ +* (c) Copyright 2000-2019 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ + +// Gaudi Array properties ( must be first ...) +#include "GaudiKernel/ParsersFactory.h" +#include "GaudiKernel/StdArrayAsProperty.h" + +// Rich Kernel +#include "RichFutureKernel/RichAlgBase.h" + +// Gaudi Functional +#include "LHCbAlgs/Transformer.h" + +// Rich Utils +#include "RichFutureUtils/RichDecodedData.h" +#include "RichUtils/RichException.h" +#include "RichUtils/RichHashMap.h" +#include "RichUtils/RichMap.h" +#include "RichUtils/RichSmartIDSorter.h" + +// RICH DAQ +#include "RichFutureDAQ/RichPDMDBDecodeMapping.h" +#include "RichFutureDAQ/RichPackedFrameSizes.h" + +// Dumper +#include "Dumper.h" +#include + +namespace { + struct RichPDMDBMapping { + RichPDMDBMapping() = default; + RichPDMDBMapping(std::vector& data, const Rich::Future::DAQ::PDMDBDecodeMapping& tel40Maps) + { + Rich::Future::DAQ::Allen::PDMDBDecodeMapping allenTel40Maps {tel40Maps}; + DumpUtils::Writer output {}; + output.write(allenTel40Maps); + data = output.buffer(); + } + }; +} // namespace + +/** + * @brief Dump cable mapping for the RICH detector. + */ +class DumpRichPDMDBMapping final + : public Allen::Dumpers::Dumper> { +public: + DumpRichPDMDBMapping(const std::string& name, ISvcLocator* svcLoc); + + void operator()(const RichPDMDBMapping&) const override; + + StatusCode initialize() override; + +private: + std::vector m_data; +}; + +DECLARE_COMPONENT(DumpRichPDMDBMapping) + +DumpRichPDMDBMapping::DumpRichPDMDBMapping(const std::string& name, ISvcLocator* svcLoc) : + Dumper(name, svcLoc, {KeyValue {"RichPDMDBMappingLocation", location(name, "pdmdbmapping")}}) +{} + +StatusCode DumpRichPDMDBMapping::initialize() +{ + return Dumper::initialize().andThen([&, this] { + register_producer(Allen::NonEventData::RichPDMDBMapping::id, "rich_pdmdbmaps", m_data); + Rich::Future::DAQ::PDMDBDecodeMapping::addConditionDerivation( + this, Rich::Future::DAQ::PDMDBDecodeMapping::DefaultConditionKey); + + addConditionDerivation( + {Rich::Future::DAQ::PDMDBDecodeMapping::DefaultConditionKey}, + inputLocation(), + [&](const Rich::Future::DAQ::PDMDBDecodeMapping& det) { + RichPDMDBMapping mapping {m_data, det}; + dump(); + return mapping; + }); + }); +} + +void DumpRichPDMDBMapping::operator()(const RichPDMDBMapping&) const {} diff --git a/Dumpers/BinaryDumpers/src/TransposeRawBanks.cpp b/Dumpers/BinaryDumpers/src/TransposeRawBanks.cpp index a4b88eb10308d86bef98f19d95b5c7056ad339d8..9d179a530e7b8af832c260284951ac4a220d1fa0 100644 --- a/Dumpers/BinaryDumpers/src/TransposeRawBanks.cpp +++ b/Dumpers/BinaryDumpers/src/TransposeRawBanks.cpp @@ -72,7 +72,8 @@ private: LHCb::RawBank::Calo, LHCb::RawBank::Muon, LHCb::RawBank::ODIN, - LHCb::RawBank::Plume}}; + LHCb::RawBank::Plume, + LHCb::RawBank::Rich}}; std::array m_histos; }; diff --git a/Rec/Allen/CMakeLists.txt b/Rec/Allen/CMakeLists.txt index 45fb5a1ac53846c1ebe847ffbd6e6a6db37dd3c4..9cc472ea35d25aefccd81c0b093c2b5bfa84a4a4 100755 --- a/Rec/Allen/CMakeLists.txt +++ b/Rec/Allen/CMakeLists.txt @@ -29,6 +29,7 @@ gaudi_add_module(AllenWrapper src/CompareRecAllenFTClusters.cpp src/CompareRecAllenVPHits.cpp src/TestVeloClusters.cpp + src/TestAllenRichPixels.cpp LINK AllenLib Backend @@ -51,6 +52,7 @@ gaudi_add_module(AllenWrapper MuonCommon EventModel Gear + Rich Gaudi::GaudiAlgLib LHCb::DAQEventLib LHCb::RecEvent @@ -60,7 +62,7 @@ gaudi_add_module(AllenWrapper LHCb::CaloDetLib LHCb::FTDetLib LHCb::MCEvent - LHCb::RecEvent + LHCb::RecEvent Rec::PrKernel) target_include_directories(AllenWrapper PRIVATE ${PROJECT_BINARY_DIR}/code_generation) @@ -98,6 +100,7 @@ gaudi_add_module(AllenAlgorithms Utils WrapperInterface mdf + Rich Gaudi::GaudiAlgLib LHCb::DAQEventLib LHCb::VPDetLib diff --git a/Rec/Allen/python/Allen/config.py b/Rec/Allen/python/Allen/config.py index 03c42d642976178e9dde8418fed1496ae3a00aa5..9c14c5934b9e54a453cf893051faeca285264774 100755 --- a/Rec/Allen/python/Allen/config.py +++ b/Rec/Allen/python/Allen/config.py @@ -20,7 +20,8 @@ from PyConf.application import configure_input, configure from PyConf.Algorithms import ( DumpBeamline, DumpCaloGeometry, DumpMagneticField, DumpVPGeometry, DumpFTGeometry, DumpUTGeometry, DumpUTLookupTables, DumpMuonGeometry, - DumpMuonTable, AllenODINProducer) + DumpMuonTable, AllenODINProducer, DumpRichPDMDBMapping, + DumpRichCableMapping) from DDDB.CheckDD4Hep import UseDD4Hep @@ -109,15 +110,25 @@ def setup_allen_non_event_data_service(allen_event_loop=False, data (geometries etc.) """ converter_types = { - 'VP': [(DumpBeamline, 'DeviceBeamline', 'beamline'), - (DumpVPGeometry, 'DeviceVPGeometry', 'velo_geometry')], - 'UT': [(DumpUTGeometry, 'DeviceUTGeometry', 'ut_geometry'), - (DumpUTLookupTables, 'DeviceUTLookupTables', 'ut_tables')], - 'ECal': [(DumpCaloGeometry, 'DeviceCaloGeometry', 'ecal_geometry')], - 'Magnet': [(DumpMagneticField, 'DeviceMagneticField', 'polarity')], - 'FTCluster': [(DumpFTGeometry, 'DeviceFTGeometry', 'scifi_geometry')], - 'Muon': [(DumpMuonGeometry, 'DeviceMuonGeometry', 'muon_geometry'), - (DumpMuonTable, 'DeviceMuonTable', 'muon_tables')] + frozenset({'VP'}): [(DumpBeamline, 'DeviceBeamline', 'beamline'), + (DumpVPGeometry, 'DeviceVPGeometry', + 'velo_geometry')], + frozenset({'UT'}): + [(DumpUTGeometry, 'DeviceUTGeometry', 'ut_geometry'), + (DumpUTLookupTables, 'DeviceUTLookupTables', 'ut_tables')], + frozenset({'ECal'}): [(DumpCaloGeometry, 'DeviceCaloGeometry', + 'ecal_geometry')], + frozenset({'Magnet'}): [(DumpMagneticField, 'DeviceMagneticField', + 'polarity')], + frozenset({'FTCluster'}): [(DumpFTGeometry, 'DeviceFTGeometry', + 'scifi_geometry')], + frozenset({'Muon'}): [(DumpMuonGeometry, 'DeviceMuonGeometry', + 'muon_geometry'), + (DumpMuonTable, 'DeviceMuonTable', + 'muon_tables')], + frozenset({'Rich1', 'Rich2'}): + [(DumpRichPDMDBMapping, 'DeviceRichPDMDBMapping', 'rich_pdmdbmaps'), + (DumpRichCableMapping, 'DeviceRichCableMapping', 'rich_tel40maps')], } detector_names = { @@ -132,7 +143,9 @@ def setup_allen_non_event_data_service(allen_event_loop=False, if type(bank_types) == list: bank_types = set(bank_types) elif bank_types is None: - bank_types = set(converter_types.keys()) + bank_types = set() + for ibt in converter_types.keys(): + bank_types.update(set(ibt)) if 'VPRetinaCluster' in bank_types: bank_types.remove('VPRetinaCluster') @@ -167,8 +180,9 @@ def setup_allen_non_event_data_service(allen_event_loop=False, if allen_event_loop: algorithm_converters.append(AllenODINProducer()) - converters = [(bt, t, tn, f) for bt, convs in converter_types.items() - for t, tn, f in convs if bt in bank_types] + converters = {(bts, t, tn, f) + for bts, convs in converter_types.items() + for t, tn, f in convs if bts.intersection(bank_types)} for bt, converter_type, converter_name, filename in converters: converter = converter_type( name=converter_name, diff --git a/Rec/Allen/src/TestAllenRichPixels.cpp b/Rec/Allen/src/TestAllenRichPixels.cpp new file mode 100644 index 0000000000000000000000000000000000000000..b1c8f990df1b95421af3f963105572675092f931 --- /dev/null +++ b/Rec/Allen/src/TestAllenRichPixels.cpp @@ -0,0 +1,65 @@ +/***************************************************************************** \ + * (c) Copyright 2000-2018 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +// Gaudi +#include "GaudiAlg/Consumer.h" + +// LHCb +#include "Event/PrHits.h" +#include "RichFutureUtils/RichDecodedData.h" + +// Allen +#include "RichDecoding.cuh" +#include "RichPDMDBDecodeMapping.cuh" + +class TestAllenRichPixels final + : public Gaudi::Functional::Consumer< + void(const std::vector&, const Rich::Future::DAQ::DecodedData&)> { + +public: + /// Standard constructor + TestAllenRichPixels(const std::string& name, ISvcLocator* pSvcLocator); + + /// Algorithm execution + void operator()(const std::vector&, const Rich::Future::DAQ::DecodedData&) const override; +}; + +DECLARE_COMPONENT(TestAllenRichPixels) + +TestAllenRichPixels::TestAllenRichPixels(const std::string& name, ISvcLocator* pSvcLocator) : + Consumer( + name, + pSvcLocator, + {KeyValue {"rich_smart_ids", ""}, KeyValue {"RichDecodedData", Rich::Future::DAQ::DecodedDataLocation::Default}}) +{} + +void TestAllenRichPixels::operator()( + const std::vector& allen_rich_smart_ids, + const Rich::Future::DAQ::DecodedData& rec_rich_pixels) const +{ + std::vector recIDs; + for (const auto& rD : rec_rich_pixels) { + for (const auto& pD : rD) { + for (const auto& mD : pD) { + for (const auto& pd : mD) { + const auto& IDs = pd.smartIDs(); + recIDs.insert(recIDs.end(), IDs.begin(), IDs.end()); + } + } + } + } + + if (recIDs.size() != allen_rich_smart_ids.size()) { + error() << "Allen and Rec Rich Smart ID containers are not the same size" << endmsg; + } + + for (const auto smart_id : recIDs) { + const auto allen_smart_id = Allen::RichSmartID {smart_id.key()}; + if ( + smart_id.pixelDataAreValid() && + std::find(allen_rich_smart_ids.begin(), allen_rich_smart_ids.end(), allen_smart_id) == + allen_rich_smart_ids.end()) { + error() << "ID " << allen_smart_id << " not present in Allen Rich Smart ID container" << endmsg; + } + } +} diff --git a/configuration/python/AllenConf/HLT1.py b/configuration/python/AllenConf/HLT1.py index d2052e3c0c9796d3d8309e00b2ed91b36909c5e3..7e77d307c0e935689ec0ae6263ae007729a71bfb 100644 --- a/configuration/python/AllenConf/HLT1.py +++ b/configuration/python/AllenConf/HLT1.py @@ -630,6 +630,7 @@ def setup_hlt1_node(enablePhysics=True, with_calo=True, with_muon=True, with_v0s=True, + with_rich=False, enableBGI=False, velo_open=False, enableDownstream=False, @@ -645,7 +646,8 @@ def setup_hlt1_node(enablePhysics=True, with_muon=with_muon, enableDownstream=enableDownstream, tracking_type=tracking_type, - velo_open=velo_open) + velo_open=velo_open, + with_rich=with_rich) hlt1_config['reconstruction'] = reconstructed_objects @@ -852,6 +854,15 @@ def setup_hlt1_node(enablePhysics=True, NodeLogic.NONLAZY_AND, force_order=False) + if with_rich: + hlt1_node = CompositeNode( + "AllenWithRich", [ + hlt1_node, + reconstructed_objects["decoded_rich"]["dev_smart_ids"].producer + ], + NodeLogic.NONLAZY_AND, + force_order=True) + if enableRateValidator: hlt1_node = CompositeNode( "AllenRateValidation", [ diff --git a/configuration/python/AllenConf/hlt1_reconstruction.py b/configuration/python/AllenConf/hlt1_reconstruction.py index e41d039b6e06f65c353334de4e0984cf786d9699..04915d8f5421f7b90b4a6018fa71552e37b3c4b3 100644 --- a/configuration/python/AllenConf/hlt1_reconstruction.py +++ b/configuration/python/AllenConf/hlt1_reconstruction.py @@ -28,7 +28,8 @@ def hlt1_reconstruction(algorithm_name='', with_ut=True, with_muon=True, velo_open=False, - enableDownstream=False): + enableDownstream=False, + with_rich=False): decoded_velo = decode_velo() decoded_scifi = decode_scifi() velo_tracks = make_velo_tracks(decoded_velo) @@ -208,6 +209,10 @@ def hlt1_reconstruction(algorithm_name='', "v0_pairs": v0_pairs }) + if with_rich: + from AllenConf.rich_reconstruction import decode_rich + output.update({"decoded_rich": decode_rich()}) + return output diff --git a/configuration/python/AllenConf/rich_reconstruction.py b/configuration/python/AllenConf/rich_reconstruction.py new file mode 100644 index 0000000000000000000000000000000000000000..a5474f82a656f4b24fa00da1724007a8f3337086 --- /dev/null +++ b/configuration/python/AllenConf/rich_reconstruction.py @@ -0,0 +1,29 @@ +############################################################################### +# (c) Copyright 2021 CERN for the benefit of the LHCb Collaboration # +############################################################################### +from AllenCore.algorithms import (data_provider_t, rich_decoding_t) +from AllenConf.utils import initialize_number_of_events +from AllenCore.generator import make_algorithm + + +def decode_rich(): + number_of_events = initialize_number_of_events() + + rich_banks = make_algorithm( + data_provider_t, name=f"rich_banks", bank_type="Rich1") + + rich_decoding = make_algorithm( + rich_decoding_t, + name=f"rich_decoding", + host_number_of_events_t=number_of_events["host_number_of_events"], + host_raw_bank_version_t=rich_banks.host_raw_bank_version_t, + dev_rich_raw_input_t=rich_banks.dev_raw_banks_t, + dev_rich_raw_input_offsets_t=rich_banks.dev_raw_offsets_t, + dev_rich_raw_input_sizes_t=rich_banks.dev_raw_sizes_t, + dev_rich_raw_input_types_t=rich_banks.dev_raw_types_t, + block_dim_x=128) + + return { + "dev_smart_ids": rich_decoding.dev_smart_ids_t, + "dev_rich_hit_offsets": rich_decoding.dev_rich_hit_offsets_t + } diff --git a/configuration/python/AllenCore/gaudi_allen_generator.py b/configuration/python/AllenCore/gaudi_allen_generator.py index 35ecf2430411038041152ba1b6d0c5d952a1879d..a5bc08b467296f6d2a391b777eade285bd289693 100644 --- a/configuration/python/AllenCore/gaudi_allen_generator.py +++ b/configuration/python/AllenCore/gaudi_allen_generator.py @@ -66,6 +66,8 @@ def make_algorithm(algorithm, name, *args, **kwargs): rawbank_list = ["Calo", "EcalPacked"] elif bank_type == "VP": rawbank_list = ["VP", "VPRetinaCluster"] + elif "Rich" in bank_type: + rawbank_list = ["Rich"] elif bank_type: rawbank_list = [bank_type] diff --git a/configuration/python/AllenSequences/hlt1_pp_rich_no_ut.py b/configuration/python/AllenSequences/hlt1_pp_rich_no_ut.py new file mode 100644 index 0000000000000000000000000000000000000000..a5d3cdd7f30e35fdfe8bfbe31724059349f1cf99 --- /dev/null +++ b/configuration/python/AllenSequences/hlt1_pp_rich_no_ut.py @@ -0,0 +1,10 @@ +############################################################################### +# (c) Copyright 2021 CERN for the benefit of the LHCb Collaboration # +############################################################################### +from AllenConf.HLT1 import setup_hlt1_node +from AllenCore.generator import generate +from AllenConf.rich_reconstruction import decode_rich +from PyConf.control_flow import NodeLogic, CompositeNode + +hlt1_node = setup_hlt1_node(with_ut=False, with_rich=True) +generate(hlt1_node) diff --git a/configuration/python/AllenSequences/rich.py b/configuration/python/AllenSequences/rich.py new file mode 100644 index 0000000000000000000000000000000000000000..a393e3b270c54c2c7abe1dd6ed41abe3f314c2b7 --- /dev/null +++ b/configuration/python/AllenSequences/rich.py @@ -0,0 +1,13 @@ +############################################################################### +# (c) Copyright 2021 CERN for the benefit of the LHCb Collaboration # +############################################################################### +from AllenConf.rich_reconstruction import decode_rich +from AllenCore.generator import generate +from PyConf.control_flow import NodeLogic, CompositeNode + +rich_decoding = CompositeNode( + "RichDecoding", [decode_rich()["dev_smart_ids"].producer], + NodeLogic.NONLAZY_AND, + force_order=True) + +generate(rich_decoding) diff --git a/device/CMakeLists.txt b/device/CMakeLists.txt index 85d8c0e5db0c7b52b5d41b7e63dd34fb77b4557c..e70e3d38ab3dc7d0943ddbc2d69f02e5e036cf2e 100644 --- a/device/CMakeLists.txt +++ b/device/CMakeLists.txt @@ -22,3 +22,4 @@ add_subdirectory(lumi) add_subdirectory(combiners) add_subdirectory(plume) add_subdirectory(downstream) +add_subdirectory(rich) diff --git a/device/rich/CMakeLists.txt b/device/rich/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..49e836d48a6ef7ff3be589d485f97dfaf13c79fa --- /dev/null +++ b/device/rich/CMakeLists.txt @@ -0,0 +1,13 @@ +############################################################################### +# (c) Copyright 2018-2020 CERN for the benefit of the LHCb Collaboration # +############################################################################### +file(GLOB rich_decoding "decoding/src/*cu") + +allen_add_device_library(Rich STATIC + ${rich_decoding} +) + +target_link_libraries(Rich PRIVATE Backend HostEventModel EventModel Utils HostPrefixSum) + +target_include_directories(Rich PUBLIC + $) diff --git a/device/rich/decoding/include/RichDecoding.cuh b/device/rich/decoding/include/RichDecoding.cuh new file mode 100644 index 0000000000000000000000000000000000000000..67c60dd7d25d29522ff1daa30716f7755d7a4899 --- /dev/null +++ b/device/rich/decoding/include/RichDecoding.cuh @@ -0,0 +1,38 @@ +/*****************************************************************************\ +* (c) Copyright 2018-2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include + +namespace rich_decoding { + struct Parameters { + HOST_INPUT(host_number_of_events_t, uint) host_number_of_events; + HOST_INPUT(host_raw_bank_version_t, int) host_raw_bank_version; + MASK_INPUT(dev_event_list_t) dev_event_list; + DEVICE_INPUT(dev_rich_raw_input_t, char) dev_rich_raw_input; + DEVICE_INPUT(dev_rich_raw_input_offsets_t, uint) dev_rich_raw_input_offsets; + DEVICE_INPUT(dev_rich_raw_input_sizes_t, uint) dev_rich_raw_input_sizes; + DEVICE_INPUT(dev_rich_raw_input_types_t, uint) dev_rich_raw_input_types; + DEVICE_OUTPUT(dev_rich_number_of_hits_t, unsigned) dev_rich_number_of_hits; + DEVICE_OUTPUT(dev_rich_hit_offsets_t, unsigned) dev_rich_hit_offsets; + HOST_OUTPUT(host_rich_hit_offsets_t, unsigned) host_rich_hit_offsets; + HOST_OUTPUT(host_rich_total_number_of_hits_t, unsigned) host_rich_total_number_of_hits; + DEVICE_OUTPUT(dev_smart_ids_t, Allen::RichSmartID) dev_smart_ids; + PROPERTY(block_dim_x_t, "block_dim_x", "block dimension x", unsigned) block_dim_x; + }; + + struct rich_decoding_t : public DeviceAlgorithm, Parameters { + void set_arguments_size(ArgumentReferences, const RuntimeOptions&, const Constants&) const; + + void operator()( + const ArgumentReferences&, + const RuntimeOptions&, + const Constants&, + const Allen::Context&) const; + + private: + Property m_block_dim_x {this, 64}; + }; +} // namespace rich_decoding diff --git a/device/rich/decoding/include/RichDefinitions.cuh b/device/rich/decoding/include/RichDefinitions.cuh new file mode 100644 index 0000000000000000000000000000000000000000..6aaef71ee324aee8567a921346852968fc04e7de --- /dev/null +++ b/device/rich/decoding/include/RichDefinitions.cuh @@ -0,0 +1,159 @@ +/*****************************************************************************\ +* (c) Copyright 2023 CERN for the benefit of the LHCb Collaboration * +* * +* This software is distributed under the terms of the Apache License * +* version 2 (Apache-2.0), copied verbatim in the file "COPYING". * +* * +* In applying this licence, CERN does not waive the privileges and immunities * +* granted to it by virtue of its status as an Intergovernmental Organization * +* or submit itself to any jurisdiction. * +\*****************************************************************************/ +#pragma once + +#include +#include +#include "BackendCommon.h" + +namespace Allen { + class RichSmartID { + uint64_t m_key; + + public: + __host__ __device__ RichSmartID() = default; + + __host__ __device__ RichSmartID(uint64_t key) : m_key(key) {} + + // Number of bits for each data field in the word + static constexpr const unsigned BitsPixelCol = 3; ///< Number of bits for MaPMT pixel column + static constexpr const unsigned BitsPixelRow = 3; ///< Number of bits for MaPMT pixel row + static constexpr const unsigned BitsPDNumInMod = 4; ///< Number of bits for MaPMT 'number in module' + static constexpr const unsigned BitsPDMod = 9; ///< Number of bits for MaPMT module + static constexpr const unsigned BitsPanel = 1; ///< Number of bits for MaPMT panel + static constexpr const unsigned BitsRich = 1; ///< Number of bits for RICH detector + static constexpr const unsigned BitsPixelSubRowIsSet = 1; + static constexpr const unsigned BitsPixelColIsSet = 1; + static constexpr const unsigned BitsPixelRowIsSet = 1; + static constexpr const unsigned BitsPDIsSet = 1; + static constexpr const unsigned BitsPanelIsSet = 1; + static constexpr const unsigned BitsRichIsSet = 1; + static constexpr const unsigned BitsLargePixel = 1; + + // The shifts + static constexpr const unsigned ShiftPixelCol = 0; + static constexpr const unsigned ShiftPixelRow = ShiftPixelCol + BitsPixelCol; + static constexpr const unsigned ShiftPDNumInMod = ShiftPixelRow + BitsPixelRow; + static constexpr const unsigned ShiftPDMod = ShiftPDNumInMod + BitsPDNumInMod; + static constexpr const unsigned ShiftPanel = ShiftPDMod + BitsPDMod; + static constexpr const unsigned ShiftRich = ShiftPanel + BitsPanel; + static constexpr const unsigned ShiftPixelSubRowIsSet = ShiftRich + BitsRich; + static constexpr const unsigned ShiftPixelColIsSet = ShiftPixelSubRowIsSet + BitsPixelSubRowIsSet; + static constexpr const unsigned ShiftPixelRowIsSet = ShiftPixelColIsSet + BitsPixelColIsSet; + static constexpr const unsigned ShiftPDIsSet = ShiftPixelRowIsSet + BitsPixelRowIsSet; + static constexpr const unsigned ShiftPanelIsSet = ShiftPDIsSet + BitsPDIsSet; + static constexpr const unsigned ShiftRichIsSet = ShiftPanelIsSet + BitsPanelIsSet; + static constexpr const unsigned ShiftLargePixel = ShiftRichIsSet + BitsRichIsSet; + + // The masks + static constexpr const unsigned MaskPixelCol = (unsigned) ((1 << BitsPixelCol) - 1) << ShiftPixelCol; + static constexpr const unsigned MaskPixelRow = (unsigned) ((1 << BitsPixelRow) - 1) << ShiftPixelRow; + static constexpr const unsigned MaskPDNumInMod = (unsigned) ((1 << BitsPDNumInMod) - 1) << ShiftPDNumInMod; + static constexpr const unsigned MaskPDMod = (unsigned) ((1 << BitsPDMod) - 1) << ShiftPDMod; + static constexpr const unsigned MaskPanel = (unsigned) ((1 << BitsPanel) - 1) << ShiftPanel; + static constexpr const unsigned MaskRich = (unsigned) ((1 << BitsRich) - 1) << ShiftRich; + static constexpr const unsigned MaskPixelSubRowIsSet = (unsigned) ((1 << BitsPixelSubRowIsSet) - 1) + << ShiftPixelSubRowIsSet; + static constexpr const unsigned MaskPixelColIsSet = (unsigned) ((1 << BitsPixelColIsSet) - 1) << ShiftPixelColIsSet; + static constexpr const unsigned MaskPixelRowIsSet = (unsigned) ((1 << BitsPixelRowIsSet) - 1) << ShiftPixelRowIsSet; + static constexpr const unsigned MaskPDIsSet = (unsigned) ((1 << BitsPDIsSet) - 1) << ShiftPDIsSet; + static constexpr const unsigned MaskPanelIsSet = (unsigned) ((1 << BitsPanelIsSet) - 1) << ShiftPanelIsSet; + static constexpr const unsigned MaskRichIsSet = (unsigned) ((1 << BitsRichIsSet) - 1) << ShiftRichIsSet; + static constexpr const unsigned MaskLargePixel = (unsigned) ((1 << BitsLargePixel) - 1) << ShiftLargePixel; + + // Max Values + static constexpr const unsigned MaxPixelCol = (unsigned) (1 << BitsPixelCol) - 1; + static constexpr const unsigned MaxPixelRow = (unsigned) (1 << BitsPixelRow) - 1; + static constexpr const unsigned MaxPDNumInMod = (unsigned) (1 << BitsPDNumInMod) - 1; + static constexpr const unsigned MaxPDMod = (unsigned) (1 << BitsPDMod) - 1; + static constexpr const unsigned MaxPanel = (unsigned) (1 << BitsPanel) - 1; + static constexpr const unsigned MaxRich = (unsigned) (1 << BitsRich) - 1; + + __host__ __device__ constexpr inline void + setData(const unsigned value, const unsigned shift, const unsigned mask) noexcept + { + m_key = ((static_cast(value) << shift) & mask) | (m_key & ~mask); + } + + __host__ __device__ constexpr inline void setData( + const unsigned value, // + const unsigned shift, // + const unsigned mask, // + const unsigned okMask) noexcept + { + m_key = ((static_cast(value) << shift) & mask) | (m_key & ~mask) | okMask; + } + + __host__ __device__ constexpr inline uint64_t getData(const unsigned shift, const unsigned mask) const noexcept + { + return (m_key & mask) >> shift; + } + + __host__ __device__ constexpr inline auto key() const noexcept { return m_key; } + + __host__ __device__ constexpr inline bool operator==(const RichSmartID& other) const noexcept + { + return m_key == other.key(); + } + + /// ostream operator + __host__ friend std::ostream& operator<<(std::ostream& str, const RichSmartID& id) { return str << id.key(); } + }; // namespace RichSmartID +} // namespace Allen + +namespace Rich::Future::DAQ { + enum DetectorType : std::int8_t { + InvalidDetector = -1, ///< Unspecified Detector + Rich1 = 0, ///< RICH1 detector + Rich2 = 1, ///< RICH2 detector + Rich = 1 ///< Single RICH detector + }; + + class PackedFrameSizes final { + public: + /// Packed type + using IntType = std::uint8_t; + + private: + // Bits for each Size + static const IntType Bits0 = 4; + static const IntType Bits1 = 4; + // shifts + static const IntType Shift0 = 0; + static const IntType Shift1 = Shift0 + Bits0; + // masks + static const IntType Mask0 = (IntType)((1 << Bits0) - 1) << Shift0; + static const IntType Mask1 = (IntType)((1 << Bits1) - 1) << Shift1; + // max values + static const IntType Max0 = (1 << Bits0) - 1; + static const IntType Max1 = (1 << Bits1) - 1; + + public: + /// Contructor from a single word + __host__ __device__ explicit PackedFrameSizes(const IntType d) : m_data(d) {} + + /// Get the overall data + __host__ __device__ inline IntType data() const noexcept { return m_data; } + + /// Get first size word + __host__ __device__ inline IntType size0() const noexcept { return ((data() & Mask0) >> Shift0); } + + /// Get second size word + __host__ __device__ inline IntType size1() const noexcept { return ((data() & Mask1) >> Shift1); } + + /// Get the total size + __host__ __device__ inline auto totalSize() const noexcept { return size0() + size1(); } + + private: + /// The data word + IntType m_data {0}; + }; +} // namespace Rich::Future::DAQ \ No newline at end of file diff --git a/device/rich/decoding/include/RichPDMDBDecodeMapping.cuh b/device/rich/decoding/include/RichPDMDBDecodeMapping.cuh new file mode 100644 index 0000000000000000000000000000000000000000..0235452a8ff1a5885e0f4057724573d5eba22114 --- /dev/null +++ b/device/rich/decoding/include/RichPDMDBDecodeMapping.cuh @@ -0,0 +1,133 @@ +/*****************************************************************************\ +* (c) Copyright 2023 CERN for the benefit of the LHCb Collaboration * +* * +* This software is distributed under the terms of the Apache License * +* version 2 (Apache-2.0), copied verbatim in the file "COPYING". * +* * +* In applying this licence, CERN does not waive the privileges and immunities * +* granted to it by virtue of its status as an Intergovernmental Organization * +* or submit itself to any jurisdiction. * +\*****************************************************************************/ +#pragma once + +#include +#include + +// STL +#include +#include +#include +#include +#include +#include + +namespace Rich::Future::DAQ::Allen { + + /// Helper class for RICH PDMDB readout mapping + class PDMDBDecodeMapping final { + public: + // data types + + /// The data for each anode + class BitData final { + public: + /// The EC number (0-3) + int8_t ec; + /// The PMT number in EC + int8_t pmtInEC; + /// The Anode index (0-63) + int8_t anode; + + public: + /// Default constructor + __host__ __device__ BitData() = default; + /// Constructor from values + __host__ __device__ BitData( + const int8_t _ec, // + const int8_t _pmt, // + const int8_t _anode) : + ec(_ec), + pmtInEC(_pmt), anode(_anode) + {} + }; + + private: + // defines + + /// Max Number of frames per PDMDB + static constexpr const unsigned FramesPerPDMDB = 6; + + /// Number of PDMDBs per module + static constexpr const unsigned PDMDBPerModule = 2; + + /// Max Number of frames per PDM + static constexpr const unsigned FramesPerPDM = PDMDBPerModule * FramesPerPDMDB; + + /// Number of bits per data frame + static constexpr const unsigned BitsPerFrame = 86; + + /// Array of Bit Data structs per frame + using FrameData = std::array; + + /// Data for each PDMDB + using PDMDBData = std::array; + + /// Data for each PDM + using PDMData = std::array; + + /// R-Type Module data for each RICH + using RTypeRichData = std::array; + + private: + // methods + + /// Get the PDMDB data for given RICH, PDMDB and frame + __host__ __device__ inline const auto& getFrameData( + const int8_t rich, // + const int8_t pdmdb, // + const int8_t link, // + const bool isHType) const + { + // Note as this is called many times from the decoding, avoid runtime range checking + // in optimised builds, as once OK it should never be invalidated. + // This is though tested in the debug builds via the asserts. + if (!isHType) { + // R type PMT + return m_pdmDataR[rich][pdmdb][link]; + } + else { + return m_pdmDataH[pdmdb][link]; + } + } + + public: + /// mapping version + __host__ __device__ inline auto version() const { return m_mappingVer; } + + /// Access the initialisation state + __host__ __device__ inline auto isInitialised() const { return m_isInitialised; } + + /// Get PDMDB data for given Tel40 data + __host__ __device__ inline const auto& getFrameData(const Tel40CableMapping::Tel40LinkData& cData) const + { + const auto rich = cData.smartID.getData(::Allen::RichSmartID::ShiftRich, ::Allen::RichSmartID::MaskRich); + return getFrameData(rich, cData.pdmdbNum, cData.linkNum, cData.isHType); + } + + private: + // data + + /// R type data + RTypeRichData m_pdmDataR; // std::array + + /// H type data + PDMData m_pdmDataH; // std::array + + /// Flag to indicate initialisation status + bool m_isInitialised {false}; + + /// Mapping version + int m_mappingVer {-1}; + }; + +} // namespace Rich::Future::DAQ::Allen diff --git a/device/rich/decoding/include/RichTel40CableMapping.cuh b/device/rich/decoding/include/RichTel40CableMapping.cuh new file mode 100644 index 0000000000000000000000000000000000000000..303d6051224f8fe38cfd6c21a269ae7e3b363a81 --- /dev/null +++ b/device/rich/decoding/include/RichTel40CableMapping.cuh @@ -0,0 +1,139 @@ +/*****************************************************************************\ +* (c) Copyright 2023 CERN for the benefit of the LHCb Collaboration * +* * +* This software is distributed under the terms of the Apache License * +* version 2 (Apache-2.0), copied verbatim in the file "COPYING". * +* * +* In applying this licence, CERN does not waive the privileges and immunities * +* granted to it by virtue of its status as an Intergovernmental Organization * +* or submit itself to any jurisdiction. * +\*****************************************************************************/ +#pragma once + +#include +#include + +namespace Rich::Future::DAQ::Allen { + /// Helper class for RICH PMT data format encoding + class Tel40CableMapping final { + + public: + /// Struct for storing data for each Tel40 Link + class Tel40LinkData final { + public: + /// RICH SmartID + ::Allen::RichSmartID smartID; + /// Module Number + int32_t moduleNum; + /// Source ID + int16_t sourceID; + /// Tel40 connector + int8_t connector; + /// PDMDB number (0,1) + int8_t pdmdbNum; + /// Link number + int8_t linkNum; + /// PMT type + bool isHType {false}; + /// Is Link Active + bool isActive {false}; + + public: + /// Default constructor + Tel40LinkData() = default; + }; + + class Tel40MetaData final { + public: + uint32_t nActiveLinks; + bool hasInactiveLinks; + }; + + /// Max number of links(frames) per PDMDB + static constexpr const uint64_t MaxLinksPerPDMDB = 6; + + /// Number of PDMDBs per module + static constexpr const uint64_t PDMDBPerModule = 2; + + /// Number of Tel40 connections per MPO + static constexpr const uint64_t ConnectionsPerTel40MPO = 12; + + /// Maximum number of active Tel40 MPOs per Source ID + static constexpr const uint64_t MaxNumberMPOsPerSourceID = 2; + + /// Maximum Number of connections per Tel40 + static constexpr const uint64_t MaxConnectionsPerTel40 = MaxNumberMPOsPerSourceID * ConnectionsPerTel40MPO; + + /// Array of Tel40 for each link in a PDMDB + using PDMDBLinkData = std::array; + + /// Array of LinkData for each PDMDB in a module + using PDMDBData = std::array; + + /// Tel40 data for each Module + using ModuleTel40Data = std::array; // 300 is totalmodules + + using Tel40SourceIDs = std::array, 164>, 2>, 2>; + using Tel40SourceMetas = std::array, 2>, 2>; + + public: + // accessors + + /// Access the initialisation state + inline bool isInitialised() const { return m_isInitialised; } + + /// Access the Tel40 Link data for given channel ID + const auto& tel40Data( + const uint32_t id, // PD ID + const int8_t pdmdb, // PDMDB ID + const int8_t frame // PDMDB Frame + ) const + { + // module number + const auto modN = (id >> 6) & 0xF; // TODO: RichSmartID.h:912 + + // return tel40 data + const auto& data = m_tel40ModuleData[modN][pdmdb][frame]; + // finally return + return data; + } + + /// Access the Tel40 connection data for a given SourceID + __device__ const auto& tel40Data(const int16_t sID) const + { + const auto payload = sID & 0x3FF; + const auto side = (sID >> 10) & 0x1; + const auto rich = (sID >> 11) == 9; // Rich1 is 4, Rich2 is 9. Then, Rich1 is identified with 0, Rich2 is 1. + return m_tel40ConnData[rich][side][payload]; + } + + // TODO: Use this when the int / boolean info is needed + __device__ const auto& tel40Meta(const int16_t sID) const + { + const auto payload = sID & 0x3FF; + const auto side = (sID >> 10) & 0x1; + const auto rich = (sID >> 11) == 9; // Rich1 is 4, Rich2 is 9. Then, Rich1 is identified with 0, Rich2 is 1. + return m_tel40ConnMeta[rich][side][payload]; + } + + /// mapping version + inline auto version() const { return m_mappingVer; } + + private: + // data + + /// Tel40 connection mapping data + Tel40SourceIDs m_tel40ConnData; + Tel40SourceMetas m_tel40ConnMeta; + + /// Tel40 Module Mapping data + ModuleTel40Data m_tel40ModuleData; + + /// Flag to indicate initialisation status + bool m_isInitialised {false}; + + /// Mapping version + int m_mappingVer {-1}; + }; + +} // namespace Rich::Future::DAQ::Allen \ No newline at end of file diff --git a/device/rich/decoding/src/RichDecoding.cu b/device/rich/decoding/src/RichDecoding.cu new file mode 100644 index 0000000000000000000000000000000000000000..1b7837782d2bab3182323ac788adba52554f25aa --- /dev/null +++ b/device/rich/decoding/src/RichDecoding.cu @@ -0,0 +1,474 @@ +/*****************************************************************************\ +* (c) Copyright 2018-2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include +#include +#include +#include +#include + +INSTANTIATE_ALGORITHM(rich_decoding::rich_decoding_t) + +__device__ unsigned rich_calculate_number_of_hits_in_raw_bank( + const Allen::RawBank& bank, + const Rich::Future::DAQ::Allen::Tel40CableMapping* cable_mapping, + const Rich::Future::DAQ::Allen::PDMDBDecodeMapping* pdmdb_mapping) +{ + unsigned number_of_hits = 0; + std::array< + Rich::Future::DAQ::PackedFrameSizes::IntType, + Rich::Future::DAQ::Allen::Tel40CableMapping::MaxConnectionsPerTel40> + connSizes {}; + + auto tel40ID = bank.source_id; + const auto& connMeta = cable_mapping->tel40Meta(tel40ID); + const auto nSizeWords = connMeta.nActiveLinks; + const auto nPackedSizeW = (nSizeWords / 2) + (nSizeWords % 2); + + auto dataW = bank.data; + auto bankEnd = bank.data + bank.size; + auto iPayloadWord = 0u; + auto iWord = 0u; + + for (; iWord < nPackedSizeW && dataW != bankEnd; ++dataW, ++iWord, iPayloadWord += 2) { + // Extract the sizes from the packed word + const Rich::Future::DAQ::PackedFrameSizes sizes(*dataW); + // extract sizes for each packed value + connSizes[iPayloadWord] = sizes.size1(); + connSizes[iPayloadWord + 1] = sizes.size0(); + } + + const auto& connData = cable_mapping->tel40Data(tel40ID); + if (connMeta.hasInactiveLinks) { + for (unsigned iL = 0; iL < connData.size(); ++iL) { + if (!connData[iL].isActive) { + for (auto i = connData.size() - 1; i > iL; --i) { + connSizes[i] = connSizes[i - 1]; + } + connSizes[iL] = 0; + } + } + } + + // finally loop over payload words and decode hits + // note iterator starts from where the above header loop ended... + unsigned iLink {0}; + while (dataW != bankEnd && iLink < connSizes.size()) { + // Do we have any words to decode for this link + if (connSizes[iLink] > 0) { + // Get the Tel40 Data for this connection + const auto& cData = connData[iLink]; + + // get the PDMDB data + const auto& frameData = pdmdb_mapping->getFrameData(cData); + + // Loop over the words for this link + uint16_t iW = 0; + while (iW < connSizes[iLink] && dataW != bankEnd) { + + // check MSB for this word + const auto isNZS = (0x80 & *dataW) != 0; + + if (!isNZS) { + // ZS decoding... word is bit index + + // load the anode data for this bit + if ((unsigned) (*dataW) < frameData.size()) { + const auto& aData = frameData[*dataW]; + + // Data 'could' be invalid, e.g. radiation-induced-upsets + // so cannot make this a hard error + if (aData.ec != -1 && aData.pmtInEC != -1 && aData.anode != -1) { + number_of_hits++; + } + } + + // move to next word + ++iW; + ++dataW; + } + else { + // NZS decoding... + + // which half of the payload are we in ? + const bool firstHalf = (0 == iW && connSizes[iLink] > 5); + // Number of words to decode depends on which half of the payload we are in + const auto nNZSwords = (firstHalf ? 6 : 5); + // bit offset per half + const auto halfBitOffset = (firstHalf ? 39 : 0); + + // look forward last NZS word and read backwards to match frame bit order + for (auto iNZS = nNZSwords - 1; iNZS >= 0; --iNZS) { + + // read the NZS word + auto nzsW = *(dataW + iNZS); + // if word zero clear MSB as this is the NZS flag + if (0 == iNZS) { + nzsW &= 0x7F; + } + + // does this word hold any active bits ? + if (nzsW > 0) { + + // Bit offset for this word + const auto bitOffset = halfBitOffset + (8 * (nNZSwords - 1 - iNZS)); + + // word has data so loop over bits to extract + for (auto iLB = 0; iLB < 8; ++iLB) { + // is bit on ? + // if ( isBitOn( nzsW, iLB ) ) { + if ((nzsW & (1 << iLB)) != 0) { + + // form frame bit value + const auto bit = iLB + bitOffset; + + // load the anode data for this bit + if ((size_t)(bit) < frameData.size()) { + const auto& aData = frameData[bit]; + + // Data 'could' be invalid, e.g. radiation-induced-upsets + // so cannot make this a hard error + if (aData.ec != -1 && aData.pmtInEC != -1 && aData.anode != -1) { + number_of_hits++; + } + } + + } // bit is on + } // loop over word bits + + } // word has any data + + } // loop over all NZS words + + // Finally skip the read NZS words + iW += nNZSwords; + dataW += nNZSwords; + } + } + + } // no data for this link, so just move on + + // move to next Tel40 link + ++iLink; + } // data word loop + + return number_of_hits; +} + +template +__global__ void rich_calculate_number_of_hits( + rich_decoding::Parameters parameters, + const unsigned event_start, + const Rich::Future::DAQ::Allen::Tel40CableMapping* cable_mapping, + const Rich::Future::DAQ::Allen::PDMDBDecodeMapping* pdmdb_mapping) +{ + const auto event_number = parameters.dev_event_list[blockIdx.x]; + + // Read raw event + const auto raw_event = Allen::RawEvent {parameters.dev_rich_raw_input, + parameters.dev_rich_raw_input_offsets, + parameters.dev_rich_raw_input_sizes, + parameters.dev_rich_raw_input_types, + event_number + event_start}; + + for (unsigned bank_number = threadIdx.x; bank_number < raw_event.number_of_raw_banks; bank_number += blockDim.x) { + const auto bank = raw_event.raw_bank(bank_number); + const auto number_of_hits_in_raw_bank = + rich_calculate_number_of_hits_in_raw_bank(bank, cable_mapping, pdmdb_mapping); + if (number_of_hits_in_raw_bank > 0) { + atomicAdd(parameters.dev_rich_number_of_hits + event_number, number_of_hits_in_raw_bank); + } + } +} + +__device__ void rich_decode_bank( + const Allen::RawBank& bank, + const Rich::Future::DAQ::Allen::Tel40CableMapping* cable_mapping, + const Rich::Future::DAQ::Allen::PDMDBDecodeMapping* pdmdb_mapping, + unsigned* event_inserted_hits, + Allen::RichSmartID* event_smart_ids) +{ + std::array< + Rich::Future::DAQ::PackedFrameSizes::IntType, + Rich::Future::DAQ::Allen::Tel40CableMapping::MaxConnectionsPerTel40> + connSizes {}; + + auto tel40ID = bank.source_id; + const auto& connMeta = cable_mapping->tel40Meta(tel40ID); + const auto nSizeWords = connMeta.nActiveLinks; + const auto nPackedSizeW = (nSizeWords / 2) + (nSizeWords % 2); + + auto dataW = bank.data; + auto bankEnd = bank.data + bank.size; + auto iPayloadWord = 0u; + auto iWord = 0u; + + for (; iWord < nPackedSizeW && dataW != bankEnd; ++dataW, ++iWord, iPayloadWord += 2) { + // Extract the sizes from the packed word + const Rich::Future::DAQ::PackedFrameSizes sizes(*dataW); + // extract sizes for each packed value + connSizes[iPayloadWord] = sizes.size1(); + connSizes[iPayloadWord + 1] = sizes.size0(); + } + + const auto& connData = cable_mapping->tel40Data(tel40ID); + if (connMeta.hasInactiveLinks) { + for (unsigned iL = 0; iL < connData.size(); ++iL) { + if (!connData[iL].isActive) { + for (auto i = connData.size() - 1; i > iL; --i) { + connSizes[i] = connSizes[i - 1]; + } + connSizes[iL] = 0; + } + } + } + + // finally loop over payload words and decode hits + // note iterator starts from where the above header loop ended... + unsigned iLink {0}; + while (dataW != bankEnd && iLink < connSizes.size()) { + // Do we have any words to decode for this link + if (connSizes[iLink] > 0) { + // Get the Tel40 Data for this connection + const auto& cData = connData[iLink]; + + // get the PDMDB data + const auto& frameData = pdmdb_mapping->getFrameData(cData); + + // Loop over the words for this link + uint16_t iW = 0; + while (iW < connSizes[iLink] && dataW != bankEnd) { + + // check MSB for this word + const auto isNZS = (0x80 & *dataW) != 0; + + if (!isNZS) { + // ZS decoding... word is bit index + + // load the anode data for this bit + if ((unsigned) (*dataW) < frameData.size()) { + + const auto& aData = frameData[*dataW]; + + // Data 'could' be invalid, e.g. radiation-induced-upsets + // so cannot make this a hard error + if (aData.ec != -1 && aData.pmtInEC != -1 && aData.anode != -1) { + // make a smart ID + auto hitID = Allen::RichSmartID {cData.smartID}; // sets RICH, side, module and PMT type + + // Add the PMT and pixel info + const auto nInMod = (aData.ec * (0 != ((hitID.key() >> 27) & 0x1) ? 1 : 4)) + aData.pmtInEC; + hitID.setData( + cData.moduleNum, + Allen::RichSmartID::ShiftPDMod, + Allen::RichSmartID::MaskPDMod, + Allen::RichSmartID::MaskPDIsSet); + hitID.setData(nInMod, Allen::RichSmartID::ShiftPDNumInMod, Allen::RichSmartID::MaskPDNumInMod); + + const auto row = aData.anode / 8; + const auto col = 8 - 1 - (aData.anode % 8); + hitID.setData( + row, + Allen::RichSmartID::ShiftPixelRow, + Allen::RichSmartID::MaskPixelRow, + Allen::RichSmartID::MaskPixelRowIsSet); + hitID.setData( + col, + Allen::RichSmartID::ShiftPixelCol, + Allen::RichSmartID::MaskPixelCol, + Allen::RichSmartID::MaskPixelColIsSet); + + const auto insert_index = atomicAdd(event_inserted_hits, 1); + event_smart_ids[insert_index] = hitID; + } + } + + // move to next word + ++iW; + ++dataW; + } + else { + // NZS decoding... + + // which half of the payload are we in ? + const bool firstHalf = (0 == iW && connSizes[iLink] > 5); + // Number of words to decode depends on which half of the payload we are in + const auto nNZSwords = (firstHalf ? 6 : 5); + // bit offset per half + const auto halfBitOffset = (firstHalf ? 39 : 0); + + // look forward last NZS word and read backwards to match frame bit order + for (auto iNZS = nNZSwords - 1; iNZS >= 0; --iNZS) { + + // read the NZS word + auto nzsW = *(dataW + iNZS); + // if word zero clear MSB as this is the NZS flag + if (0 == iNZS) { + nzsW &= 0x7F; + } + + // does this word hold any active bits ? + if (nzsW > 0) { + + // Bit offset for this word + const auto bitOffset = halfBitOffset + (8 * (nNZSwords - 1 - iNZS)); + + // word has data so loop over bits to extract + for (auto iLB = 0; iLB < 8; ++iLB) { + // is bit on ? + // if ( isBitOn( nzsW, iLB ) ) { + if ((nzsW & (1 << iLB)) != 0) { + + // form frame bit value + const auto bit = iLB + bitOffset; + + // load the anode data for this bit + if ((size_t)(bit) < frameData.size()) { + const auto& aData = frameData[bit]; + // Data 'could' be invalid, e.g. radiation-induced-upsets + // so cannot make this a hard error + if (aData.ec != -1 && aData.pmtInEC != -1 && aData.anode != -1) { + + // make a smart ID + auto hitID = Allen::RichSmartID {cData.smartID}; // sets RICH, side, module and PMT type + + // Add the PMT and pixel info + const auto nInMod = (aData.ec * (0 != ((hitID.key() >> 27) & 0x1) ? 1 : 4)) + aData.pmtInEC; + hitID.setData( + cData.moduleNum, + Allen::RichSmartID::ShiftPDMod, + Allen::RichSmartID::MaskPDMod, + Allen::RichSmartID::MaskPDIsSet); + hitID.setData(nInMod, Allen::RichSmartID::ShiftPDNumInMod, Allen::RichSmartID::MaskPDNumInMod); + + const auto row = aData.anode / 8; + const auto col = 8 - 1 - (aData.anode % 8); + hitID.setData( + row, + Allen::RichSmartID::ShiftPixelRow, + Allen::RichSmartID::MaskPixelRow, + Allen::RichSmartID::MaskPixelRowIsSet); + hitID.setData( + col, + Allen::RichSmartID::ShiftPixelCol, + Allen::RichSmartID::MaskPixelCol, + Allen::RichSmartID::MaskPixelColIsSet); + + const auto insert_index = atomicAdd(event_inserted_hits, 1); + event_smart_ids[insert_index] = hitID; + } + } + + } // bit is on + } // loop over word bits + + } // word has any data + + } // loop over all NZS words + + // Finally skip the read NZS words + iW += nNZSwords; + dataW += nNZSwords; + } + } + + } // no data for this link, so just move on + + // move to next Tel40 link + ++iLink; + } // data word loop +} + +template +__global__ void rich_decoding_kernel( + rich_decoding::Parameters parameters, + const unsigned event_start, + const Rich::Future::DAQ::Allen::Tel40CableMapping* cable_mapping, + const Rich::Future::DAQ::Allen::PDMDBDecodeMapping* pdmdb_mapping, + unsigned* dev_rich_number_of_inserted_hits) +{ + const auto event_number = parameters.dev_event_list[blockIdx.x]; + + auto event_inserted_hits = dev_rich_number_of_inserted_hits + event_number; + auto event_smart_ids = parameters.dev_smart_ids + parameters.dev_rich_hit_offsets[event_number]; + + // Read raw event + const auto raw_event = Allen::RawEvent {parameters.dev_rich_raw_input, + parameters.dev_rich_raw_input_offsets, + parameters.dev_rich_raw_input_sizes, + parameters.dev_rich_raw_input_types, + event_number + event_start}; + + for (unsigned bank_number = threadIdx.x; bank_number < raw_event.number_of_raw_banks; bank_number += blockDim.x) { + const auto bank = raw_event.raw_bank(bank_number); + rich_decode_bank(bank, cable_mapping, pdmdb_mapping, event_inserted_hits, event_smart_ids); + } +} + +void rich_decoding::rich_decoding_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&) const +{ + set_size(arguments, size(arguments)); + set_size(arguments, size(arguments) + 1); + set_size(arguments, size(arguments) + 1); + set_size(arguments, 1); +} + +void rich_decoding::rich_decoding_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + const Allen::Context& context) const +{ + const auto bank_version = first(arguments); + if (bank_version != 10) { + throw StrException("Rich bank version not supported (" + std::to_string(bank_version) + ")"); + } + + auto cable_mapping = reinterpret_cast(constants.dev_rich_cable_mapping); + auto pdmdb_mapping = + reinterpret_cast(constants.dev_rich_pdmdb_mapping); + + // Calculate number of hits into dev_rich_number_of_hits_t + Allen::memset_async(arguments, 0, context); + global_function( + runtime_options.mep_layout ? rich_calculate_number_of_hits : rich_calculate_number_of_hits)( + dim3(size(arguments)), property(), context)( + arguments, std::get<0>(runtime_options.event_interval), cable_mapping, pdmdb_mapping); + + // Copy to host + data(arguments)[0] = 0; + Allen::copy( + arguments, context, size(arguments), 1, 0); + + // Prefix sum + host_prefix_sum::host_prefix_sum_impl( + data(arguments), + size(arguments), + data(arguments)); + + // Copy prefix summed container to device + Allen::copy_async(arguments, context); + + // Decode RICH hits + auto dev_rich_number_of_inserted_hits = make_device_buffer(arguments, size(arguments)); + Allen::memset_async( + dev_rich_number_of_inserted_hits.data(), 0, dev_rich_number_of_inserted_hits.size_bytes(), context); + resize(arguments, first(arguments)); + + global_function(runtime_options.mep_layout ? rich_decoding_kernel : rich_decoding_kernel)( + dim3(size(arguments)), property(), context)( + arguments, + std::get<0>(runtime_options.event_interval), + cable_mapping, + pdmdb_mapping, + dev_rich_number_of_inserted_hits.data()); + + if (property() >= logger::debug) { + // Print output + print(arguments); + print(arguments); + } +} diff --git a/input/detector_configuration/rich_pdmdbmaps.bin b/input/detector_configuration/rich_pdmdbmaps.bin new file mode 100644 index 0000000000000000000000000000000000000000..f13aba7b331e39b34e1a128b67a4571097f15285 Binary files /dev/null and b/input/detector_configuration/rich_pdmdbmaps.bin differ diff --git a/input/detector_configuration/rich_tel40maps.bin b/input/detector_configuration/rich_tel40maps.bin new file mode 100644 index 0000000000000000000000000000000000000000..467738102542693be65565d1726d0b849c91803f Binary files /dev/null and b/input/detector_configuration/rich_tel40maps.bin differ diff --git a/integration/non_event_data/src/Updater.cpp b/integration/non_event_data/src/Updater.cpp index d5f28bdf69128b2e2acf33f4e73ce78631425c36..3ad84319ac1c7fa4b942596f5d48a8e43fbf30e6 100644 --- a/integration/non_event_data/src/Updater.cpp +++ b/integration/non_event_data/src/Updater.cpp @@ -55,7 +55,9 @@ namespace Allen { tuple {NonEventData::SciFiGeometry {}, std::string("scifi_geometry.bin")}, tuple {NonEventData::ECalGeometry {}, std::string("ecal_geometry.bin")}, tuple {NonEventData::MuonGeometry {}, std::string("muon_geometry.bin")}, - tuple {NonEventData::MuonLookupTables {}, std::string("muon_tables.bin")}}; + tuple {NonEventData::MuonLookupTables {}, std::string("muon_tables.bin")}, + tuple {NonEventData::RichPDMDBMapping {}, std::string("rich_pdmdbmaps.bin")}, + tuple {NonEventData::RichCableMapping {}, std::string("rich_tel40maps.bin")}}; for_each(producers, [this, &geometry_producer](const auto& p) { using id_t = typename std::remove_reference_t(p))>; diff --git a/main/include/BankMapping.h b/main/include/BankMapping.h index 912e2d64ab4ce403d40d6d5697211a927e7523f8..b41f173edaedb436ea1263830a35fb6af23d40ff 100644 --- a/main/include/BankMapping.h +++ b/main/include/BankMapping.h @@ -10,20 +10,21 @@ #include "BankTypes.h" namespace Allen { - const std::unordered_map bank_mapping = { - {LHCb::RawBank::VP, BankTypes::VP}, - {LHCb::RawBank::Velo, BankTypes::VP}, - {LHCb::RawBank::VPRetinaCluster, BankTypes::VP}, - {LHCb::RawBank::UT, BankTypes::UT}, - {LHCb::RawBank::FTCluster, BankTypes::FT}, - {LHCb::RawBank::Muon, BankTypes::MUON}, - {LHCb::RawBank::ODIN, BankTypes::ODIN}, - {LHCb::RawBank::HcalPacked, BankTypes::HCal}, - {LHCb::RawBank::EcalPacked, BankTypes::ECal}, - {LHCb::RawBank::Calo, BankTypes::ECal}, - {LHCb::RawBank::OTError, BankTypes::MCVertices}, // used for PV MC info - {LHCb::RawBank::OTRaw, BankTypes::MCTracks}, - {LHCb::RawBank::Plume, BankTypes::Plume}}; // used for track MC info + const std::unordered_map> bank_mapping = { + {LHCb::RawBank::VP, {BankTypes::VP}}, + {LHCb::RawBank::Velo, {BankTypes::VP}}, + {LHCb::RawBank::VPRetinaCluster, {BankTypes::VP}}, + {LHCb::RawBank::UT, {BankTypes::UT}}, + {LHCb::RawBank::FTCluster, {BankTypes::FT}}, + {LHCb::RawBank::Muon, {BankTypes::MUON}}, + {LHCb::RawBank::ODIN, {BankTypes::ODIN}}, + {LHCb::RawBank::HcalPacked, {BankTypes::HCal}}, + {LHCb::RawBank::EcalPacked, {BankTypes::ECal}}, + {LHCb::RawBank::Calo, {BankTypes::ECal}}, + {LHCb::RawBank::Rich, {BankTypes::Rich1, BankTypes::Rich2}}, + {LHCb::RawBank::OTError, {BankTypes::MCVertices}}, // used for PV MC info + {LHCb::RawBank::OTRaw, {BankTypes::MCTracks}}, + {LHCb::RawBank::Plume, {BankTypes::Plume}}}; // used for track MC info const std::unordered_map subdetectors = {{SourceIdSys::SourceIdSys_ODIN, BankTypes::ODIN}, {SourceIdSys::SourceIdSys_VELO_A, BankTypes::VP}, diff --git a/main/include/MEPTools.h b/main/include/MEPTools.h index 372eabde89e07b447e7f83e31f5fe5c037b24aff..cf70332edc7f4fb6f21ccd5b504ec6957a8c9d14 100644 --- a/main/include/MEPTools.h +++ b/main/include/MEPTools.h @@ -220,3 +220,64 @@ namespace MEP { } }; } // namespace MEP + +namespace Allen { + struct RawBank { + uint32_t source_id = 0; + uint16_t size = 0; + uint8_t const* data = nullptr; + uint8_t const type; + + // For Allen format + __device__ __host__ RawBank(const char* raw_bank, const uint16_t s, const uint8_t t) : + RawBank {*reinterpret_cast(raw_bank), raw_bank + sizeof(uint32_t), s, t} + {} + + // For MEP format + __device__ __host__ RawBank(const uint32_t sid, const char* fragment, const uint16_t s, const uint8_t t) : + source_id {sid}, size {s}, data {reinterpret_cast(fragment)}, type {t} + {} + }; + + template + struct RawEvent { + + uint32_t number_of_raw_banks = 0; + const char* data = nullptr; + const uint32_t* offsets = nullptr; + typename std::conditional_t* sizes = nullptr; + typename std::conditional_t* types = nullptr; + const unsigned event = 0; + + // For Allen format + __device__ __host__ + RawEvent(char const* d, uint32_t const* o, uint32_t const* s, uint32_t const* t, unsigned const event_number) : + offsets {o}, + event {event_number} + { + if constexpr (mep_layout) { + data = d; + number_of_raw_banks = MEP::number_of_banks(o); + sizes = s; + types = t; + } + else { + data = d + offsets[event]; + number_of_raw_banks = reinterpret_cast(data)[0]; + sizes = Allen::bank_sizes(s, event); + types = Allen::bank_types(t, event); + } + } + + __device__ __host__ RawBank raw_bank(unsigned const n) const + { + if constexpr (mep_layout) { + return MEP::raw_bank(data, offsets, sizes, types, event, n); + } + else { + uint32_t const* bank_offsets = reinterpret_cast(data) + 1; + return RawBank {data + (number_of_raw_banks + 2) * sizeof(uint32_t) + bank_offsets[n], sizes[n], types[n]}; + } + } + }; +} // namespace Allen diff --git a/main/src/RegisterConsumers.cpp b/main/src/RegisterConsumers.cpp index 265f5340fe01b9e1b20202b28e8a2b36280b4b01..fe4752da70d257ec4f3b4bf3d5e9760627a2549c 100644 --- a/main/src/RegisterConsumers.cpp +++ b/main/src/RegisterConsumers.cpp @@ -68,7 +68,35 @@ void register_consumers( return std::make_unique( constants.host_muon_lookup_tables_raw, constants.dev_muon_lookup_tables_raw, constants.dev_muon_tables); }, - BankTypes::MUON)); + BankTypes::MUON), + std::make_tuple( + Allen::NonEventData::RichPDMDBMapping {}, + [&constants]() { + return std::make_unique( + constants.host_rich_pdmdb_mapping, constants.dev_rich_pdmdb_mapping); + }, + BankTypes::Rich1), + std::make_tuple( + Allen::NonEventData::RichCableMapping {}, + [&constants]() { + return std::make_unique( + constants.host_rich_cable_mapping, constants.dev_rich_cable_mapping); + }, + BankTypes::Rich1), + std::make_tuple( + Allen::NonEventData::RichPDMDBMapping {}, + [&constants]() { + return std::make_unique( + constants.host_rich_pdmdb_mapping, constants.dev_rich_pdmdb_mapping); + }, + BankTypes::Rich2), + std::make_tuple( + Allen::NonEventData::RichCableMapping {}, + [&constants]() { + return std::make_unique( + constants.host_rich_cable_mapping, constants.dev_rich_cable_mapping); + }, + BankTypes::Rich2)); const auto unconditional_consumers = std::make_tuple(std::make_tuple(Allen::NonEventData::MagneticField {}, [&constants]() { diff --git a/main/src/Transpose.cpp b/main/src/Transpose.cpp index dc9cea7f3abe5fba7297e69487bd1343cceb7260..63bf98961877c84aad46de1a1216eb6f68a11bb2 100644 --- a/main/src/Transpose.cpp +++ b/main/src/Transpose.cpp @@ -17,7 +17,14 @@ std::array Allen::bank_ids() std::array ids; for (auto bt : LHCb::RawBank::types()) { auto it = Allen::bank_mapping.find(bt); - ids[bt] = (it != Allen::bank_mapping.end() ? to_integral(it->second) : -1); + if (it != Allen::bank_mapping.end()) { + for (auto allen_bt : it->second) { + ids[bt] = static_cast(allen_bt); + } + } + else { + ids[bt] = -1; + } } return ids; } diff --git a/mdf/test/test_mep_banks.cpp b/mdf/test/test_mep_banks.cpp index 85c3e91dee43d51314a7a78ff63323da8ac4b2b2..f9565d99bba1ee58c8047a61188e1b26c3bfbc1f 100644 --- a/mdf/test/test_mep_banks.cpp +++ b/mdf/test/test_mep_banks.cpp @@ -24,7 +24,6 @@ #include #include #include -#include #include #include @@ -241,6 +240,44 @@ int main(int argc, char* argv[]) template struct compare { + void operator()( + const int, + gsl::span mep_fragments, + gsl::span mep_offsets, + gsl::span mep_sizes, + gsl::span mep_types, + gsl::span allen_banks, + gsl::span allen_offsets, + gsl::span allen_sizes, + gsl::span allen_types, + unsigned const i_event) + { + + const auto allen_raw_event = + Allen::RawEvent(allen_banks.data(), allen_offsets.data(), allen_sizes.data(), allen_types.data(), i_event); + const auto mep_raw_event = Allen::RawEvent( + mep_fragments.data(), mep_offsets.data(), mep_sizes.data(), mep_types.data(), i_event); + auto const mep_n_banks = mep_raw_event.number_of_raw_banks; + + REQUIRE(mep_n_banks == allen_raw_event.number_of_raw_banks); + + for (unsigned bank = 0; bank < mep_n_banks; ++bank) { + // Read raw bank + auto const mep_bank = mep_raw_event.raw_bank(bank); + auto const allen_bank = allen_raw_event.raw_bank(bank); + auto mep_len = mep_bank.size; + auto allen_len = allen_bank.size; + REQUIRE(mep_len == allen_len); + + REQUIRE(mep_bank.type == allen_bank.type); + + auto top5_mask = (allen_bank.source_id >> 11 == 0) ? 0x7FF : 0xFFFF; + REQUIRE((mep_bank.source_id & top5_mask) == allen_bank.source_id); + for (long j = 0; j < mep_len; ++j) { + REQUIRE(allen_bank.data[j] == mep_bank.data[j]); + } + } + } }; template @@ -456,48 +493,6 @@ struct compare { } }; -template -struct compare { - void operator()( - const int, - gsl::span mep_fragments, - gsl::span mep_offsets, - gsl::span mep_sizes, - gsl::span mep_types, - gsl::span allen_banks, - gsl::span allen_offsets, - gsl::span allen_sizes, - gsl::span allen_types, - unsigned const i_event) - { - - const auto allen_raw_event = - Calo::RawEvent(allen_banks.data(), allen_offsets.data(), allen_sizes.data(), allen_types.data(), i_event); - const auto mep_raw_event = Calo::RawEvent( - mep_fragments.data(), mep_offsets.data(), mep_sizes.data(), mep_types.data(), i_event); - auto const mep_n_banks = mep_raw_event.number_of_raw_banks; - - REQUIRE(mep_n_banks == allen_raw_event.number_of_raw_banks); - - for (unsigned bank = 0; bank < mep_n_banks; ++bank) { - // Read raw bank - auto const mep_bank = mep_raw_event.raw_bank(bank); - auto const allen_bank = allen_raw_event.raw_bank(bank); - auto mep_len = mep_bank.end - mep_bank.data; - auto allen_len = allen_bank.end - allen_bank.data; - REQUIRE(mep_len == allen_len); - - REQUIRE(mep_bank.type == allen_bank.type); - - auto top5_mask = (allen_bank.source_id >> 11 == 0) ? 0x7FF : 0xFFFF; - REQUIRE((mep_bank.source_id & top5_mask) == allen_bank.source_id); - for (long j = 0; j < mep_len; ++j) { - REQUIRE(allen_bank.data[j] == mep_bank.data[j]); - } - } - } -}; - template struct BTTag { inline static const BankTypes BT = BT_; @@ -509,6 +504,8 @@ using SciFiTag = BTTag; using UTTag = BTTag; using MuonTag = BTTag; using ECalTag = BTTag; +using Rich1Tag = BTTag; +using Rich2Tag = BTTag; /** * @brief Check banks @@ -576,7 +573,7 @@ void check_banks(BanksAndOffsets const& mep_data, BanksAndOffsets const& allen_d // Main test case, multiple bank types are checked // VeloTag, UTTag, SciFiTag, -TEMPLATE_TEST_CASE("MEP vs MDF", "[MEP MDF]", ECalTag, MuonTag, VeloTag, SciFiTag, ODINTag) +TEMPLATE_TEST_CASE("MEP vs MDF", "[MEP MDF]", ECalTag, MuonTag, VeloTag, SciFiTag, ODINTag, Rich1Tag, Rich2Tag) { if (!s_config.run) return; diff --git a/scripts/ci/test_config.yaml b/scripts/ci/test_config.yaml index 7b7356813f5ffbe9509ca80369c884842f8555c4..9c7d187cb7c690af843963a806d936ed7bebb140 100644 --- a/scripts/ci/test_config.yaml +++ b/scripts/ci/test_config.yaml @@ -211,6 +211,13 @@ full: dataset: "Beam6800GeV-expected-2024-MagDown-nu7.6" geometry: "geometry_dddb-20231017_sim-20231017-vc-md100" + #throughput test with RICH decoding -> medium/long term development + - type: "throughput" + sequence: + - hlt1_pp_rich_no_ut + dataset: "Run_0000248711_HLT20840_20221011-113809-426" + geometry: "MiniBrunel_2018_MinBias_FTv4_DIGI_ecalv4_scifiv7_muonv3_RICH" + # run physics efficiency for downstream reconstruction - type: "efficiency" sequence: diff --git a/stream/CMakeLists.txt b/stream/CMakeLists.txt index 1f3478cd293fd486a7225de5bfcd361377b2700e..764b898ab687f505f13fff90b2059b2d53247e62 100644 --- a/stream/CMakeLists.txt +++ b/stream/CMakeLists.txt @@ -41,7 +41,10 @@ target_link_libraries(Stream Gear track_matching MuonCommon - ) + Rich + PUBLIC + Utils + Selections) if(STANDALONE) add_dependencies(Stream checkout_lhcb checkout_gaudi) diff --git a/stream/sequence/include/Constants.cuh b/stream/sequence/include/Constants.cuh index fb4a9af333cbad1d0d9ba98989ed3ddb1fa77514..6f1d72e37edf254b5352e8a763fc05789b6b4ffd 100644 --- a/stream/sequence/include/Constants.cuh +++ b/stream/sequence/include/Constants.cuh @@ -36,6 +36,10 @@ namespace MatchUpstreamMuon { namespace TrackMatchingConsts { struct MagnetParametrization; } +namespace Rich::Future::DAQ::Allen { + class PDMDBDecodeMapping; + class Tel40CableMapping; +} // namespace Rich::Future::DAQ::Allen /** * @brief Struct intended as a singleton with constants defined on GPU. @@ -47,7 +51,6 @@ namespace TrackMatchingConsts { * The pointers are hard-coded. Feel free to write more as needed. */ struct Constants { - gsl::span dev_velo_candidate_ks; gsl::span dev_velo_sp_patterns; gsl::span dev_velo_sp_fx; @@ -137,6 +140,12 @@ struct Constants { // Kalman filter ParKalmanFilter::KalmanParametrizations* dev_kalman_params = nullptr; + // Rich + std::vector host_rich_pdmdb_mapping; + std::vector host_rich_cable_mapping; + char* dev_rich_pdmdb_mapping; + char* dev_rich_cable_mapping; + /** * @brief Reserves and initializes constants. */