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