From e88f256479d8955e45ab27ca2e4ecfe6900a0a87 Mon Sep 17 00:00:00 2001
From: thboettc <boettcts@ucmail.uc.edu>
Date: Mon, 25 Jul 2022 12:42:24 -0400
Subject: [PATCH 1/6] Remove limits on selected tracks and vertices in
 selreport writer

---
 .../selections/Hlt1/include/DecReporter.cuh   |  1 +
 .../Hlt1/include/MakeSelectedObjectLists.cuh  | 11 ++-
 .../selections/Hlt1/include/MakeSubBanks.cuh  | 10 +--
 device/selections/Hlt1/src/DecReporter.cu     |  8 ++
 .../Hlt1/src/MakeSelectedObjectLists.cu       | 84 +++++++++----------
 device/selections/Hlt1/src/MakeSubBanks.cu    | 32 +++----
 6 files changed, 77 insertions(+), 69 deletions(-)

diff --git a/device/selections/Hlt1/include/DecReporter.cuh b/device/selections/Hlt1/include/DecReporter.cuh
index 073d1e45d67..449ba299f0c 100644
--- a/device/selections/Hlt1/include/DecReporter.cuh
+++ b/device/selections/Hlt1/include/DecReporter.cuh
@@ -12,6 +12,7 @@ namespace dec_reporter {
     DEVICE_INPUT(dev_number_of_active_lines_t, unsigned) dev_number_of_active_lines;
     DEVICE_INPUT(dev_selections_t, bool) dev_selections;
     DEVICE_INPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
+    DEVICE_OUTPUT(dev_selected_candidates_counts_t, unsigned) dev_selected_candidates_counts;
     DEVICE_OUTPUT(dev_dec_reports_t, unsigned) dev_dec_reports;
     HOST_OUTPUT(host_dec_reports_t, unsigned) host_dec_reports;
     PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
diff --git a/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh b/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
index d1aeca46974..34fb4cfc4c0 100644
--- a/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
+++ b/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
@@ -17,6 +17,7 @@ namespace make_selected_object_lists {
   struct Parameters {
     HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
     HOST_INPUT(host_number_of_active_lines_t, unsigned) host_number_of_active_lines;
+    HOST_INPUT(host_max_objects_t, unsigned) host_max_objects;
     DEVICE_INPUT(dev_dec_reports_t, unsigned) dev_dec_reports;
     DEVICE_INPUT(dev_number_of_active_lines_t, unsigned) dev_number_of_active_lines;
     MASK_INPUT(dev_event_list_t) dev_event_list;
@@ -25,6 +26,7 @@ namespace make_selected_object_lists {
     dev_multi_event_particle_containers;
     DEVICE_INPUT(dev_selections_t, bool) dev_selections;
     DEVICE_INPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
+    DEVICE_INPUT(dev_max_objects_offsets_t, unsigned) dev_max_objects_offsets;
     DEVICE_OUTPUT(dev_candidate_count_t, unsigned) dev_candidate_count;
     DEVICE_OUTPUT(dev_sel_track_count_t, unsigned) dev_sel_track_count;
     DEVICE_OUTPUT(dev_sel_track_indices_t, unsigned) dev_sel_track_indices;
@@ -56,10 +58,8 @@ namespace make_selected_object_lists {
       DEPENDENCIES(dev_multi_event_particle_containers_t),
       Allen::Views::Physics::CompositeParticle*)
     dev_selected_composite_particle_ptrs;
-    PROPERTY(max_selected_tracks_t, "max_selected_tracks", "Maximum number of selected tracks per event.", unsigned)
-    max_selected_tracks;
-    PROPERTY(max_selected_svs_t, "max_selected_svs", "Maximum number of selected SVs per event.", unsigned)
-    max_selected_svs;
+    PROPERTY(max_children_per_object_t, "max_children_per_object", "Maximum number of children per selected object", unsigned)
+    max_children_per_object;
     PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
   };
 
@@ -83,8 +83,7 @@ namespace make_selected_object_lists {
 
   private:
     Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
-    Property<max_selected_tracks_t> m_max_selected_tracks {this, 100};
-    Property<max_selected_svs_t> m_max_selected_svs {this, 100};
+    Property<max_children_per_object_t> m_max_children_per_object {this, 4};
   };
 
 } // namespace make_selected_object_lists
\ No newline at end of file
diff --git a/device/selections/Hlt1/include/MakeSubBanks.cuh b/device/selections/Hlt1/include/MakeSubBanks.cuh
index f6df32ea69f..930131fbbe3 100644
--- a/device/selections/Hlt1/include/MakeSubBanks.cuh
+++ b/device/selections/Hlt1/include/MakeSubBanks.cuh
@@ -24,6 +24,7 @@ namespace make_subbanks {
     DEVICE_INPUT(dev_dec_reports_t, unsigned) dev_dec_reports;
     DEVICE_INPUT(dev_selections_t, bool) dev_selections;
     DEVICE_INPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
+    DEVICE_INPUT(dev_max_objects_offsets_t, unsigned) dev_max_objects_offsets;
     DEVICE_INPUT(dev_sel_count_t, unsigned) dev_sel_count;
     DEVICE_INPUT(dev_sel_list_t, unsigned) dev_sel_list;
     DEVICE_INPUT(dev_candidate_count_t, unsigned) dev_candidate_count;
@@ -52,10 +53,8 @@ namespace make_subbanks {
     DEVICE_OUTPUT(dev_rb_stdinfo_t, unsigned) dev_rb_stdinfo;
     // TODO: This needs to be the same as the properties in
     // MakeSelectedObjectLists. These should be saved as constants somewhere.
-    PROPERTY(max_selected_tracks_t, "max_selected_tracks", "Maximum number of selected tracks per event.", unsigned)
-    max_selected_tracks;
-    PROPERTY(max_selected_svs_t, "max_selected_svs", "Maximum number of selected SVs per event.", unsigned)
-    max_selected_svs;
+    PROPERTY(max_children_per_object_t, "max_children_per_object", "Maximum number of children per selected object", unsigned)
+    max_children_per_object;
     PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
   };
 
@@ -79,7 +78,6 @@ namespace make_subbanks {
 
   private:
     Property<block_dim_t> m_block_dim {this, {{64, 1, 1}}};
-    Property<max_selected_tracks_t> m_max_selected_tracks {this, 100};
-    Property<max_selected_svs_t> m_max_selected_svs {this, 100};
+    Property<max_children_per_object_t> m_max_children_per_object {this, 4};
   };
 } // namespace make_subbanks
\ No newline at end of file
diff --git a/device/selections/Hlt1/src/DecReporter.cu b/device/selections/Hlt1/src/DecReporter.cu
index 2baefee5c28..6705bc7c9f9 100644
--- a/device/selections/Hlt1/src/DecReporter.cu
+++ b/device/selections/Hlt1/src/DecReporter.cu
@@ -17,6 +17,8 @@ void dec_reporter::dec_reporter_t::set_arguments_size(
     arguments, (2 + first<host_number_of_active_lines_t>(arguments)) * first<host_number_of_events_t>(arguments));
   set_size<host_dec_reports_t>(
     arguments, (2 + first<host_number_of_active_lines_t>(arguments)) * first<host_number_of_events_t>(arguments));
+  set_size<dev_selected_candidates_counts_t>(
+    arguments, first<host_number_of_active_lines_t>(arguments)*first<host_number_of_events_t>(arguments));
 }
 
 void dec_reporter::dec_reporter_t::operator()(
@@ -27,6 +29,7 @@ void dec_reporter::dec_reporter_t::operator()(
   const Allen::Context& context) const
 {
   Allen::memset_async<host_dec_reports_t>(arguments, 0, context);
+  Allen::memset_async<dev_selected_candidates_counts_t>(arguments, 0, context);
 
   global_function(dec_reporter)(dim3(first<host_number_of_events_t>(arguments)), property<block_dim_t>(), context)(
     arguments);
@@ -52,6 +55,8 @@ __global__ void dec_reporter::dec_reporter(dec_reporter::Parameters parameters)
 
   uint32_t* event_dec_reports =
     parameters.dev_dec_reports + (2 + parameters.dev_number_of_active_lines[0]) * event_index;
+  unsigned* event_selected_candidates_counts =
+    parameters.dev_selected_candidates_counts + event_index * parameters.dev_number_of_active_lines[0];
 
   if (threadIdx.x == 0) {
     // Set TCK and taskID for each event dec report
@@ -68,6 +73,9 @@ __global__ void dec_reporter::dec_reporter(dec_reporter::Parameters parameters)
     auto decs = selections.get_span(line_index, event_index);
     for (unsigned i = 0; i < decs.size(); ++i) {
       final_decision |= decs[i];
+      if (decs[i]) {
+        event_selected_candidates_counts[i]++;
+      }
     }
 
     HltDecReport dec_report;
diff --git a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
index 708e8c9be2b..9264298d957 100644
--- a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
+++ b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
@@ -32,30 +32,25 @@ void make_selected_object_lists::make_selected_object_lists_t::set_arguments_siz
   set_size<dev_sel_sv_count_t>(arguments, first<host_number_of_events_t>(arguments));
   // These are effectively 3D arrays. Use the convention: X = candidate, Y = event, Z = line.
   set_size<dev_sel_track_indices_t>(
-    arguments,
-    property<max_selected_tracks_t>() * first<host_number_of_events_t>(arguments) *
-      first<host_number_of_active_lines_t>(arguments));
+    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
   set_size<dev_sel_sv_indices_t>(
-    arguments,
-    property<max_selected_svs_t>() * first<host_number_of_events_t>(arguments) *
-      first<host_number_of_active_lines_t>(arguments));
+    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
 
   // For saving selected candidates.
   // We could have multiple track and SV containers, so we can either set these
   // sizes arbitrarily, or create an algorithm to calculate them.
-  set_size<dev_selected_basic_particle_ptrs_t>(
-    arguments, first<host_number_of_events_t>(arguments) * property<max_selected_tracks_t>());
-  set_size<dev_selected_composite_particle_ptrs_t>(
-    arguments, first<host_number_of_events_t>(arguments) * property<max_selected_svs_t>());
+  set_size<dev_selected_basic_particle_ptrs_t>(arguments, 4*first<host_max_objects_t>(arguments));
+  set_size<dev_selected_composite_particle_ptrs_t>(arguments, 4*first<host_max_objects_t>(arguments));
 
   // For removing duplicates.
   set_size<dev_track_duplicate_map_t>(
-    arguments, first<host_number_of_events_t>(arguments) * property<max_selected_tracks_t>());
+    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
   set_size<dev_sv_duplicate_map_t>(
-    arguments, first<host_number_of_events_t>(arguments) * property<max_selected_svs_t>());
+    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
   set_size<dev_unique_track_list_t>(
-    arguments, first<host_number_of_events_t>(arguments) * property<max_selected_tracks_t>());
-  set_size<dev_unique_sv_list_t>(arguments, first<host_number_of_events_t>(arguments) * property<max_selected_svs_t>());
+    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
+  set_size<dev_unique_sv_list_t>(
+    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
   set_size<dev_unique_track_count_t>(arguments, first<host_number_of_events_t>(arguments));
   set_size<dev_unique_sv_count_t>(arguments, first<host_number_of_events_t>(arguments));
 
@@ -107,9 +102,11 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
   const unsigned total_events)
 {
   const auto event_number = blockIdx.x;
-  const unsigned selected_track_offset = parameters.max_selected_tracks * event_number;
-  const unsigned selected_sv_offset = parameters.max_selected_svs * event_number;
   const unsigned n_lines = parameters.dev_number_of_active_lines[0];
+  const unsigned n_children = parameters.max_children_per_object;
+  const unsigned* line_selected_object_offsets = 
+    parameters.dev_max_objects_offsets + n_lines * event_number;
+  const unsigned selected_object_offset = n_children*line_selected_object_offsets[0];
   const uint32_t* event_dec_reports =
     parameters.dev_dec_reports + (2 + parameters.dev_number_of_active_lines[0]) * event_number;
   unsigned* event_candidate_count =
@@ -138,9 +135,9 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
           const unsigned track_candidate_index = atomicAdd(event_candidate_count + line_index, 1);
           const unsigned track_insert_index = atomicAdd(parameters.dev_sel_track_count + event_number, 1);
           parameters.dev_sel_track_indices
-            [(event_number * n_lines + line_index) * parameters.max_selected_tracks + track_candidate_index] =
+            [n_children*line_selected_object_offsets[line_index] + track_candidate_index] =
             track_insert_index;
-          parameters.dev_selected_basic_particle_ptrs[selected_track_offset + track_insert_index] =
+          parameters.dev_selected_basic_particle_ptrs[selected_object_offset + track_insert_index] =
             const_cast<Allen::Views::Physics::BasicParticle*>(event_tracks.particle_pointer(track_index));
         }
       }
@@ -157,9 +154,9 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
           const unsigned sv_candidate_index = atomicAdd(event_candidate_count + line_index, 1);
           const unsigned sv_insert_index = atomicAdd(parameters.dev_sel_sv_count + event_number, 1);
           parameters.dev_sel_sv_indices
-            [(event_number * n_lines + line_index) * parameters.max_selected_svs + sv_candidate_index] =
+            [n_children*line_selected_object_offsets[line_index] + sv_candidate_index] =
             sv_insert_index;
-          parameters.dev_selected_composite_particle_ptrs[selected_sv_offset + sv_insert_index] =
+          parameters.dev_selected_composite_particle_ptrs[selected_object_offset + sv_insert_index] =
             const_cast<Allen::Views::Physics::CompositeParticle*>(&event_svs.particle(sv_index));
 
           // Parse the substructure.
@@ -171,13 +168,13 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
             if (substr->type_id() == Allen::TypeIDs::BasicParticle) { // Handle track substructures.
               const unsigned track_insert_index = atomicAdd(parameters.dev_sel_track_count + event_number, 1);
               const auto basic_substr = static_cast<const Allen::Views::Physics::BasicParticle*>(substr);
-              parameters.dev_selected_basic_particle_ptrs[selected_track_offset + track_insert_index] =
+              parameters.dev_selected_basic_particle_ptrs[selected_object_offset + track_insert_index] =
                 const_cast<Allen::Views::Physics::BasicParticle*>(basic_substr);
             }
             else { // Handle composite substructures.
               const unsigned sv_insert_index = atomicAdd(parameters.dev_sel_sv_count + event_number, 1);
               const auto composite_substr = static_cast<const Allen::Views::Physics::CompositeParticle*>(substr);
-              parameters.dev_selected_composite_particle_ptrs[selected_sv_offset + sv_insert_index] =
+              parameters.dev_selected_composite_particle_ptrs[selected_object_offset + sv_insert_index] =
                 const_cast<Allen::Views::Physics::CompositeParticle*>(composite_substr);
 
               // Manually handle sub-substructure to avoid recursion.
@@ -187,7 +184,7 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
                 // Assume all sub-substructures are BasicParticles.
                 const auto basic_subsubstr =
                   static_cast<const Allen::Views::Physics::BasicParticle*>(composite_substr->child(i_subsubstr));
-                parameters.dev_selected_basic_particle_ptrs[selected_sv_offset + track_insert_index] =
+                parameters.dev_selected_basic_particle_ptrs[selected_object_offset + track_insert_index] =
                   const_cast<Allen::Views::Physics::BasicParticle*>(basic_subsubstr);
               } // End sub-substructure loop.
             }
@@ -207,15 +204,15 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
     const auto n_selected_tracks = parameters.dev_sel_track_count[event_number];
     for (unsigned i_track = 0; i_track < n_selected_tracks; i_track += 1) {
       // Skip tracks that are already marked as duplicates.
-      if (parameters.dev_track_duplicate_map[selected_track_offset + i_track] >= 0) continue;
+      if (parameters.dev_track_duplicate_map[selected_object_offset + i_track] >= 0) continue;
       const unsigned track_insert = atomicAdd(parameters.dev_unique_track_count + event_number, 1);
-      parameters.dev_unique_track_list[selected_track_offset + track_insert] = i_track;
-      const auto trackA = parameters.dev_selected_basic_particle_ptrs[selected_track_offset + i_track];
+      parameters.dev_unique_track_list[selected_object_offset + track_insert] = i_track;
+      const auto trackA = parameters.dev_selected_basic_particle_ptrs[selected_object_offset + i_track];
       // Check for duplicate tracks.
       for (unsigned j_track = i_track + 1; j_track < n_selected_tracks; j_track++) {
-        const auto trackB = parameters.dev_selected_basic_particle_ptrs[selected_track_offset + j_track];
+        const auto trackB = parameters.dev_selected_basic_particle_ptrs[selected_object_offset + j_track];
         if (trackA == trackB) {
-          parameters.dev_track_duplicate_map[selected_track_offset + j_track] = i_track;
+          parameters.dev_track_duplicate_map[selected_object_offset + j_track] = i_track;
         }
       }
     }
@@ -223,15 +220,15 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
     const auto n_selected_svs = parameters.dev_sel_sv_count[event_number];
     for (unsigned i_sv = 0; i_sv < n_selected_svs; i_sv += 1) {
       // Skip SVs that are already marked as duplicates.
-      if (parameters.dev_sv_duplicate_map[selected_sv_offset + i_sv] >= 0) continue;
+      if (parameters.dev_sv_duplicate_map[selected_object_offset + i_sv] >= 0) continue;
       const unsigned sv_insert = atomicAdd(parameters.dev_unique_sv_count + event_number, 1);
-      parameters.dev_unique_sv_list[selected_sv_offset + sv_insert] = i_sv;
-      const auto svA = parameters.dev_selected_composite_particle_ptrs[selected_sv_offset + i_sv];
+      parameters.dev_unique_sv_list[selected_object_offset + sv_insert] = i_sv;
+      const auto svA = parameters.dev_selected_composite_particle_ptrs[selected_object_offset + i_sv];
       // Check for duplicate SVs.
       for (unsigned j_sv = i_sv + 1; j_sv < n_selected_svs; j_sv++) {
-        const auto svB = parameters.dev_selected_composite_particle_ptrs[selected_sv_offset + j_sv];
+        const auto svB = parameters.dev_selected_composite_particle_ptrs[selected_object_offset + j_sv];
         if (svA == svB) {
-          parameters.dev_sv_duplicate_map[selected_sv_offset + j_sv] = i_sv;
+          parameters.dev_sv_duplicate_map[selected_object_offset + j_sv] = i_sv;
         }
       }
     }
@@ -241,19 +238,22 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
 __global__ void make_selected_object_lists::calc_rb_sizes(make_selected_object_lists::Parameters parameters)
 {
   const auto event_number = blockIdx.x;
-  const unsigned selected_track_offset = parameters.max_selected_tracks * event_number;
-  const unsigned selected_sv_offset = parameters.max_selected_svs * event_number;
-  const auto event_track_ptrs = parameters.dev_selected_basic_particle_ptrs + selected_track_offset;
-  const auto event_unique_track_list = parameters.dev_unique_track_list + selected_track_offset;
-  const auto event_sv_ptrs = parameters.dev_selected_composite_particle_ptrs + selected_sv_offset;
-  const auto event_unique_sv_list = parameters.dev_unique_sv_list + selected_sv_offset;
+  const unsigned n_children = parameters.max_children_per_object;
+  const unsigned n_lines = parameters.dev_number_of_active_lines[0];
+  const unsigned* line_selected_object_offsets = 
+    parameters.dev_max_objects_offsets + n_lines * event_number;
+  const unsigned selected_object_offset = n_children*line_selected_object_offsets[0];
+  const auto event_track_ptrs = parameters.dev_selected_basic_particle_ptrs + selected_object_offset;
+  const auto event_unique_track_list = parameters.dev_unique_track_list + selected_object_offset;
+  const auto event_sv_ptrs = parameters.dev_selected_composite_particle_ptrs + selected_object_offset;
+  const auto event_unique_sv_list = parameters.dev_unique_sv_list + selected_object_offset;
   const auto n_selected_tracks = parameters.dev_unique_track_count[event_number];
   const auto n_selected_svs = parameters.dev_unique_sv_count[event_number];
   const uint32_t* event_dec_reports =
-    parameters.dev_dec_reports + (2 + parameters.dev_number_of_active_lines[0]) * event_number;
+    parameters.dev_dec_reports + (2 + n_lines) * event_number;
   unsigned* event_candidate_count =
-    parameters.dev_candidate_count + event_number * parameters.dev_number_of_active_lines[0];
-  unsigned* event_sel_list = parameters.dev_sel_list + event_number * parameters.dev_number_of_active_lines[0];
+    parameters.dev_candidate_count + event_number * n_lines;
+  unsigned* event_sel_list = parameters.dev_sel_list + event_number * n_lines;
 
   // Calculate the size of the hits bank.
   for (unsigned i_track = threadIdx.x; i_track < n_selected_tracks; i_track += blockDim.x) {
diff --git a/device/selections/Hlt1/src/MakeSubBanks.cu b/device/selections/Hlt1/src/MakeSubBanks.cu
index 636658336d3..206e986ffa0 100644
--- a/device/selections/Hlt1/src/MakeSubBanks.cu
+++ b/device/selections/Hlt1/src/MakeSubBanks.cu
@@ -48,21 +48,22 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
     unsigned* event_rb_substr = parameters.dev_rb_substr + parameters.dev_rb_substr_offsets[event_number];
     const unsigned event_rb_substr_size =
       parameters.dev_rb_substr_offsets[event_number + 1] - parameters.dev_rb_substr_offsets[event_number];
-    const unsigned sv_offset = parameters.max_selected_svs * event_number;
-    const unsigned track_offset = parameters.max_selected_tracks * event_number;
+    const unsigned n_lines = parameters.dev_number_of_active_lines[0];
+    const unsigned* line_object_offsets = parameters.dev_max_objects_offsets + n_lines*event_number;
+    const unsigned selected_object_offset = line_object_offsets[0];
     const unsigned n_tracks = parameters.dev_unique_track_count[event_number];
     const unsigned n_svs = parameters.dev_unique_sv_count[event_number];
     const unsigned n_sels = parameters.dev_sel_count[event_number];
-    const unsigned n_lines = parameters.dev_number_of_active_lines[0];
+    
 
     const unsigned sels_start_short = 2;
     const unsigned svs_start_short = sels_start_short + parameters.dev_substr_sel_size[event_number];
     const unsigned tracks_start_short = svs_start_short + parameters.dev_substr_sv_size[event_number];
 
-    const auto event_track_ptrs = parameters.dev_basic_particle_ptrs + track_offset;
-    const auto event_sv_ptrs = parameters.dev_composite_particle_ptrs + sv_offset;
-    const unsigned* event_unique_track_list = parameters.dev_unique_track_list + track_offset;
-    const unsigned* event_unique_sv_list = parameters.dev_unique_sv_list + sv_offset;
+    const auto event_track_ptrs = parameters.dev_basic_particle_ptrs + selected_object_offset;
+    const auto event_sv_ptrs = parameters.dev_composite_particle_ptrs + selected_object_offset;
+    const unsigned* event_unique_track_list = parameters.dev_unique_track_list + selected_object_offset;
+    const unsigned* event_unique_sv_list = parameters.dev_unique_sv_list + selected_object_offset;
 
     // Add the track substructures.
     // Each track substructure has one pointer to a sequence of LHCbIDs.
@@ -159,7 +160,7 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
       // Handle lines that select BasicParticles.
       if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventBasicParticles*>(mec)) {
         const unsigned* line_candidate_indices =
-          parameters.dev_sel_track_indices + (event_number * n_lines + line_id) * parameters.max_selected_tracks;
+          parameters.dev_sel_track_indices + line_object_offsets[i_line];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -170,13 +171,13 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
         insert_short++;
         for (unsigned i_cand = 0; i_cand < n_cand; i_cand++) {
           const unsigned i_track = line_candidate_indices[i_cand];
-          const unsigned track_index = parameters.dev_track_duplicate_map[track_offset + i_track] >= 0 ?
-                                         parameters.dev_track_duplicate_map[track_offset + i_track] :
+          const unsigned track_index = parameters.dev_track_duplicate_map[selected_object_offset + i_track] >= 0 ?
+                                         parameters.dev_track_duplicate_map[selected_object_offset + i_track] :
                                          i_track;
           unsigned obj_index = 0;
           // if (track_index < 0) track_index = i_track;
           for (unsigned j_track = 0; j_track < n_tracks; j_track++) {
-            const unsigned test_index = parameters.dev_unique_track_list[track_offset + j_track];
+            const unsigned test_index = parameters.dev_unique_track_list[selected_object_offset + j_track];
             if (track_index == test_index) {
               obj_index = n_sels + n_svs + j_track;
               break;
@@ -193,7 +194,7 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
       // Handle lines that select CompositeParticles.
       else if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventCompositeParticles*>(mec)) {
         const unsigned* line_candidate_indices =
-          parameters.dev_sel_sv_indices + (event_number * n_lines + line_id) * parameters.max_selected_svs;
+          parameters.dev_sel_sv_indices + line_object_offsets[line_id];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -204,8 +205,8 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
         insert_short++;
         for (unsigned i_cand = 0; i_cand < n_cand; i_cand++) {
           const unsigned i_sv = line_candidate_indices[i_cand];
-          const unsigned sv_index = parameters.dev_sv_duplicate_map[sv_offset + i_sv] >= 0 ?
-                                      parameters.dev_sv_duplicate_map[sv_offset + i_sv] :
+          const unsigned sv_index = parameters.dev_sv_duplicate_map[selected_object_offset + i_sv] >= 0 ?
+                                      parameters.dev_sv_duplicate_map[selected_object_offset + i_sv] :
                                       i_sv;
           unsigned obj_index = 0;
           // if (sv_index < 0) sv_index = i_sv;
@@ -355,7 +356,8 @@ __global__ void make_subbanks::make_rb_hits(make_subbanks::Parameters parameters
   const unsigned n_hit_sequences = parameters.dev_unique_track_count[event_number];
   unsigned* event_rb_hits = parameters.dev_rb_hits + parameters.dev_rb_hits_offsets[event_number];
   const unsigned bank_info_size = 1 + (n_hit_sequences / 2);
-  const unsigned track_offset = parameters.max_selected_tracks * event_number;
+  const unsigned n_lines = parameters.dev_number_of_active_lines[0];
+  const unsigned track_offset = parameters.dev_max_objects_offsets[event_number*n_lines];
 
   // Run sequentially over tracks and in parallel over hits. There will usually
   // only be ~1 selected track anyway.
-- 
GitLab


From 0967aad1e9a0853985d1d34f8651a238718c09ea Mon Sep 17 00:00:00 2001
From: thboettc <boettcts@ucmail.uc.edu>
Date: Mon, 25 Jul 2022 12:56:53 -0400
Subject: [PATCH 2/6] Update persistency configuration

---
 configuration/python/AllenConf/persistency.py | 10 +++++++++-
 1 file changed, 9 insertions(+), 1 deletion(-)

diff --git a/configuration/python/AllenConf/persistency.py b/configuration/python/AllenConf/persistency.py
index 2c8d7884507..bbc67d6e750 100644
--- a/configuration/python/AllenConf/persistency.py
+++ b/configuration/python/AllenConf/persistency.py
@@ -72,12 +72,18 @@ def make_sel_report_writer(lines, forward_tracks, secondary_vertices):
     dec_reporter = make_dec_reporter(lines)
     number_of_events = initialize_number_of_events()
 
+    prefix_sum_max_objects = make_algorithm(
+        host_prefix_sum_t,
+        name="prefix_sum_max_objects",
+        dev_input_buffer_t=dec_reporter.dev_selected_candidates_counts_t)
+
     make_selected_object_lists = make_algorithm(
         make_selected_object_lists_t,
         name="make_selected_object_lists",
         host_number_of_events_t=number_of_events["host_number_of_events"],
         host_number_of_active_lines_t=gather_selections.
         host_number_of_active_lines_t,
+        host_max_objects_t=prefix_sum_max_objects.host_total_sum_holder_t,
         dev_dec_reports_t=dec_reporter.dev_dec_reports_t,
         dev_number_of_active_lines_t=gather_selections.
         dev_number_of_active_lines_t,
@@ -85,7 +91,8 @@ def make_sel_report_writer(lines, forward_tracks, secondary_vertices):
         dev_multi_event_particle_containers_t=gather_selections.
         dev_particle_containers_t,
         dev_selections_t=gather_selections.dev_selections_t,
-        dev_selections_offsets_t=gather_selections.dev_selections_offsets_t)
+        dev_selections_offsets_t=gather_selections.dev_selections_offsets_t,
+        dev_max_objects_offsets_t=prefix_sum_max_objects.dev_output_buffer_t)
 
     prefix_sum_hits_size = make_algorithm(
         host_prefix_sum_t,
@@ -131,6 +138,7 @@ def make_sel_report_writer(lines, forward_tracks, secondary_vertices):
         dev_dec_reports_t=dec_reporter.dev_dec_reports_t,
         dev_selections_t=gather_selections.dev_selections_t,
         dev_selections_offsets_t=gather_selections.dev_selections_offsets_t,
+        dev_max_objects_offsets_t=prefix_sum_max_objects.dev_output_buffer_t,
         dev_sel_count_t=make_selected_object_lists.dev_sel_count_t,
         dev_sel_list_t=make_selected_object_lists.dev_sel_list_t,
         dev_candidate_count_t=make_selected_object_lists.dev_candidate_count_t,
-- 
GitLab


From 18c6a8e58ce4c28de03b7ec4e33c0cb366e1ce59 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Mon, 25 Jul 2022 16:59:55 +0000
Subject: [PATCH 3/6] Fixed formatting

patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/23524861
---
 .../Hlt1/include/MakeSelectedObjectLists.cuh  |  6 ++-
 .../selections/Hlt1/include/MakeSubBanks.cuh  |  6 ++-
 device/selections/Hlt1/src/DecReporter.cu     |  2 +-
 .../Hlt1/src/MakeSelectedObjectLists.cu       | 39 ++++++++-----------
 device/selections/Hlt1/src/MakeSubBanks.cu    | 11 ++----
 5 files changed, 32 insertions(+), 32 deletions(-)

diff --git a/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh b/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
index 34fb4cfc4c0..767e3e718a3 100644
--- a/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
+++ b/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
@@ -58,7 +58,11 @@ namespace make_selected_object_lists {
       DEPENDENCIES(dev_multi_event_particle_containers_t),
       Allen::Views::Physics::CompositeParticle*)
     dev_selected_composite_particle_ptrs;
-    PROPERTY(max_children_per_object_t, "max_children_per_object", "Maximum number of children per selected object", unsigned)
+    PROPERTY(
+      max_children_per_object_t,
+      "max_children_per_object",
+      "Maximum number of children per selected object",
+      unsigned)
     max_children_per_object;
     PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
   };
diff --git a/device/selections/Hlt1/include/MakeSubBanks.cuh b/device/selections/Hlt1/include/MakeSubBanks.cuh
index 930131fbbe3..83eeb5766f1 100644
--- a/device/selections/Hlt1/include/MakeSubBanks.cuh
+++ b/device/selections/Hlt1/include/MakeSubBanks.cuh
@@ -53,7 +53,11 @@ namespace make_subbanks {
     DEVICE_OUTPUT(dev_rb_stdinfo_t, unsigned) dev_rb_stdinfo;
     // TODO: This needs to be the same as the properties in
     // MakeSelectedObjectLists. These should be saved as constants somewhere.
-    PROPERTY(max_children_per_object_t, "max_children_per_object", "Maximum number of children per selected object", unsigned)
+    PROPERTY(
+      max_children_per_object_t,
+      "max_children_per_object",
+      "Maximum number of children per selected object",
+      unsigned)
     max_children_per_object;
     PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim;
   };
diff --git a/device/selections/Hlt1/src/DecReporter.cu b/device/selections/Hlt1/src/DecReporter.cu
index 6705bc7c9f9..2164a54db4a 100644
--- a/device/selections/Hlt1/src/DecReporter.cu
+++ b/device/selections/Hlt1/src/DecReporter.cu
@@ -18,7 +18,7 @@ void dec_reporter::dec_reporter_t::set_arguments_size(
   set_size<host_dec_reports_t>(
     arguments, (2 + first<host_number_of_active_lines_t>(arguments)) * first<host_number_of_events_t>(arguments));
   set_size<dev_selected_candidates_counts_t>(
-    arguments, first<host_number_of_active_lines_t>(arguments)*first<host_number_of_events_t>(arguments));
+    arguments, first<host_number_of_active_lines_t>(arguments) * first<host_number_of_events_t>(arguments));
 }
 
 void dec_reporter::dec_reporter_t::operator()(
diff --git a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
index 9264298d957..30d7abd145f 100644
--- a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
+++ b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
@@ -32,25 +32,25 @@ void make_selected_object_lists::make_selected_object_lists_t::set_arguments_siz
   set_size<dev_sel_sv_count_t>(arguments, first<host_number_of_events_t>(arguments));
   // These are effectively 3D arrays. Use the convention: X = candidate, Y = event, Z = line.
   set_size<dev_sel_track_indices_t>(
-    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
   set_size<dev_sel_sv_indices_t>(
-    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
 
   // For saving selected candidates.
   // We could have multiple track and SV containers, so we can either set these
   // sizes arbitrarily, or create an algorithm to calculate them.
-  set_size<dev_selected_basic_particle_ptrs_t>(arguments, 4*first<host_max_objects_t>(arguments));
-  set_size<dev_selected_composite_particle_ptrs_t>(arguments, 4*first<host_max_objects_t>(arguments));
+  set_size<dev_selected_basic_particle_ptrs_t>(arguments, 4 * first<host_max_objects_t>(arguments));
+  set_size<dev_selected_composite_particle_ptrs_t>(arguments, 4 * first<host_max_objects_t>(arguments));
 
   // For removing duplicates.
   set_size<dev_track_duplicate_map_t>(
-    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
   set_size<dev_sv_duplicate_map_t>(
-    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
   set_size<dev_unique_track_list_t>(
-    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
   set_size<dev_unique_sv_list_t>(
-    arguments, property<max_children_per_object_t>()*first<host_max_objects_t>(arguments));
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
   set_size<dev_unique_track_count_t>(arguments, first<host_number_of_events_t>(arguments));
   set_size<dev_unique_sv_count_t>(arguments, first<host_number_of_events_t>(arguments));
 
@@ -104,9 +104,8 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
   const auto event_number = blockIdx.x;
   const unsigned n_lines = parameters.dev_number_of_active_lines[0];
   const unsigned n_children = parameters.max_children_per_object;
-  const unsigned* line_selected_object_offsets = 
-    parameters.dev_max_objects_offsets + n_lines * event_number;
-  const unsigned selected_object_offset = n_children*line_selected_object_offsets[0];
+  const unsigned* line_selected_object_offsets = parameters.dev_max_objects_offsets + n_lines * event_number;
+  const unsigned selected_object_offset = n_children * line_selected_object_offsets[0];
   const uint32_t* event_dec_reports =
     parameters.dev_dec_reports + (2 + parameters.dev_number_of_active_lines[0]) * event_number;
   unsigned* event_candidate_count =
@@ -134,8 +133,8 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
         if (decs[track_index]) {
           const unsigned track_candidate_index = atomicAdd(event_candidate_count + line_index, 1);
           const unsigned track_insert_index = atomicAdd(parameters.dev_sel_track_count + event_number, 1);
-          parameters.dev_sel_track_indices
-            [n_children*line_selected_object_offsets[line_index] + track_candidate_index] =
+          parameters
+            .dev_sel_track_indices[n_children * line_selected_object_offsets[line_index] + track_candidate_index] =
             track_insert_index;
           parameters.dev_selected_basic_particle_ptrs[selected_object_offset + track_insert_index] =
             const_cast<Allen::Views::Physics::BasicParticle*>(event_tracks.particle_pointer(track_index));
@@ -153,8 +152,7 @@ __global__ void make_selected_object_lists::make_selected_object_lists(
         if (decs[sv_index]) {
           const unsigned sv_candidate_index = atomicAdd(event_candidate_count + line_index, 1);
           const unsigned sv_insert_index = atomicAdd(parameters.dev_sel_sv_count + event_number, 1);
-          parameters.dev_sel_sv_indices
-            [n_children*line_selected_object_offsets[line_index] + sv_candidate_index] =
+          parameters.dev_sel_sv_indices[n_children * line_selected_object_offsets[line_index] + sv_candidate_index] =
             sv_insert_index;
           parameters.dev_selected_composite_particle_ptrs[selected_object_offset + sv_insert_index] =
             const_cast<Allen::Views::Physics::CompositeParticle*>(&event_svs.particle(sv_index));
@@ -240,19 +238,16 @@ __global__ void make_selected_object_lists::calc_rb_sizes(make_selected_object_l
   const auto event_number = blockIdx.x;
   const unsigned n_children = parameters.max_children_per_object;
   const unsigned n_lines = parameters.dev_number_of_active_lines[0];
-  const unsigned* line_selected_object_offsets = 
-    parameters.dev_max_objects_offsets + n_lines * event_number;
-  const unsigned selected_object_offset = n_children*line_selected_object_offsets[0];
+  const unsigned* line_selected_object_offsets = parameters.dev_max_objects_offsets + n_lines * event_number;
+  const unsigned selected_object_offset = n_children * line_selected_object_offsets[0];
   const auto event_track_ptrs = parameters.dev_selected_basic_particle_ptrs + selected_object_offset;
   const auto event_unique_track_list = parameters.dev_unique_track_list + selected_object_offset;
   const auto event_sv_ptrs = parameters.dev_selected_composite_particle_ptrs + selected_object_offset;
   const auto event_unique_sv_list = parameters.dev_unique_sv_list + selected_object_offset;
   const auto n_selected_tracks = parameters.dev_unique_track_count[event_number];
   const auto n_selected_svs = parameters.dev_unique_sv_count[event_number];
-  const uint32_t* event_dec_reports =
-    parameters.dev_dec_reports + (2 + n_lines) * event_number;
-  unsigned* event_candidate_count =
-    parameters.dev_candidate_count + event_number * n_lines;
+  const uint32_t* event_dec_reports = parameters.dev_dec_reports + (2 + n_lines) * event_number;
+  unsigned* event_candidate_count = parameters.dev_candidate_count + event_number * n_lines;
   unsigned* event_sel_list = parameters.dev_sel_list + event_number * n_lines;
 
   // Calculate the size of the hits bank.
diff --git a/device/selections/Hlt1/src/MakeSubBanks.cu b/device/selections/Hlt1/src/MakeSubBanks.cu
index 206e986ffa0..f95502618de 100644
--- a/device/selections/Hlt1/src/MakeSubBanks.cu
+++ b/device/selections/Hlt1/src/MakeSubBanks.cu
@@ -49,12 +49,11 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
     const unsigned event_rb_substr_size =
       parameters.dev_rb_substr_offsets[event_number + 1] - parameters.dev_rb_substr_offsets[event_number];
     const unsigned n_lines = parameters.dev_number_of_active_lines[0];
-    const unsigned* line_object_offsets = parameters.dev_max_objects_offsets + n_lines*event_number;
+    const unsigned* line_object_offsets = parameters.dev_max_objects_offsets + n_lines * event_number;
     const unsigned selected_object_offset = line_object_offsets[0];
     const unsigned n_tracks = parameters.dev_unique_track_count[event_number];
     const unsigned n_svs = parameters.dev_unique_sv_count[event_number];
     const unsigned n_sels = parameters.dev_sel_count[event_number];
-    
 
     const unsigned sels_start_short = 2;
     const unsigned svs_start_short = sels_start_short + parameters.dev_substr_sel_size[event_number];
@@ -159,8 +158,7 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
 
       // Handle lines that select BasicParticles.
       if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventBasicParticles*>(mec)) {
-        const unsigned* line_candidate_indices =
-          parameters.dev_sel_track_indices + line_object_offsets[i_line];
+        const unsigned* line_candidate_indices = parameters.dev_sel_track_indices + line_object_offsets[i_line];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -193,8 +191,7 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
       }
       // Handle lines that select CompositeParticles.
       else if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventCompositeParticles*>(mec)) {
-        const unsigned* line_candidate_indices =
-          parameters.dev_sel_sv_indices + line_object_offsets[line_id];
+        const unsigned* line_candidate_indices = parameters.dev_sel_sv_indices + line_object_offsets[line_id];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -357,7 +354,7 @@ __global__ void make_subbanks::make_rb_hits(make_subbanks::Parameters parameters
   unsigned* event_rb_hits = parameters.dev_rb_hits + parameters.dev_rb_hits_offsets[event_number];
   const unsigned bank_info_size = 1 + (n_hit_sequences / 2);
   const unsigned n_lines = parameters.dev_number_of_active_lines[0];
-  const unsigned track_offset = parameters.dev_max_objects_offsets[event_number*n_lines];
+  const unsigned track_offset = parameters.dev_max_objects_offsets[event_number * n_lines];
 
   // Run sequentially over tracks and in parallel over hits. There will usually
   // only be ~1 selected track anyway.
-- 
GitLab


From deb7011fb1cbd65e65825fc8c9631b0f22ba6e76 Mon Sep 17 00:00:00 2001
From: thboettc <boettcts@ucmail.uc.edu>
Date: Mon, 25 Jul 2022 18:46:57 -0400
Subject: [PATCH 4/6] Fix selected object offsets

---
 device/selections/Hlt1/src/MakeSelectedObjectLists.cu |  4 ++--
 device/selections/Hlt1/src/MakeSubBanks.cu            | 10 +++++-----
 2 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
index 30d7abd145f..d193e4cb34b 100644
--- a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
+++ b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
@@ -39,8 +39,8 @@ void make_selected_object_lists::make_selected_object_lists_t::set_arguments_siz
   // For saving selected candidates.
   // We could have multiple track and SV containers, so we can either set these
   // sizes arbitrarily, or create an algorithm to calculate them.
-  set_size<dev_selected_basic_particle_ptrs_t>(arguments, 4 * first<host_max_objects_t>(arguments));
-  set_size<dev_selected_composite_particle_ptrs_t>(arguments, 4 * first<host_max_objects_t>(arguments));
+  set_size<dev_selected_basic_particle_ptrs_t>(arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
+  set_size<dev_selected_composite_particle_ptrs_t>(arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
 
   // For removing duplicates.
   set_size<dev_track_duplicate_map_t>(
diff --git a/device/selections/Hlt1/src/MakeSubBanks.cu b/device/selections/Hlt1/src/MakeSubBanks.cu
index f95502618de..4172b365eda 100644
--- a/device/selections/Hlt1/src/MakeSubBanks.cu
+++ b/device/selections/Hlt1/src/MakeSubBanks.cu
@@ -50,7 +50,8 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
       parameters.dev_rb_substr_offsets[event_number + 1] - parameters.dev_rb_substr_offsets[event_number];
     const unsigned n_lines = parameters.dev_number_of_active_lines[0];
     const unsigned* line_object_offsets = parameters.dev_max_objects_offsets + n_lines * event_number;
-    const unsigned selected_object_offset = line_object_offsets[0];
+    const unsigned n_children = parameters.max_children_per_object;
+    const unsigned selected_object_offset = n_children*line_object_offsets[0];
     const unsigned n_tracks = parameters.dev_unique_track_count[event_number];
     const unsigned n_svs = parameters.dev_unique_sv_count[event_number];
     const unsigned n_sels = parameters.dev_sel_count[event_number];
@@ -105,7 +106,6 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
         event_rb_substr[i_word] = (event_rb_substr[i_word] & ~(mask << bits)) | (sv_struct << bits);
       }
       for (unsigned i_substr = 0; i_substr < n_substr; i_substr++) {
-
         // Find the location of the substructure in the bank.
         const auto substr = sv->child(i_substr);
         const auto basic_substr = Allen::dyn_cast<const Allen::Views::Physics::BasicParticle*>(substr);
@@ -158,7 +158,7 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
 
       // Handle lines that select BasicParticles.
       if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventBasicParticles*>(mec)) {
-        const unsigned* line_candidate_indices = parameters.dev_sel_track_indices + line_object_offsets[i_line];
+        const unsigned* line_candidate_indices = parameters.dev_sel_track_indices + n_children*line_object_offsets[line_id];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -191,7 +191,7 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
       }
       // Handle lines that select CompositeParticles.
       else if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventCompositeParticles*>(mec)) {
-        const unsigned* line_candidate_indices = parameters.dev_sel_sv_indices + line_object_offsets[line_id];
+        const unsigned* line_candidate_indices = parameters.dev_sel_sv_indices + n_children*line_object_offsets[line_id];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -354,7 +354,7 @@ __global__ void make_subbanks::make_rb_hits(make_subbanks::Parameters parameters
   unsigned* event_rb_hits = parameters.dev_rb_hits + parameters.dev_rb_hits_offsets[event_number];
   const unsigned bank_info_size = 1 + (n_hit_sequences / 2);
   const unsigned n_lines = parameters.dev_number_of_active_lines[0];
-  const unsigned track_offset = parameters.dev_max_objects_offsets[event_number * n_lines];
+  const unsigned track_offset = parameters.max_children_per_object*parameters.dev_max_objects_offsets[event_number * n_lines];
 
   // Run sequentially over tracks and in parallel over hits. There will usually
   // only be ~1 selected track anyway.
-- 
GitLab


From cef37f388aac7ee87df47a8e5c1ebbadb9ebe568 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Mon, 25 Jul 2022 22:48:01 +0000
Subject: [PATCH 5/6] Fixed formatting

patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/23530339
---
 device/selections/Hlt1/src/MakeSelectedObjectLists.cu |  6 ++++--
 device/selections/Hlt1/src/MakeSubBanks.cu            | 11 +++++++----
 2 files changed, 11 insertions(+), 6 deletions(-)

diff --git a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
index d193e4cb34b..5a90fe33cfe 100644
--- a/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
+++ b/device/selections/Hlt1/src/MakeSelectedObjectLists.cu
@@ -39,8 +39,10 @@ void make_selected_object_lists::make_selected_object_lists_t::set_arguments_siz
   // For saving selected candidates.
   // We could have multiple track and SV containers, so we can either set these
   // sizes arbitrarily, or create an algorithm to calculate them.
-  set_size<dev_selected_basic_particle_ptrs_t>(arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
-  set_size<dev_selected_composite_particle_ptrs_t>(arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
+  set_size<dev_selected_basic_particle_ptrs_t>(
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
+  set_size<dev_selected_composite_particle_ptrs_t>(
+    arguments, property<max_children_per_object_t>() * first<host_max_objects_t>(arguments));
 
   // For removing duplicates.
   set_size<dev_track_duplicate_map_t>(
diff --git a/device/selections/Hlt1/src/MakeSubBanks.cu b/device/selections/Hlt1/src/MakeSubBanks.cu
index 4172b365eda..2ac3a22fb98 100644
--- a/device/selections/Hlt1/src/MakeSubBanks.cu
+++ b/device/selections/Hlt1/src/MakeSubBanks.cu
@@ -51,7 +51,7 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
     const unsigned n_lines = parameters.dev_number_of_active_lines[0];
     const unsigned* line_object_offsets = parameters.dev_max_objects_offsets + n_lines * event_number;
     const unsigned n_children = parameters.max_children_per_object;
-    const unsigned selected_object_offset = n_children*line_object_offsets[0];
+    const unsigned selected_object_offset = n_children * line_object_offsets[0];
     const unsigned n_tracks = parameters.dev_unique_track_count[event_number];
     const unsigned n_svs = parameters.dev_unique_sv_count[event_number];
     const unsigned n_sels = parameters.dev_sel_count[event_number];
@@ -158,7 +158,8 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
 
       // Handle lines that select BasicParticles.
       if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventBasicParticles*>(mec)) {
-        const unsigned* line_candidate_indices = parameters.dev_sel_track_indices + n_children*line_object_offsets[line_id];
+        const unsigned* line_candidate_indices =
+          parameters.dev_sel_track_indices + n_children * line_object_offsets[line_id];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -191,7 +192,8 @@ __global__ void make_subbanks::make_rb_substr(make_subbanks::Parameters paramete
       }
       // Handle lines that select CompositeParticles.
       else if (Allen::dyn_cast<const Allen::Views::Physics::MultiEventCompositeParticles*>(mec)) {
-        const unsigned* line_candidate_indices = parameters.dev_sel_sv_indices + n_children*line_object_offsets[line_id];
+        const unsigned* line_candidate_indices =
+          parameters.dev_sel_sv_indices + n_children * line_object_offsets[line_id];
         unsigned n_cand = event_candidate_offsets[line_id + 1] - event_candidate_offsets[line_id];
         unsigned i_word = insert_short / 2;
         unsigned i_part = insert_short % 2;
@@ -354,7 +356,8 @@ __global__ void make_subbanks::make_rb_hits(make_subbanks::Parameters parameters
   unsigned* event_rb_hits = parameters.dev_rb_hits + parameters.dev_rb_hits_offsets[event_number];
   const unsigned bank_info_size = 1 + (n_hit_sequences / 2);
   const unsigned n_lines = parameters.dev_number_of_active_lines[0];
-  const unsigned track_offset = parameters.max_children_per_object*parameters.dev_max_objects_offsets[event_number * n_lines];
+  const unsigned track_offset =
+    parameters.max_children_per_object * parameters.dev_max_objects_offsets[event_number * n_lines];
 
   // Run sequentially over tracks and in parallel over hits. There will usually
   // only be ~1 selected track anyway.
-- 
GitLab


From da2cb49b81e113a37e1a5400a35f3a00cb8bb75a Mon Sep 17 00:00:00 2001
From: thboettc <boettcts@ucmail.uc.edu>
Date: Mon, 25 Jul 2022 21:32:09 -0400
Subject: [PATCH 6/6] Fix line index in dec reporter

---
 device/selections/Hlt1/src/DecReporter.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/device/selections/Hlt1/src/DecReporter.cu b/device/selections/Hlt1/src/DecReporter.cu
index 2164a54db4a..1bc97faca6b 100644
--- a/device/selections/Hlt1/src/DecReporter.cu
+++ b/device/selections/Hlt1/src/DecReporter.cu
@@ -74,7 +74,7 @@ __global__ void dec_reporter::dec_reporter(dec_reporter::Parameters parameters)
     for (unsigned i = 0; i < decs.size(); ++i) {
       final_decision |= decs[i];
       if (decs[i]) {
-        event_selected_candidates_counts[i]++;
+        event_selected_candidates_counts[line_index]++;
       }
     }
 
-- 
GitLab