From 96ceeed879abfe8f1f76f940c5f570412bb36260 Mon Sep 17 00:00:00 2001 From: Adrian Casais Vidal Date: Tue, 27 Sep 2022 11:18:46 +0200 Subject: [PATCH 1/6] Added StandaloneMuon, VelMuon and FitMuon (to be removed) with their validators. (Help from @dcampora through extreme programming) --- CMakeLists.txt | 4 + checker/CMakeLists.txt | 1 + .../plotting/muon/muon_id_efficiency_plots.py | 28 +- .../plotting/tracking/momentum_resolution.py | 7 +- checker/tracking/include/TrackChecker.h | 3 + checker/tracking/src/PrepareTracks.cpp | 47 +- configuration/python/AllenConf/HLT1.py | 17 +- .../python/AllenConf/hlt1_muon_lines.py | 54 ++- .../python/AllenConf/hlt1_reconstruction.py | 27 +- .../python/AllenConf/muon_reconstruction.py | 137 +++++- configuration/python/AllenConf/validators.py | 30 +- .../AllenSequences/hlt1_complex_validation.py | 8 +- .../trackmatching_veloscifi_validation.py | 4 +- .../common/include/CopyTrackParameters.cuh | 55 +++ .../muon/include/MuonDefinitions.cuh | 203 ++++++++ device/muon/CMakeLists.txt | 10 +- .../include/ConsolidateMuon.cuh | 58 +++ .../include/ConsolidateVeloMuon.cuh | 58 +++ .../match_velo_muon/include/FindMuonHits.cuh | 71 +++ .../muon/match_velo_muon/include/FitMuon.cuh | 57 +++ .../match_velo_muon/include/MatchVeloMuon.cuh | 70 +++ .../match_velo_muon/src/ConsolidateMuon.cu | 52 ++ .../src/ConsolidateVeloMuon.cu | 69 +++ .../muon/match_velo_muon/src/FindMuonHits.cu | 456 ++++++++++++++++++ device/muon/match_velo_muon/src/FitMuon.cu | 178 +++++++ .../muon/match_velo_muon/src/MatchVeloMuon.cu | 287 +++++++++++ .../lines/muon/include/OneMuonTrackLine.cuh | 85 ++++ .../muon/include/OneVeloMuonTrackLine.cuh | 87 ++++ .../lines/muon/src/OneMuonTrackLine.cu | 25 + .../lines/muon/src/OneVeloMuonTrackLine.cu | 30 ++ .../validators/include/LongTrackValidator.cuh | 9 + device/validators/include/MuonValidator.cuh | 5 + device/validators/src/LongTrackValidator.cu | 37 +- device/validators/src/MuonValidator.cu | 25 +- host/validators/include/HostVeloValidator.h | 1 + host/validators/src/HostVeloValidator.cpp | 1 + stream/sequence/include/Constants.cuh | 6 +- stream/sequence/src/Constants.cpp | 13 + 38 files changed, 2270 insertions(+), 45 deletions(-) create mode 100644 device/muon/match_velo_muon/include/ConsolidateMuon.cuh create mode 100644 device/muon/match_velo_muon/include/ConsolidateVeloMuon.cuh create mode 100644 device/muon/match_velo_muon/include/FindMuonHits.cuh create mode 100644 device/muon/match_velo_muon/include/FitMuon.cuh create mode 100644 device/muon/match_velo_muon/include/MatchVeloMuon.cuh create mode 100644 device/muon/match_velo_muon/src/ConsolidateMuon.cu create mode 100644 device/muon/match_velo_muon/src/ConsolidateVeloMuon.cu create mode 100644 device/muon/match_velo_muon/src/FindMuonHits.cu create mode 100644 device/muon/match_velo_muon/src/FitMuon.cu create mode 100644 device/muon/match_velo_muon/src/MatchVeloMuon.cu create mode 100644 device/selections/lines/muon/include/OneMuonTrackLine.cuh create mode 100644 device/selections/lines/muon/include/OneVeloMuonTrackLine.cuh create mode 100644 device/selections/lines/muon/src/OneMuonTrackLine.cu create mode 100644 device/selections/lines/muon/src/OneVeloMuonTrackLine.cu diff --git a/CMakeLists.txt b/CMakeLists.txt index 5af62d9cb69..facbc2401b9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -726,6 +726,7 @@ if (NOT STANDALONE) lhcb_finalize_configuration() endif() + # Download external files if in STANDALONE if (STANDALONE) file(MAKE_DIRECTORY "${PROJECT_BINARY_DIR}/external") @@ -748,3 +749,6 @@ if (STANDALONE) message(STATUS "PARAMFILESROOT set to ${PARAMFILESROOT}") endif() endif() + +file(MAKE_DIRECTORY ${CMAKE_BINARY_DIR}/plotsfornote) +file(MAKE_DIRECTORY ${CMAKE_BINARY_DIR}/plotsfornote_root) \ No newline at end of file diff --git a/checker/CMakeLists.txt b/checker/CMakeLists.txt index e945a4b5634..301949babbf 100644 --- a/checker/CMakeLists.txt +++ b/checker/CMakeLists.txt @@ -5,3 +5,4 @@ add_subdirectory(clustering) add_subdirectory(tracking) add_subdirectory(pv) add_subdirectory(selections) +file(COPY plotting DESTINATION ./) \ No newline at end of file diff --git a/checker/plotting/muon/muon_id_efficiency_plots.py b/checker/plotting/muon/muon_id_efficiency_plots.py index 0c3623090c1..5a0e4f0451b 100644 --- a/checker/plotting/muon/muon_id_efficiency_plots.py +++ b/checker/plotting/muon/muon_id_efficiency_plots.py @@ -28,6 +28,14 @@ from common.LHCbStyle import * from common.Legend import * from common.ConfigHistos import * +PLACES = [ + (start, stop - y_width, start + x_width, stop), # top left opt + (start, start, start + x_width, start + y_width), # bottom left opt + (stop - x_width, stop - y_width, stop, stop), # top right opt + (stop - x_width, start, stop, start + y_width), # bottom right opt + (stop - x_width, 0.6 - y_width / 2, stop, 0.6 + y_width / 2), # right + (start, 0.5 - y_width / 2, start + x_width, 0.5 + y_width / 2) +] # left def getEfficiencyHistoNames(): @@ -65,7 +73,6 @@ def muonCategoryDict(): # ROOT.TFile.Open("../../../output/Z2MuMu/PrCheckerPLots-Z2MuMu.root", "read"), # ] - f = [ROOT.TFile.Open("../../../output/PrCheckerPlots.root", "read")] outputfile = ROOT.TFile( "../../../plotsfornote_root/muon_id_efficiency_plots.root", "recreate") @@ -87,14 +94,14 @@ for category in muonCategories: title = muonCatDict[category]["title"] + " vs. " + histo canvas = ROOT.TCanvas(title, title) ROOT.gPad.SetTicks() - numeratorName = "Forward/" + muonCatDict[category][ + numeratorName = "velo_tracks muon/" + muonCatDict[category][ "numerator"] + efficiencyHistoDict[histo][ "variable"] + "_reconstructed" print("Opening " + numeratorName) numerator = f[0].Get(numeratorName) for infile in f[1:]: numerator.Add(infile.Get(numeratorName)) - denominatorName = "Forward/" + muonCatDict[category][ + denominatorName = "velo_tracks muon/" + muonCatDict[category][ "denominator"] + efficiencyHistoDict[histo][ "variable"] + "_reconstructible" denominator = f[0].Get(denominatorName) @@ -127,7 +134,7 @@ for category in muonCategories: numerator.Draw("hist bar same") if (category == "matched_isMuon"): - place = find_place(canvas, 3) + place = PLACES[0] else: place = PLACES[0] #place = find_place(canvas, 0) @@ -143,14 +150,15 @@ for category in muonCategories: # Draw second y axis low = 0 high = 1.05 - axis = ROOT.TGaxis(gPad.GetUxmax(), gPad.GetUymin(),gPad.GetUxmax(),gPad.GetUymax(),low,high,510,"+L") + axis = ROOT.TGaxis(gPad.GetUxmax(), gPad.GetUymin(), gPad.GetUxmax(), + gPad.GetUymax(), low, high, 510, "+L") axis.SetTitleFont(132) axis.SetTitleSize(0.06) axis.SetTitleOffset(0.55) axis.SetTitle("Number of events [a.u.]") axis.SetLabelSize(0) axis.Draw() - + canvas.Write() cleantitle = muonCatDict[category]["title"].replace(" ", "").replace( ",", "_").replace("<", "_") @@ -162,9 +170,9 @@ for histo in ghostHistos: title = "muon ID in ghost tracks vs. " + histo canvas = ROOT.TCanvas(title, title) ROOT.gPad.SetTicks() - numeratorName = "Forward/ghost_isMuon_" + efficiencyHistoDict[histo][ - "variable"] + "_reconstructed" - denominatorName = "Forward/" + histo + "_Ghosts" + numeratorName = "velo_tracks muon/ghost_isMuon_" + efficiencyHistoDict[ + histo]["variable"] + "_reconstructed" + denominatorName = "velo_tracks muon/" + histo + "_Ghosts" print("Opening " + numeratorName) print("Opening " + denominatorName) @@ -194,7 +202,7 @@ for histo in ghostHistos: numerator.SetLineColor(ROOT.kWhite) numerator.Draw("hist bar same") - place = find_place(canvas, 0) + place = PLACES[0] legend = TLegend(place[0], place[1], place[2], place[3]) legend.AddEntry(g_efficiency, "muon ID in ghost tracks", "ep") legend.AddEntry(numerator, diff --git a/checker/plotting/tracking/momentum_resolution.py b/checker/plotting/tracking/momentum_resolution.py index 7e91547c9d6..24866e349f4 100644 --- a/checker/plotting/tracking/momentum_resolution.py +++ b/checker/plotting/tracking/momentum_resolution.py @@ -53,7 +53,7 @@ def getHistos(): def getTrackers(): - return ["Upstream", "Forward"] + return ["veloUT_validator", "forward_validator", "velo_muon_tracks"] def getResolutionInSlices(histo2D, var, var_dict): @@ -68,9 +68,10 @@ def getResolutionInSlices(histo2D, var, var_dict): histo1D = histo2D.ProjectionY("_py", i, i, "") if histo1D.GetEntries() >= 100: # fit Gaussian - if tracker == "Forward": + g1 = ROOT.TF1("g1", "gaus", -0.5, 0.5) + if tracker == "forward_validator": g1 = ROOT.TF1("g1", "gaus", -0.05, 0.05) - elif tracker == "Upstream": + elif tracker == "veloUT_validator": g1 = ROOT.TF1("g1", "gaus", -0.5, 0.5) #histo1D.GetYaxis().SetTitle("Entries") #histo1D.GetXaxis().SetTitle("Resolution (%/100)") diff --git a/checker/tracking/include/TrackChecker.h b/checker/tracking/include/TrackChecker.h index 105ac5b8a29..9e049f4020f 100644 --- a/checker/tracking/include/TrackChecker.h +++ b/checker/tracking/include/TrackChecker.h @@ -418,6 +418,9 @@ public: if (std::is_same_v && mcp.hasVelo && mcp.hasUT) { m_histos->fillMomentumResolutionHisto(mcp, track.p, track.qop); } + if (std::is_same_v && mcp.hasVelo && track.p > 0) { + m_histos->fillMomentumResolutionHisto(mcp, track.p, track.qop); + } } for (auto& category : m_categories) { diff --git a/checker/tracking/src/PrepareTracks.cpp b/checker/tracking/src/PrepareTracks.cpp index 5452c26b24c..befff05b539 100644 --- a/checker/tracking/src/PrepareTracks.cpp +++ b/checker/tracking/src/PrepareTracks.cpp @@ -16,6 +16,7 @@ #include "CopyTrackParameters.cuh" #include "ROOTHeaders.h" #include +#include "MuonDefinitions.cuh" std::vector prepareVeloTracks( const unsigned number_of_events, @@ -28,10 +29,8 @@ std::vector prepareVeloTracks( std::vector checker_tracks(event_list.size()); for (unsigned i = 0; i < event_list.size(); i++) { const auto event_number = event_list[i]; - // Tracks of this event auto& tracks = checker_tracks[i]; - Velo::Consolidated::ConstTracks velo_tracks { track_atomics.data(), track_hit_number.data(), event_number, number_of_events}; const unsigned number_of_tracks_event = velo_tracks.number_of_tracks(event_number); @@ -40,7 +39,6 @@ std::vector prepareVeloTracks( for (unsigned i_track = 0; i_track < number_of_tracks_event; i_track++) { auto& t = tracks[i_track]; t.p = 0.f; - const auto velo_lhcb_ids = velo_tracks.get_lhcbids_for_track(track_hits.data(), i_track); for (const auto id : velo_lhcb_ids) { t.addId(id); @@ -51,6 +49,49 @@ std::vector prepareVeloTracks( return checker_tracks; } +std::vector prepareVeloMuonTracks( + const unsigned number_of_events, + gsl::span track_atomics, + gsl::span track_hit_number, + gsl::span track_hits, + gsl::span velomuon_atomics, + // gsl::span velomuon_offsets, + gsl::span velomuon_qop, + gsl::span velomuon_velo_indices, + gsl::span>, // velomuon_muon_id, + gsl::span event_list) +{ + /* Tracks to be checked, save in format for checker */ + std::vector checker_tracks(event_list.size()); + for (unsigned i = 0; i < event_list.size(); i++) { + const auto event_number = event_list[i]; + + // VeloMuon Tracks + const auto event_velomuon_offset = event_number * Muon::Constants::max_number_of_tracks; + const auto event_number_of_velomuon_tracks = velomuon_atomics[event_number]; + // Tracks of this event + auto& tracks = checker_tracks[i]; + Velo::Consolidated::ConstTracks velo_tracks { + track_atomics.data(), track_hit_number.data(), event_number, number_of_events}; + tracks.resize(event_number_of_velomuon_tracks); + for (unsigned i_track = 0; i_track < event_number_of_velomuon_tracks; i_track++) { + auto& t = tracks[i_track]; + const auto i_velo_track = velomuon_velo_indices[event_velomuon_offset + i_track]; + const auto velo_lhcb_ids = velo_tracks.get_lhcbids_for_track(track_hits.data(), i_velo_track); + for (const auto id : velo_lhcb_ids) { + t.addId(id); + } + // momentum + const float qop = velomuon_qop[event_velomuon_offset + i_track]; + + t.p = 1.f / std::abs(qop); + t.qop = qop; + } // tracks + } + + return checker_tracks; +}; + std::vector prepareUTTracks( const unsigned number_of_events, gsl::span velo_track_atomics, diff --git a/configuration/python/AllenConf/HLT1.py b/configuration/python/AllenConf/HLT1.py index c4fa54a5777..a8572f6dec7 100644 --- a/configuration/python/AllenConf/HLT1.py +++ b/configuration/python/AllenConf/HLT1.py @@ -7,7 +7,7 @@ from AllenConf.hlt1_reconstruction import hlt1_reconstruction, validator_node from AllenConf.hlt1_inclusive_hadron_lines import make_track_mva_line, make_two_track_mva_line, make_kstopipi_line, make_two_track_line_ks from AllenConf.hlt1_charm_lines import make_d2kk_line, make_d2pipi_line, make_two_track_mva_charm_xsec_line from AllenConf.hlt1_calibration_lines import make_d2kpi_line, make_passthrough_line, make_rich_1_line, make_rich_2_line, make_displaced_dimuon_mass_line, make_di_muon_mass_align_line -from AllenConf.hlt1_muon_lines import make_single_high_pt_muon_line, make_single_high_pt_muon_no_muid_line, make_low_pt_muon_line, make_di_muon_mass_line, make_di_muon_soft_line, make_low_pt_di_muon_line, make_track_muon_mva_line, make_di_muon_no_ip_line +from AllenConf.hlt1_muon_lines import make_one_velomuon_track_line, make_one_muon_track_line, make_single_high_pt_muon_line, make_single_high_pt_muon_no_muid_line, make_low_pt_muon_line, make_di_muon_mass_line, make_di_muon_soft_line, make_low_pt_di_muon_line, make_track_muon_mva_line, make_di_muon_no_ip_line from AllenConf.hlt1_electron_lines import make_track_electron_mva_line, make_single_high_pt_electron_line, make_lowmass_noip_dielectron_line, make_displaced_dielectron_line, make_displaced_leptons_line, make_single_high_et_line from AllenConf.hlt1_monitoring_lines import make_beam_line, make_velo_micro_bias_line, make_odin_event_type_line, make_beam_gas_line from AllenConf.hlt1_smog2_lines import ( @@ -19,15 +19,16 @@ from AllenConf.validators import rate_validation, routingbits_validation from PyConf.control_flow import NodeLogic, CompositeNode from PyConf.tonic import configurable from AllenConf.lumi_reconstruction import lumi_reconstruction -from AllenCore.generator import is_allen_standalone def default_physics_lines(reconstructed_objects, with_calo, with_muon): + velo_tracks = reconstructed_objects["velo_tracks"] long_tracks = reconstructed_objects["long_tracks"] long_track_particles = reconstructed_objects["long_track_particles"] secondary_vertices = reconstructed_objects["secondary_vertices"] + velo_muon_objects = reconstructed_objects['velo_muon'] lines = [ make_two_track_mva_charm_xsec_line( @@ -47,6 +48,16 @@ def default_physics_lines(reconstructed_objects, with_calo, with_muon): if with_muon: lines += [ + make_one_velomuon_track_line( + velo_muon_objects['match_velo_muon'], + velo_muon_objects['consolidate_velo_muon'], + velo_muon_objects['prefix_sum_velomuon_tracks'], + name="Hlt1OneVeloMuonTrackLine"), + make_one_muon_track_line( + velo_muon_objects['find_muon_hits'], + velo_muon_objects['consolidate_muon'], + velo_muon_objects['prefix_sum_muon_tracks'], + name = "Hlt1OneMuonTrackLine"), make_single_high_pt_muon_line( long_tracks, long_track_particles, name="Hlt1SingleHighPtMuon"), @@ -85,6 +96,8 @@ def default_physics_lines(reconstructed_objects, with_calo, with_muon): pre_scaler=.1) ] + + if with_calo: ecal_clusters = reconstructed_objects["ecal_clusters"] calo_matching_objects = reconstructed_objects["calo_matching_objects"] diff --git a/configuration/python/AllenConf/hlt1_muon_lines.py b/configuration/python/AllenConf/hlt1_muon_lines.py index 576e6d98fcf..12790a1ee11 100644 --- a/configuration/python/AllenConf/hlt1_muon_lines.py +++ b/configuration/python/AllenConf/hlt1_muon_lines.py @@ -4,11 +4,63 @@ from AllenAlgorithms.algorithms import ( single_high_pt_muon_line_t, single_high_pt_muon_no_muid_line_t, low_pt_muon_line_t, di_muon_mass_line_t, di_muon_soft_line_t, - low_pt_di_muon_line_t, track_muon_mva_line_t, di_muon_no_ip_line_t) + low_pt_di_muon_line_t, track_muon_mva_line_t, di_muon_no_ip_line_t, + one_muon_track_line_t, one_velomuon_track_line_t) from AllenConf.utils import initialize_number_of_events, mep_layout from AllenCore.generator import make_algorithm +def make_one_velomuon_track_line(velomuon_tracks_algo, + consolidate_tracks_algo, + velomuon_prefix_sum, + name="Hlt1OneVeloMuonTrack", + pre_scaler_hash_string=None, + post_scaler_hash_string=None): + number_of_events = initialize_number_of_events() + velomuon_qop = consolidate_tracks_algo.dev_velomuon_tracks_qop_output_t + dev_number_of_velomuon_tracks = velomuon_tracks_algo.dev_velomuon_number_of_tracks_t + dev_velomuon_tracks_offsets = velomuon_prefix_sum.dev_output_buffer_t + dev_velomuon_tracks = consolidate_tracks_algo.dev_velomuon_tracks_output_t + host_velomuon_total_number_of_tracks = velomuon_prefix_sum.host_total_sum_holder_t + return make_algorithm( + one_velomuon_track_line_t, + name=name, + host_number_of_events_t=number_of_events["host_number_of_events"], + dev_number_of_events_t=number_of_events["dev_number_of_events"], + pre_scaler_hash_string=pre_scaler_hash_string or name + "_pre", + post_scaler_hash_string=post_scaler_hash_string or name + "_post", + dev_velomuon_tracks_t=dev_velomuon_tracks, + dev_velomuon_tracks_qop_t=velomuon_qop, + dev_velomuon_number_of_tracks_t=dev_number_of_velomuon_tracks, + host_velomuon_total_number_of_tracks_t= + host_velomuon_total_number_of_tracks, + dev_velomuon_tracks_offsets_t=dev_velomuon_tracks_offsets) + + +def make_one_muon_track_line(standalone_muon_tracks_algo, + consolidate_tracks_algo, + standalone_muon_prefix_sum, + name="Hlt1OneMuonTrack", + pre_scaler_hash_string=None, + post_scaler_hash_string=None): + number_of_events = initialize_number_of_events() + muon_tracks = consolidate_tracks_algo.dev_muon_tracks_output_t + dev_number_of_muon_tracks = standalone_muon_tracks_algo.dev_muon_number_of_tracks_t + dev_muon_tracks_offsets = standalone_muon_prefix_sum.dev_output_buffer_t + host_muon_total_number_of_tracks = standalone_muon_prefix_sum.host_total_sum_holder_t + return make_algorithm( + one_muon_track_line_t, + name=name, + host_number_of_events_t=number_of_events["host_number_of_events"], + dev_number_of_events_t=number_of_events["dev_number_of_events"], + pre_scaler_hash_string=pre_scaler_hash_string or name + "_pre", + post_scaler_hash_string=post_scaler_hash_string or name + "_post", + dev_muon_tracks_t=muon_tracks, + dev_muon_number_of_tracks_t=dev_number_of_muon_tracks, + host_muon_total_number_of_tracks_t=host_muon_total_number_of_tracks, + dev_muon_tracks_offsets_t=dev_muon_tracks_offsets) + + def make_single_high_pt_muon_line(long_tracks, long_track_particles, name="Hlt1SingleHighPtMuon", diff --git a/configuration/python/AllenConf/hlt1_reconstruction.py b/configuration/python/AllenConf/hlt1_reconstruction.py index b733bc575df..7cc4b1ee6a3 100755 --- a/configuration/python/AllenConf/hlt1_reconstruction.py +++ b/configuration/python/AllenConf/hlt1_reconstruction.py @@ -5,7 +5,7 @@ from AllenConf.velo_reconstruction import decode_velo, make_velo_tracks, run_vel from AllenConf.ut_reconstruction import decode_ut, make_ut_tracks from AllenConf.scifi_reconstruction import decode_scifi, make_forward_tracks, make_seeding_XZ_tracks, make_seeding_tracks from AllenConf.matching_reconstruction import make_velo_scifi_matches -from AllenConf.muon_reconstruction import decode_muon, is_muon, fake_muon_id +from AllenConf.muon_reconstruction import decode_muon, is_muon,fake_muon_id, make_velo_muon from AllenConf.calo_reconstruction import decode_calo, make_track_matching, make_ecal_clusters from AllenConf.primary_vertex_reconstruction import make_pvs from AllenConf.secondary_vertex_reconstruction import make_kalman_velo_only, make_basic_particles, fit_secondary_vertices @@ -27,11 +27,13 @@ def hlt1_reconstruction(matching=False, velo_tracks = make_velo_tracks(decoded_velo) velo_states = run_velo_kalman_filter(velo_tracks) pvs = make_pvs(velo_tracks) - + velo_muon = make_velo_muon() + output = { "velo_tracks": velo_tracks, "velo_states": velo_states, - "pvs": pvs + "pvs": pvs, + 'velo_muon':velo_muon } if matching: @@ -86,10 +88,11 @@ def hlt1_reconstruction(matching=False, secondary_vertices = fit_secondary_vertices( long_tracks, pvs, kalman_velo_only, long_track_particles) - + velo_muon = make_velo_muon() output.update({ "long_track_particles": long_track_particles, - "secondary_vertices": secondary_vertices + "secondary_vertices": secondary_vertices, + "velo_muon": velo_muon }) return output @@ -137,18 +140,22 @@ def validator_node(reconstructed_objects, line_algorithms, matching, with_ut, validators += [ make_composite_node_with_gec( "long_validation", - long_validation(reconstructed_objects["long_tracks"]), + long_validation(reconstructed_objects["long_tracks"], + reconstructed_objects["velo_muon"]), with_scifi=True, with_ut=with_ut) ] + if with_muon: + validators += make_composite_node_with_gec( - "muon_validation", - muon_validation(reconstructed_objects["muonID"]), - with_scifi=True, - with_ut=with_ut), + "muon_validation", + muon_validation(reconstructed_objects["muonID"], + reconstructed_objects['velo_muon']), + with_scifi=True, + with_ut=with_ut), validators += [ make_composite_node_with_gec( "pv_validation", diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py index 1ed0b169d30..cfae519179b 100644 --- a/configuration/python/AllenConf/muon_reconstruction.py +++ b/configuration/python/AllenConf/muon_reconstruction.py @@ -4,7 +4,8 @@ from AllenAlgorithms.algorithms import ( data_provider_t, muon_calculate_srq_size_t, host_prefix_sum_t, muon_populate_tile_and_tdc_t, muon_add_coords_crossing_maps_t, - muon_populate_hits_t, is_muon_t, empty_lepton_id_t) + muon_populate_hits_t, is_muon_t, empty_lepton_id_t, find_muon_hits_t, fit_muon_t, + ut_select_velo_tracks_t, match_velo_muon_t, consolidate_muon_t, consolidate_velo_muon_t) from AllenConf.utils import initialize_number_of_events from AllenCore.generator import make_algorithm @@ -165,3 +166,137 @@ def muon_id(): muonID = is_muon(decoded_muon, long_tracks) alg = muonID["dev_is_muon"].producer return alg + + +def make_velo_muon(): + from AllenConf.velo_reconstruction import decode_velo, make_velo_tracks + from AllenConf.ut_reconstruction import decode_ut, make_ut_tracks + from AllenConf.velo_reconstruction import run_velo_kalman_filter + + number_of_events = initialize_number_of_events() + host_number_of_events = number_of_events["host_number_of_events"] + dev_number_of_events = number_of_events["dev_number_of_events"] + decoded_velo = decode_velo() + velo_tracks = make_velo_tracks(decoded_velo) + decoded_ut = decode_ut() + decoded_muon = decode_muon() + ut_tracks = make_ut_tracks(decoded_ut, velo_tracks) + velo_tracks = ut_tracks["velo_tracks"] + velo_states = ut_tracks["velo_states"] + + velo_kalman_filter = run_velo_kalman_filter(velo_tracks) + + ut_select_velo_tracks = make_algorithm( + ut_select_velo_tracks_t, + name="ut_select_velo_tracks", + host_number_of_events_t=number_of_events["host_number_of_events"], + host_number_of_reconstructed_velo_tracks_t=velo_tracks[ + "host_number_of_reconstructed_velo_tracks"], + dev_velo_tracks_view_t=velo_tracks["dev_velo_tracks_view"], + dev_velo_states_view_t=velo_states[ + "dev_velo_kalman_beamline_states_view"], + dev_accepted_velo_tracks_t=velo_tracks["dev_accepted_velo_tracks"]) + + find_muon_hits = make_algorithm( + find_muon_hits_t, + name='find_muon_hits', + host_number_of_events_t=number_of_events["host_number_of_events"], + dev_number_of_events_t=number_of_events["dev_number_of_events"], + host_number_of_reconstructed_ut_tracks_t=ut_tracks[ + "host_number_of_reconstructed_ut_tracks"], + dev_offsets_ut_tracks_t=ut_tracks["dev_offsets_ut_tracks"], + dev_offsets_ut_track_hit_number_t=ut_tracks[ + "dev_offsets_ut_track_hit_number"], + dev_ut_qop_t=ut_tracks["dev_ut_qop"], + dev_ut_track_velo_indices_t=ut_tracks["dev_ut_track_velo_indices"], + dev_offsets_all_velo_tracks_t=velo_tracks[ + "dev_offsets_all_velo_tracks"], + dev_offsets_velo_track_hit_number_t=velo_tracks[ + "dev_offsets_velo_track_hit_number"], + dev_velo_kalman_beamline_states_t=velo_kalman_filter[ + "dev_velo_kalman_beamline_states"], + dev_station_ocurrences_offset_t=decoded_muon[ + "dev_station_ocurrences_offset"], + dev_muon_hits_t=decoded_muon["dev_muon_hits"], + ) + prefix_sum_muon_tracks = make_algorithm( + host_prefix_sum_t, + name="prefix_sum_muon_tracks_find_hits", + dev_input_buffer_t=find_muon_hits.dev_muon_number_of_tracks_t) + + consolidate_muon = make_algorithm( + consolidate_muon_t, + name='consolidate_muon_t', + host_number_of_events_t=number_of_events["host_number_of_events"], + dev_number_of_events_t=number_of_events["dev_number_of_events"], + dev_muon_tracks_input_t=find_muon_hits.dev_muon_tracks_t, + # dev_muon_tracks_tx_input_t=find_muon_hits.dev_muon_tracks_tx_t, + # dev_muon_tracks_ty_input_t=find_muon_hits.dev_muon_tracks_ty_t, + # dev_muon_tracks_state_muon_index_input_t=find_muon_hits. + # dev_muon_tracks_states_t, + dev_muon_number_of_tracks_t=find_muon_hits.dev_muon_number_of_tracks_t, + dev_muon_tracks_offsets_t=prefix_sum_muon_tracks.dev_output_buffer_t, + host_muon_total_number_of_tracks_t=prefix_sum_muon_tracks. + host_total_sum_holder_t, + ) + + match_velo_muon = make_algorithm( + match_velo_muon_t, + name='match_velo_muon', + host_number_of_events_t=number_of_events["host_number_of_events"], + dev_number_of_events_t=number_of_events["dev_number_of_events"], + dev_muon_tracks_t=consolidate_muon.dev_muon_tracks_output_t, + dev_muon_number_of_tracks_t=find_muon_hits.dev_muon_number_of_tracks_t, + host_muon_total_number_of_tracks_t=prefix_sum_muon_tracks. + host_total_sum_holder_t, + host_number_of_reconstructed_velo_tracks_t=velo_tracks[ + "host_number_of_reconstructed_velo_tracks"], + dev_muon_tracks_offsets_t=prefix_sum_muon_tracks.dev_output_buffer_t, + dev_station_ocurrences_offset_t=decoded_muon[ + "dev_station_ocurrences_offset"], + dev_muon_hits_t=decoded_muon["dev_muon_hits"], + dev_ut_number_of_selected_velo_tracks_t=ut_select_velo_tracks. + dev_ut_number_of_selected_velo_tracks_t, + dev_ut_selected_velo_tracks_t=ut_select_velo_tracks. + dev_ut_selected_velo_tracks_t, + dev_velo_tracks_view_t=velo_tracks["dev_velo_tracks_view"], + dev_velo_states_view_t=velo_states[ + "dev_velo_kalman_endvelo_states_view"], + # dev_muon_tracks_tx_t=consolidate_muon.dev_muon_tracks_tx_output_t, + # dev_muon_tracks_ty_t=consolidate_muon.dev_muon_tracks_ty_output_t, + # dev_muon_tracks_state_muon_index_t=consolidate_muon. + # dev_muon_tracks_state_muon_index_output_t) + ) + + prefix_sum_velomuon_tracks = make_algorithm( + host_prefix_sum_t, + name="prefix_sum_muon_tracks", + dev_input_buffer_t=match_velo_muon.dev_velomuon_number_of_tracks_t) + + consolidate_velo_muon = make_algorithm( + consolidate_velo_muon_t, + name="consolidate_velo_muon_t", + host_number_of_events_t=number_of_events["host_number_of_events"], + dev_number_of_events_t=number_of_events["dev_number_of_events"], + dev_velomuon_number_of_tracks_t=match_velo_muon. + dev_velomuon_number_of_tracks_t, + dev_velomuon_tracks_offsets_t=prefix_sum_velomuon_tracks. + dev_output_buffer_t, + host_velomuon_total_number_of_tracks_t=prefix_sum_velomuon_tracks. + host_total_sum_holder_t, + dev_velomuon_tracks_velo_indices_input_t=match_velo_muon. + dev_velomuon_tracks_velo_indices_t, + dev_velomuon_tracks_muon_indices_input_t=match_velo_muon. + dev_velomuon_tracks_muon_indices_t, + dev_velomuon_tracks_qop_input_t=match_velo_muon. + dev_velomuon_tracks_qop_t, + dev_velomuon_tracks_input_t=match_velo_muon.dev_velomuon_tracks_t) + + return { + 'find_muon_hits': find_muon_hits, + 'consolidate_muon': consolidate_muon, + 'match_velo_muon': match_velo_muon, + 'prefix_sum_muon_tracks': prefix_sum_muon_tracks, + 'prefix_sum_velomuon_tracks': prefix_sum_velomuon_tracks, + 'consolidate_velo_muon': consolidate_velo_muon + } diff --git a/configuration/python/AllenConf/validators.py b/configuration/python/AllenConf/validators.py index 82fedde9d1e..63b4fcf5172 100644 --- a/configuration/python/AllenConf/validators.py +++ b/configuration/python/AllenConf/validators.py @@ -14,6 +14,8 @@ from AllenAlgorithms.algorithms import (host_prefix_sum_t, seeding_copy_trackXZ_hit_number_t) from AllenConf.scifi_reconstruction import decode_scifi, make_seeding_XZ_tracks +from AllenConf.velo_reconstruction import decode_velo, make_velo_tracks, run_velo_kalman_filter + def mc_data_provider(): host_mc_particle_banks = make_algorithm( @@ -78,10 +80,11 @@ def veloUT_validation(veloUT_tracks, name="veloUT_validator"): dev_ut_qop_t=veloUT_tracks["dev_ut_qop"]) -def long_validation(long_tracks, name="long_validator"): +def long_validation(long_tracks, velo_muon, name="long_validator"): mc_events = mc_data_provider() number_of_events = initialize_number_of_events() - + decoded_velo = decode_velo() + velo_tracks = make_velo_tracks(decoded_velo) velo_kalman_filter = long_tracks["velo_kalman_filter"] return make_algorithm( @@ -91,11 +94,20 @@ def long_validation(long_tracks, name="long_validator"): host_mc_events_t=mc_events.host_mc_events_t, host_number_of_reconstructed_long_tracks_t=long_tracks[ "host_number_of_reconstructed_scifi_tracks"], + host_number_of_reconstructed_velomuon_tracks_t=velo_muon[ + 'prefix_sum_velomuon_tracks'].host_total_sum_holder_t, dev_velo_states_view_t=velo_kalman_filter[ "dev_velo_kalman_endvelo_states_view"], + dev_velo_tracks_view_t=velo_tracks['dev_velo_tracks_view'], dev_multi_event_long_tracks_view_t=long_tracks[ "dev_multi_event_long_tracks_view"], - dev_offsets_long_tracks_t=long_tracks["dev_offsets_long_tracks"]) + dev_offsets_long_tracks_t=long_tracks["dev_offsets_long_tracks"], + dev_offsets_velomuon_tracks_t=velo_muon["prefix_sum_velomuon_tracks"]. + dev_output_buffer_t, + dev_velomuon_tracks_qop_t=velo_muon['consolidate_velo_muon']. + dev_velomuon_tracks_qop_output_t, + dev_velomuon_tracks_velo_indices_t=velo_muon['consolidate_velo_muon']. + dev_velomuon_tracks_velo_indices_output_t) def seeding_xz_validation(name="seed_xz_validator"): @@ -191,11 +203,13 @@ def velo_scifi_dump(matched_tracks, name="veloscifi_dump"): dev_seeding_states_t=seeding_tracks["dev_seeding_states"]) -def muon_validation(muonID, name="muon_validator"): +def muon_validation(muonID, velo_muon, name="muon_validator"): mc_events = mc_data_provider() number_of_events = initialize_number_of_events() long_tracks = muonID["long_tracks"] + decoded_velo = decode_velo() + velo_tracks = make_velo_tracks(decoded_velo) velo_kalman_filter = long_tracks["velo_kalman_filter"] return make_algorithm( @@ -205,12 +219,18 @@ def muon_validation(muonID, name="muon_validator"): host_mc_events_t=mc_events.host_mc_events_t, host_number_of_reconstructed_long_tracks_t=long_tracks[ "host_number_of_reconstructed_scifi_tracks"], + host_number_of_reconstructed_velo_tracks_t=velo_tracks[ + "host_number_of_reconstructed_velo_tracks"], dev_velo_states_view_t=velo_kalman_filter[ "dev_velo_kalman_endvelo_states_view"], dev_multi_event_long_tracks_view_t=long_tracks[ "dev_multi_event_long_tracks_view"], + dev_velo_tracks_view_t=velo_tracks['dev_velo_tracks_view'], dev_offsets_long_tracks_t=long_tracks["dev_offsets_long_tracks"], - dev_is_muon_t=muonID["dev_is_muon"]) + dev_offsets_velo_tracks_t=velo_tracks["dev_offsets_all_velo_tracks"], + dev_is_muon_t=muonID["dev_is_muon"], + dev_match_velo_muon_t=velo_muon['match_velo_muon']. + dev_velomuon_muon_id_t) def pv_validation(pvs, name="pv_validator"): diff --git a/configuration/python/AllenSequences/hlt1_complex_validation.py b/configuration/python/AllenSequences/hlt1_complex_validation.py index 45cf3371156..70c8ff26b5f 100644 --- a/configuration/python/AllenSequences/hlt1_complex_validation.py +++ b/configuration/python/AllenSequences/hlt1_complex_validation.py @@ -12,6 +12,7 @@ from AllenConf.validators import ( from PyConf.control_flow import NodeLogic, CompositeNode from AllenCore.generator import generate +from AllenConf.muon_reconstruction import make_velo_muon # Reconstructed objects with make_ut_tracks.bind(restricted=False): @@ -19,6 +20,7 @@ with make_ut_tracks.bind(restricted=False): restricted_hlt1_reconstruction = hlt1_reconstruction() gec = make_gec(count_scifi=True, count_ut=True) +velo_muon = make_velo_muon() lines = [] with line_maker.bind(prefilter=gec): @@ -87,25 +89,27 @@ validators_leaf = CompositeNode( make_composite_node_with_gec( "restricted_long_validator", long_validation(restricted_hlt1_reconstruction["long_tracks"], + restricted_hlt1_reconstruction['velo_muon'], "restricted_long_validator"), with_scifi=True, with_ut=True), make_composite_node_with_gec( "non-restricted_long_validator", long_validation(non_restricted_hlt1_reconstruction["long_tracks"], + non_restricted_hlt1_reconstruction['velo_muon'], "non-restricted_long_validator"), with_scifi=True, with_ut=True), make_composite_node_with_gec( "restricted_muon_validation", muon_validation(restricted_hlt1_reconstruction["muonID"], - "restricted_muon_validation"), + velo_muon, "restricted_muon_validation"), with_scifi=True, with_ut=True), make_composite_node_with_gec( "non-restricted_muon_validation", muon_validation(non_restricted_hlt1_reconstruction["muonID"], - "non-restricted_muon_validation"), + velo_muon, "non-restricted_muon_validation"), with_scifi=True, with_ut=True), make_composite_node_with_gec( diff --git a/configuration/python/AllenSequences/trackmatching_veloscifi_validation.py b/configuration/python/AllenSequences/trackmatching_veloscifi_validation.py index b7d6110afbc..b2786206cae 100644 --- a/configuration/python/AllenSequences/trackmatching_veloscifi_validation.py +++ b/configuration/python/AllenSequences/trackmatching_veloscifi_validation.py @@ -3,6 +3,7 @@ ############################################################################### from AllenConf.scifi_reconstruction import decode_scifi, seeding_xz, make_seeding_XZ_tracks, make_seeding_tracks from AllenConf.matching_reconstruction import make_velo_scifi_matches +from AllenConf.muon_reconstruction import make_velo_muon from AllenConf.hlt1_reconstruction import make_composite_node_with_gec from AllenConf.validators import velo_validation, seeding_validation, seeding_xz_validation, long_validation, velo_scifi_dump from AllenConf.velo_reconstruction import decode_velo, make_velo_tracks, run_velo_kalman_filter @@ -20,7 +21,8 @@ seed = seeding_validation(seeding_tracks) seed_xz = seeding_xz_validation() matched_tracks = make_velo_scifi_matches(velo_tracks, velo_states, seeding_tracks) -velo_scifi = long_validation(matched_tracks) +velo_muon = make_velo_muon() +velo_scifi = long_validation(matched_tracks, velo_muon) velo_scifi_matching_sequence = CompositeNode( "Validators", [ make_composite_node_with_gec( diff --git a/device/event_model/common/include/CopyTrackParameters.cuh b/device/event_model/common/include/CopyTrackParameters.cuh index 94c4071c014..dafdce9bbbf 100644 --- a/device/event_model/common/include/CopyTrackParameters.cuh +++ b/device/event_model/common/include/CopyTrackParameters.cuh @@ -7,6 +7,61 @@ #include "PV_Definitions.cuh" #include "patPV_Definitions.cuh" +__device__ inline void prepare_velo_tracks( + const Allen::Views::Velo::Consolidated::Tracks event_velo_tracks, + Checker::Track* velo_checker_tracks) +{ + const unsigned number_of_tracks_event = event_velo_tracks.size(); + for (unsigned i_track = 0; i_track < number_of_tracks_event; i_track++) { + Checker::Track t; + const auto velo_track = event_velo_tracks.track(i_track); + + const auto total_number_of_hits = velo_track.number_of_hits(); + for (unsigned ihit = 0; ihit < total_number_of_hits; ihit++) { + const auto hit = velo_track.hit(ihit); + t.addId(hit.id()); + } + velo_checker_tracks[i_track] = t; + } +} + +__device__ inline void prepare_velomuon_tracks( + const unsigned number_of_tracks_event, + const Allen::Views::Velo::Consolidated::Tracks event_velo_tracks, + const Allen::Views::Physics::KalmanStates endvelo_states, + const float* velomuon_qop, + const unsigned* velomuon_velo_indices, + Checker::Track* velomuon_checker_tracks) +{ + for (unsigned i_track = 0; i_track < number_of_tracks_event; i_track++) { + Checker::Track t; + const auto velo_track = event_velo_tracks.track(velomuon_velo_indices[i_track]); + const auto velo_track_index = velo_track.track_index(); + const auto velo_state = endvelo_states.state(velo_track_index); + t.velo_track_index = velo_track_index; + // momentum + const auto qop = velomuon_qop[i_track]; + t.p = 1.f / std::abs(qop); + t.qop = qop; + // direction at first state -> velo state of track + const float tx = velo_state.tx(); + const float ty = velo_state.ty(); + const float slope2 = tx * tx + ty * ty; + t.pt = std::sqrt(slope2 / (1.0f + slope2)) / std::fabs(qop); + // pseudorapidity + const float rho = std::sqrt(slope2); + t.rho = rho; + + // add all hits + const auto total_number_of_hits = velo_track.number_of_hits(); + for (unsigned ihit = 0; ihit < total_number_of_hits; ihit++) { + const auto hit = velo_track.hit(ihit); + t.addId(hit.id()); + } + velomuon_checker_tracks[i_track] = t; + } +} + __device__ inline void prepare_long_tracks( const Allen::Views::Physics::LongTracks event_long_tracks, const Allen::Views::Physics::KalmanStates endvelo_states, diff --git a/device/event_model/muon/include/MuonDefinitions.cuh b/device/event_model/muon/include/MuonDefinitions.cuh index 4a0c40116e2..84a81082d11 100644 --- a/device/event_model/muon/include/MuonDefinitions.cuh +++ b/device/event_model/muon/include/MuonDefinitions.cuh @@ -95,9 +95,212 @@ namespace Muon { __host__ __device__ const float* params_begin_const() const { return reinterpret_cast(m_params); } }; + struct MatchVeloWindows { + float Xmax[16] = { + // R1 R2 R3 R4 + 100., + 200., + 300., + 400., // M2 + 100., + 200., + 300., + 400., // M3 + 400., + 400., + 400., + 400., // M4 + 400., + 400., + 400., + 400.}; // M5 + + float Ymax[16] = { + // R1 R2 R3 R4 + 60., + 120., + 180., + 240., // M2 + 60., + 120., + 240., + 240., // M3 + 60., + 120., + 240., + 480., // M4 + // 60., + // 120., + // 240., + // 480., // M5 + 60., + 120., + 240., + 480., // M5 + + }; + + float tolForRegion[4] = {2.0, 4.0, 8.0, 10.0}; + float z[4] {15205.f, 16400.f, 17550.f, 18850.f}; + float xcale[4] {0.06f, 0.1f, 0.15f, 0.15f}; + float distance_cut[5] {30 * 30, 60 * 60, 110 * 110, 200 * 200}; + }; + static constexpr unsigned max_number_of_tracks = 120; + struct MomParam { + float C0[10] = { + 113.21271, + -115.8049, + -657.5307, + 60.049296, + -491.3313, + 2381.3909, + 2684.9605, + 12571.073, + 45041.933, + -85422.44, + }; + float C1[14] = { + -0.154380, + -119.347058, + -841.363037, + 16543.416488, + -170652.099599, + 28.940644, + -2313.431503, + 45068.764734, + 3439.026842, + -143227.239737, + 1796054.751602, + 596.408641, + -226695.997197, + 471660.297116, + }; + float C2[10] = {-48.940097, + -213.741769, + -3006.193908, + 135.633447, + -822.544939, + 12407.568991, + 18899.246135, + 40529.981272, + 184896.154097, + -376756.125511}; + float C3[14] = { + 0.215404, + -122.816177, + -902.317011, + 25227.824861, + -206136.914140, + -134.229687, + 12320.745080, + -247916.687297, + 6915.080393, + -232643.972656, + 2009303.058464, + -46719.185323, + 2276730.013659, + 177694.558072, + }; + float C4[10] = { + -23.473752, + -10.128032, + -889.713080, + 468.227962, + -11156.964015, + -2248.650393, + 11765.688620, + 59454.044613, + 10303.576287, + 137177.920462, + }; + }; } // namespace Constants + namespace TrackMatchingConsts { + constexpr float z_match = 5240.0f; // FIXME_GEOMETRY_HARDCODING + constexpr float zMatchY = 10000.f; // in mm + + constexpr float dxTol = 8.f; // in mm // to scan + constexpr float dxTolSlope = 80.f; // in mm // to scan + + constexpr float dyTol = 6.f; // in mm // to scan + constexpr float dyTolSlope = 300.f; // in mm // to scan + constexpr float fastYTol = 250.f; // to scan + // The main cut values + // constexpr float maxChi2 = 15.f; // to scan + constexpr float maxChi2 = 2.5f; // to scan + constexpr float minNN = 0.25f; // not used so far + + // Magnetic parametrization from Velo + SciFi tracks + struct MagnetParametrization { + float zMagnetParamsMatch[5] {5287.6f, -7.98878f, 317.683f, 0.0119379f, -1418.42f}; + float bendYParams[2] {-347.801f, -42663.6f}; + }; + } // namespace TrackMatchingConsts + + struct MatchingResult { + float dSlopeX; + float dSlopeY; + float distX; + float distY; + float zForX; + float chi2; + }; } // namespace Muon +struct MuonTrack { + int m_hits[4] {-1, -1, -1, -1}; + int m_velo_index = -1; + uint8_t m_number_of_hits = 0; + float m_tx; + float m_ty; + float m_ax; + float m_ay; + float m_chi2x; + float m_chi2y; + float m_p; + float m_chi2match; + float m_state_muon_index; + + __host__ __device__ MuonTrack() {} + + __host__ __device__ void add_hit_to_station(const unsigned hit_index, const int station_index) + { + ++m_number_of_hits; + m_hits[station_index] = hit_index; + } + + __host__ __device__ int hit(const int station_index) const { return m_hits[station_index]; } + + __host__ __device__ uint8_t number_of_hits() const { return m_number_of_hits; } + + // __host__ __device__ void set_velo_index(int velo_index){ m_velo_index = velo_index;} + + __host__ __device__ void set_tx(float tx) { m_tx = tx; } + __host__ __device__ void set_ty(float ty) { m_ty = ty; } + __host__ __device__ void set_ax(float ax) { m_ax = ax; } + __host__ __device__ void set_ay(float ay) { m_ay = ay; } + __host__ __device__ void set_chi2x(float chi2x) { m_chi2x = chi2x; } + __host__ __device__ void set_chi2y(float chi2y) { m_chi2y = chi2y; } + __host__ __device__ void set_p(float p) { m_p = p; } + __host__ __device__ void set_chi2match(float chi2) { m_chi2match = chi2; } + __host__ __device__ void set_state(int state) { m_state_muon_index = state; } + __host__ __device__ void set_velo_index(int idx) { m_velo_index = idx; } + + __host__ __device__ float get_tx() const { return m_tx; } + __host__ __device__ float get_ty() const { return m_ty; } + __host__ __device__ float get_ax() const { return m_ax; } + __host__ __device__ float get_ay() const { return m_ay; } + __host__ __device__ float get_chi2x() const { return m_chi2x; } + __host__ __device__ float get_chi2y() const { return m_chi2y; } + __host__ __device__ float get_p() const { return m_p; } + __host__ __device__ float get_chi2match() const { return m_chi2match; } + __host__ __device__ int get_state() const { return m_state_muon_index; } + __host__ __device__ int get_velo_index() const { return m_velo_index; } + __host__ __device__ float state_x() const { return m_tx * 15205.f + m_ax; } + __host__ __device__ float state_y() const { return m_ty * 15205.f + m_ay; } + __host__ __device__ float state_z() const { return 15205.f; } +}; + namespace MatchUpstreamMuon { static constexpr float kickOffset = 338.92f * Gaudi::Units::MeV; // KickOffset static constexpr float kickScale = 1218.62f * Gaudi::Units::MeV; // KickScale diff --git a/device/muon/CMakeLists.txt b/device/muon/CMakeLists.txt index 80373cd7b22..a6f216f3720 100644 --- a/device/muon/CMakeLists.txt +++ b/device/muon/CMakeLists.txt @@ -6,7 +6,11 @@ file(GLOB muon_classification "classification/src/*cu") file(GLOB muon_filtering "is_muon/src/*cu") file(GLOB muon_decoding "decoding/src/*cu") file(GLOB muon_upstream_filtering "match_upstream_muon/src/*cu") +file(GLOB muon_velo_filtering "match_velo_muon/src/*cu") file(GLOB muon_filter "muon_filter/src/*cu") +file(GLOB consolidate_muon "consolidate_muon/src/*cu") + + allen_add_device_library(Muon STATIC ${muon_filtering} @@ -14,7 +18,9 @@ allen_add_device_library(Muon STATIC ${muon_classification} ${muon_decoding} ${muon_upstream_filtering} + ${muon_velo_filtering} ${muon_filter} + ${consolidate_muon} ) add_library(MuonCommon INTERFACE) @@ -32,4 +38,6 @@ target_include_directories(Muon PUBLIC $ $ $ - $) + $ + $ + ) diff --git a/device/muon/match_velo_muon/include/ConsolidateMuon.cuh b/device/muon/match_velo_muon/include/ConsolidateMuon.cuh new file mode 100644 index 00000000000..7af0e366dfa --- /dev/null +++ b/device/muon/match_velo_muon/include/ConsolidateMuon.cuh @@ -0,0 +1,58 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include "MuonDefinitions.cuh" +#include "UTConsolidated.cuh" +#include "VeloConsolidated.cuh" +#include "States.cuh" + +static constexpr int M2 {0}, M3 {1}, M4 {2}, M5 {3}; + +namespace consolidate_muon { + struct Parameters { + HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; + MASK_INPUT(dev_event_list_t) dev_event_list; + DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; + // DEVICE_INPUT(dev_station_ocurrences_offset_t, unsigned) dev_station_ocurrences_offset; + // DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; + DEVICE_INPUT(dev_muon_tracks_input_t, MuonTrack) dev_muon_tracks_input; + // DEVICE_INPUT(dev_muon_tracks_tx_input_t, float) dev_muon_tracks_tx_input; + // DEVICE_INPUT(dev_muon_tracks_ty_input_t, float) dev_muon_tracks_ty_input; + // DEVICE_INPUT(dev_muon_tracks_state_muon_index_input_t, int) dev_muon_tracks_state_muon_index_input; + DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; + + DEVICE_INPUT(dev_muon_tracks_offsets_t, unsigned) dev_muon_tracks_offsets; + HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; + + DEVICE_OUTPUT(dev_muon_tracks_output_t, MuonTrack) dev_muon_tracks_output; + // DEVICE_OUTPUT(dev_muon_tracks_tx_output_t, float) dev_muon_tracks_tx_output; + // DEVICE_OUTPUT(dev_muon_tracks_ty_output_t, float) dev_muon_tracks_ty_output; + // DEVICE_OUTPUT(dev_muon_tracks_state_muon_index_output_t, int) dev_muon_tracks_state_muon_index_output; + + PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; + }; + + __global__ void consolidate_muon(Parameters); + + struct consolidate_muon_t : public DeviceAlgorithm, Parameters { + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const; + + void operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + HostBuffers& host_buffers, + const Allen::Context& context) const; + + private: + Property m_block_dim {this, {{64, 1, 1}}}; + }; + +} // namespace consolidate_muon \ No newline at end of file diff --git a/device/muon/match_velo_muon/include/ConsolidateVeloMuon.cuh b/device/muon/match_velo_muon/include/ConsolidateVeloMuon.cuh new file mode 100644 index 00000000000..ed325c8154e --- /dev/null +++ b/device/muon/match_velo_muon/include/ConsolidateVeloMuon.cuh @@ -0,0 +1,58 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include "MuonDefinitions.cuh" +#include "UTConsolidated.cuh" +#include "VeloConsolidated.cuh" +#include "States.cuh" + +static constexpr int M2 {0}, M3 {1}, M4 {2}, M5 {3}; + +namespace consolidate_velo_muon { + struct Parameters { + HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; + MASK_INPUT(dev_event_list_t) dev_event_list; + DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; + + DEVICE_INPUT(dev_velomuon_number_of_tracks_t, unsigned) dev_velomuon_number_of_tracks; + + DEVICE_INPUT(dev_velomuon_tracks_offsets_t, unsigned) dev_velomuon_tracks_offsets; + HOST_INPUT(host_velomuon_total_number_of_tracks_t, unsigned) host_velomuon_total_number_of_tracks; + + DEVICE_INPUT(dev_velomuon_tracks_input_t, MuonTrack) dev_velomuon_tracks_input; + DEVICE_INPUT(dev_velomuon_tracks_velo_indices_input_t, unsigned) dev_velomuon_tracks_velo_indices_input; + DEVICE_INPUT(dev_velomuon_tracks_muon_indices_input_t, unsigned) dev_velomuon_tracks_muon_indices_input; + DEVICE_INPUT(dev_velomuon_tracks_qop_input_t, float) dev_velomuon_tracks_qop_input; + + DEVICE_OUTPUT(dev_velomuon_tracks_output_t, MuonTrack) dev_velomuon_tracks_output; + DEVICE_OUTPUT(dev_velomuon_tracks_velo_indices_output_t, unsigned) dev_velomuon_tracks_velo_indices_output; + DEVICE_OUTPUT(dev_velomuon_tracks_muon_indices_output_t, unsigned) dev_velomuon_tracks_muon_indices_output; + DEVICE_OUTPUT(dev_velomuon_tracks_qop_output_t, float) dev_velomuon_tracks_qop_output; + + PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; + }; + + __global__ void consolidate_velo_muon(Parameters); + + struct consolidate_velo_muon_t : public DeviceAlgorithm, Parameters { + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const; + + void operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + HostBuffers& host_buffers, + const Allen::Context& context) const; + + private: + Property m_block_dim {this, {{64, 1, 1}}}; + }; + +} // namespace consolidate_velo_muon \ No newline at end of file diff --git a/device/muon/match_velo_muon/include/FindMuonHits.cuh b/device/muon/match_velo_muon/include/FindMuonHits.cuh new file mode 100644 index 00000000000..854aea20efd --- /dev/null +++ b/device/muon/match_velo_muon/include/FindMuonHits.cuh @@ -0,0 +1,71 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include "MuonDefinitions.cuh" +#include "UTConsolidated.cuh" +#include "VeloConsolidated.cuh" +#include "States.cuh" +#include "ROOTService.h" + +static constexpr int M2 {0}, M3 {1}, M4 {2}, M5 {3}; + +namespace find_muon_hits { + struct Parameters { + HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; + HOST_INPUT(host_number_of_reconstructed_ut_tracks_t, unsigned) host_number_of_reconstructed_ut_tracks; + MASK_INPUT(dev_event_list_t) dev_event_list; + DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; + DEVICE_INPUT(dev_offsets_all_velo_tracks_t, unsigned) dev_atomics_velo; + DEVICE_INPUT(dev_offsets_velo_track_hit_number_t, unsigned) dev_velo_track_hit_number; + DEVICE_INPUT(dev_velo_kalman_beamline_states_t, char) dev_kalmanvelo_states; + DEVICE_INPUT(dev_offsets_ut_tracks_t, unsigned) dev_atomics_ut; + DEVICE_INPUT(dev_offsets_ut_track_hit_number_t, unsigned) dev_ut_track_hit_number; + DEVICE_INPUT(dev_ut_qop_t, float) dev_ut_qop; + DEVICE_INPUT(dev_ut_track_velo_indices_t, unsigned) dev_ut_track_velo_indices; + DEVICE_INPUT(dev_station_ocurrences_offset_t, unsigned) dev_station_ocurrences_offset; + HOST_OUTPUT(host_station_ocurrences_offset_t, unsigned) host_station_ocurrences_offset; + DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; + HOST_OUTPUT(host_muon_hits_t, char) host_muon_hits; + + DEVICE_OUTPUT(dev_muon_tracks_t, MuonTrack) dev_muon_tracks; + HOST_OUTPUT(host_muon_tracks_t, MuonTrack) host_velomuon_tracks; + + DEVICE_OUTPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; + HOST_OUTPUT(host_muon_number_of_tracks_t, unsigned) host_muon_number_of_tracks; + // DEVICE_OUTPUT(dev_muon_tracks_tx_t, float) dev_muon_tracks_tx; + // DEVICE_OUTPUT(dev_muon_tracks_ty_t, float) dev_muon_tracks_ty; + // DEVICE_OUTPUT(dev_muon_tracks_chi2x_t, float) dev_muon_tracks_chi2x; + // DEVICE_OUTPUT(dev_muon_tracks_chi2y_t, float) dev_muon_tracks_chi2y; + // DEVICE_OUTPUT(dev_muon_tracks_states_t, int) dev_muon_tracks_states; + PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; + }; + + __global__ void find_muon_hits(Parameters, const Muon::Constants::MatchVeloWindows* dev_match_velo_windows); + + struct find_muon_hits_t : public DeviceAlgorithm, Parameters { + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const; + + void operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + HostBuffers& host_buffers, + const Allen::Context& context) const; + + void output_monitor( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Allen::Context& context) const; + + private: + Property m_block_dim {this, {{64, 1, 1}}}; + }; + +} // namespace find_muon_hits \ No newline at end of file diff --git a/device/muon/match_velo_muon/include/FitMuon.cuh b/device/muon/match_velo_muon/include/FitMuon.cuh new file mode 100644 index 00000000000..76d393e0bc3 --- /dev/null +++ b/device/muon/match_velo_muon/include/FitMuon.cuh @@ -0,0 +1,57 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include "MuonDefinitions.cuh" +#include "UTConsolidated.cuh" +#include "VeloConsolidated.cuh" +#include "States.cuh" + +static constexpr int M2 {0}, M3 {1}, M4 {2}, M5 {3}; + +namespace fit_muon { + struct Parameters { + HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; + DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; + MASK_INPUT(dev_event_list_t) dev_event_list; + DEVICE_INPUT(dev_muon_tracks_t, MuonTrack) dev_muon_tracks; + DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; + DEVICE_INPUT(dev_station_ocurrences_offset_t, unsigned) dev_station_ocurrences_offset; + DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; + + DEVICE_INPUT(dev_muon_tracks_offsets_t, unsigned) dev_muon_tracks_offsets; + HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; + + DEVICE_OUTPUT(dev_muon_tracks_tx_t, float) dev_muon_tracks_tx; + DEVICE_OUTPUT(dev_muon_tracks_ty_t, float) dev_muon_tracks_ty; + DEVICE_OUTPUT(dev_muon_tracks_state_muon_index_t, unsigned) dev_muon_tracks_state_muon_index; + DEVICE_OUTPUT(dev_muon_tracks_chi2x_t, float) dev_muon_tracks_chi2x; + DEVICE_OUTPUT(dev_muon_tracks_chi2y_t, float) dev_muon_tracks_chi2y; + DEVICE_OUTPUT(dev_muon_number_of_fitted_tracks_t, unsigned) dev_muon_number_of_fitted_tracks; + // DEVICE_OUTPUT(dev_output_muon_tracks); + PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; + }; + + __global__ void fit_muon(Parameters); + + struct fit_muon_t : public DeviceAlgorithm, Parameters { + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const; + + void operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + HostBuffers& host_buffers, + const Allen::Context& context) const; + + private: + Property m_block_dim {this, {{64, 1, 1}}}; + }; + +} // namespace fit_muon \ No newline at end of file diff --git a/device/muon/match_velo_muon/include/MatchVeloMuon.cuh b/device/muon/match_velo_muon/include/MatchVeloMuon.cuh new file mode 100644 index 00000000000..add81726d80 --- /dev/null +++ b/device/muon/match_velo_muon/include/MatchVeloMuon.cuh @@ -0,0 +1,70 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include "MuonDefinitions.cuh" +#include "UTConsolidated.cuh" +#include "VeloConsolidated.cuh" +#include "States.cuh" + +static constexpr int M2 {0}, M3 {1}, M4 {2}, M5 {3}; + +namespace match_velo_muon { + struct Parameters { + HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; + MASK_INPUT(dev_event_list_t) dev_event_list; + DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; + DEVICE_INPUT(dev_station_ocurrences_offset_t, unsigned) dev_station_ocurrences_offset; + DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; + HOST_INPUT(host_number_of_reconstructed_velo_tracks_t, unsigned) host_number_of_reconstructed_velo_tracks; + DEVICE_INPUT(dev_muon_tracks_t, MuonTrack) dev_muon_tracks; + DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; + // DEVICE_INPUT(dev_muon_tracks_tx_t, float) dev_muon_tracks_tx; + // DEVICE_INPUT(dev_muon_tracks_ty_t, float) dev_muon_tracks_ty; + // DEVICE_INPUT(dev_muon_tracks_state_muon_index_t, unsigned) dev_muon_tracks_state_muon_index; + + DEVICE_INPUT(dev_velo_tracks_view_t, Allen::Views::Velo::Consolidated::Tracks) dev_velo_tracks_view; + DEVICE_INPUT(dev_velo_states_view_t, Allen::Views::Physics::KalmanStates) dev_velo_states_view; + + DEVICE_INPUT(dev_ut_number_of_selected_velo_tracks_t, unsigned) dev_ut_number_of_selected_velo_tracks; + DEVICE_INPUT(dev_ut_selected_velo_tracks_t, unsigned) dev_ut_selected_velo_tracks; + + DEVICE_INPUT(dev_muon_tracks_offsets_t, unsigned) dev_muon_tracks_offsets; + HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; + + DEVICE_OUTPUT(dev_velomuon_tracks_t, MuonTrack) dev_velomuon_tracks; + + DEVICE_OUTPUT(dev_velomuon_tracks_velo_indices_t, unsigned) dev_velomuon_tracks_velo_indices; + DEVICE_OUTPUT(dev_velomuon_tracks_muon_indices_t, unsigned) dev_velomuon_tracks_muon_indices; + DEVICE_OUTPUT(dev_velomuon_tracks_qop_t, float) dev_velomuon_tracks_qop; + DEVICE_OUTPUT(dev_velomuon_number_of_tracks_t, unsigned) dev_velomuon_number_of_tracks; + DEVICE_OUTPUT(dev_velomuon_muon_id_t, bool) dev_velomuon_muon_id; + + PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; + }; + + __global__ void + match_velo_muon(Parameters, const float* magnet_polarity, const Muon::Constants::MomParam* muon_mom_param); + + struct match_velo_muon_t : public DeviceAlgorithm, Parameters { + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const; + + void operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + HostBuffers& host_buffers, + const Allen::Context& context) const; + + private: + Property m_block_dim {this, {{64, 1, 1}}}; + }; + float xcale[4] {0.06f, 0.1f, 0.15f, 0.15f}; + float distance_cut[5] {30 * 30, 60 * 60, 110 * 110, 200 * 200}; +} // namespace match_velo_muon \ No newline at end of file diff --git a/device/muon/match_velo_muon/src/ConsolidateMuon.cu b/device/muon/match_velo_muon/src/ConsolidateMuon.cu new file mode 100644 index 00000000000..c8c099eafac --- /dev/null +++ b/device/muon/match_velo_muon/src/ConsolidateMuon.cu @@ -0,0 +1,52 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include "ConsolidateMuon.cuh" + +#include "Common.h" +#include "VeloDefinitions.cuh" +#include "VeloEventModel.cuh" +#include + +INSTANTIATE_ALGORITHM(consolidate_muon::consolidate_muon_t) + +void consolidate_muon::consolidate_muon_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const +{ + set_size(arguments, first(arguments)); + // set_size(arguments, first(arguments)); + // set_size(arguments, first(arguments)); + // set_size(arguments, + // first(arguments)) +} + +void consolidate_muon::consolidate_muon_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions&, + const Constants&, + HostBuffers&, + const Allen::Context& context) const +{ + global_function(consolidate_muon)(dim3(size(arguments)), property(), context)( + arguments); +} + +__global__ void consolidate_muon::consolidate_muon(consolidate_muon::Parameters parameters) +{ + const unsigned event_number = parameters.dev_event_list[blockIdx.x]; + + // Input + auto event_muon_tracks_input = + parameters.dev_muon_tracks_input + event_number * Muon::Constants::max_number_of_tracks; + + // Output + auto event_muon_tracks_output = parameters.dev_muon_tracks_output + parameters.dev_muon_tracks_offsets[event_number]; + + for (unsigned i_muon_track = threadIdx.x; i_muon_track < parameters.dev_muon_number_of_tracks[event_number]; + i_muon_track += blockDim.x) { + event_muon_tracks_output[i_muon_track] = event_muon_tracks_input[i_muon_track]; + } +} diff --git a/device/muon/match_velo_muon/src/ConsolidateVeloMuon.cu b/device/muon/match_velo_muon/src/ConsolidateVeloMuon.cu new file mode 100644 index 00000000000..6225b6d0317 --- /dev/null +++ b/device/muon/match_velo_muon/src/ConsolidateVeloMuon.cu @@ -0,0 +1,69 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include "ConsolidateVeloMuon.cuh" + +#include "Common.h" +#include "VeloDefinitions.cuh" +#include "VeloEventModel.cuh" +#include + +INSTANTIATE_ALGORITHM(consolidate_velo_muon::consolidate_velo_muon_t) + +void consolidate_velo_muon::consolidate_velo_muon_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const +{ + set_size(arguments, first(arguments)); + set_size( + arguments, first(arguments)); + set_size( + arguments, first(arguments)); + set_size(arguments, first(arguments)); +} + +void consolidate_velo_muon::consolidate_velo_muon_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions&, + const Constants& constants, + HostBuffers&, + const Allen::Context& context) const +{ + global_function(consolidate_velo_muon)(dim3(size(arguments)), property(), context)( + arguments); +} + +__global__ void consolidate_velo_muon::consolidate_velo_muon(consolidate_velo_muon::Parameters parameters) +{ + const unsigned event_number = parameters.dev_event_list[blockIdx.x]; + + // Input + auto event_velomuon_tracks_velo_indices_input = + parameters.dev_velomuon_tracks_velo_indices_input + event_number * Muon::Constants::max_number_of_tracks; + auto event_velomuon_tracks_muon_indices_input = + parameters.dev_velomuon_tracks_muon_indices_input + event_number * Muon::Constants::max_number_of_tracks; + auto event_velomuon_tracks_qop_input = + parameters.dev_velomuon_tracks_qop_input + event_number * Muon::Constants::max_number_of_tracks; + auto event_velomuon_tracks_input = + parameters.dev_velomuon_tracks_input + event_number * Muon::Constants::max_number_of_tracks; + + // Output + auto event_velomuon_tracks_velo_indices_output = + parameters.dev_velomuon_tracks_velo_indices_output + parameters.dev_velomuon_tracks_offsets[event_number]; + auto event_velomuon_tracks_muon_indices_output = + parameters.dev_velomuon_tracks_muon_indices_output + parameters.dev_velomuon_tracks_offsets[event_number]; + auto event_velomuon_tracks_qop_output = + parameters.dev_velomuon_tracks_qop_output + parameters.dev_velomuon_tracks_offsets[event_number]; + auto event_velomuon_tracks_output = + parameters.dev_velomuon_tracks_output + parameters.dev_velomuon_tracks_offsets[event_number]; + + for (unsigned i_muon_track = threadIdx.x; i_muon_track < parameters.dev_velomuon_number_of_tracks[event_number]; + i_muon_track += blockDim.x) { + event_velomuon_tracks_velo_indices_output[i_muon_track] = event_velomuon_tracks_velo_indices_input[i_muon_track]; + event_velomuon_tracks_muon_indices_output[i_muon_track] = event_velomuon_tracks_muon_indices_input[i_muon_track]; + event_velomuon_tracks_qop_output[i_muon_track] = event_velomuon_tracks_qop_input[i_muon_track]; + event_velomuon_tracks_output[i_muon_track] = event_velomuon_tracks_input[i_muon_track]; + } +} diff --git a/device/muon/match_velo_muon/src/FindMuonHits.cu b/device/muon/match_velo_muon/src/FindMuonHits.cu new file mode 100644 index 00000000000..a0e13e8e685 --- /dev/null +++ b/device/muon/match_velo_muon/src/FindMuonHits.cu @@ -0,0 +1,456 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include "FindMuonHits.cuh" + +#include "Common.h" +#include "VeloDefinitions.cuh" +#include "VeloEventModel.cuh" +#include + +INSTANTIATE_ALGORITHM(find_muon_hits::find_muon_hits_t) + +void find_muon_hits::find_muon_hits_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const +{ + set_size(arguments, 4 * first(arguments) + 1); + set_size(arguments, first(arguments) * 500 * Muon::Hits::element_size); + set_size( + arguments, Muon::Constants::max_number_of_tracks * first(arguments)); + set_size( + arguments, Muon::Constants::max_number_of_tracks * first(arguments)); + set_size(arguments, first(arguments)); + set_size(arguments, first(arguments)); +} + +void find_muon_hits::find_muon_hits_t::output_monitor( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Allen::Context& context) const +{ +#ifdef WITH_ROOT + auto handler = runtime_options.root_service->handle(name()); + auto tree = handler.tree("monitor_tree"); + + float chi2x, chi2y, tx, ty, ax, ay, x0, x1, x2, x3, y0, y1, y2, y3, z0, z1, z2, z3; + handler.branch(tree, "chi2y", chi2y); + handler.branch(tree, "chi2x", chi2x); + handler.branch(tree, "ty", ty); + handler.branch(tree, "tx", tx); + handler.branch(tree, "ay", ay); + handler.branch(tree, "ax", ax); + handler.branch(tree, "x0", x0); + handler.branch(tree, "x1", x1); + handler.branch(tree, "x2", x2); + handler.branch(tree, "x3", x3); + handler.branch(tree, "y0", y0); + handler.branch(tree, "y1", y1); + handler.branch(tree, "y2", y2); + handler.branch(tree, "y3", y3); + handler.branch(tree, "z0", z0); + handler.branch(tree, "z1", z1); + handler.branch(tree, "z2", z2); + handler.branch(tree, "z3", z3); + + Allen::copy(arguments, context); + Allen::copy(arguments, context); + + Allen::copy(arguments, context); + Allen::copy(arguments, context); + Allen::synchronize(context); + const auto n_tracks = data(arguments); + const auto tracks = data(arguments); + const auto n_events = first(arguments); + const auto muon_hits = data(arguments); + const auto muon_offsets = data(arguments); + + const auto muon_total_number_of_hits = muon_offsets[n_events * Muon::Constants::n_stations]; + const auto event_muon_hits = Muon::ConstHits {muon_hits, muon_total_number_of_hits}; + for (unsigned event_number = 0; event_number < n_events; event_number++) { + for (unsigned i_track = 0; i_track < n_tracks[event_number]; i_track++) { + const MuonTrack track = tracks[event_number * Muon::Constants::max_number_of_tracks + i_track]; + std::array x {-99999.f, -99999.f, -99999.f, -99999.f}; + std::array y {-99999.f, -99999.f, -99999.f, -99999.f}; + std::array z {-99999.f, -99999.f, -99999.f, -99999.f}; + + for (unsigned muon_st = 0; muon_st < 4; muon_st++) { + const int muon_idx = track.hit(muon_st); + if (muon_idx != -1) { + x[muon_st] = event_muon_hits.x(muon_idx); + y[muon_st] = event_muon_hits.y(muon_idx); + z[muon_st] = event_muon_hits.z(muon_idx); + } + } + + std::tie(x0, x1, x2, x3) = std::make_tuple(x[0], x[1], x[2], x[3]); + std::tie(y0, y1, y2, y3) = std::make_tuple(y[0], y[1], y[2], y[3]); + std::tie(z0, z1, z2, z3) = std::make_tuple(z[0], z[1], z[2], z[3]); + std::tie(chi2x, chi2y, tx, ty, ax, ay) = std::make_tuple( + track.get_chi2x(), track.get_chi2y(), track.get_tx(), track.get_ty(), track.get_ax(), track.get_ay()); + tree->Fill(); + } + } +#endif +} + +void find_muon_hits::find_muon_hits_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + HostBuffers&, + const Allen::Context& context) const +{ + Allen::memset_async(arguments, 0, context); + + global_function(find_muon_hits)(dim3(size(arguments)), property(), context)( + arguments, constants.dev_match_velo_windows); + Allen::synchronize(context); + output_monitor(arguments, runtime_options, context); + Allen::synchronize(context); +} +__device__ void applyWeightedFit(MuonTrack& muon_track, Muon::ConstHits& muon_hits) +{ + const auto n_hits_track = muon_track.number_of_hits(); + float szx2, szy2, sz_x, sz_y, sx0, sxz, sx, sx2; + szx2 = sz_x = szy2 = sz_y = sx0 = sxz = sx = sx2 = 0.f; + float sy0, syz, sy, sy2; + sy0 = syz = sy = sy2 = 0.f; + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + float x = muon_hits.x(muon_track.hit(i_hit)); + float y = muon_hits.y(muon_track.hit(i_hit)); + float z = muon_hits.z(muon_track.hit(i_hit)); + + float xerr = muon_hits.dx(muon_track.hit(i_hit)); + float yerr = muon_hits.dy(muon_track.hit(i_hit)); + + szx2 += z * z / xerr / xerr; + sz_x += z / xerr / xerr; + sx0 += 1.f / xerr / xerr; + sxz += z * x / xerr / xerr; + sx += x / xerr / xerr; + sx2 += x * x / xerr / xerr; + + szy2 += z * z / yerr / yerr; + sz_y += z / yerr / yerr; + sy0 += 1.f / yerr / yerr; + syz += z * y / yerr / yerr; + sy += y / yerr / yerr; + sy2 += y * y / yerr / yerr; + } + float xdet = szx2 * sx0 - sz_x * sz_x; + float ydet = szy2 * sy0 - sz_y * sz_y; + float tx, ty, ax, ay; + tx = ty = ax = ay = 9999.f; + float chi2xndof = 99999.f; + float chi2yndof = 99999.f; + + if (xdet != 0 && ydet != 0) { + tx = (sxz * sx0 - sx * sz_x) / xdet; + ax = (sx * szx2 - sxz * sz_x) / xdet; + + chi2xndof = (sx2 + tx * tx * szx2 + ax * ax * sx0 - 2.f * tx * sxz - 2.f * ax * sx + 2.f * tx * ax * sz_x) / + (n_hits_track - 2); + + ty = (syz * sy0 - sy * sz_y) / ydet; + ay = (sy * szy2 - syz * sz_y) / ydet; + + chi2yndof = (sy2 + ty * ty * szy2 + ay * ay * sy0 - 2.f * ty * syz - 2.f * ay * sy + 2.f * ty * ay * sz_y) / + (n_hits_track - 2); + } + muon_track.set_tx(tx); + muon_track.set_ty(ty); + muon_track.set_ax(ax); + muon_track.set_ay(ay); + muon_track.set_chi2x(chi2xndof); + muon_track.set_chi2y(chi2yndof); + muon_track.set_state(muon_track.hit(M2)); +} + +__device__ void applyFit(MuonTrack& muon_track, Muon::ConstHits& muon_hits) +{ + const auto n_hits_track = muon_track.number_of_hits(); + float xz = 0; + float yz = 0; + float x = 0; + float y = 0; + float z = 0; + float x2 = 0; + float y2 = 0; + float z2 = 0; + float chi2x = 0; + float chi2y = 0; + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + + xz += muon_hits.x(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + yz += muon_hits.y(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + x += muon_hits.x(muon_track.hit(i_hit)) / n_hits_track; + y += muon_hits.y(muon_track.hit(i_hit)) / n_hits_track; + z += muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + x2 += muon_hits.x(muon_track.hit(i_hit)) * muon_hits.x(muon_track.hit(i_hit)) / n_hits_track; + y2 += muon_hits.y(muon_track.hit(i_hit)) * muon_hits.y(muon_track.hit(i_hit)) / n_hits_track; + z2 += muon_hits.z(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + } + float tx = (xz - x * z) / (z2 - z * z); + float ax = (x - tx * z); + float ty = (yz - y * z) / (z2 - z * z); + float ay = (y - ty * z); + + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + chi2x += (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) * + (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) / 1.f; + chi2y += (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) * + (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) / 1.f; + } + muon_track.set_tx(tx); + muon_track.set_ty(ty); + muon_track.set_ax(ax); + muon_track.set_ay(ay); + auto chi2xndof = chi2x / (n_hits_track - 2); + auto chi2yndof = chi2y / (n_hits_track - 2); + muon_track.set_chi2x(chi2xndof); + muon_track.set_chi2y(chi2yndof); + muon_track.set_state(muon_track.hit(M2)); +} + +__device__ int find_compatible_hit_in_station( + float x, + float y, + unsigned region_seed, + unsigned i_station, + const unsigned* station_ocurrences_offset, + const Muon::ConstHits muon_hits, + const float* Xmax, + const float* Ymax, + const float* tolForRegion) +{ + const auto station_offset = station_ocurrences_offset[i_station]; + const auto nhits_i_station = station_ocurrences_offset[i_station + 1] - station_offset; + + // TODO: Understand logic here... + // https://gitlab.cern.ch/lhcb/Rec/-/blob/master/Tr/TrackTools/src/StandaloneMuonRec.cpp#L200 + const auto xmax = Xmax[i_station * 4 + region_seed]; + const auto ymax = Ymax[i_station * 4 + region_seed]; + + const auto tol = tolForRegion[region_seed]; + + int track_index = -1; + float deltaXmin, deltaYmin; + + for (unsigned i_hit = 0; i_hit < nhits_i_station; i_hit++) { + const auto idx = station_offset + i_hit; + const auto deltaX = fabsf(x - muon_hits.x(idx)); + const auto deltaY = fabsf(y - muon_hits.y(idx)); + + // NOTE: This (previous) implementation depends on the order in which hits were decoded! + // if ( + // deltaX < xmax && deltaY < ymax && + // (track_index == -1 || + // (deltaY < deltaYmin - tol || + // (deltaY < deltaYmin + tol && (deltaX < deltaXmin - tol || fabsf(deltaXmin - deltaX) < 0.1f))))) { + + if (deltaX < xmax && deltaY < ymax && (track_index == -1 || deltaY < deltaYmin && deltaX < deltaXmin)) { + deltaXmin = deltaX; + deltaYmin = deltaY; + track_index = static_cast(i_hit); + } + } + + return track_index; +} + +__device__ std::pair extrapolate( + const Muon::ConstHits muon_hits, + const float* z, + const MuonTrack& muon_track, + const unsigned& previous_index, + const unsigned& current_index, + const unsigned& following_index) +{ + const auto x_current_station = muon_hits.x(muon_track.hit(current_index)); + const auto x_previous_station = muon_hits.x(muon_track.hit(previous_index)); + const auto y_current_station = muon_hits.y(muon_track.hit(current_index)); + const auto y_previous_station = muon_hits.y(muon_track.hit(previous_index)); + const auto z_previous_station = muon_hits.z(muon_track.hit(previous_index)); + const auto z_following_station = z[following_index]; + const auto z_current_station = muon_hits.z(muon_track.hit(current_index)); + auto x = x_current_station + (x_current_station - x_previous_station) / (z_current_station - z_previous_station) * + (z_following_station - z_current_station); + auto y = y_current_station + (y_current_station - y_previous_station) / (z_current_station - z_previous_station) * + (z_following_station - z_current_station); + return std::pair {x, y}; +} + +__device__ void seedAndFind( + Muon::ConstHits muon_hits, + const unsigned* station_ocurrences_offset, + const int* st_order, + unsigned n_stations, + Muon::Constants::MatchVeloWindows match_velo_windows, + int required_station, + unsigned* number_of_muon_tracks_atomic, + MuonTrack* muon_tracks) +{ + + const auto first_st_ocurrences_offset = station_ocurrences_offset[st_order[0]]; + const auto number_of_hits_first_st = station_ocurrences_offset[st_order[0] + 1] - first_st_ocurrences_offset; + for (unsigned seed = threadIdx.x; seed < number_of_hits_first_st; seed += blockDim.x) { + const unsigned first_st_abs_idx = first_st_ocurrences_offset + seed; + const auto xseed = muon_hits.x(first_st_abs_idx); + const auto yseed = muon_hits.y(first_st_abs_idx); + const auto zseed_inverse = 1.f / muon_hits.z(first_st_abs_idx); + auto region_seed = muon_hits.region(first_st_abs_idx); + + MuonTrack muon_track; + muon_track.add_hit_to_station(first_st_abs_idx, st_order[0]); + + float x, y; + std::array last_two_stations_found {st_order[0], -1}; + + for (unsigned i_station = 1; i_station < n_stations; i_station++) { + const auto station = st_order[i_station]; + if (station == st_order[3] && muon_track.number_of_hits() < 2) { + // If we are on the last station and we have less than 2 hits, early exit + break; + } + else if (muon_track.number_of_hits() >= 2) { + // Update x, y parameters by extrapolating the building track + const auto xy = extrapolate( + muon_hits, match_velo_windows.z, muon_track, last_two_stations_found[0], last_two_stations_found[1], station); + x = xy.first; + y = xy.second; + } + else { + // Create x, y parameters by linear extrapolation from origin + const auto slope = match_velo_windows.z[station] * zseed_inverse; + x = xseed * slope; + y = yseed * slope; + } + + const auto hit_index = find_compatible_hit_in_station( + x, + y, + region_seed, + station, + station_ocurrences_offset, + muon_hits, + match_velo_windows.Xmax, + match_velo_windows.Ymax, + match_velo_windows.tolForRegion); + + if (hit_index != -1) { + muon_track.add_hit_to_station(station_ocurrences_offset[station] + hit_index, station); + last_two_stations_found[1] = last_two_stations_found[0]; + last_two_stations_found[0] = station; + region_seed = muon_hits.region(station_ocurrences_offset[station] + hit_index); + } + } + + if (muon_track.number_of_hits() > 3 && muon_track.hit(required_station) != -1) { + applyWeightedFit(muon_track, muon_hits); + const auto insert_index = atomicAdd(number_of_muon_tracks_atomic, 1); + muon_tracks[insert_index] = muon_track; + } + } +} + +__global__ void find_muon_hits::find_muon_hits( + find_muon_hits::Parameters parameters, + const Muon::Constants::MatchVeloWindows* dev_match_velo_windows) +{ + const unsigned event_number = parameters.dev_event_list[blockIdx.x]; + const unsigned number_of_events = parameters.dev_number_of_events[0]; + + const auto muon_total_number_of_hits = + parameters.dev_station_ocurrences_offset[number_of_events * Muon::Constants::n_stations]; + const auto station_ocurrences_offset = + parameters.dev_station_ocurrences_offset + event_number * Muon::Constants::n_stations; + const auto muon_hits = Muon::ConstHits {parameters.dev_muon_hits, muon_total_number_of_hits}; + + // TODO: Change into container taking into consideration actual sizes of M4 and M5 (use prefix sum) + + // Output + auto tracks_offset = event_number * Muon::Constants::max_number_of_tracks; + auto event_muon_tracks = parameters.dev_muon_tracks + tracks_offset; + + auto event_number_of_tracks = parameters.dev_muon_number_of_tracks + event_number; + + // Station processing order + constexpr std::array st_order {M5, M4, M3, M2}; + constexpr std::array st_order2 {M4, M3, M2}; + + const auto match_velo_windows = dev_match_velo_windows[0]; + + // TODO: Test speed removing + __shared__ float muon_tracks_shared_container[Muon::Constants::max_number_of_tracks * sizeof(MuonTrack)]; + MuonTrack* muon_tracks = reinterpret_cast(muon_tracks_shared_container); + __shared__ unsigned number_of_muon_tracks_atomic[1]; + if (threadIdx.x == 0) number_of_muon_tracks_atomic[0] = 0; + + __syncthreads(); + seedAndFind( + muon_hits, + station_ocurrences_offset, + st_order.data(), + 4, + match_velo_windows, + M2, + number_of_muon_tracks_atomic, + muon_tracks); + __syncthreads(); + + // need to tune track making in order to make available second loop: take care of not repeating tracks! + + // seedAndFind( + // muon_hits, + // station_ocurrences_offset, + // st_order2.data(), + // 3, + // match_velo_windows, + // M2, + // number_of_muon_tracks_atomic, + // muon_tracks); + + // Clone killing + const auto is_clone_of = [&](const MuonTrack& track_a, const MuonTrack& track_b) { + if ( + track_a.hit(M2) == track_b.hit(M2) || track_a.hit(M3) == track_b.hit(M3) || + ((track_a.hit(M4) != -1 && track_a.hit(M4) == track_b.hit(M4)) || + (track_a.hit(M5) != -1 && track_a.hit(M5) == track_b.hit(M5)))) { + + auto chi2_ax = track_a.m_chi2x; + auto chi2_ay = track_a.m_chi2y; + + auto chi2_bx = track_b.m_chi2x; + auto chi2_by = track_b.m_chi2y; + + return (((chi2_ax + chi2_ay) > (chi2_bx + chi2_by)) - ((chi2_ax + chi2_ay) < (chi2_bx + chi2_by))); + } + return -1; + }; + + for (unsigned track = threadIdx.x; track < *number_of_muon_tracks_atomic; track += blockDim.x) { + const auto muon_track = muon_tracks[track]; + bool is_clone = false; + + for (unsigned other_track = 0; other_track < number_of_muon_tracks_atomic[0]; ++other_track) { + if (track != other_track) { + const int is_clone_of_value = is_clone_of(muon_track, muon_tracks[other_track]); + if (is_clone_of_value == 1 || (is_clone_of_value == 0 && track > other_track)) { + is_clone = true; + break; + } + } + } + float threshold = 0.25; + if (!is_clone && muon_track.get_chi2x() < 2.f * threshold && muon_track.get_chi2y() < threshold) { + const auto insert_index = atomicAdd(event_number_of_tracks, 1); + event_muon_tracks[insert_index] = muon_track; + } + } +} \ No newline at end of file diff --git a/device/muon/match_velo_muon/src/FitMuon.cu b/device/muon/match_velo_muon/src/FitMuon.cu new file mode 100644 index 00000000000..28595ca7183 --- /dev/null +++ b/device/muon/match_velo_muon/src/FitMuon.cu @@ -0,0 +1,178 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include "FitMuon.cuh" + +#include "Common.h" +#include "VeloDefinitions.cuh" +#include "VeloEventModel.cuh" +#include + +INSTANTIATE_ALGORITHM(fit_muon::fit_muon_t) + +void fit_muon::fit_muon_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const +{ + set_size(arguments, first(arguments)); + set_size(arguments, first(arguments)); + + set_size(arguments, first(arguments)); + + set_size(arguments, first(arguments)); + set_size(arguments, first(arguments)); + + set_size(arguments, first(arguments)); +} + +void fit_muon::fit_muon_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions&, + const Constants&, + HostBuffers&, + const Allen::Context& context) const +{ + // initialize(arguments, 0, context); + // initialize(arguments, 0, context); + // print(arguments); + Allen::memset_async(arguments, 0, context); + global_function(fit_muon)(dim3(size(arguments)), property(), context)(arguments); +} + +__global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) +{ + const unsigned event_number = parameters.dev_event_list[blockIdx.x]; + const unsigned number_of_events = parameters.dev_number_of_events[0]; + + const auto muon_total_number_of_hits = + parameters.dev_station_ocurrences_offset[number_of_events * Muon::Constants::n_stations]; + // const auto station_ocurrences_offset = + // parameters.dev_station_ocurrences_offset + event_number * Muon::Constants::n_stations; + const auto muon_hits = Muon::ConstHits {parameters.dev_muon_hits, muon_total_number_of_hits}; + + // TODO: Change into container taking into consideration actual sizes of M4 and M5 (use prefix sum) + const auto event_muon_tracks = parameters.dev_muon_tracks + event_number * Muon::Constants::max_number_of_tracks; + const auto event_number_of_tracks = parameters.dev_muon_number_of_tracks[event_number]; + const auto event_muon_tracks_offset = parameters.dev_muon_tracks_offsets[event_number]; + + auto event_muon_tx = parameters.dev_muon_tracks_tx + event_muon_tracks_offset; + auto event_muon_ty = parameters.dev_muon_tracks_ty + event_muon_tracks_offset; + auto event_muon_states = parameters.dev_muon_tracks_state_muon_index + event_muon_tracks_offset; + + const auto event_muon_chi2x = parameters.dev_muon_tracks_chi2x + event_muon_tracks_offset; + const auto event_muon_chi2y = parameters.dev_muon_tracks_chi2y + event_muon_tracks_offset; + + auto event_number_of_fitted_tracks = parameters.dev_muon_number_of_fitted_tracks + event_number; + bool normal_fit = true; + for (unsigned idx = threadIdx.x; idx < event_number_of_tracks; idx += blockDim.x) { + const auto muon_track = event_muon_tracks[idx]; + const auto n_hits_track = muon_track.number_of_hits(); + const auto ndof = n_hits_track - 2; + if (normal_fit) { + // printf("Entering normal fit\n"); + float x_mean = 0; + float y_mean = 0; + float z_mean = 0; + float chi2x = 0; + float chi2y = 0; + + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + x_mean += muon_hits.x(muon_track.hit(i_hit)); + y_mean += muon_hits.y(muon_track.hit(i_hit)); + z_mean += muon_hits.z(muon_track.hit(i_hit)); + } + + x_mean = x_mean / n_hits_track; + y_mean = y_mean / n_hits_track; + z_mean = z_mean / n_hits_track; + + float szx = 0; + float szy = 0; + float sz2 = 0; + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + szx += (muon_hits.z(muon_track.hit(i_hit)) - z_mean) * (muon_hits.x(muon_track.hit(i_hit)) - x_mean); + szy += (muon_hits.z(muon_track.hit(i_hit)) - z_mean) * (muon_hits.y(muon_track.hit(i_hit)) - y_mean); + sz2 += (muon_hits.z(muon_track.hit(i_hit)) - z_mean) * (muon_hits.z(muon_track.hit(i_hit)) - z_mean); + } + float tx = szx / sz2; + float ax = x_mean - tx * z_mean; + float ty = szy / sz2; + float ay = y_mean - ty * z_mean; + event_muon_tx[idx] = tx; + event_muon_ty[idx] = ty; + // muon_track.set_tx(tx); + // muon_track.set_ty(ty); + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + chi2x += (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) * + (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) / + muon_hits.x(muon_track.hit(i_hit)); + chi2y += (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) * + (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) / + muon_hits.y(muon_track.hit(i_hit)); + } + + if (chi2x / ndof < 100.f && chi2y / ndof < 100.f) { + // if (true) { + int state_idx = muon_track.hit(M5) != -1.f ? M5 : M4; + state_idx = M2; + const auto insert_idx = atomicAdd(event_number_of_fitted_tracks, 1); + event_muon_chi2x[insert_idx] = chi2x; + event_muon_chi2y[insert_idx] = chi2y; + event_muon_tx[insert_idx] = tx; + event_muon_ty[insert_idx] = ty; + event_muon_states[insert_idx] = muon_track.hit(state_idx); + // __syncthreads(); + } + } + else { + // printf("Entering alternate fit\n"); + float xz = 0; + float yz = 0; + float x = 0; + float y = 0; + float z = 0; + float x2 = 0; + float y2 = 0; + float z2 = 0; + float chi2x = 0; + float chi2y = 0; + // printf("n_hits_track=%i\n",n_hits_track); + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + // printf("summing variables\n"); + xz += muon_hits.x(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + yz += muon_hits.y(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + x += muon_hits.x(muon_track.hit(i_hit)) / n_hits_track; + y += muon_hits.y(muon_track.hit(i_hit)) / n_hits_track; + z += muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + x2 += muon_hits.x(muon_track.hit(i_hit)) * muon_hits.x(muon_track.hit(i_hit)) / n_hits_track; + y2 += muon_hits.y(muon_track.hit(i_hit)) * muon_hits.y(muon_track.hit(i_hit)) / n_hits_track; + z2 += muon_hits.z(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; + } + float tx = (xz - x * z) / (z2 - z * z); + float ax = (x - tx * z); + float ty = (yz - y * z) / (z2 - z * z); + float ay = (y - ty * z); + event_muon_tx[idx] = tx; + event_muon_ty[idx] = ty; + + for (unsigned i_hit = 0; i_hit < 4; i_hit++) { + if (muon_track.hit(i_hit) == -1) continue; + chi2x += (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) * + (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))); + chi2y += (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) * + (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))); + } + // printf("offset=%i\n",event_muon_tracks_offset); + event_muon_chi2x[idx] = chi2x; + event_muon_chi2y[idx] = chi2y; + } + // printf("tx = %f, ty = %f, chi2x/nodf = %f, chi2y/ndof = + // %f\n",event_muon_tx[idx],event_muon_ty[idx],event_muon_chi2x[idx]/2.f,event_muon_chi2y[idx]/2.f); + } +} diff --git a/device/muon/match_velo_muon/src/MatchVeloMuon.cu b/device/muon/match_velo_muon/src/MatchVeloMuon.cu new file mode 100644 index 00000000000..5914254e7cb --- /dev/null +++ b/device/muon/match_velo_muon/src/MatchVeloMuon.cu @@ -0,0 +1,287 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include "MatchVeloMuon.cuh" + +#include "Common.h" +#include "VeloDefinitions.cuh" +#include "VeloEventModel.cuh" +#include + +INSTANTIATE_ALGORITHM(match_velo_muon::match_velo_muon_t) + +void match_velo_muon::match_velo_muon_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions&, + const Constants&, + const HostBuffers&) const +{ + set_size( + arguments, Muon::Constants::max_number_of_tracks * first(arguments)); + set_size( + arguments, Muon::Constants::max_number_of_tracks * first(arguments)); + set_size( + arguments, Muon::Constants::max_number_of_tracks * first(arguments)); + set_size( + arguments, Muon::Constants::max_number_of_tracks * first(arguments)); + set_size(arguments, first(arguments)); + set_size(arguments, first(arguments)); +} + +void match_velo_muon::match_velo_muon_t::operator()( + const ArgumentReferences& arguments, + const RuntimeOptions&, + const Constants& constants, + HostBuffers&, + const Allen::Context& context) const +{ + Allen::memset_async(arguments, 0, context); + Allen::memset_async(arguments, 0, context); + global_function(match_velo_muon)(dim3(size(arguments)), property(), context)( + arguments, constants.dev_magnet_polarity.data(), constants.dev_muon_mom_param); +} + +__device__ float qop_calculation( + Muon::Constants::MomParam const* muon_mom_param, + float const magSign, + float const z0SciFi, + float const x0SciFi, + float const y0SciFi, + float const xVelo, + float const yVelo, + float const zVelo, + float const txO, + float const tyO, + float const txSciFi, + float const tySciFi) +{ + const auto zMatch = (x0SciFi - xVelo + txO * zVelo - txSciFi * z0SciFi) / (txO - txSciFi); + const auto xMatch = xVelo + txO * (zMatch - zVelo); + const auto yMatch = yVelo + tyO * (zMatch - zVelo); + + const auto xVelo_at0 = xVelo - txO * zVelo; + const auto yVelo_at0 = yVelo - tyO * zVelo; + const auto FLIGHTPATH_MAGNET_SCI_SQ = (x0SciFi - xMatch) * (x0SciFi - xMatch) + + (y0SciFi - yMatch) * (y0SciFi - yMatch) + + (z0SciFi - zMatch) * (z0SciFi - zMatch); + const auto FLIGHTPATH_VELO_MAGNET_SQ = + (xVelo_at0 - xMatch) * (xVelo_at0 - xMatch) + (yVelo_at0 - yMatch) * (yVelo_at0 - yMatch) + zMatch * zMatch; + const auto FLIGHTPATH = 0.001f * sqrtf(FLIGHTPATH_MAGNET_SCI_SQ + FLIGHTPATH_VELO_MAGNET_SQ); + const auto MAGFIELD = FLIGHTPATH * cosf(asinf(tyO)); + const auto DSLOPE = + txSciFi / (sqrtf(1.f + txSciFi * txSciFi + tySciFi * tySciFi)) - txO / (sqrtf(1.f + txO * txO + tyO * tyO)); + + const auto txO2 = txO * txO; + const auto txO3 = txO * txO * txO; + const auto txO4 = txO * txO * txO * txO; + const auto txO5 = txO * txO * txO * txO * txO; + const auto txO6 = txO * txO * txO * txO * txO * txO; + const auto txO7 = txO * txO * txO * txO * txO * txO * txO; + const auto tyO2 = tyO * tyO; + const auto tyO4 = tyO * tyO * tyO * tyO; + const auto tyO5 = tyO * tyO * tyO * tyO * tyO; + const auto tyO6 = tyO * tyO * tyO * tyO * tyO * tyO; + + const auto C0 = muon_mom_param->C0[0] + muon_mom_param->C0[1] * txO2 + muon_mom_param->C0[2] * txO4 + + muon_mom_param->C0[3] * tyO2 + muon_mom_param->C0[4] * tyO4 + muon_mom_param->C0[5] * txO2 * tyO2 + + muon_mom_param->C0[6] * txO6 + muon_mom_param->C0[7] * tyO6 + muon_mom_param->C0[8] * txO4 * tyO2 + + muon_mom_param->C0[9] * txO2 * tyO4; + const auto C1 = muon_mom_param->C1[0] + muon_mom_param->C1[1] * txO + muon_mom_param->C1[2] * txO3 + + muon_mom_param->C1[3] * txO5 + muon_mom_param->C1[4] * txO7 + muon_mom_param->C1[5] * tyO2 + + muon_mom_param->C1[6] * tyO4 + muon_mom_param->C1[7] * tyO6 + muon_mom_param->C1[8] * txO * tyO2 + + muon_mom_param->C1[9] * txO * tyO4 + muon_mom_param->C1[10] * txO * tyO6 + + muon_mom_param->C1[11] * txO3 * tyO2 + muon_mom_param->C1[12] * txO3 * tyO4 + + muon_mom_param->C1[13] * txO5 * tyO2; + const auto C2 = muon_mom_param->C2[0] + muon_mom_param->C2[1] * txO2 + muon_mom_param->C2[2] * txO4 + + muon_mom_param->C2[3] * tyO2 + muon_mom_param->C2[4] * tyO4 + muon_mom_param->C2[5] * txO2 * tyO2 + + muon_mom_param->C2[6] * txO6 + muon_mom_param->C2[7] * tyO6 + muon_mom_param->C2[8] * txO4 * tyO2 + + muon_mom_param->C2[9] * txO2 * tyO4; + const auto C3 = muon_mom_param->C3[0] + muon_mom_param->C3[1] * txO + muon_mom_param->C3[2] * txO3 + + muon_mom_param->C3[3] * txO5 + muon_mom_param->C3[4] * txO7 + muon_mom_param->C3[5] * tyO2 + + muon_mom_param->C3[6] * tyO4 + muon_mom_param->C3[7] * tyO6 + muon_mom_param->C3[8] * txO * tyO2 + + muon_mom_param->C3[9] * txO * tyO4 + muon_mom_param->C3[10] * txO * tyO6 + + muon_mom_param->C3[11] * txO3 * tyO2 + muon_mom_param->C3[12] * txO3 * tyO4 + + muon_mom_param->C3[13] * txO5 * tyO2; + const auto C4 = muon_mom_param->C4[0] + muon_mom_param->C4[1] * txO2 + muon_mom_param->C4[2] * txO4 + + muon_mom_param->C4[3] * tyO2 + muon_mom_param->C4[4] * tyO4 + muon_mom_param->C4[5] * txO2 * tyO2 + + muon_mom_param->C4[6] * txO6 + muon_mom_param->C4[7] * tyO6 + muon_mom_param->C4[8] * txO4 * tyO2 + + muon_mom_param->C4[9] * txO2 * tyO4; + + const auto MAGFIELD_updated = + MAGFIELD * magSign * + (C0 + C1 * DSLOPE + C2 * DSLOPE * DSLOPE + C3 * DSLOPE * DSLOPE * DSLOPE + C4 * DSLOPE * DSLOPE * DSLOPE * DSLOPE); + const auto qop = DSLOPE / MAGFIELD_updated; + return qop; +} +__device__ Muon::MatchingResult getChi2Match( + const MiniState velo_state, + float tx_muon, + float ty_muon, + float xpos_muon, + float ypos_muon, + float zpos_muon, + int region) +{ + const float xpos_velo = velo_state.x, ypos_velo = velo_state.y, zpos_velo = velo_state.z, tx_velo = velo_state.tx, + ty_velo = velo_state.ty; + + if (fabsf(tx_velo) > 0.4f || fabsf(ty_velo) > 0.3f) return {-1.f, -1.f, -1.f, -1.f, -1.f, 99999.f}; + const float dSlopeX = tx_velo - tx_muon; + // if (std::abs(dSlopeX) > 0.6f) + // return {-1.f, -1.f, -1.f, -1.f, -1.f, 99999.f}; // matching the UT/muon slopes in X (bending -> large tolerance) + + const float dSlopeY = ty_velo - ty_muon; + if (std::abs(dSlopeY) > 0.05f) + return {-1.f, -1.f, -1.f, -1.f, -1.f, 99999.f}; // matching the UT/muon slopes in Y (no bending) + + const float zForX = 1.f * (MatchUpstreamMuon::za + MatchUpstreamMuon::zb * velo_state.tx * velo_state.tx); + const float zForY = 17500.f; + const float xV = xpos_velo + (zForX - zpos_velo) * tx_velo; + // -- This is the function that calculates the 'bending' in y-direction + // -- The parametrisation can be derived with the MatchFitParams package + const float yV = (ypos_velo + (zForY - zpos_velo) * ty_velo); + + const float xS = xpos_muon + (zForX - zpos_muon) * tx_muon; + const float yS = ypos_muon + (zForY - zpos_muon) * ty_muon; + const float distX = xS - xV; + if (std::abs(distX) > 100.f) return {-1.f, -1.f, -1.f, -1.f, -1.f, 99999.f}; // to scan + const float distY = yS - yV; + // if (std::abs(distY) > 150.f) return {-1.f, -1.f, -1.f, -1.f, -1.f, 99999.f}; // to scan + float fracDistSlope = 0.5; + float fracChi2xy = 0.2f; + float fracChi2xyslope = 0.8f; + + float chi2 = 0.f; + chi2 += fracDistSlope * (fracChi2xy) *distY * distY; + chi2 += fracDistSlope * (1 - fracChi2xy) * distX * distX; + chi2 += (1 - fracDistSlope) * (fracChi2xyslope) *dSlopeY * dSlopeY; + chi2 += (1 - fracDistSlope) * (1 - fracChi2xyslope) * dSlopeX * dSlopeX; + + return {dSlopeX, dSlopeY, distX, distY, zForX, chi2}; +} + +__global__ void match_velo_muon::match_velo_muon( + match_velo_muon::Parameters parameters, + const float* magnet_polarity, + const Muon::Constants::MomParam* muon_mom_param) +{ + const unsigned event_number = parameters.dev_event_list[blockIdx.x]; + const unsigned number_of_events = parameters.dev_number_of_events[0]; + + // Velo views + const auto velo_tracks = parameters.dev_velo_tracks_view[event_number]; + const auto velo_states = parameters.dev_velo_states_view[event_number]; + + const unsigned event_velo_seeds_offset = velo_tracks.offset(); + + // filtered velo tracks + const auto ut_number_of_selected_tracks = parameters.dev_ut_number_of_selected_velo_tracks[event_number]; + const auto ut_selected_velo_tracks = parameters.dev_ut_selected_velo_tracks + event_velo_seeds_offset; + + const auto muon_total_number_of_hits = + parameters.dev_station_ocurrences_offset[number_of_events * Muon::Constants::n_stations]; + const auto muon_hits = Muon::ConstHits {parameters.dev_muon_hits, muon_total_number_of_hits}; + + // TODO: Change into container taking into consideration actual sizes of M4 and M5 (use prefix sum) + const auto event_velomuon_tracks_offset = event_number * Muon::Constants::max_number_of_tracks; + const auto event_muon_tracks_offset = parameters.dev_muon_tracks_offsets[event_number]; + const auto event_number_of_tracks = parameters.dev_muon_number_of_tracks[event_number]; + + const auto event_muon_tracks = parameters.dev_muon_tracks + event_muon_tracks_offset; + + // Output + auto event_velomuon_tracks = parameters.dev_velomuon_tracks + event_velomuon_tracks_offset; + auto event_velomuon_velo_indices = parameters.dev_velomuon_tracks_velo_indices + event_velomuon_tracks_offset; + auto event_velomuon_muon_indices = parameters.dev_velomuon_tracks_muon_indices + event_velomuon_tracks_offset; + auto event_velomuon_qop = parameters.dev_velomuon_tracks_qop + event_velomuon_tracks_offset; + auto event_velomuon_number_of_tracks = parameters.dev_velomuon_number_of_tracks + event_number; + auto event_velomuon_muon_id = parameters.dev_velomuon_muon_id + event_velo_seeds_offset; + + __shared__ unsigned number_of_matched_velo[1]; + const auto number = Muon::Constants::max_number_of_tracks; + + __shared__ unsigned event_matched_muon_indices_with_clones[number]; + __shared__ unsigned event_muon_velo_indices_with_clones[number]; + __shared__ float event_match_chi2[number]; + __shared__ float event_qop_with_clones[number]; + + if (threadIdx.x == 0) number_of_matched_velo[0] = 0; + __syncthreads(); + for (unsigned i_muon_track = threadIdx.x; i_muon_track < event_number_of_tracks; i_muon_track += blockDim.x) { + auto muon_track = event_muon_tracks[i_muon_track]; + auto tx_muon = muon_track.get_tx(); + auto ty_muon = muon_track.get_ty(); + + auto muon_idx = muon_track.get_state(); + auto xpos_muon = muon_track.state_x(); + auto ypos_muon = muon_track.state_y(); + auto zpos_muon = muon_track.state_z(); + auto region = muon_hits.region(muon_idx); + + int matched_velo_idx = -1; + float min_chi2 = 99999.f; + + for (unsigned ivelo = 0; ivelo < ut_number_of_selected_tracks; ivelo++) { + const auto velo_track_index = ut_selected_velo_tracks[ivelo]; + const auto endvelo_state = velo_states.state(velo_track_index); + + auto matching_result = getChi2Match(endvelo_state, tx_muon, ty_muon, xpos_muon, ypos_muon, zpos_muon, region); + auto chi2 = matching_result.chi2; + if (chi2 < min_chi2 || matched_velo_idx == -1) { + matched_velo_idx = velo_track_index; + min_chi2 = chi2; + } + } + if (min_chi2 >= 99999.f) continue; + const auto insert_index = atomicAdd(number_of_matched_velo, 1); + event_matched_muon_indices_with_clones[insert_index] = i_muon_track; + event_muon_velo_indices_with_clones[insert_index] = matched_velo_idx; + const auto velo_state = velo_states.state(matched_velo_idx); + event_match_chi2[insert_index] = min_chi2; + float poq = MatchUpstreamMuon::kickScale / fabsf(tx_muon - velo_states.state(matched_velo_idx).tx()) + + MatchUpstreamMuon::kickOffset; + float qop = qop_calculation( + muon_mom_param, + magnet_polarity[0], + zpos_muon, + xpos_muon, + ypos_muon, + velo_state.x(), + velo_state.y(), + velo_state.z(), + velo_state.tx(), + velo_state.ty(), + tx_muon, + ty_muon); + event_qop_with_clones[insert_index] = qop; + } + __syncthreads(); + + for (unsigned n_track_1 = threadIdx.x; n_track_1 < number_of_matched_velo[0]; n_track_1 += blockDim.x) { + int velo_index_1 = static_cast(event_muon_velo_indices_with_clones[n_track_1]); + int velo_index_2 = 0; + bool shared_seeds = false; + float minimum_weight = 99999.f; + unsigned track_2_id = 0; + for (unsigned n_track_2 = 0; n_track_2 < number_of_matched_velo[0]; n_track_2++) { + if (n_track_1 == n_track_2) continue; + velo_index_2 = static_cast(event_muon_velo_indices_with_clones[n_track_2]); + if (velo_index_1 == velo_index_2 && minimum_weight > event_match_chi2[n_track_2]) { + shared_seeds = true; + track_2_id = n_track_2; + minimum_weight = event_match_chi2[n_track_2]; + } + } + if (!(shared_seeds && (event_match_chi2[track_2_id] <= event_match_chi2[n_track_1]))) { + const auto insert_index = atomicAdd(event_velomuon_number_of_tracks, 1); + auto velomuon_track = MuonTrack(event_muon_tracks[event_matched_muon_indices_with_clones[n_track_1]]); + velomuon_track.set_p(1 / event_qop_with_clones[n_track_1]); + velomuon_track.set_velo_index(velo_index_1); + event_velomuon_tracks[insert_index] = velomuon_track; + event_velomuon_velo_indices[insert_index] = velo_index_1; + event_velomuon_muon_indices[insert_index] = event_matched_muon_indices_with_clones[n_track_1]; + event_velomuon_qop[insert_index] = event_qop_with_clones[n_track_1]; + event_velomuon_muon_id[velo_index_1] = true; + } + } +} diff --git a/device/selections/lines/muon/include/OneMuonTrackLine.cuh b/device/selections/lines/muon/include/OneMuonTrackLine.cuh new file mode 100644 index 00000000000..faf3fc26111 --- /dev/null +++ b/device/selections/lines/muon/include/OneMuonTrackLine.cuh @@ -0,0 +1,85 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include "Line.cuh" +#include "MuonDefinitions.cuh" + +namespace one_muon_track_line { + struct Parameters { + // Commonly required inputs, outputs and properties + HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; + DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; + MASK_INPUT(dev_event_list_t) dev_event_list; + HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size; + HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler; + HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash; + HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters; + // Line-specific inputs and properties + DEVICE_INPUT(dev_muon_tracks_t, MuonTrack) dev_muon_tracks; + DEVICE_INPUT(dev_muon_tracks_offsets_t, unsigned) dev_muon_tracks_offsets; + DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; + HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; + + // PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler; + PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler; + PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler; + PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string); + PROPERTY(post_scaler_hash_string_t, "post_scaler_hash_string", "Post-scaling hash string", std::string); + }; + + // SelectionAlgorithm definition + struct one_muon_track_line_t : public SelectionAlgorithm, Parameters, Line { + + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + const HostBuffers& host_buffers) const; + + // Offset function + __device__ static unsigned offset(const Parameters& parameters, const unsigned event_number) + { + return parameters.dev_muon_tracks_offsets[event_number]; + } + + // Get decision size function + static unsigned get_decisions_size(const ArgumentReferences& arguments) + { + return first(arguments); + } + + // Get input function + __device__ static std::tuple + get_input(const Parameters& parameters, const unsigned event_number, const unsigned i) + { + // Get the number of events + // const unsigned number_of_events = parameters.dev_number_of_events[0]; + const auto muon_tracks_offsets = parameters.dev_muon_tracks_offsets; + const auto muon_tracks = parameters.dev_muon_tracks; + + // Get the ith velo track + const unsigned track_index = i + muon_tracks_offsets[event_number]; + + return std::forward_as_tuple(muon_tracks[track_index]); + } + + __device__ static unsigned input_size(const Parameters& parameters, const unsigned event_number) + { + return parameters.dev_muon_number_of_tracks[event_number]; + } + + // Selection function + __device__ static bool select(const Parameters& parameters, std::tuple input); + + private: + // Commonly required properties + Property m_pre_scaler {this, 1.f}; + Property m_post_scaler {this, 1.f}; + Property m_pre_scaler_hash_string {this, ""}; + Property m_post_scaler_hash_string {this, ""}; + // Line-specific properties + }; +} // namespace one_muon_track_line \ No newline at end of file diff --git a/device/selections/lines/muon/include/OneVeloMuonTrackLine.cuh b/device/selections/lines/muon/include/OneVeloMuonTrackLine.cuh new file mode 100644 index 00000000000..e7a8ad8bd08 --- /dev/null +++ b/device/selections/lines/muon/include/OneVeloMuonTrackLine.cuh @@ -0,0 +1,87 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#pragma once + +#include "AlgorithmTypes.cuh" +#include "Line.cuh" +#include "MuonDefinitions.cuh" + +namespace one_velomuon_track_line { + struct Parameters { + // Commonly required inputs, outputs and properties + HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; + DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; + MASK_INPUT(dev_event_list_t) dev_event_list; + HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size; + HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler; + HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash; + HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters; + // Line-specific inputs and properties + DEVICE_INPUT(dev_velomuon_tracks_t, MuonTrack) dev_velomuon_tracks; + DEVICE_INPUT(dev_velomuon_tracks_qop_t, float) dev_velomuon_tracks_qop; + DEVICE_INPUT(dev_velomuon_tracks_offsets_t, unsigned) dev_velomuon_tracks_offsets; + DEVICE_INPUT(dev_velomuon_number_of_tracks_t, unsigned) dev_velomuon_number_of_tracks; + HOST_INPUT(host_velomuon_total_number_of_tracks_t, unsigned) host_velomuon_total_number_of_tracks; + + PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler; + PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler; + PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string); + PROPERTY(post_scaler_hash_string_t, "post_scaler_hash_string", "Post-scaling hash string", std::string); + }; + + // SelectionAlgorithm definition + struct one_velomuon_track_line_t : public SelectionAlgorithm, + Parameters, + Line { + + void set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + const HostBuffers& host_buffers) const; + + // Offset function + __device__ static unsigned offset(const Parameters& parameters, const unsigned event_number) + { + return parameters.dev_velomuon_tracks_offsets[event_number]; + } + + // Get decision size function + static unsigned get_decisions_size(const ArgumentReferences& arguments) + { + return first(arguments); + } + + // Get input function + __device__ static std::tuple + get_input(const Parameters& parameters, const unsigned event_number, const unsigned i) + { + // Get the number of events + // const unsigned number_of_events = parameters.dev_number_of_events[0]; + const auto velomuon_tracks_offsets = parameters.dev_velomuon_tracks_offsets; + const auto velomuon_track = parameters.dev_velomuon_tracks; + + // Get the ith velo track + const unsigned track_index = i + velomuon_tracks_offsets[event_number]; + + return std::forward_as_tuple(velomuon_track[track_index]); + } + + __device__ static unsigned input_size(const Parameters& parameters, const unsigned event_number) + { + return parameters.dev_velomuon_number_of_tracks[event_number]; + } + + // Selection function + __device__ static bool select(const Parameters& parameters, std::tuple input); + + private: + // Commonly required properties + Property m_pre_scaler {this, 1.f}; + Property m_post_scaler {this, 1.f}; + Property m_pre_scaler_hash_string {this, ""}; + Property m_post_scaler_hash_string {this, ""}; + // Line-specific properties + }; +} // namespace one_velomuon_track_line \ No newline at end of file diff --git a/device/selections/lines/muon/src/OneMuonTrackLine.cu b/device/selections/lines/muon/src/OneMuonTrackLine.cu new file mode 100644 index 00000000000..7df2ab7cee0 --- /dev/null +++ b/device/selections/lines/muon/src/OneMuonTrackLine.cu @@ -0,0 +1,25 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include "OneMuonTrackLine.cuh" + +// Explicit instantiation of the line +INSTANTIATE_LINE(one_muon_track_line::one_muon_track_line_t, one_muon_track_line::Parameters) + +void one_muon_track_line::one_muon_track_line_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + const HostBuffers& host_buffers) const +{ + static_cast(this)->set_arguments_size(arguments, runtime_options, constants, host_buffers); +} +// Selection function +__device__ bool one_muon_track_line::one_muon_track_line_t::select(const Parameters&, std::tuple input) +{ + const auto& muon_track = std::get<0>(input); + const bool decision = + (fabsf(muon_track.get_tx()) > 0.2f && muon_track.get_chi2x() < 0.1f && muon_track.get_chi2y() < 0.1f); + + return decision; +} \ No newline at end of file diff --git a/device/selections/lines/muon/src/OneVeloMuonTrackLine.cu b/device/selections/lines/muon/src/OneVeloMuonTrackLine.cu new file mode 100644 index 00000000000..1a96ff7c7c7 --- /dev/null +++ b/device/selections/lines/muon/src/OneVeloMuonTrackLine.cu @@ -0,0 +1,30 @@ +/*****************************************************************************\ +* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * +\*****************************************************************************/ +#include "OneVeloMuonTrackLine.cuh" + +// Explicit instantiation of the line +INSTANTIATE_LINE(one_velomuon_track_line::one_velomuon_track_line_t, one_velomuon_track_line::Parameters) + +void one_velomuon_track_line::one_velomuon_track_line_t::set_arguments_size( + ArgumentReferences arguments, + const RuntimeOptions& runtime_options, + const Constants& constants, + const HostBuffers& host_buffers) const +{ + static_cast(this)->set_arguments_size(arguments, runtime_options, constants, host_buffers); +} +// Selection function +__device__ bool one_velomuon_track_line::one_velomuon_track_line_t::select( + const Parameters&, + std::tuple input) +{ + const auto& velomuon_track = std::get<0>(input); + const auto tx = velomuon_track.get_tx(); + const auto ty = velomuon_track.get_ty(); + const auto p = velomuon_track.get_p(); + const auto pt = sqrtf((tx * tx + ty * ty) / (1 + tx * tx + ty * ty)) * p; + const bool decision = (p > 5000 && pt > 500); + + return decision; +} \ No newline at end of file diff --git a/device/validators/include/LongTrackValidator.cuh b/device/validators/include/LongTrackValidator.cuh index d2ea081d800..685e41982cb 100644 --- a/device/validators/include/LongTrackValidator.cuh +++ b/device/validators/include/LongTrackValidator.cuh @@ -16,13 +16,22 @@ namespace long_track_validator { struct Parameters { HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; HOST_INPUT(host_number_of_reconstructed_long_tracks_t, unsigned) host_number_of_reconstructed_long_tracks; + HOST_INPUT(host_number_of_reconstructed_velomuon_tracks_t, unsigned) host_number_of_reconstructed_velomuon_tracks; HOST_INPUT(host_mc_events_t, const MCEvents*) host_mc_events; DEVICE_INPUT(dev_velo_states_view_t, Allen::Views::Physics::KalmanStates) dev_velo_states_view; + DEVICE_INPUT(dev_velo_tracks_view_t, Allen::Views::Velo::Consolidated::Tracks) dev_velo_tracks_view; MASK_INPUT(dev_event_list_t) dev_event_list; DEVICE_INPUT(dev_multi_event_long_tracks_view_t, Allen::Views::Physics::MultiEventLongTracks) dev_multi_event_long_tracks_view; DEVICE_INPUT(dev_offsets_long_tracks_t, unsigned) dev_offsets_long_tracks; + DEVICE_OUTPUT(dev_long_checker_tracks_t, Checker::Track) dev_long_checker_tracks; + DEVICE_OUTPUT(dev_velomuon_checker_tracks_t, Checker::Track) dev_velomuon_checker_tracks; + + DEVICE_INPUT(dev_velomuon_tracks_velo_indices_t, unsigned) dev_velomuon_tracks_velo_indices; + DEVICE_INPUT(dev_velomuon_tracks_qop_t, float) dev_velomuon_tracks_qop; + DEVICE_INPUT(dev_offsets_velomuon_tracks_t, unsigned) dev_offsets_velomuon_tracks; + PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; PROPERTY(root_output_filename_t, "root_output_filename", "root output filename", std::string); }; diff --git a/device/validators/include/MuonValidator.cuh b/device/validators/include/MuonValidator.cuh index 7487ea80eee..bd49e7ef9ad 100644 --- a/device/validators/include/MuonValidator.cuh +++ b/device/validators/include/MuonValidator.cuh @@ -16,14 +16,19 @@ namespace muon_validator { struct Parameters { HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; HOST_INPUT(host_number_of_reconstructed_long_tracks_t, unsigned) host_number_of_reconstructed_long_tracks; + HOST_INPUT(host_number_of_reconstructed_velo_tracks_t, unsigned) host_number_of_reconstructed_velo_tracks; HOST_INPUT(host_mc_events_t, const MCEvents*) host_mc_events; DEVICE_INPUT(dev_velo_states_view_t, Allen::Views::Physics::KalmanStates) dev_velo_states_view; MASK_INPUT(dev_event_list_t) dev_event_list; DEVICE_INPUT(dev_multi_event_long_tracks_view_t, Allen::Views::Physics::MultiEventLongTracks) dev_multi_event_long_tracks_view; + DEVICE_INPUT(dev_velo_tracks_view_t, Allen::Views::Velo::Consolidated::Tracks) dev_velo_tracks_view; + DEVICE_INPUT(dev_offsets_velo_tracks_t, unsigned) dev_offsets_velo_tracks; DEVICE_INPUT(dev_offsets_long_tracks_t, unsigned) dev_offsets_long_tracks; DEVICE_INPUT(dev_is_muon_t, bool) dev_is_muon; + DEVICE_INPUT(dev_match_velo_muon_t, bool) dev_match_velo_muon; DEVICE_OUTPUT(dev_muon_checker_tracks_t, Checker::Track) dev_muon_checker_tracks; + DEVICE_OUTPUT(dev_muon_checker_velotracks_t, Checker::Track) dev_muon_checker_velotracks; PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; PROPERTY(root_output_filename_t, "root_output_filename", "root output filename", std::string); }; diff --git a/device/validators/src/LongTrackValidator.cu b/device/validators/src/LongTrackValidator.cu index 369b93260c0..a3fc7177e0e 100644 --- a/device/validators/src/LongTrackValidator.cu +++ b/device/validators/src/LongTrackValidator.cu @@ -10,11 +10,25 @@ __global__ void long_track_validator::long_track_validator(long_track_validator: { const unsigned event_number = blockIdx.x; const auto event_long_tracks = parameters.dev_multi_event_long_tracks_view->container(event_number); + const auto event_velo_tracks = parameters.dev_velo_tracks_view[event_number]; const auto endvelo_states = parameters.dev_velo_states_view[event_number]; const unsigned offset_long_tracks = event_long_tracks.offset(); - Checker::Track* long_checker_tracks_event = parameters.dev_long_checker_tracks + offset_long_tracks; + const unsigned offset_velomuon_tracks = parameters.dev_offsets_velomuon_tracks[event_number]; + const unsigned number_of_velomuon_tracks = + parameters.dev_offsets_velomuon_tracks[event_number + 1] - parameters.dev_offsets_velomuon_tracks[event_number]; + const float* event_velomuon_qop = parameters.dev_velomuon_tracks_qop + offset_velomuon_tracks; + const unsigned* event_velomuon_velo_indices = parameters.dev_velomuon_tracks_velo_indices + offset_velomuon_tracks; + Checker::Track* long_checker_tracks_event = parameters.dev_long_checker_tracks + offset_long_tracks; + Checker::Track* velomuon_checker_tracks_event = parameters.dev_velomuon_checker_tracks + offset_velomuon_tracks; prepare_long_tracks(event_long_tracks, endvelo_states, long_checker_tracks_event); + prepare_velomuon_tracks( + number_of_velomuon_tracks, + event_velo_tracks, + endvelo_states, + event_velomuon_qop, + event_velomuon_velo_indices, + velomuon_checker_tracks_event); } void long_track_validator::long_track_validator_t::set_arguments_size( @@ -24,6 +38,7 @@ void long_track_validator::long_track_validator_t::set_arguments_size( const HostBuffers&) const { set_size(arguments, first(arguments)); + set_size(arguments, first(arguments)); } void long_track_validator::long_track_validator_t::operator()( @@ -38,18 +53,34 @@ void long_track_validator::long_track_validator_t::operator()( const auto event_list = make_host_buffer(arguments, context); const auto long_tracks_for_checker = make_host_buffer(arguments, context); const auto event_tracks_offsets = make_host_buffer(arguments, context); + + const auto velomuon_tracks_for_checker = make_host_buffer(arguments, context); + const auto event_velomuon_offsets = make_host_buffer(arguments, context); std::vector tracks; + std::vector velomuon_tracks; + tracks.resize(event_list.size()); + velomuon_tracks.resize(event_list.size()); for (size_t i = 0; i < event_list.size(); ++i) { const auto evnum = event_list[i]; const auto event_offset = event_tracks_offsets[evnum]; + const auto event_velomuon_offset = event_velomuon_offsets[evnum]; const auto n_tracks = event_tracks_offsets[evnum + 1] - event_offset; - std::vector event_trakcs = {long_tracks_for_checker.begin() + event_offset, + const auto n_velomuon_tracks = event_velomuon_offsets[evnum + 1] - event_velomuon_offset; + std::vector event_tracks = {long_tracks_for_checker.begin() + event_offset, long_tracks_for_checker.begin() + event_offset + n_tracks}; - tracks[i] = event_trakcs; + + std::vector event_velomuon_tracks = {velomuon_tracks_for_checker.begin() + event_velomuon_offset, + velomuon_tracks_for_checker.begin() + event_velomuon_offset + + n_velomuon_tracks}; + tracks[i] = event_tracks; + velomuon_tracks[i] = event_velomuon_tracks; } auto& checker = runtime_options.checker_invoker->checker(name(), property()); checker.accumulate(*first(arguments), tracks, event_list); + auto& velomuon_checker = + runtime_options.checker_invoker->checker("velomuon tracks", property()); + velomuon_checker.accumulate(*first(arguments), velomuon_tracks, event_list); } \ No newline at end of file diff --git a/device/validators/src/MuonValidator.cu b/device/validators/src/MuonValidator.cu index 769842037d0..c886a867a97 100644 --- a/device/validators/src/MuonValidator.cu +++ b/device/validators/src/MuonValidator.cu @@ -10,14 +10,19 @@ __global__ void muon_validator::muon_validator(muon_validator::Parameters parame { const unsigned event_number = blockIdx.x; const auto event_long_tracks = parameters.dev_multi_event_long_tracks_view->container(event_number); + const auto event_velo_tracks = parameters.dev_velo_tracks_view[event_number]; const auto endvelo_states = parameters.dev_velo_states_view[event_number]; const unsigned offset_long_tracks = event_long_tracks.offset(); const bool* is_muon = parameters.dev_is_muon + offset_long_tracks; + const bool* match_velo_muon = parameters.dev_match_velo_muon + event_velo_tracks.offset(); Checker::Track* muon_checker_tracks_event = parameters.dev_muon_checker_tracks + offset_long_tracks; + Checker::Track* muon_checker_velotracks_event = parameters.dev_muon_checker_velotracks + event_velo_tracks.offset(); prepare_long_tracks(event_long_tracks, endvelo_states, muon_checker_tracks_event); + prepare_velo_tracks(event_velo_tracks, muon_checker_velotracks_event); prepare_muons(event_long_tracks.size(), muon_checker_tracks_event, is_muon); + prepare_muons(event_velo_tracks.size(), muon_checker_velotracks_event, match_velo_muon); } void muon_validator::muon_validator_t::set_arguments_size( @@ -27,6 +32,7 @@ void muon_validator::muon_validator_t::set_arguments_size( const HostBuffers&) const { set_size(arguments, first(arguments)); + set_size(arguments, first(arguments)); } void muon_validator::muon_validator_t::operator()( @@ -41,18 +47,33 @@ void muon_validator::muon_validator_t::operator()( const auto event_list = make_host_buffer(arguments, context); const auto muon_tracks_for_checker = make_host_buffer(arguments, context); const auto event_tracks_offsets = make_host_buffer(arguments, context); + + const auto muon_velotracks_for_checker = make_host_buffer(arguments, context); + const auto event_velotracks_offsets = make_host_buffer(arguments, context); std::vector tracks; + std::vector velotracks; tracks.resize(event_list.size()); + velotracks.resize(event_list.size()); for (size_t i = 0; i < event_list.size(); ++i) { const auto evnum = event_list[i]; const auto event_offset = event_tracks_offsets[evnum]; + const auto event_velooffset = event_velotracks_offsets[evnum]; const auto n_tracks = event_tracks_offsets[evnum + 1] - event_offset; - std::vector event_trakcs = {muon_tracks_for_checker.begin() + event_offset, + const auto n_velotracks = event_velotracks_offsets[evnum + 1] - event_velooffset; + std::vector event_tracks = {muon_tracks_for_checker.begin() + event_offset, muon_tracks_for_checker.begin() + event_offset + n_tracks}; - tracks[i] = event_trakcs; + + std::vector event_velotracks = {muon_velotracks_for_checker.begin() + event_velooffset, + muon_velotracks_for_checker.begin() + event_velooffset + + n_velotracks}; + tracks[i] = event_tracks; + velotracks[i] = event_velotracks; } auto& checker = runtime_options.checker_invoker->checker(name(), property()); checker.accumulate(*first(arguments), tracks, event_list); + auto& checker_velo = + runtime_options.checker_invoker->checker("velo muon checker", property()); + checker_velo.accumulate(*first(arguments), velotracks, event_list); } \ No newline at end of file diff --git a/host/validators/include/HostVeloValidator.h b/host/validators/include/HostVeloValidator.h index 45347566b58..347080db0cd 100644 --- a/host/validators/include/HostVeloValidator.h +++ b/host/validators/include/HostVeloValidator.h @@ -12,6 +12,7 @@ namespace host_velo_validator { DEVICE_INPUT(dev_offsets_all_velo_tracks_t, unsigned) dev_offsets_all_velo_tracks; DEVICE_INPUT(dev_offsets_velo_track_hit_number_t, unsigned) dev_offsets_velo_track_hit_number; DEVICE_INPUT(dev_velo_track_hits_t, char) dev_velo_track_hits; + // DEVICE_INPUT(dev_velomuon_muon_id_t, bool) dev_velomuon_muon_id; MASK_INPUT(dev_event_list_t) dev_event_list; HOST_INPUT(host_mc_events_t, const MCEvents*) host_mc_events; PROPERTY(root_output_filename_t, "root_output_filename", "root output filename", std::string); diff --git a/host/validators/src/HostVeloValidator.cpp b/host/validators/src/HostVeloValidator.cpp index 11ffccb59b3..4c4084c7a81 100644 --- a/host/validators/src/HostVeloValidator.cpp +++ b/host/validators/src/HostVeloValidator.cpp @@ -23,6 +23,7 @@ void host_velo_validator::host_velo_validator_t::operator()( offsets_all_velo_tracks, offsets_velo_track_hit_number, velo_track_hits, + // velomuon_id, event_list); auto& checker = diff --git a/stream/sequence/include/Constants.cuh b/stream/sequence/include/Constants.cuh index 78f05ab8334..7de5fa927d0 100644 --- a/stream/sequence/include/Constants.cuh +++ b/stream/sequence/include/Constants.cuh @@ -20,7 +20,9 @@ namespace Muon { class MuonTables; namespace Constants { struct FieldOfInterest; - } + struct MatchVeloWindows; + struct MomParam; + } // namespace Constants } // namespace Muon namespace LookingForward { struct Constants; @@ -100,6 +102,8 @@ struct Constants { std::vector host_muon_lookup_tables_raw; Muon::MuonGeometry* dev_muon_geometry = nullptr; Muon::MuonTables* dev_muon_tables = nullptr; + Muon::Constants::MatchVeloWindows* dev_match_velo_windows = nullptr; + Muon::Constants::MomParam* dev_muon_mom_param = nullptr; // Velo-UT-muon MatchUpstreamMuon::MuonChambers* dev_muonmatch_search_muon_chambers = nullptr; diff --git a/stream/sequence/src/Constants.cpp b/stream/sequence/src/Constants.cpp index d752d9e9d2d..9e78154cb00 100644 --- a/stream/sequence/src/Constants.cpp +++ b/stream/sequence/src/Constants.cpp @@ -22,6 +22,8 @@ void Constants::reserve_constants() Allen::malloc((void**) &dev_muon_momentum_cuts, 3 * sizeof(float)); Allen::malloc((void**) &dev_muonmatch_search_muon_chambers, sizeof(MatchUpstreamMuon::MuonChambers)); Allen::malloc((void**) &dev_muonmatch_search_windows, sizeof(MatchUpstreamMuon::SearchWindows)); + Allen::malloc((void**) &dev_match_velo_windows, sizeof(Muon::Constants::MatchVeloWindows)); + Allen::malloc((void**) &dev_muon_mom_param, sizeof(Muon::Constants::MomParam)); host_ut_region_offsets.resize(UT::Constants::n_layers * UT::Constants::n_regions_in_layer + 1); host_ut_dxDy.resize(UT::Constants::n_layers); @@ -70,6 +72,17 @@ void Constants::initialize_constants( Allen::memcpy(dev_muon_momentum_cuts, &Muon::Constants::momentum_cuts, 3 * sizeof(float), Allen::memcpyHostToDevice); Allen::memcpy(dev_muon_foi, &host_muon_foi, sizeof(Muon::Constants::FieldOfInterest), Allen::memcpyHostToDevice); + // VeloMuon + Muon::Constants::MatchVeloWindows host_match_velo_windows; + Allen::memcpy( + dev_match_velo_windows, + &host_match_velo_windows, + sizeof(Muon::Constants::MatchVeloWindows), + Allen::memcpyHostToDevice); + + Muon::Constants::MomParam host_muon_mom_param; + Allen::memcpy(dev_muon_mom_param, &host_muon_mom_param, sizeof(Muon::Constants::MomParam), Allen::memcpyHostToDevice); + // Velo-UT-muon MatchUpstreamMuon::MuonChambers host_muonmatch_search_muon_chambers; MatchUpstreamMuon::SearchWindows host_muonmatch_search_windows; -- GitLab From caec21de992eaa7d5258d16513b3880a6056a774 Mon Sep 17 00:00:00 2001 From: Gitlab CI Date: Tue, 27 Sep 2022 15:59:32 +0000 Subject: [PATCH 2/6] Fixed formatting patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/24817858 --- configuration/python/AllenConf/HLT1.py | 5 +---- .../python/AllenConf/hlt1_reconstruction.py | 18 ++++++++---------- .../python/AllenConf/muon_reconstruction.py | 5 +++-- 3 files changed, 12 insertions(+), 16 deletions(-) diff --git a/configuration/python/AllenConf/HLT1.py b/configuration/python/AllenConf/HLT1.py index a8572f6dec7..a680b08cd92 100644 --- a/configuration/python/AllenConf/HLT1.py +++ b/configuration/python/AllenConf/HLT1.py @@ -22,7 +22,6 @@ from AllenConf.lumi_reconstruction import lumi_reconstruction def default_physics_lines(reconstructed_objects, with_calo, with_muon): - velo_tracks = reconstructed_objects["velo_tracks"] long_tracks = reconstructed_objects["long_tracks"] @@ -57,7 +56,7 @@ def default_physics_lines(reconstructed_objects, with_calo, with_muon): velo_muon_objects['find_muon_hits'], velo_muon_objects['consolidate_muon'], velo_muon_objects['prefix_sum_muon_tracks'], - name = "Hlt1OneMuonTrackLine"), + name="Hlt1OneMuonTrackLine"), make_single_high_pt_muon_line( long_tracks, long_track_particles, name="Hlt1SingleHighPtMuon"), @@ -96,8 +95,6 @@ def default_physics_lines(reconstructed_objects, with_calo, with_muon): pre_scaler=.1) ] - - if with_calo: ecal_clusters = reconstructed_objects["ecal_clusters"] calo_matching_objects = reconstructed_objects["calo_matching_objects"] diff --git a/configuration/python/AllenConf/hlt1_reconstruction.py b/configuration/python/AllenConf/hlt1_reconstruction.py index 7cc4b1ee6a3..526aa50845b 100755 --- a/configuration/python/AllenConf/hlt1_reconstruction.py +++ b/configuration/python/AllenConf/hlt1_reconstruction.py @@ -5,7 +5,7 @@ from AllenConf.velo_reconstruction import decode_velo, make_velo_tracks, run_vel from AllenConf.ut_reconstruction import decode_ut, make_ut_tracks from AllenConf.scifi_reconstruction import decode_scifi, make_forward_tracks, make_seeding_XZ_tracks, make_seeding_tracks from AllenConf.matching_reconstruction import make_velo_scifi_matches -from AllenConf.muon_reconstruction import decode_muon, is_muon,fake_muon_id, make_velo_muon +from AllenConf.muon_reconstruction import decode_muon, is_muon, fake_muon_id, make_velo_muon from AllenConf.calo_reconstruction import decode_calo, make_track_matching, make_ecal_clusters from AllenConf.primary_vertex_reconstruction import make_pvs from AllenConf.secondary_vertex_reconstruction import make_kalman_velo_only, make_basic_particles, fit_secondary_vertices @@ -28,12 +28,12 @@ def hlt1_reconstruction(matching=False, velo_states = run_velo_kalman_filter(velo_tracks) pvs = make_pvs(velo_tracks) velo_muon = make_velo_muon() - + output = { "velo_tracks": velo_tracks, "velo_states": velo_states, "pvs": pvs, - 'velo_muon':velo_muon + 'velo_muon': velo_muon } if matching: @@ -146,16 +146,14 @@ def validator_node(reconstructed_objects, line_algorithms, matching, with_ut, with_ut=with_ut) ] - if with_muon: validators += make_composite_node_with_gec( - - "muon_validation", - muon_validation(reconstructed_objects["muonID"], - reconstructed_objects['velo_muon']), - with_scifi=True, - with_ut=with_ut), + "muon_validation", + muon_validation(reconstructed_objects["muonID"], + reconstructed_objects['velo_muon']), + with_scifi=True, + with_ut=with_ut), validators += [ make_composite_node_with_gec( "pv_validation", diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py index cfae519179b..903c658e4d1 100644 --- a/configuration/python/AllenConf/muon_reconstruction.py +++ b/configuration/python/AllenConf/muon_reconstruction.py @@ -4,8 +4,9 @@ from AllenAlgorithms.algorithms import ( data_provider_t, muon_calculate_srq_size_t, host_prefix_sum_t, muon_populate_tile_and_tdc_t, muon_add_coords_crossing_maps_t, - muon_populate_hits_t, is_muon_t, empty_lepton_id_t, find_muon_hits_t, fit_muon_t, - ut_select_velo_tracks_t, match_velo_muon_t, consolidate_muon_t, consolidate_velo_muon_t) + muon_populate_hits_t, is_muon_t, empty_lepton_id_t, find_muon_hits_t, + fit_muon_t, ut_select_velo_tracks_t, match_velo_muon_t, consolidate_muon_t, + consolidate_velo_muon_t) from AllenConf.utils import initialize_number_of_events from AllenCore.generator import make_algorithm -- GitLab From 4b6d776c0e2a6d995d753220312a39a89e5f5ab8 Mon Sep 17 00:00:00 2001 From: Adrian Casais Vidal Date: Tue, 27 Sep 2022 23:42:39 +0200 Subject: [PATCH 3/6] Removed comments --- .../python/AllenConf/muon_reconstruction.py | 9 --------- .../event_model/muon/include/MuonDefinitions.cuh | 2 -- .../match_velo_muon/include/ConsolidateMuon.cuh | 11 ----------- .../muon/match_velo_muon/include/FindMuonHits.cuh | 7 ------- device/muon/match_velo_muon/include/FitMuon.cuh | 3 --- .../muon/match_velo_muon/src/ConsolidateMuon.cu | 4 ---- device/muon/match_velo_muon/src/FitMuon.cu | 15 --------------- .../lines/muon/include/OneMuonTrackLine.cuh | 1 - host/validators/include/HostVeloValidator.h | 1 - host/validators/src/HostVeloValidator.cpp | 1 - 10 files changed, 54 deletions(-) diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py index 903c658e4d1..b03d94f8414 100644 --- a/configuration/python/AllenConf/muon_reconstruction.py +++ b/configuration/python/AllenConf/muon_reconstruction.py @@ -231,10 +231,6 @@ def make_velo_muon(): host_number_of_events_t=number_of_events["host_number_of_events"], dev_number_of_events_t=number_of_events["dev_number_of_events"], dev_muon_tracks_input_t=find_muon_hits.dev_muon_tracks_t, - # dev_muon_tracks_tx_input_t=find_muon_hits.dev_muon_tracks_tx_t, - # dev_muon_tracks_ty_input_t=find_muon_hits.dev_muon_tracks_ty_t, - # dev_muon_tracks_state_muon_index_input_t=find_muon_hits. - # dev_muon_tracks_states_t, dev_muon_number_of_tracks_t=find_muon_hits.dev_muon_number_of_tracks_t, dev_muon_tracks_offsets_t=prefix_sum_muon_tracks.dev_output_buffer_t, host_muon_total_number_of_tracks_t=prefix_sum_muon_tracks. @@ -263,11 +259,6 @@ def make_velo_muon(): dev_velo_tracks_view_t=velo_tracks["dev_velo_tracks_view"], dev_velo_states_view_t=velo_states[ "dev_velo_kalman_endvelo_states_view"], - # dev_muon_tracks_tx_t=consolidate_muon.dev_muon_tracks_tx_output_t, - # dev_muon_tracks_ty_t=consolidate_muon.dev_muon_tracks_ty_output_t, - # dev_muon_tracks_state_muon_index_t=consolidate_muon. - # dev_muon_tracks_state_muon_index_output_t) - ) prefix_sum_velomuon_tracks = make_algorithm( host_prefix_sum_t, diff --git a/device/event_model/muon/include/MuonDefinitions.cuh b/device/event_model/muon/include/MuonDefinitions.cuh index 84a81082d11..69103f3f929 100644 --- a/device/event_model/muon/include/MuonDefinitions.cuh +++ b/device/event_model/muon/include/MuonDefinitions.cuh @@ -273,8 +273,6 @@ struct MuonTrack { __host__ __device__ uint8_t number_of_hits() const { return m_number_of_hits; } - // __host__ __device__ void set_velo_index(int velo_index){ m_velo_index = velo_index;} - __host__ __device__ void set_tx(float tx) { m_tx = tx; } __host__ __device__ void set_ty(float ty) { m_ty = ty; } __host__ __device__ void set_ax(float ax) { m_ax = ax; } diff --git a/device/muon/match_velo_muon/include/ConsolidateMuon.cuh b/device/muon/match_velo_muon/include/ConsolidateMuon.cuh index 7af0e366dfa..8be36b4f6fa 100644 --- a/device/muon/match_velo_muon/include/ConsolidateMuon.cuh +++ b/device/muon/match_velo_muon/include/ConsolidateMuon.cuh @@ -16,22 +16,11 @@ namespace consolidate_muon { HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; MASK_INPUT(dev_event_list_t) dev_event_list; DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; - // DEVICE_INPUT(dev_station_ocurrences_offset_t, unsigned) dev_station_ocurrences_offset; - // DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; DEVICE_INPUT(dev_muon_tracks_input_t, MuonTrack) dev_muon_tracks_input; - // DEVICE_INPUT(dev_muon_tracks_tx_input_t, float) dev_muon_tracks_tx_input; - // DEVICE_INPUT(dev_muon_tracks_ty_input_t, float) dev_muon_tracks_ty_input; - // DEVICE_INPUT(dev_muon_tracks_state_muon_index_input_t, int) dev_muon_tracks_state_muon_index_input; DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; - DEVICE_INPUT(dev_muon_tracks_offsets_t, unsigned) dev_muon_tracks_offsets; HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; - DEVICE_OUTPUT(dev_muon_tracks_output_t, MuonTrack) dev_muon_tracks_output; - // DEVICE_OUTPUT(dev_muon_tracks_tx_output_t, float) dev_muon_tracks_tx_output; - // DEVICE_OUTPUT(dev_muon_tracks_ty_output_t, float) dev_muon_tracks_ty_output; - // DEVICE_OUTPUT(dev_muon_tracks_state_muon_index_output_t, int) dev_muon_tracks_state_muon_index_output; - PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; }; diff --git a/device/muon/match_velo_muon/include/FindMuonHits.cuh b/device/muon/match_velo_muon/include/FindMuonHits.cuh index 854aea20efd..3a186226745 100644 --- a/device/muon/match_velo_muon/include/FindMuonHits.cuh +++ b/device/muon/match_velo_muon/include/FindMuonHits.cuh @@ -29,17 +29,10 @@ namespace find_muon_hits { HOST_OUTPUT(host_station_ocurrences_offset_t, unsigned) host_station_ocurrences_offset; DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; HOST_OUTPUT(host_muon_hits_t, char) host_muon_hits; - DEVICE_OUTPUT(dev_muon_tracks_t, MuonTrack) dev_muon_tracks; HOST_OUTPUT(host_muon_tracks_t, MuonTrack) host_velomuon_tracks; - DEVICE_OUTPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; HOST_OUTPUT(host_muon_number_of_tracks_t, unsigned) host_muon_number_of_tracks; - // DEVICE_OUTPUT(dev_muon_tracks_tx_t, float) dev_muon_tracks_tx; - // DEVICE_OUTPUT(dev_muon_tracks_ty_t, float) dev_muon_tracks_ty; - // DEVICE_OUTPUT(dev_muon_tracks_chi2x_t, float) dev_muon_tracks_chi2x; - // DEVICE_OUTPUT(dev_muon_tracks_chi2y_t, float) dev_muon_tracks_chi2y; - // DEVICE_OUTPUT(dev_muon_tracks_states_t, int) dev_muon_tracks_states; PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; }; diff --git a/device/muon/match_velo_muon/include/FitMuon.cuh b/device/muon/match_velo_muon/include/FitMuon.cuh index 76d393e0bc3..3bb41225dcf 100644 --- a/device/muon/match_velo_muon/include/FitMuon.cuh +++ b/device/muon/match_velo_muon/include/FitMuon.cuh @@ -20,17 +20,14 @@ namespace fit_muon { DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; DEVICE_INPUT(dev_station_ocurrences_offset_t, unsigned) dev_station_ocurrences_offset; DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; - DEVICE_INPUT(dev_muon_tracks_offsets_t, unsigned) dev_muon_tracks_offsets; HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; - DEVICE_OUTPUT(dev_muon_tracks_tx_t, float) dev_muon_tracks_tx; DEVICE_OUTPUT(dev_muon_tracks_ty_t, float) dev_muon_tracks_ty; DEVICE_OUTPUT(dev_muon_tracks_state_muon_index_t, unsigned) dev_muon_tracks_state_muon_index; DEVICE_OUTPUT(dev_muon_tracks_chi2x_t, float) dev_muon_tracks_chi2x; DEVICE_OUTPUT(dev_muon_tracks_chi2y_t, float) dev_muon_tracks_chi2y; DEVICE_OUTPUT(dev_muon_number_of_fitted_tracks_t, unsigned) dev_muon_number_of_fitted_tracks; - // DEVICE_OUTPUT(dev_output_muon_tracks); PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; }; diff --git a/device/muon/match_velo_muon/src/ConsolidateMuon.cu b/device/muon/match_velo_muon/src/ConsolidateMuon.cu index c8c099eafac..a22e534c3f0 100644 --- a/device/muon/match_velo_muon/src/ConsolidateMuon.cu +++ b/device/muon/match_velo_muon/src/ConsolidateMuon.cu @@ -17,10 +17,6 @@ void consolidate_muon::consolidate_muon_t::set_arguments_size( const HostBuffers&) const { set_size(arguments, first(arguments)); - // set_size(arguments, first(arguments)); - // set_size(arguments, first(arguments)); - // set_size(arguments, - // first(arguments)) } void consolidate_muon::consolidate_muon_t::operator()( diff --git a/device/muon/match_velo_muon/src/FitMuon.cu b/device/muon/match_velo_muon/src/FitMuon.cu index 28595ca7183..cd4c83c742a 100644 --- a/device/muon/match_velo_muon/src/FitMuon.cu +++ b/device/muon/match_velo_muon/src/FitMuon.cu @@ -34,9 +34,6 @@ void fit_muon::fit_muon_t::operator()( HostBuffers&, const Allen::Context& context) const { - // initialize(arguments, 0, context); - // initialize(arguments, 0, context); - // print(arguments); Allen::memset_async(arguments, 0, context); global_function(fit_muon)(dim3(size(arguments)), property(), context)(arguments); } @@ -48,8 +45,6 @@ __global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) const auto muon_total_number_of_hits = parameters.dev_station_ocurrences_offset[number_of_events * Muon::Constants::n_stations]; - // const auto station_ocurrences_offset = - // parameters.dev_station_ocurrences_offset + event_number * Muon::Constants::n_stations; const auto muon_hits = Muon::ConstHits {parameters.dev_muon_hits, muon_total_number_of_hits}; // TODO: Change into container taking into consideration actual sizes of M4 and M5 (use prefix sum) @@ -71,7 +66,6 @@ __global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) const auto n_hits_track = muon_track.number_of_hits(); const auto ndof = n_hits_track - 2; if (normal_fit) { - // printf("Entering normal fit\n"); float x_mean = 0; float y_mean = 0; float z_mean = 0; @@ -104,8 +98,6 @@ __global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) float ay = y_mean - ty * z_mean; event_muon_tx[idx] = tx; event_muon_ty[idx] = ty; - // muon_track.set_tx(tx); - // muon_track.set_ty(ty); for (unsigned i_hit = 0; i_hit < 4; i_hit++) { if (muon_track.hit(i_hit) == -1) continue; chi2x += (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) * @@ -126,11 +118,9 @@ __global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) event_muon_tx[insert_idx] = tx; event_muon_ty[insert_idx] = ty; event_muon_states[insert_idx] = muon_track.hit(state_idx); - // __syncthreads(); } } else { - // printf("Entering alternate fit\n"); float xz = 0; float yz = 0; float x = 0; @@ -141,10 +131,8 @@ __global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) float z2 = 0; float chi2x = 0; float chi2y = 0; - // printf("n_hits_track=%i\n",n_hits_track); for (unsigned i_hit = 0; i_hit < 4; i_hit++) { if (muon_track.hit(i_hit) == -1) continue; - // printf("summing variables\n"); xz += muon_hits.x(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; yz += muon_hits.y(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; x += muon_hits.x(muon_track.hit(i_hit)) / n_hits_track; @@ -168,11 +156,8 @@ __global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) chi2y += (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) * (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))); } - // printf("offset=%i\n",event_muon_tracks_offset); event_muon_chi2x[idx] = chi2x; event_muon_chi2y[idx] = chi2y; } - // printf("tx = %f, ty = %f, chi2x/nodf = %f, chi2y/ndof = - // %f\n",event_muon_tx[idx],event_muon_ty[idx],event_muon_chi2x[idx]/2.f,event_muon_chi2y[idx]/2.f); } } diff --git a/device/selections/lines/muon/include/OneMuonTrackLine.cuh b/device/selections/lines/muon/include/OneMuonTrackLine.cuh index faf3fc26111..135d8538a47 100644 --- a/device/selections/lines/muon/include/OneMuonTrackLine.cuh +++ b/device/selections/lines/muon/include/OneMuonTrackLine.cuh @@ -23,7 +23,6 @@ namespace one_muon_track_line { DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; - // PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler; PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler; PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler; PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string); diff --git a/host/validators/include/HostVeloValidator.h b/host/validators/include/HostVeloValidator.h index 347080db0cd..45347566b58 100644 --- a/host/validators/include/HostVeloValidator.h +++ b/host/validators/include/HostVeloValidator.h @@ -12,7 +12,6 @@ namespace host_velo_validator { DEVICE_INPUT(dev_offsets_all_velo_tracks_t, unsigned) dev_offsets_all_velo_tracks; DEVICE_INPUT(dev_offsets_velo_track_hit_number_t, unsigned) dev_offsets_velo_track_hit_number; DEVICE_INPUT(dev_velo_track_hits_t, char) dev_velo_track_hits; - // DEVICE_INPUT(dev_velomuon_muon_id_t, bool) dev_velomuon_muon_id; MASK_INPUT(dev_event_list_t) dev_event_list; HOST_INPUT(host_mc_events_t, const MCEvents*) host_mc_events; PROPERTY(root_output_filename_t, "root_output_filename", "root output filename", std::string); diff --git a/host/validators/src/HostVeloValidator.cpp b/host/validators/src/HostVeloValidator.cpp index 4c4084c7a81..11ffccb59b3 100644 --- a/host/validators/src/HostVeloValidator.cpp +++ b/host/validators/src/HostVeloValidator.cpp @@ -23,7 +23,6 @@ void host_velo_validator::host_velo_validator_t::operator()( offsets_all_velo_tracks, offsets_velo_track_hit_number, velo_track_hits, - // velomuon_id, event_list); auto& checker = -- GitLab From f52fb169f72ec8fbe0952e9916818a42bdaca65f Mon Sep 17 00:00:00 2001 From: Adrian Casais Vidal Date: Wed, 28 Sep 2022 12:30:00 +0200 Subject: [PATCH 4/6] Resolving threads --- CMakeLists.txt | 5 +---- configuration/python/AllenConf/muon_reconstruction.py | 1 + device/muon/match_velo_muon/include/MatchVeloMuon.cuh | 4 +++- device/muon/match_velo_muon/src/MatchVeloMuon.cu | 9 ++++++++- 4 files changed, 13 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index facbc2401b9..64b2eae7257 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -748,7 +748,4 @@ if (STANDALONE) add_compile_definitions(PARAMFILESROOTPATH=${PARAMFILESROOT}) message(STATUS "PARAMFILESROOT set to ${PARAMFILESROOT}") endif() -endif() - -file(MAKE_DIRECTORY ${CMAKE_BINARY_DIR}/plotsfornote) -file(MAKE_DIRECTORY ${CMAKE_BINARY_DIR}/plotsfornote_root) \ No newline at end of file +endif() \ No newline at end of file diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py index b03d94f8414..e5d7dc84f70 100644 --- a/configuration/python/AllenConf/muon_reconstruction.py +++ b/configuration/python/AllenConf/muon_reconstruction.py @@ -259,6 +259,7 @@ def make_velo_muon(): dev_velo_tracks_view_t=velo_tracks["dev_velo_tracks_view"], dev_velo_states_view_t=velo_states[ "dev_velo_kalman_endvelo_states_view"], + create_plot_dirs=True) prefix_sum_velomuon_tracks = make_algorithm( host_prefix_sum_t, diff --git a/device/muon/match_velo_muon/include/MatchVeloMuon.cuh b/device/muon/match_velo_muon/include/MatchVeloMuon.cuh index add81726d80..7794d818243 100644 --- a/device/muon/match_velo_muon/include/MatchVeloMuon.cuh +++ b/device/muon/match_velo_muon/include/MatchVeloMuon.cuh @@ -41,7 +41,8 @@ namespace match_velo_muon { DEVICE_OUTPUT(dev_velomuon_tracks_qop_t, float) dev_velomuon_tracks_qop; DEVICE_OUTPUT(dev_velomuon_number_of_tracks_t, unsigned) dev_velomuon_number_of_tracks; DEVICE_OUTPUT(dev_velomuon_muon_id_t, bool) dev_velomuon_muon_id; - + + PROPERTY(create_plot_dirs_t, "create_plot_dirs", "create plot dirs", bool) create_plot_dirs; PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; }; @@ -64,6 +65,7 @@ namespace match_velo_muon { private: Property m_block_dim {this, {{64, 1, 1}}}; + Property m_create_plot_dir {this, false}; }; float xcale[4] {0.06f, 0.1f, 0.15f, 0.15f}; float distance_cut[5] {30 * 30, 60 * 60, 110 * 110, 200 * 200}; diff --git a/device/muon/match_velo_muon/src/MatchVeloMuon.cu b/device/muon/match_velo_muon/src/MatchVeloMuon.cu index 5914254e7cb..005fd203538 100644 --- a/device/muon/match_velo_muon/src/MatchVeloMuon.cu +++ b/device/muon/match_velo_muon/src/MatchVeloMuon.cu @@ -8,6 +8,9 @@ #include "VeloEventModel.cuh" #include +#include +namespace fs = std::filesystem; + INSTANTIATE_ALGORITHM(match_velo_muon::match_velo_muon_t) void match_velo_muon::match_velo_muon_t::set_arguments_size( @@ -34,7 +37,11 @@ void match_velo_muon::match_velo_muon_t::operator()( const Constants& constants, HostBuffers&, const Allen::Context& context) const -{ +{ + if (property()){ + fs::create_directory("plotsfornote"); + fs::create_directory("plotsfornote_root"); +} Allen::memset_async(arguments, 0, context); Allen::memset_async(arguments, 0, context); global_function(match_velo_muon)(dim3(size(arguments)), property(), context)( -- GitLab From 1b42d646b18942dbc4ffd13babaf0b67c926252d Mon Sep 17 00:00:00 2001 From: Gitlab CI Date: Wed, 28 Sep 2022 10:31:04 +0000 Subject: [PATCH 5/6] Fixed formatting patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/24834060 --- device/muon/match_velo_muon/include/MatchVeloMuon.cuh | 4 ++-- device/muon/match_velo_muon/src/MatchVeloMuon.cu | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/device/muon/match_velo_muon/include/MatchVeloMuon.cuh b/device/muon/match_velo_muon/include/MatchVeloMuon.cuh index 7794d818243..4c8fe7d9213 100644 --- a/device/muon/match_velo_muon/include/MatchVeloMuon.cuh +++ b/device/muon/match_velo_muon/include/MatchVeloMuon.cuh @@ -41,7 +41,7 @@ namespace match_velo_muon { DEVICE_OUTPUT(dev_velomuon_tracks_qop_t, float) dev_velomuon_tracks_qop; DEVICE_OUTPUT(dev_velomuon_number_of_tracks_t, unsigned) dev_velomuon_number_of_tracks; DEVICE_OUTPUT(dev_velomuon_muon_id_t, bool) dev_velomuon_muon_id; - + PROPERTY(create_plot_dirs_t, "create_plot_dirs", "create plot dirs", bool) create_plot_dirs; PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; }; @@ -65,7 +65,7 @@ namespace match_velo_muon { private: Property m_block_dim {this, {{64, 1, 1}}}; - Property m_create_plot_dir {this, false}; + Property m_create_plot_dir {this, false}; }; float xcale[4] {0.06f, 0.1f, 0.15f, 0.15f}; float distance_cut[5] {30 * 30, 60 * 60, 110 * 110, 200 * 200}; diff --git a/device/muon/match_velo_muon/src/MatchVeloMuon.cu b/device/muon/match_velo_muon/src/MatchVeloMuon.cu index 005fd203538..bc0f246a4b1 100644 --- a/device/muon/match_velo_muon/src/MatchVeloMuon.cu +++ b/device/muon/match_velo_muon/src/MatchVeloMuon.cu @@ -37,11 +37,11 @@ void match_velo_muon::match_velo_muon_t::operator()( const Constants& constants, HostBuffers&, const Allen::Context& context) const -{ - if (property()){ +{ + if (property()) { fs::create_directory("plotsfornote"); fs::create_directory("plotsfornote_root"); -} + } Allen::memset_async(arguments, 0, context); Allen::memset_async(arguments, 0, context); global_function(match_velo_muon)(dim3(size(arguments)), property(), context)( -- GitLab From c1488a0dc95f246923e035755e0fbcb0f9ee9b4f Mon Sep 17 00:00:00 2001 From: Adrian Casais Vidal Date: Wed, 28 Sep 2022 15:53:13 +0200 Subject: [PATCH 6/6] remove fitmuon --- .../python/AllenConf/muon_reconstruction.py | 2 +- .../muon/match_velo_muon/include/FitMuon.cuh | 54 ------ device/muon/match_velo_muon/src/FitMuon.cu | 163 ------------------ 3 files changed, 1 insertion(+), 218 deletions(-) delete mode 100644 device/muon/match_velo_muon/include/FitMuon.cuh delete mode 100644 device/muon/match_velo_muon/src/FitMuon.cu diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py index e5d7dc84f70..79a26caf504 100644 --- a/configuration/python/AllenConf/muon_reconstruction.py +++ b/configuration/python/AllenConf/muon_reconstruction.py @@ -5,7 +5,7 @@ from AllenAlgorithms.algorithms import ( data_provider_t, muon_calculate_srq_size_t, host_prefix_sum_t, muon_populate_tile_and_tdc_t, muon_add_coords_crossing_maps_t, muon_populate_hits_t, is_muon_t, empty_lepton_id_t, find_muon_hits_t, - fit_muon_t, ut_select_velo_tracks_t, match_velo_muon_t, consolidate_muon_t, + ut_select_velo_tracks_t, match_velo_muon_t, consolidate_muon_t, consolidate_velo_muon_t) from AllenConf.utils import initialize_number_of_events from AllenCore.generator import make_algorithm diff --git a/device/muon/match_velo_muon/include/FitMuon.cuh b/device/muon/match_velo_muon/include/FitMuon.cuh deleted file mode 100644 index 3bb41225dcf..00000000000 --- a/device/muon/match_velo_muon/include/FitMuon.cuh +++ /dev/null @@ -1,54 +0,0 @@ -/*****************************************************************************\ -* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * -\*****************************************************************************/ -#pragma once - -#include "AlgorithmTypes.cuh" -#include "MuonDefinitions.cuh" -#include "UTConsolidated.cuh" -#include "VeloConsolidated.cuh" -#include "States.cuh" - -static constexpr int M2 {0}, M3 {1}, M4 {2}, M5 {3}; - -namespace fit_muon { - struct Parameters { - HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events; - DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events; - MASK_INPUT(dev_event_list_t) dev_event_list; - DEVICE_INPUT(dev_muon_tracks_t, MuonTrack) dev_muon_tracks; - DEVICE_INPUT(dev_muon_number_of_tracks_t, unsigned) dev_muon_number_of_tracks; - DEVICE_INPUT(dev_station_ocurrences_offset_t, unsigned) dev_station_ocurrences_offset; - DEVICE_INPUT(dev_muon_hits_t, char) dev_muon_hits; - DEVICE_INPUT(dev_muon_tracks_offsets_t, unsigned) dev_muon_tracks_offsets; - HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks; - DEVICE_OUTPUT(dev_muon_tracks_tx_t, float) dev_muon_tracks_tx; - DEVICE_OUTPUT(dev_muon_tracks_ty_t, float) dev_muon_tracks_ty; - DEVICE_OUTPUT(dev_muon_tracks_state_muon_index_t, unsigned) dev_muon_tracks_state_muon_index; - DEVICE_OUTPUT(dev_muon_tracks_chi2x_t, float) dev_muon_tracks_chi2x; - DEVICE_OUTPUT(dev_muon_tracks_chi2y_t, float) dev_muon_tracks_chi2y; - DEVICE_OUTPUT(dev_muon_number_of_fitted_tracks_t, unsigned) dev_muon_number_of_fitted_tracks; - PROPERTY(block_dim_t, "block_dim", "block dimensions", DeviceDimensions) block_dim; - }; - - __global__ void fit_muon(Parameters); - - struct fit_muon_t : public DeviceAlgorithm, Parameters { - void set_arguments_size( - ArgumentReferences arguments, - const RuntimeOptions&, - const Constants&, - const HostBuffers&) const; - - void operator()( - const ArgumentReferences& arguments, - const RuntimeOptions& runtime_options, - const Constants& constants, - HostBuffers& host_buffers, - const Allen::Context& context) const; - - private: - Property m_block_dim {this, {{64, 1, 1}}}; - }; - -} // namespace fit_muon \ No newline at end of file diff --git a/device/muon/match_velo_muon/src/FitMuon.cu b/device/muon/match_velo_muon/src/FitMuon.cu deleted file mode 100644 index cd4c83c742a..00000000000 --- a/device/muon/match_velo_muon/src/FitMuon.cu +++ /dev/null @@ -1,163 +0,0 @@ -/*****************************************************************************\ -* (c) Copyright 2020 CERN for the benefit of the LHCb Collaboration * -\*****************************************************************************/ -#include "FitMuon.cuh" - -#include "Common.h" -#include "VeloDefinitions.cuh" -#include "VeloEventModel.cuh" -#include - -INSTANTIATE_ALGORITHM(fit_muon::fit_muon_t) - -void fit_muon::fit_muon_t::set_arguments_size( - ArgumentReferences arguments, - const RuntimeOptions&, - const Constants&, - const HostBuffers&) const -{ - set_size(arguments, first(arguments)); - set_size(arguments, first(arguments)); - - set_size(arguments, first(arguments)); - - set_size(arguments, first(arguments)); - set_size(arguments, first(arguments)); - - set_size(arguments, first(arguments)); -} - -void fit_muon::fit_muon_t::operator()( - const ArgumentReferences& arguments, - const RuntimeOptions&, - const Constants&, - HostBuffers&, - const Allen::Context& context) const -{ - Allen::memset_async(arguments, 0, context); - global_function(fit_muon)(dim3(size(arguments)), property(), context)(arguments); -} - -__global__ void fit_muon::fit_muon(fit_muon::Parameters parameters) -{ - const unsigned event_number = parameters.dev_event_list[blockIdx.x]; - const unsigned number_of_events = parameters.dev_number_of_events[0]; - - const auto muon_total_number_of_hits = - parameters.dev_station_ocurrences_offset[number_of_events * Muon::Constants::n_stations]; - const auto muon_hits = Muon::ConstHits {parameters.dev_muon_hits, muon_total_number_of_hits}; - - // TODO: Change into container taking into consideration actual sizes of M4 and M5 (use prefix sum) - const auto event_muon_tracks = parameters.dev_muon_tracks + event_number * Muon::Constants::max_number_of_tracks; - const auto event_number_of_tracks = parameters.dev_muon_number_of_tracks[event_number]; - const auto event_muon_tracks_offset = parameters.dev_muon_tracks_offsets[event_number]; - - auto event_muon_tx = parameters.dev_muon_tracks_tx + event_muon_tracks_offset; - auto event_muon_ty = parameters.dev_muon_tracks_ty + event_muon_tracks_offset; - auto event_muon_states = parameters.dev_muon_tracks_state_muon_index + event_muon_tracks_offset; - - const auto event_muon_chi2x = parameters.dev_muon_tracks_chi2x + event_muon_tracks_offset; - const auto event_muon_chi2y = parameters.dev_muon_tracks_chi2y + event_muon_tracks_offset; - - auto event_number_of_fitted_tracks = parameters.dev_muon_number_of_fitted_tracks + event_number; - bool normal_fit = true; - for (unsigned idx = threadIdx.x; idx < event_number_of_tracks; idx += blockDim.x) { - const auto muon_track = event_muon_tracks[idx]; - const auto n_hits_track = muon_track.number_of_hits(); - const auto ndof = n_hits_track - 2; - if (normal_fit) { - float x_mean = 0; - float y_mean = 0; - float z_mean = 0; - float chi2x = 0; - float chi2y = 0; - - for (unsigned i_hit = 0; i_hit < 4; i_hit++) { - if (muon_track.hit(i_hit) == -1) continue; - x_mean += muon_hits.x(muon_track.hit(i_hit)); - y_mean += muon_hits.y(muon_track.hit(i_hit)); - z_mean += muon_hits.z(muon_track.hit(i_hit)); - } - - x_mean = x_mean / n_hits_track; - y_mean = y_mean / n_hits_track; - z_mean = z_mean / n_hits_track; - - float szx = 0; - float szy = 0; - float sz2 = 0; - for (unsigned i_hit = 0; i_hit < 4; i_hit++) { - if (muon_track.hit(i_hit) == -1) continue; - szx += (muon_hits.z(muon_track.hit(i_hit)) - z_mean) * (muon_hits.x(muon_track.hit(i_hit)) - x_mean); - szy += (muon_hits.z(muon_track.hit(i_hit)) - z_mean) * (muon_hits.y(muon_track.hit(i_hit)) - y_mean); - sz2 += (muon_hits.z(muon_track.hit(i_hit)) - z_mean) * (muon_hits.z(muon_track.hit(i_hit)) - z_mean); - } - float tx = szx / sz2; - float ax = x_mean - tx * z_mean; - float ty = szy / sz2; - float ay = y_mean - ty * z_mean; - event_muon_tx[idx] = tx; - event_muon_ty[idx] = ty; - for (unsigned i_hit = 0; i_hit < 4; i_hit++) { - if (muon_track.hit(i_hit) == -1) continue; - chi2x += (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) * - (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) / - muon_hits.x(muon_track.hit(i_hit)); - chi2y += (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) * - (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) / - muon_hits.y(muon_track.hit(i_hit)); - } - - if (chi2x / ndof < 100.f && chi2y / ndof < 100.f) { - // if (true) { - int state_idx = muon_track.hit(M5) != -1.f ? M5 : M4; - state_idx = M2; - const auto insert_idx = atomicAdd(event_number_of_fitted_tracks, 1); - event_muon_chi2x[insert_idx] = chi2x; - event_muon_chi2y[insert_idx] = chi2y; - event_muon_tx[insert_idx] = tx; - event_muon_ty[insert_idx] = ty; - event_muon_states[insert_idx] = muon_track.hit(state_idx); - } - } - else { - float xz = 0; - float yz = 0; - float x = 0; - float y = 0; - float z = 0; - float x2 = 0; - float y2 = 0; - float z2 = 0; - float chi2x = 0; - float chi2y = 0; - for (unsigned i_hit = 0; i_hit < 4; i_hit++) { - if (muon_track.hit(i_hit) == -1) continue; - xz += muon_hits.x(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; - yz += muon_hits.y(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; - x += muon_hits.x(muon_track.hit(i_hit)) / n_hits_track; - y += muon_hits.y(muon_track.hit(i_hit)) / n_hits_track; - z += muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; - x2 += muon_hits.x(muon_track.hit(i_hit)) * muon_hits.x(muon_track.hit(i_hit)) / n_hits_track; - y2 += muon_hits.y(muon_track.hit(i_hit)) * muon_hits.y(muon_track.hit(i_hit)) / n_hits_track; - z2 += muon_hits.z(muon_track.hit(i_hit)) * muon_hits.z(muon_track.hit(i_hit)) / n_hits_track; - } - float tx = (xz - x * z) / (z2 - z * z); - float ax = (x - tx * z); - float ty = (yz - y * z) / (z2 - z * z); - float ay = (y - ty * z); - event_muon_tx[idx] = tx; - event_muon_ty[idx] = ty; - - for (unsigned i_hit = 0; i_hit < 4; i_hit++) { - if (muon_track.hit(i_hit) == -1) continue; - chi2x += (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))) * - (tx * muon_hits.z(muon_track.hit(i_hit)) + ax - muon_hits.x(muon_track.hit(i_hit))); - chi2y += (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))) * - (ty * muon_hits.z(muon_track.hit(i_hit)) + ay - muon_hits.y(muon_track.hit(i_hit))); - } - event_muon_chi2x[idx] = chi2x; - event_muon_chi2y[idx] = chi2y; - } - } -} -- GitLab