creating new HLT1 line for error banks (with raw event model)

configuration/python/AllenConf/hlt1_calibration_lines.py

@@ -2,7 +2,7 @@
 # (c) Copyright 2021 CERN for the benefit of the LHCb Collaboration           #
 ###############################################################################
 from AllenConf.algorithms import (d2kpi_line_t, passthrough_line_t,
-                                  rich_1_line_t, rich_2_line_t)
+                                  rich_1_line_t, rich_2_line_t, error_banks_line_t, bank_types_provider_t)
 from AllenConf.utils import initialize_number_of_events, mep_layout
 from AllenCore.generator import make_algorithm
 from PyConf.tonic import configurable
@@ -47,6 +47,29 @@ def make_passthrough_line(pre_scaler_hash_string="passthrough_line_pre",
         pre_scaler_hash_string=pre_scaler_hash_string,
         post_scaler_hash_string=post_scaler_hash_string)
 
+def make_error_banks_line(pre_scaler_hash_string="error_banks_line_pre",
+                          post_scaler_hash_string="error_banks_line_post",
+                          name="Hlt1ErrorBanks"):
+    number_of_events = initialize_number_of_events()
+    odin = decode_odin()
+    # layout = mep_layout()
+    ecal_banks = make_algorithm(
+        bank_types_provider_t, name="ecal_banks", bank_type="ECal")
+
+    return make_algorithm(
+        error_banks_line_t,
+        name=name,
+        host_number_of_events_t=number_of_events["host_number_of_events"],
+        dev_number_of_events_t=number_of_events["dev_number_of_events"],
+        dev_odin_data_t=odin["dev_odin_data"],
+        # dev_mep_layout_t=layout["dev_mep_layout"],
+        # dev_ecal_raw_input_t=ecal_banks.dev_raw_banks_t,
+        dev_ecal_raw_input_offsets_t=ecal_banks.dev_raw_offsets_t,
+        # dev_ecal_raw_input_sizes_t=ecal_banks.dev_raw_sizes_t,
+        dev_ecal_raw_input_types_t=ecal_banks.dev_raw_types_t,
+        pre_scaler_hash_string=pre_scaler_hash_string,
+        post_scaler_hash_string=post_scaler_hash_string)
+
 
 def make_rich_line(line_type, forward_tracks, long_track_particles,
                    pre_scaler_hash_string, post_scaler_hash_string, name):

configuration/sequences/select_error_banks.py

+###############################################################################
+# (c) Copyright 2021 CERN for the benefit of the LHCb Collaboration           #
+###############################################################################
+from PyConf.control_flow import NodeLogic, CompositeNode
+from AllenCore.generator import generate, make_algorithm
+from AllenConf.hlt1_calibration_lines import make_passthrough_line, make_error_banks_line
+from AllenConf.persistency import make_global_decision
+from AllenConf.odin import decode_odin
+from AllenConf.algorithms import data_provider_t
+from AllenConf.HLT1 import line_maker
+from AllenConf.validators import rate_validation
+
+bank_providers = [decode_odin()['dev_odin_data'].producer]
+
+# To test memory traffic when copying to the device, add the following
+# for det, bt in (("velo", "VP"), ("ut", "UT"), ("scifi", "FT"),
+#                 ("muon", "Muon")):
+#     bank_providers.append(
+#         make_algorithm(data_provider_t, name=det + "_banks", bank_type=bt))
+
+passthrough_line = line_maker(
+    "Hlt1Passthrough",
+    make_passthrough_line(
+        name="Hlt1Passthrough",
+        pre_scaler_hash_string="passthrough_line_pre",
+        post_scaler_hash_string="passthrough_line_post"),
+    enableGEC=False)
+
+error_banks_line = line_maker(
+    "Hlt1ErrorBanks",
+    make_error_banks_line(
+        name="Hlt1ErrorBanks",
+        pre_scaler_hash_string="error_banks_line_pre",
+        post_scaler_hash_string="error_banks_line_post"),
+    enableGEC=False)
+
+# line_algorithms = [passthrough_line[0],error_banks_line[0]]
+line_algorithms = [error_banks_line[0]]
+
+# line_nodes = [passthrough_line[1],error_banks_line[1]]
+line_nodes = [error_banks_line[1]]
+
+
+global_decision = make_global_decision(lines=line_algorithms)
+
+providers = CompositeNode(
+    "Providers", bank_providers, NodeLogic.NONLAZY_AND, force_order=False)
+
+lines = CompositeNode(
+    "AllLines", line_nodes, NodeLogic.NONLAZY_OR, force_order=False)
+
+select_error_banks_sequence = CompositeNode(
+    "SelectErrorBanks",
+    [lines, global_decision,
+     rate_validation(lines=line_algorithms)],
+    NodeLogic.NONLAZY_AND,
+    force_order=True)
+
+generate(select_error_banks_sequence)

device/event_model/selections/include/RawBankTypesRawEvent.cuh

