diff --git a/configuration/python/AllenConf/lumi_reconstruction.py b/configuration/python/AllenConf/lumi_reconstruction.py
index 3eb1e70ddcb9432b940ec4b7715b3bb1d0127191..cbe273cedd2783825c8b7101b15c77f992be652d 100644
--- a/configuration/python/AllenConf/lumi_reconstruction.py
+++ b/configuration/python/AllenConf/lumi_reconstruction.py
@@ -180,7 +180,18 @@ def lumi_reconstruction(
("VeloClustersInnerS22", 3000), ("VeloClustersOuterS22", 3000),
("VeloClustersInnerS23", 3000), ("VeloClustersOuterS23", 3000),
("VeloClustersInnerS24", 3000), ("VeloClustersOuterS24", 3000),
- ("VeloClustersInnerS25", 3000), ("VeloClustersOuterS25", 3000)
+ ("VeloClustersInnerS25", 3000), ("VeloClustersOuterS25", 3000),
+ ("MuonHitsTell01", 400), ("MuonHitsTell02", 400),
+ ("MuonHitsTell03", 400), ("MuonHitsTell04", 400),
+ ("MuonHitsTell05", 400), ("MuonHitsTell06", 400),
+ ("MuonHitsTell07", 400), ("MuonHitsTell08", 400),
+ ("MuonHitsTell09", 400), ("MuonHitsTell10", 400),
+ ("MuonHitsTell11", 400), ("MuonHitsTell12", 400),
+ ("MuonHitsTell13", 400), ("MuonHitsTell14", 400),
+ ("MuonHitsTell15", 400), ("MuonHitsTell16", 400),
+ ("MuonHitsTell17", 400), ("MuonHitsTell18", 400),
+ ("MuonHitsTell19", 400), ("MuonHitsTell20", 400),
+ ("MuonHitsTell21", 400), ("MuonHitsTell22", 400)
],
counterFactors={
"ECalET": (0x10000, 0.2),
@@ -333,12 +344,14 @@ def lumi_reconstruction(
host_number_of_events_t=number_of_events["host_number_of_events"],
host_lumi_summaries_count_t=prefix_sum_lumi_present.
host_total_sum_holder_t,
+ host_raw_bank_version_t=decoded_muon["host_raw_bank_version"],
dev_lumi_event_indices_t=prefix_sum_lumi_present.
dev_output_buffer_t,
dev_storage_station_region_quarter_offsets_t=decoded_muon[
"dev_storage_station_region_quarter_offsets"],
dev_muon_number_of_tracks_t=muon_stubs[
"dev_muon_number_of_tracks"],
+ dev_muon_tell_number_t=decoded_muon["dev_muon_tell_number"],
lumi_counter_schema=schema_for_algorithms,
lumi_counter_shifts_and_scales=shifts_and_scales_for_algorithms)
diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py
index 066b26471c1f556e80add54bd4ec37edc0862add..0988a24475a240660c1b8727f5a37ae01d1fb69c 100644
--- a/configuration/python/AllenConf/muon_reconstruction.py
+++ b/configuration/python/AllenConf/muon_reconstruction.py
@@ -100,7 +100,11 @@ def decode_muon(empty_banks=False):
"dev_muon_hits":
muon_populate_hits.dev_muon_hits_t,
"dev_station_ocurrences_offset":
- muon_station_ocurrence_prefix_sum.dev_output_buffer_t
+ muon_station_ocurrence_prefix_sum.dev_output_buffer_t,
+ "host_raw_bank_version":
+ muon_banks.host_raw_bank_version_t,
+ "dev_muon_tell_number":
+ muon_populate_tile_and_tdc.dev_muon_tell_number_t
}
diff --git a/device/event_model/lumi/include/LumiDefinitions.cuh b/device/event_model/lumi/include/LumiDefinitions.cuh
index 4eef0891c162f1cc7ff3155efdd2261d2b1e86af..e1a1b0ba9ad6c8f9406451348665c4fea15342aa 100644
--- a/device/event_model/lumi/include/LumiDefinitions.cuh
+++ b/device/event_model/lumi/include/LumiDefinitions.cuh
@@ -62,7 +62,8 @@ namespace Lumi {
static constexpr unsigned n_pv_counters = 5u;
static constexpr unsigned n_scifi_counters = 38u;
static constexpr unsigned n_calo_counters = 8u;
- static constexpr unsigned n_muon_counters = 13u;
+ // 1u for muon tracks
+ static constexpr unsigned n_muon_counters = n_muon_station_regions + 1u + Muon::Constants::maxTell40Number;
static constexpr unsigned n_plume_counters = 47u;
// number of velo eta bins edges
@@ -166,19 +167,13 @@ namespace Lumi {
"ECalETOuterBottom",
"ECalETMiddleBottom",
"ECalETInnerBottom"};
- const std::array<std::string, n_muon_counters> muon_counter_names = {"MuonHitsM2R1",
- "MuonHitsM2R2",
- "MuonHitsM2R3",
- "MuonHitsM2R4",
- "MuonHitsM3R1",
- "MuonHitsM3R2",
- "MuonHitsM3R3",
- "MuonHitsM3R4",
- "MuonHitsM4R1",
- "MuonHitsM4R2",
- "MuonHitsM4R3",
- "MuonHitsM4R4",
- "MuonTracks"};
+ const std::array<std::string, n_muon_counters> muon_counter_names = {
+ "MuonHitsM2R1", "MuonHitsM2R2", "MuonHitsM2R3", "MuonHitsM2R4", "MuonHitsM3R1", "MuonHitsM3R2",
+ "MuonHitsM3R3", "MuonHitsM3R4", "MuonHitsM4R1", "MuonHitsM4R2", "MuonHitsM4R3", "MuonHitsM4R4",
+ "MuonTracks", "MuonHitsTell01", "MuonHitsTell02", "MuonHitsTell03", "MuonHitsTell04", "MuonHitsTell05",
+ "MuonHitsTell06", "MuonHitsTell07", "MuonHitsTell08", "MuonHitsTell09", "MuonHitsTell10", "MuonHitsTell11",
+ "MuonHitsTell12", "MuonHitsTell13", "MuonHitsTell14", "MuonHitsTell15", "MuonHitsTell16", "MuonHitsTell17",
+ "MuonHitsTell18", "MuonHitsTell19", "MuonHitsTell20", "MuonHitsTell21", "MuonHitsTell22"};
const std::array<std::string, n_plume_counters> plume_counter_names = {
"PlumeAvgLumiADC", "PlumeLumiOverthrLow", "PlumeLumiOverthrHigh", "PlumeLumiADC00", "PlumeLumiADC01",
"PlumeLumiADC02", "PlumeLumiADC03", "PlumeLumiADC04", "PlumeLumiADC05", "PlumeLumiADC06",
diff --git a/device/lumi/include/MuonLumiCounters.cuh b/device/lumi/include/MuonLumiCounters.cuh
index 2832d27bb7b5eeb839e5718c2c333392e3556da9..c6055e09d99374ff4c40693b4f09f480929ce3c4 100644
--- a/device/lumi/include/MuonLumiCounters.cuh
+++ b/device/lumi/include/MuonLumiCounters.cuh
@@ -16,16 +16,19 @@
#include <LumiDefinitions.cuh>
-#include <VeloConsolidated.cuh>
+#include "MuonDefinitions.cuh"
+#include "MuonEventModel.cuh"
namespace muon_lumi_counters {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
HOST_INPUT(host_lumi_summaries_count_t, unsigned) host_lumi_summaries_count;
+ HOST_INPUT(host_raw_bank_version_t, int) host_raw_bank_version;
DEVICE_INPUT(dev_lumi_event_indices_t, unsigned) dev_lumi_event_indices;
DEVICE_INPUT(dev_storage_station_region_quarter_offsets_t, unsigned) dev_storage_station_region_quarter_offsets;
DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks;
+ DEVICE_INPUT(dev_muon_tell_number_t, unsigned short) dev_muon_tell_number;
DEVICE_OUTPUT(dev_lumi_infos_t, Lumi::LumiInfo) dev_lumi_infos;
PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
PROPERTY(
@@ -49,6 +52,7 @@ namespace muon_lumi_counters {
Parameters,
const unsigned number_of_events,
const unsigned number_of_gec_events,
+ const int decoding_version,
const offsets_and_sizes_t offsets_and_sizes,
const shifts_and_scales_t shifts_and_scales);
diff --git a/device/lumi/src/MuonLumiCounters.cu b/device/lumi/src/MuonLumiCounters.cu
index e80f1fc18a122a649a18ab7b6065d6b6c9e12164..76a86a871fb0e4ee13625cd3e9b5710d517a2147 100644
--- a/device/lumi/src/MuonLumiCounters.cu
+++ b/device/lumi/src/MuonLumiCounters.cu
@@ -64,6 +64,7 @@ void muon_lumi_counters::muon_lumi_counters_t::operator()(
arguments,
first<host_number_of_events_t>(arguments),
size<dev_event_list_t>(arguments),
+ first<host_raw_bank_version_t>(arguments),
m_offsets_and_sizes,
m_shifts_and_scales);
}
@@ -72,6 +73,7 @@ __global__ void muon_lumi_counters::muon_lumi_counters(
muon_lumi_counters::Parameters parameters,
const unsigned number_of_events,
const unsigned number_of_gec_events,
+ const int decoding_version,
const offsets_and_sizes_t offsets_and_sizes,
const shifts_and_scales_t shifts_and_scales)
{
@@ -85,6 +87,7 @@ __global__ void muon_lumi_counters::muon_lumi_counters(
const auto muon_hits_offsets =
parameters.dev_storage_station_region_quarter_offsets + event_number * Lumi::Constants::MuonBankSize;
+ const auto event_number_of_hits = muon_hits_offsets[Lumi::Constants::MuonBankSize] - muon_hits_offsets[0];
unsigned info_offset = Lumi::Constants::n_muon_counters * lumi_evt_index;
@@ -111,6 +114,31 @@ __global__ void muon_lumi_counters::muon_lumi_counters(
shifts_and_scales[2 * i],
shifts_and_scales[2 * i + 1]);
}
+
+ const auto tell_number = parameters.dev_muon_tell_number + muon_hits_offsets[0];
+ std::array<unsigned, Muon::Constants::maxTell40Number> hits_per_tell = {0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u,
+ 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u};
+ unsigned firstTell = 1u;
+ // tell number starts from 0 for decoding version 2
+ // and only 10 tell numbers defined
+ if (decoding_version == 2) {
+ firstTell = 0u;
+ for (unsigned i = 10u; i < Muon::Constants::maxTell40Number; ++i)
+ hits_per_tell[i] = 0xffffffff;
+ }
+ for (unsigned i = 0; i < event_number_of_hits; ++i) {
+ assert(tell_number[i] <= Muon::Constants::maxTell40Number);
+ ++hits_per_tell[tell_number[i] - firstTell];
+ }
+ for (unsigned i = Lumi::Constants::n_muon_station_regions + 1u; i < Lumi::Constants::n_muon_counters; ++i) {
+ fillLumiInfo(
+ parameters.dev_lumi_infos[info_offset + i],
+ offsets_and_sizes[2 * i],
+ offsets_and_sizes[2 * i + 1],
+ hits_per_tell[i - Lumi::Constants::n_muon_station_regions - 1u],
+ shifts_and_scales[2 * i],
+ shifts_and_scales[2 * i + 1]);
+ }
}
// now fill counters that are only avaiable inside the GEC
diff --git a/device/muon/decoding/include/MuonPopulateTileAndTDC.cuh b/device/muon/decoding/include/MuonPopulateTileAndTDC.cuh
index 9af4582a6f5fea91b3b33ff9743e8e3d64fc7017..a8071350f5bb1f7d5c560b0f667ecb22a88a1e70 100644
--- a/device/muon/decoding/include/MuonPopulateTileAndTDC.cuh
+++ b/device/muon/decoding/include/MuonPopulateTileAndTDC.cuh
@@ -25,6 +25,7 @@ namespace muon_populate_tile_and_tdc {
DEVICE_OUTPUT(dev_atomics_muon_t, unsigned) dev_atomics_muon;
DEVICE_OUTPUT(dev_muon_tile_used_t, bool) dev_muon_tile_used;
DEVICE_OUTPUT(dev_station_ocurrences_sizes_t, unsigned) dev_station_ocurrences_sizes;
+ DEVICE_OUTPUT(dev_muon_tell_number_t, unsigned short) dev_muon_tell_number;
};
struct muon_populate_tile_and_tdc_t : public DeviceAlgorithm, Parameters {
diff --git a/device/muon/decoding/src/MuonPopulateTileAndTDC.cu b/device/muon/decoding/src/MuonPopulateTileAndTDC.cu
index 705a56cf831b198b769ca75060046340f368087e..02b9eef44fe5720f64e99196b4c5f40ba745cba1 100644
--- a/device/muon/decoding/src/MuonPopulateTileAndTDC.cu
+++ b/device/muon/decoding/src/MuonPopulateTileAndTDC.cu
@@ -84,7 +84,8 @@ __device__ void decode_muon_bank(
const unsigned* storage_station_region_quarter_offsets,
unsigned* atomics_muon,
unsigned* dev_storage_tile_id,
- unsigned* dev_storage_tdc_value)
+ unsigned* dev_storage_tdc_value,
+ unsigned short* dev_muon_tell_number)
{
if constexpr (decoding_version == 2) {
for (unsigned batch_index = threadIdx.y; batch_index < Muon::batches_per_bank; batch_index += blockDim.y) {
@@ -115,6 +116,7 @@ __device__ void decode_muon_bank(
const auto insert_index = atomicAdd(atomics_muon + storage_srq_layout, 1);
dev_storage_tile_id[storage_station_region_quarter_offsets[storage_srq_layout] + insert_index] = tileId;
dev_storage_tdc_value[storage_station_region_quarter_offsets[storage_srq_layout] + insert_index] = tdc_value;
+ dev_muon_tell_number[storage_station_region_quarter_offsets[storage_srq_layout] + insert_index] = tell_number;
}
}
}
@@ -233,6 +235,8 @@ __device__ void decode_muon_bank(
tileId;
dev_storage_tdc_value[storage_station_region_quarter_offsets[storage_srq_layout] + insert_index] =
tdc_value;
+ dev_muon_tell_number[storage_station_region_quarter_offsets[storage_srq_layout] + insert_index] =
+ tell_number;
}
}
}
@@ -275,7 +279,8 @@ __global__ void muon_populate_tile_and_tdc_kernel(
storage_station_region_quarter_offsets,
atomics_muon,
parameters.dev_storage_tile_id,
- parameters.dev_storage_tdc_value);
+ parameters.dev_storage_tdc_value,
+ parameters.dev_muon_tell_number);
}
}
@@ -293,6 +298,7 @@ void muon_populate_tile_and_tdc::muon_populate_tile_and_tdc_t::set_arguments_siz
set_size<dev_muon_tile_used_t>(arguments, first<host_muon_total_number_of_tiles_t>(arguments));
set_size<dev_station_ocurrences_sizes_t>(
arguments, first<host_number_of_events_t>(arguments) * Muon::Constants::n_stations);
+ set_size<dev_muon_tell_number_t>(arguments, first<host_muon_total_number_of_tiles_t>(arguments));
}
void muon_populate_tile_and_tdc::muon_populate_tile_and_tdc_t::operator()(
@@ -306,6 +312,7 @@ void muon_populate_tile_and_tdc::muon_populate_tile_and_tdc_t::operator()(
Allen::memset_async<dev_storage_tdc_value_t>(arguments, 0, context);
Allen::memset_async<dev_muon_tile_used_t>(arguments, 0, context);
Allen::memset_async<dev_station_ocurrences_sizes_t>(arguments, 0, context);
+ Allen::memset_async<dev_muon_tell_number_t>(arguments, 0xffff, context);
const auto bank_version = first<host_raw_bank_version_t>(arguments);
if (bank_version < 0) return; // no Muon banks present in data