Allen RICH decoding

This MR introduces RICH decoding, concretely producing RichSmartIDs from RICH1 and RICH2 raw banks.

• Created algorithm rich_decoding_t, which performs the decoding by calculating the number of hits, prefix sum and finally decoding the hits.
• Added simplified copies of RichSmartID, PackedFrameSizes, Tel40CableMapping and PDMDBDecodeMapping adjusted to the requirements of GPU computing, namely no dynamic memory allocation, no std::string elements and __host__ __device__ attributes.
• Created Allen::RawEvent and Allen::RawBank which replace CaloRawEvent and CaloRawBank. Ideally this should be used in other subdetectors as well due to the raw bank format being mostly identical.
• Added configurations rich and hlt1_pp_rich_no_ut.
• Added throughput test of hlt1_pp_rich_no_ut.
• Added two geometry files corresponding to Tel40CableMapping and PDMDBDecodeMapping, alongside constants and the necessary scaffolding.

Goes with LHCb!4001 (merged) and Moore!2171 (merged)

Done in collaboration with @caschbur .

configuration/python/AllenConf/rich_reconstruction.py

@@ -20,7 +20,8 @@ def decode_rich1():
         dev_rich_raw_input_t=rich1_banks.dev_raw_banks_t,
         dev_rich_raw_input_offsets_t=rich1_banks.dev_raw_offsets_t,
         dev_rich_raw_input_sizes_t=rich1_banks.dev_raw_sizes_t,
-        dev_rich_raw_input_types_t=rich1_banks.dev_raw_types_t)
+        dev_rich_raw_input_types_t=rich1_banks.dev_raw_types_t,
+        block_dim_x=128)
 
     return rich1_decoding
 
@@ -39,6 +40,7 @@ def decode_rich2():
         dev_rich_raw_input_t=rich2_banks.dev_raw_banks_t,
         dev_rich_raw_input_offsets_t=rich2_banks.dev_raw_offsets_t,
         dev_rich_raw_input_sizes_t=rich2_banks.dev_raw_sizes_t,
-        dev_rich_raw_input_types_t=rich2_banks.dev_raw_types_t)
+        dev_rich_raw_input_types_t=rich2_banks.dev_raw_types_t,
+        block_dim_x=128)
 
     return rich2_decoding