+/*****************************************************************************\
+* (c) Copyright 2021 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 "LICENSE".              *
+*                                                                             *
+* 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 <type_traits>
+#include "BackendCommon.h"
+#include "MEPTools.h"
+
+namespace RawBankTypes {
+
+  struct RawEvent {
+    const uint32_t* offsets = nullptr;
+    uint32_t const* types = nullptr;
+    const unsigned event = 0;
+    
+    __device__ __host__  unsigned number_of_raw_banks() const { return offsets[0]; }
+
+    __device__ __host__ unsigned char bank_type(const unsigned index) const
+    {
+      return MEP::bank_type(nullptr, types, event, index);
+    }
+  };
+} // namespace RawBankTypes

device/selections/lines/calibration/include/ErrorBanksLine.cuh

+/*****************************************************************************\
+* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration           *
+\*****************************************************************************/
+#pragma once
+
+#include "AlgorithmTypes.cuh"
+#include "Event/ODIN.h"
+#include "EventLine.cuh"
+#include "RawBankTypesRawEvent.cuh"
+#include "BankMapping.h"
+
+
+namespace error_banks_line {
+  struct Parameters {
+    HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
+    DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events;
+    MASK_INPUT(dev_event_list_t) dev_event_list;
+    MASK_OUTPUT(dev_selected_events_t) dev_selected_events;
+    HOST_OUTPUT(host_selected_events_size_t, unsigned) host_selected_events_size;
+    DEVICE_OUTPUT(dev_selected_events_size_t, unsigned) dev_selected_events_size;
+    DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
+    // DEVICE_INPUT(dev_ecal_raw_input_t, char) dev_ecal_raw_input;
+    DEVICE_INPUT(dev_ecal_raw_input_offsets_t, unsigned) dev_ecal_raw_input_offsets;
+    // DEVICE_INPUT(dev_ecal_raw_input_sizes_t, unsigned) dev_ecal_raw_input_sizes;
+    DEVICE_INPUT(dev_ecal_raw_input_types_t, unsigned) dev_ecal_raw_input_types;
+    DEVICE_OUTPUT(dev_decisions_t, bool) dev_decisions;
+    DEVICE_OUTPUT(dev_decisions_offsets_t, unsigned) dev_decisions_offsets;
+    HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
+    HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+    DEVICE_OUTPUT(dev_particle_container_ptr_t, Allen::IMultiEventContainer*)
+    dev_particle_container_ptr;
+    PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
+    PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
+    PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string);
+    PROPERTY(post_scaler_hash_string_t, "post_scaler_hash_string", "Post-scaling hash string", std::string);
+  };
+
+  
+
+  struct error_banks_line_t : public SelectionAlgorithm, Parameters, EventLine<error_banks_line_t, Parameters> {
+    __device__ static std::tuple<const unsigned> get_input(const Parameters& parameters, const unsigned event_number);
+
+    __device__ static bool select(const Parameters& parameters, std::tuple<const unsigned> input);
+
+  private:
+    Property<pre_scaler_t> m_pre_scaler {this, 1.f};
+    Property<post_scaler_t> m_post_scaler {this, 1.f};
+    Property<pre_scaler_hash_string_t> m_pre_scaler_hash_string {this, ""};
+    Property<post_scaler_hash_string_t> m_post_scaler_hash_string {this, ""};
+  };
+} // namespace error_banks_line

device/selections/lines/calibration/src/ErrorBanksLine.cu

+/*****************************************************************************\
+* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration           *
+\*****************************************************************************/
+#include "ErrorBanksLine.cuh"
+
+
+
+// Explicit instantiation
+INSTANTIATE_LINE(error_banks_line::error_banks_line_t, error_banks_line::Parameters)
+
+
+// Get input function
+__device__ std::tuple<const unsigned> error_banks_line::error_banks_line_t::get_input(const Parameters& parameters, const unsigned event_number)
+  {
+    
+    unsigned error_banks_counter = 0;
+    const auto raw_event = RawBankTypes::RawEvent {parameters.dev_ecal_raw_input_offsets,
+                                                parameters.dev_ecal_raw_input_types,
+                                                event_number};
+    unsigned number_of_raw_banks = raw_event.number_of_raw_banks();
+
+    for (unsigned raw_bank_number = threadIdx.x; raw_bank_number < number_of_raw_banks; raw_bank_number += blockDim.x) {
+      const auto bank_type = raw_event.bank_type(raw_bank_number);
+      // printf("Inside event %u raw bank number %u has type %u\n",event_number,raw_bank_number,static_cast<int>(bank_type));
+      //list of error banks
+      if (std::count(Allen::ErrorBanksList.begin(), Allen::ErrorBanksList.end(), static_cast<int>(bank_type))) {
+        error_banks_counter++;
+      }  
+    }//loop of all raw banks for specific event and specific subdetector
+    return std::forward_as_tuple(error_banks_counter);
+  }
+
+
+__device__ bool error_banks_line::error_banks_line_t::select(const Parameters&,  std::tuple<const unsigned> input)
+{
+  const unsigned number_of_error_banks = std::get<0>(input);
+  return number_of_error_banks > 0;
+}