Skip to content
Snippets Groups Projects

Compare revisions

Changes are shown as if the source revision was being merged into the target revision. Learn more about comparing revisions.

Source

Select target project
No results found

Target

Select target project
  • laolivei/Allen
  • pvfinder/inference-engine
  • rmatev/Allen
  • suali/Allen
  • mstahl/Allen
  • roneil/Allen
  • graemes/Allen
  • cburr/Allen
  • jonrob/Allen
  • bjashal/Allen-HIP
  • dcampora/MiniAllen
  • brij/Allen
  • raaij/cuda_hlt
  • bsm-fleet/cuda_hlt
  • abrearod/cuda_hlt
  • aalvesju/cuda_hlt
  • lhcb/Allen
17 results
Show changes
Commits on Source (6)
Hide whitespace changes
Inline Side-by-side
Showing
with 209 additions and 68 deletions
Rec/Allen/tests/refs/allen_event_loop_v3.ref
View file @ 9a8e05f4
......@@ -14,7 +14,6 @@ Using 3 input slices.
ApplicationMgr INFO Application Manager Started successfully
HLTControlFlowMgr INFO Will measure time between events 0 and 0 (stop might be some events later)
HLTControlFlowMgr INFO Starting loop on events
DeviceUTLookupTables.PrUTMagnetTool INFO No B field detected.
DeviceFTGeometry INFO Conditions DB is compatible with FT bank version 7 and 8.
Starting timer for throughput measurement
Input complete
......@@ -54,17 +53,17 @@ velo_consolidate_tracks INFO Number of counters : 1
| "n_velo_tracks" | 1000 | 0 | 0.0000 |
error_bank_filter INFO 1D histograms in directory "error_bank_filter" : 3
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
| ODIN_banks | "ODIN_banks" | 1000 | 0 | 0.0000 |-nan |-nan |
| Plume_banks | "Plume_banks" | 3000 | 0 | 0.0000 |-nan |-nan |
| ODIN_banks | "ODIN_banks" | 1000 | 0 | 0.0000 | 0 | 0 |
| Plume_banks | "Plume_banks" | 3000 | 0 | 0.0000 | 0 | 0 |
| n_data_banks | "n_data_banks" | 4000 | 8 | 3.4641 | -1.1547 | -0.66667 |
gather_selections INFO 1D histograms in directory "gather_selections" : 2
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
| line_passes | "line passes" | 1000 | 58.5 | 0.0000 |-nan |-nan |
| line_rates | "line rates" | 1000 | 58.5 | 0.0000 |-nan |-nan |
| line_passes | "line passes" | 1000 | 58.5 | 0.0000 | 0 | 0 |
| line_rates | "line rates" | 1000 | 58.5 | 0.0000 | 0 | 0 |
pv_beamline_cleanup INFO 1D histograms in directory "pv_beamline_cleanup" : 2
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
| n_pvs_event | "n_pvs_event" | 1000 | 0 | 0.0000 |-nan |-nan |
| n_smog2_PVs | "n_smog2_PVs" | 1000 | 0 | 0.0000 |-nan |-nan |
| n_pvs_event | "n_pvs_event" | 1000 | 0 | 0.0000 | 0 | 0 |
| n_smog2_PVs | "n_smog2_PVs" | 1000 | 0 | 0.0000 | 0 | 0 |
velo_consolidate_tracks INFO 1D histograms in directory "velo_consolidate_tracks" : 1
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
| n_velo_tracks_event | "n_velo_tracks_event" | 1000 | 0 | 0.0000 |-nan |-nan |
| n_velo_tracks_event | "n_velo_tracks_event" | 1000 | 0 | 0.0000 | 0 | 0 |
Rec/Allen/tests/refs/allen_event_loop_v4.ref
View file @ 9a8e05f4
......@@ -14,7 +14,6 @@ Using 3 input slices.
ApplicationMgr INFO Application Manager Started successfully
HLTControlFlowMgr INFO Will measure time between events 0 and 0 (stop might be some events later)
HLTControlFlowMgr INFO Starting loop on events
DeviceUTLookupTables.PrUTMagnetTool INFO No B field detected.
DeviceFTGeometry INFO Conditions DB is compatible with FT bank version 7 and 8.
Starting timer for throughput measurement
Input complete
......@@ -56,8 +55,8 @@ velo_consolidate_tracks INFO Number of counters : 1
| "n_velo_tracks" | 10000 | 0 | 0.0000 |
error_bank_filter INFO 1D histograms in directory "error_bank_filter" : 3
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
| ODIN_banks | "ODIN_banks" | 10000 | 0 | 0.0000 |-nan |-nan |
| Plume_banks | "Plume_banks" | 40000 | 0 | 0.0000 |-nan |-nan |
| ODIN_banks | "ODIN_banks" | 10000 | 0 | 0.0000 | 0 | 0 |
| Plume_banks | "Plume_banks" | 40000 | 0 | 0.0000 | 0 | 0 |
| n_data_banks | "n_data_banks" | 50000 | 8.4 | 3.2000 | -1.5 | 0.25 |
gather_selections INFO 1D histograms in directory "gather_selections" : 2
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
......@@ -65,8 +64,8 @@ gather_selections INFO 1D histograms in directory "gather_s
| line_rates | "line rates" | 10001 | 58.5 | 0.039996 | 99.99 | 9996 |
pv_beamline_cleanup INFO 1D histograms in directory "pv_beamline_cleanup" : 2
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
| n_pvs_event | "n_pvs_event" | 10000 | 0 | 0.0000 |-nan |-nan |
| n_smog2_PVs | "n_smog2_PVs" | 10000 | 0 | 0.0000 |-nan |-nan |
| n_pvs_event | "n_pvs_event" | 10000 | 0 | 0.0000 | 0 | 0 |
| n_smog2_PVs | "n_smog2_PVs" | 10000 | 0 | 0.0000 | 0 | 0 |
velo_consolidate_tracks INFO 1D histograms in directory "velo_consolidate_tracks" : 1
| ID | Title | # | Mean | RMS | Skewness | Kurtosis |
| n_velo_tracks_event | "n_velo_tracks_event" | 10000 | 0 | 0.0000 |-nan |-nan |
| n_velo_tracks_event | "n_velo_tracks_event" | 10000 | 0 | 0.0000 | 0 | 0 |
configuration/python/AllenConf/calo_reconstruction.py
View file @ 9a8e05f4
......@@ -14,7 +14,7 @@ from AllenCore.algorithms import (
momentum_brem_correction_t, calo_seed_clusters_t, calo_find_clusters_t,
calo_prefilter_clusters_t, calo_filter_clusters_t, calo_find_twoclusters_t,
total_ecal_energy_t, make_neutral_particles_t, calo_overlap_clusters_t,
electronid_nn_t)
electronid_nn_t, fake_digit_matching_t)
from AllenConf.utils import initialize_number_of_events
from AllenCore.generator import make_algorithm
from PyConf.tonic import configurable
......@@ -49,7 +49,6 @@ def decode_calo(empty_banks=False):
sum_ecal_energy = make_algorithm(
total_ecal_energy_t,
name='total_ecal_energy_{hash}',
host_number_of_events_t=number_of_events["host_number_of_events"],
host_ecal_number_of_digits_t=calo_count_digits.host_total_sum_holder_t,
dev_ecal_digits_offsets_t=calo_count_digits.dev_digits_offsets_t,
dev_ecal_digits_t=calo_decode.dev_ecal_digits_t)
......@@ -75,6 +74,8 @@ def make_is_electron(decoded_calo, host_number_of_tracks,
dev_scifi_states_t=dev_scifi_states,
# dev_long_tracks_view_t=dev_multi_event_tracks_ptr,
dev_tracks_view_t=dev_multi_event_tracks_ptr,
host_ecal_number_of_digits_t=decoded_calo[
"host_ecal_number_of_digits"],
dev_ecal_digits_t=decoded_calo["dev_ecal_digits"],
dev_ecal_digits_offsets_t=decoded_calo["dev_ecal_digits_offsets"],
dev_number_of_events_t=number_of_events["dev_number_of_events"])
......@@ -119,6 +120,8 @@ def make_track_matching(decoded_calo, velo_tracks, velo_states, long_tracks,
dev_scifi_states_t=long_tracks["dev_scifi_states"],
# dev_long_tracks_view_t=long_tracks["dev_multi_event_long_tracks_view"],
dev_tracks_view_t=long_tracks["dev_multi_event_long_tracks_ptr"],
host_ecal_number_of_digits_t=decoded_calo[
"host_ecal_number_of_digits"],
dev_ecal_digits_t=decoded_calo["dev_ecal_digits"],
dev_ecal_digits_offsets_t=decoded_calo["dev_ecal_digits_offsets"],
dev_number_of_events_t=number_of_events["dev_number_of_events"])
......@@ -134,6 +137,8 @@ def make_track_matching(decoded_calo, velo_tracks, velo_states, long_tracks,
"dev_offsets_velo_track_hit_number"],
dev_velo_kalman_beamline_states_t=velo_states[
"dev_velo_kalman_beamline_states"],
host_ecal_number_of_digits_t=decoded_calo[
"host_ecal_number_of_digits"],
dev_ecal_digits_t=decoded_calo["dev_ecal_digits"],
dev_ecal_digits_offsets_t=decoded_calo["dev_ecal_digits_offsets"],
dev_number_of_events_t=number_of_events["dev_number_of_events"])
......@@ -175,6 +180,8 @@ def make_track_matching(decoded_calo, velo_tracks, velo_states, long_tracks,
track_digit_selective_matching.dev_track_Eop3x3_t,
"dev_track_isElectron":
track_digit_selective_matching.dev_track_isElectron_t,
"dev_ecal_digits_isTrackMatched":
track_digit_selective_matching.dev_ecal_digits_isTrackMatched_t,
"dev_brem_E":
brem_recovery.dev_brem_E_t,
"dev_brem_ET":
......@@ -185,6 +192,8 @@ def make_track_matching(decoded_calo, velo_tracks, velo_states, long_tracks,
brem_recovery.dev_brem_ecal_digits_size_t,
"dev_brem_ecal_digits":
brem_recovery.dev_brem_ecal_digits_t,
"dev_ecal_digits_isBremMatched":
brem_recovery.dev_ecal_digits_isBremMatched_t,
"dev_brem_corrected_p":
momentum_brem_correction.dev_brem_corrected_p_t,
"dev_brem_corrected_pt":
......@@ -192,6 +201,23 @@ def make_track_matching(decoded_calo, velo_tracks, velo_states, long_tracks,
}
# Produce arrays with fake matching decisions, all set to false,
# to run calo_find_clusters without vetoing any seed clusters.
def fake_digit_matching(decoded_calo):
fake_digit_match = make_algorithm(
fake_digit_matching_t,
name='empty_digit_track_matching_{hash}',
host_ecal_number_of_digits_t=decoded_calo["host_ecal_number_of_digits"]
)
return {
"dev_ecal_digits_isTrackMatched":
fake_digit_match.dev_ecal_digits_isMatched_t,
"dev_ecal_digits_isBremMatched":
fake_digit_match.dev_ecal_digits_isMatched_t
}
@configurable
def make_ecal_clusters(decoded_calo,
calo_matching_objects=None,
......@@ -214,7 +240,7 @@ def make_ecal_clusters(decoded_calo,
calo_overlap_clusters = make_algorithm(
calo_overlap_clusters_t,
name="calo_overlap_clusters",
name="calo_overlap_clusters_{hash}",
host_ecal_number_of_clusters_t=calo_seed_clusters.
host_total_sum_holder_t,
dev_ecal_digits_t=decoded_calo["dev_ecal_digits"],
......@@ -224,10 +250,18 @@ def make_ecal_clusters(decoded_calo,
dev_ecal_cluster_offsets_t,
dev_ecal_digit_is_seed_t=calo_seed_clusters.dev_ecal_digit_is_seed_t)
# If calo_matching_objects is None, produce arrays with fake matching decisions,
# all set to false, to run calo_find_clusters without vetoing any seed clusters.
# Otherwise, calo_matching_objects contains the output from make_track_matching(),
# and seed clusters matched with charged tracks or brem photons are vetoed.
if calo_matching_objects is None:
calo_matching_objects = fake_digit_matching(decoded_calo)
calo_find_clusters = make_algorithm(
calo_find_clusters_t,
name=str(calo_find_clusters_name),
name=calo_find_clusters_name,
ecal_min_adc=neighbour_min_adc,
host_number_of_events_t=number_of_events["host_number_of_events"],
host_ecal_number_of_clusters_t=calo_seed_clusters.
host_total_sum_holder_t,
dev_ecal_digits_t=decoded_calo["dev_ecal_digits"],
......@@ -235,16 +269,23 @@ def make_ecal_clusters(decoded_calo,
dev_ecal_seed_clusters_t=calo_seed_clusters.dev_ecal_seed_clusters_t,
dev_ecal_cluster_offsets_t=calo_seed_clusters.
dev_ecal_cluster_offsets_t,
dev_ecal_corrections_t=calo_overlap_clusters.dev_ecal_corrections_t)
dev_ecal_corrections_t=calo_overlap_clusters.dev_ecal_corrections_t,
dev_ecal_digits_isTrackMatched_t=calo_matching_objects[
"dev_ecal_digits_isTrackMatched"],
dev_ecal_digits_isBremMatched_t=calo_matching_objects[
"dev_ecal_digits_isBremMatched"])
make_neutral_particles = make_algorithm(
make_neutral_particles_t,
name="make_neutral_particles",
name="make_neutral_particles_{hash}",
host_number_of_events_t=number_of_events["host_number_of_events"],
host_number_of_clusters_t=calo_seed_clusters.host_total_sum_holder_t,
host_number_of_neutral_clusters_t=calo_find_clusters.
host_total_sum_holder_t,
dev_number_of_events_t=number_of_events["dev_number_of_events"],
dev_ecal_cluster_offsets_t=calo_seed_clusters.
dev_ecal_cluster_offsets_t,
dev_ecal_neutral_cluster_offsets_t=calo_find_clusters.
dev_ecal_neutral_cluster_offsets_t,
dev_ecal_clusters_t=calo_find_clusters.dev_ecal_clusters_t)
calo_prefilter_clusters = make_algorithm(
......@@ -254,7 +295,7 @@ def make_ecal_clusters(decoded_calo,
minE19_clusters=min_e19,
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_ecal_number_of_clusters_t=calo_seed_clusters.
host_ecal_number_of_neutral_clusters_t=calo_find_clusters.
host_total_sum_holder_t,
dev_neutral_particles_t=make_neutral_particles.
dev_multi_event_neutral_particles_view_t)
......@@ -263,8 +304,6 @@ def make_ecal_clusters(decoded_calo,
calo_filter_clusters_t,
name='calo_filter_clusters_{hash}',
host_number_of_events_t=number_of_events["host_number_of_events"],
host_ecal_number_of_clusters_t=calo_seed_clusters.
host_total_sum_holder_t,
host_ecal_number_of_twoclusters_t=calo_prefilter_clusters.
host_total_sum_holder_t,
dev_neutral_particles_t=make_neutral_particles.
......@@ -295,10 +334,14 @@ def make_ecal_clusters(decoded_calo,
return {
"host_ecal_number_of_clusters":
calo_seed_clusters.host_total_sum_holder_t,
"host_ecal_number_of_neutral_particles":
calo_find_clusters.host_total_sum_holder_t,
"host_ecal_number_of_twoclusters":
calo_prefilter_clusters.host_total_sum_holder_t,
"dev_ecal_cluster_offsets":
calo_seed_clusters.dev_ecal_cluster_offsets_t,
"dev_ecal_neutral_particle_offsets":
calo_find_clusters.dev_ecal_neutral_cluster_offsets_t,
"dev_ecal_clusters":
calo_find_clusters.dev_ecal_clusters_t,
"dev_multi_event_neutral_particles":
......
configuration/python/AllenConf/hlt1_calibration_lines.py
View file @ 9a8e05f4
......@@ -26,7 +26,8 @@ def make_pi02gammagamma_line(calo,
name="Hlt1Pi02GammaGamma",
pre_scaler=0.05,
pre_scaler_hash_string=None,
post_scaler_hash_string=None):
post_scaler_hash_string=None,
enable_tupling=False):
number_of_events = initialize_number_of_events()
return make_algorithm(
......@@ -37,7 +38,6 @@ def make_pi02gammagamma_line(calo,
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",
host_ecal_number_of_clusters_t=calo["host_ecal_number_of_clusters"],
dev_velo_tracks_t=velo_tracks["dev_velo_tracks_view"],
dev_particle_container_t=calo["dev_multi_event_diphotons"],
dev_cluster_particle_container_t=calo[
......@@ -50,7 +50,7 @@ def make_pi02gammagamma_line(calo,
minE19_clusters=0.7,
minPtEta=200, #Pi0Pt>minPtEta*(10-Pi0Eta)
max_n_pvs=1,
enable_tupling=False)
enable_tupling=enable_tupling)
def make_dst_line(dstars,
......
configuration/python/AllenConf/hlt1_heavy_ions_lines.py
View file @ 9a8e05f4
......@@ -35,7 +35,8 @@ def make_photon_lowmult_line(calo,
host_number_of_events_t=number_of_events["host_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",
host_ecal_number_of_clusters_t=calo["host_ecal_number_of_clusters"],
host_ecal_number_of_clusters_t=calo[
"host_ecal_number_of_neutral_particles"],
dev_particle_container_t=calo["dev_multi_event_neutral_particles"],
minEt=minEt,
minAbsY_cluster=min_absY,
......@@ -69,7 +70,6 @@ def make_diphoton_lowmult_line(calo,
pre_scaler_hash_string=pre_scaler_hash_string or name + "_pre",
post_scaler_hash_string=post_scaler_hash_string or name + "_post",
host_number_of_events_t=number_of_events["host_number_of_events"],
host_ecal_number_of_clusters_t=calo["host_ecal_number_of_clusters"],
dev_number_of_events_t=number_of_events["dev_number_of_events"],
dev_velo_tracks_t=velo_tracks["dev_velo_tracks_view"],
dev_particle_container_t=calo["dev_multi_event_diphotons"],
......
configuration/python/AllenConf/hlt1_photon_lines.py
View file @ 9a8e05f4
......@@ -9,7 +9,7 @@
# or submit itself to any jurisdiction. #
###############################################################################
from AllenCore.algorithms import single_calo_cluster_line_t, two_calo_clusters_line_t
from AllenConf.utils import initialize_number_of_events, mep_layout
from AllenConf.utils import initialize_number_of_events
from AllenCore.generator import make_algorithm
......@@ -29,7 +29,8 @@ def make_single_calo_cluster_line(calo,
host_number_of_events_t=number_of_events["host_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",
host_ecal_number_of_clusters_t=calo["host_ecal_number_of_clusters"],
host_ecal_number_of_clusters_t=calo[
"host_ecal_number_of_neutral_particles"],
dev_particle_container_t=calo["dev_multi_event_neutral_particles"],
minEt=minEt,
maxEt=maxEt,
......@@ -56,7 +57,6 @@ def make_diphotonhighmass_line(calo,
pre_scaler=pre_scaler,
post_scaler=post_scaler,
host_number_of_events_t=number_of_events["host_number_of_events"],
host_ecal_number_of_clusters_t=calo["host_ecal_number_of_clusters"],
dev_number_of_events_t=number_of_events["dev_number_of_events"],
dev_velo_tracks_t=velo_tracks["dev_velo_tracks_view"],
dev_particle_container_t=calo["dev_multi_event_diphotons"],
......
configuration/python/AllenConf/hlt1_reconstruction.py
View file @ 9a8e05f4
......@@ -27,7 +27,6 @@ from AllenConf.validators import (
data_quality_validation_occupancy, data_quality_validation_pv,
data_quality_validation_velo, downstream_validation)
from PyConf.control_flow import NodeLogic, CompositeNode
from PyConf.tonic import configurable
from AllenConf.persistency import make_gather_selections, make_sel_report_writer
from AllenConf.filters import make_gec
from AllenConf.best_track_creator import best_track_creator
......@@ -158,7 +157,6 @@ def hlt1_reconstruction(algorithm_name='',
if with_calo:
decoded_calo = decode_calo()
ecal_clusters = make_ecal_clusters(decoded_calo)
calo_matching_objects = make_track_matching(decoded_calo, velo_tracks,
velo_states, long_tracks,
......@@ -167,8 +165,8 @@ def hlt1_reconstruction(algorithm_name='',
ecal_clusters = make_ecal_clusters(
decoded_calo,
calo_matching_objects,
calo_find_clusters_name='calo_find_clusters')
calo_matching_objects=calo_matching_objects,
calo_find_clusters_name=algorithm_name + 'calo_find_clusters')
long_track_particles = make_basic_particles(
kalman_velo_only,
......
configuration/python/AllenConf/jet_reconstruction.py
View file @ 9a8e05f4
......@@ -11,7 +11,6 @@
from AllenCore.algorithms import build_cone_jets_t, make_neutral_particles_t
from AllenConf.utils import initialize_number_of_events
from AllenCore.generator import make_algorithm
from AllenConf.odin import decode_odin
def make_cone_jets(long_tracks,
......@@ -27,7 +26,7 @@ def make_cone_jets(long_tracks,
host_number_of_tracks_t=long_tracks[
"host_number_of_reconstructed_scifi_tracks"],
host_number_of_neutrals_t=ecal_clusters[
"host_ecal_number_of_clusters"],
"host_ecal_number_of_neutral_particles"],
dev_number_of_events_t=number_of_events["dev_number_of_events"],
dev_long_track_particle_container_t=long_track_particles[
"dev_multi_event_basic_particles"],
......@@ -39,9 +38,11 @@ def make_cone_jets(long_tracks,
make_neutral_particles_t,
name="make_neutral_particles_{hash}",
host_number_of_events_t=number_of_events["host_number_of_events"],
host_number_of_clusters_t=build_cone_jets.host_number_of_jets_t,
host_number_of_neutral_clusters_t=build_cone_jets.
host_number_of_jets_t,
dev_number_of_events_t=number_of_events["dev_number_of_events"],
dev_ecal_cluster_offsets_t=build_cone_jets.dev_jet_offsets_t,
dev_ecal_neutral_cluster_offsets_t=build_cone_jets.dev_jet_offsets_t,
dev_ecal_clusters_t=build_cone_jets.dev_jet_clusters_t)
return {
......
configuration/python/AllenConf/validators.py
View file @ 9a8e05f4
......@@ -22,7 +22,7 @@ from AllenConf.utils import initialize_number_of_events
from AllenCore.generator import make_algorithm
from AllenConf.persistency import make_dec_reporter, make_gather_selections, make_routingbits_writer, rb_map
from AllenCore.algorithms import seeding_copy_trackXZ_hit_number_t
from AllenConf.scifi_reconstruction import decode_scifi, make_seeding_XZ_tracks, make_seeding_tracks
from AllenConf.scifi_reconstruction import decode_scifi, make_seeding_XZ_tracks
from AllenConf.primary_vertex_reconstruction import make_pvs
from AllenConf.muon_reconstruction import decode_muon
from AllenConf.velo_reconstruction import decode_velo, make_velo_tracks
......@@ -460,7 +460,8 @@ def data_quality_validation_occupancy(name="data_quality_validator"):
decoded_muon = decode_muon()
decoded_calo = decode_calo()
ecal_clusters = make_ecal_clusters(
decoded_calo, calo_find_clusters_name='calo_find_clusters')
decoded_calo,
calo_find_clusters_name='calo_find_clusters_dq_validator')
decoded_velo = decode_velo()
velo_tracks = make_velo_tracks(decoded_velo)
......
device/calo/clustering/include/CaloFilterClusters.cuh
View file @ 9a8e05f4
......@@ -22,7 +22,6 @@ namespace calo_filter_clusters {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_ecal_number_of_clusters_t, unsigned) host_ecal_number_of_clusters;
HOST_INPUT(host_ecal_number_of_twoclusters_t, unsigned) host_ecal_number_of_twoclusters;
MASK_INPUT(dev_event_list_t) dev_event_list;
......
device/calo/clustering/include/CaloFindClusters.cuh
View file @ 9a8e05f4
......@@ -20,14 +20,19 @@
namespace calo_find_clusters {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_ecal_number_of_clusters_t, unsigned) host_ecal_number_of_clusters;
MASK_INPUT(dev_event_list_t) dev_event_list;
DEVICE_INPUT(dev_ecal_digits_t, CaloDigit) dev_ecal_digits;
DEVICE_INPUT(dev_ecal_digits_offsets_t, unsigned) dev_ecal_digits_offsets;
DEVICE_INPUT(dev_ecal_seed_clusters_t, CaloSeedCluster) dev_ecal_seed_clusters;
DEVICE_INPUT(dev_ecal_cluster_offsets_t, unsigned) dev_ecal_cluster_offsets;
DEVICE_INPUT(dev_ecal_digits_isTrackMatched_t, bool) dev_ecal_digits_isTrackMatched;
DEVICE_INPUT(dev_ecal_digits_isBremMatched_t, bool) dev_ecal_digits_isBremMatched;
DEVICE_INPUT(dev_ecal_corrections_t, float) dev_ecal_corrections;
DEVICE_OUTPUT(dev_ecal_clusters_t, CaloCluster) dev_ecal_clusters;
DEVICE_OUTPUT(dev_ecal_neutral_cluster_offsets_t, unsigned) dev_ecal_neutral_cluster_offsets;
HOST_OUTPUT(host_total_sum_holder_t, unsigned) host_total_sum_holder;
};
// Global function
......@@ -54,7 +59,6 @@ namespace calo_find_clusters {
private:
Allen::Property<unsigned> m_block_dim_x {this, "block_dim_x", 64, "block dimension X"};
Allen::Property<unsigned> m_block_dim_y {this, "block_dim_y", 16, "block dimension Y"};
Allen::Property<int16_t> m_ecal_min_adc {this, "ecal_min_adc", 10, "cluster neighbors' minimum ADC"};
Allen::Monitoring::Histogram<> m_histogram_n_clusters {this, "n_ecal_clusters", "NClusters", {401u, -0.5f, 400.5f}};
......
device/calo/clustering/include/CaloPreFilterClusters.cuh
View file @ 9a8e05f4
......@@ -20,7 +20,7 @@ namespace calo_prefilter_clusters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events;
HOST_INPUT(host_ecal_number_of_clusters_t, unsigned) host_ecal_number_of_clusters;
HOST_INPUT(host_ecal_number_of_neutral_clusters_t, unsigned) host_ecal_number_neutral_of_clusters;
MASK_INPUT(dev_event_list_t) dev_event_list;
......
device/calo/clustering/include/MakeNeutralBasicParticles.cuh
View file @ 9a8e05f4
......@@ -17,10 +17,11 @@
namespace make_neutral_basic_particles {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_number_of_clusters_t, unsigned) host_number_of_clusters;
HOST_INPUT(host_number_of_neutral_clusters_t, unsigned) host_number_of_neutral_clusters;
MASK_INPUT(dev_event_list_t) dev_event_list;
DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events;
DEVICE_INPUT(dev_ecal_cluster_offsets_t, unsigned) dev_ecal_cluster_offsets;
DEVICE_INPUT(dev_ecal_neutral_cluster_offsets_t, unsigned) dev_ecal_neutral_cluster_offsets;
DEVICE_INPUT(dev_ecal_clusters_t, CaloCluster) dev_ecal_clusters;
DEVICE_OUTPUT_WITH_DEPENDENCIES(
dev_neutral_basic_particle_view_t,
......
device/calo/clustering/src/CaloFilterClusters.cu
View file @ 9a8e05f4
......@@ -54,7 +54,6 @@ __global__ void calo_filter_clusters::calo_filter_clusters(calo_filter_clusters:
unsigned* event_cluster1_idx = parameters.dev_cluster1_idx + ecal_twoclusters_offsets;
unsigned* event_cluster2_idx = parameters.dev_cluster2_idx + ecal_twoclusters_offsets;
// const unsigned ecal_cluster_offsets = parameters.dev_ecal_cluster_offsets[event_number];
const auto event_neutral_particles = parameters.dev_neutral_particles->container(event_number);
const unsigned* prefiltered_clusters_idx = parameters.dev_prefiltered_clusters_idx + event_neutral_particles.offset();
const unsigned n_prefltred_clusters = parameters.dev_num_prefiltered_clusters[event_number];
......
device/calo/clustering/src/CaloFindClusters.cu
View file @ 9a8e05f4
......@@ -10,14 +10,18 @@
\*****************************************************************************/
#include <CaloCluster.cuh>
#include <CaloFindClusters.cuh>
#include <PrefixSum.cuh>
INSTANTIATE_ALGORITHM(calo_find_clusters::calo_find_clusters_t)
__device__ void simple_clusters(
CaloDigit const* digits,
const bool* digits_isTrackMatched,
const bool* digits_isBremMatched,
CaloSeedCluster const* seed_clusters,
CaloCluster* clusters,
unsigned const num_clusters,
unsigned* num_neutral_clusters,
const CaloGeometry& calo,
const int16_t min_adc,
float const* corrections,
......@@ -27,10 +31,16 @@ __device__ void simple_clusters(
Allen::Monitoring::Histogram<>::DeviceType histo_ecal_cluster_x,
Allen::Monitoring::Histogram<>::DeviceType histo_ecal_cluster_y)
{
__shared__ unsigned num_neutral_clusters_shared;
if (threadIdx.x == 0) num_neutral_clusters_shared = 0u;
__syncthreads();
for (unsigned c = threadIdx.x; c < num_clusters; c += blockDim.x) {
auto const& seed_cluster = seed_clusters[c];
auto& cluster = clusters[c];
cluster = CaloCluster(calo, seed_cluster);
auto const seed_cluster = seed_clusters[c];
const bool isTrackMatched = digits_isTrackMatched[seed_cluster.id];
const bool isBremMatched = digits_isBremMatched[seed_cluster.id];
auto cluster = CaloCluster(calo, seed_cluster, isTrackMatched, isBremMatched);
uint16_t const* neighbors = &(calo.neighbors[seed_cluster.id * Calo::Constants::max_neighbours]);
for (uint16_t n = 0; n < Calo::Constants::max_neighbours; n++) {
......@@ -65,7 +75,16 @@ __device__ void simple_clusters(
histo_ecal_cluster_et.increment(cluster.et);
histo_ecal_cluster_x.increment(cluster.x);
histo_ecal_cluster_y.increment(cluster.y);
clusters[c] = cluster;
if (!cluster.isTrackMatched) {
atomicAdd(&num_neutral_clusters_shared, 1);
}
}
__syncthreads();
if (threadIdx.x == 0) *num_neutral_clusters = num_neutral_clusters_shared;
}
__global__ void calo_find_clusters::calo_find_clusters(
......@@ -93,9 +112,12 @@ __global__ void calo_find_clusters::calo_find_clusters(
simple_clusters(
parameters.dev_ecal_digits + ecal_digits_offset,
parameters.dev_ecal_digits_isTrackMatched + ecal_digits_offset,
parameters.dev_ecal_digits_isBremMatched + ecal_digits_offset,
parameters.dev_ecal_seed_clusters + Calo::Constants::ecal_max_index / 8 * event_number,
parameters.dev_ecal_clusters + ecal_clusters_offset,
ecal_num_clusters,
parameters.dev_ecal_neutral_cluster_offsets + event_number,
ecal_geometry,
min_adc,
parameters.dev_ecal_corrections + ecal_clusters_offset,
......@@ -111,7 +133,10 @@ void calo_find_clusters::calo_find_clusters_t::set_arguments_size(
const RuntimeOptions&,
const Constants&) const
{
auto const n_events = first<host_number_of_events_t>(arguments);
set_size<dev_ecal_clusters_t>(arguments, first<host_ecal_number_of_clusters_t>(arguments));
set_size<dev_ecal_neutral_cluster_offsets_t>(arguments, n_events + 1);
set_size<host_total_sum_holder_t>(arguments, 1);
}
__host__ void calo_find_clusters::calo_find_clusters_t::operator()(
......@@ -127,6 +152,8 @@ __host__ void calo_find_clusters::calo_find_clusters_t::operator()(
auto dev_histo_ecal_cluster_x = m_histogram_ecal_cluster_x.data(context);
auto dev_histo_ecal_cluster_y = m_histogram_ecal_cluster_y.data(context);
Allen::memset_async<dev_ecal_neutral_cluster_offsets_t>(arguments, 0, context);
// Find clusters.
global_function(calo_find_clusters)(dim3(size<dev_event_list_t>(arguments)), dim3(m_block_dim_x), context)(
arguments,
......@@ -138,4 +165,6 @@ __host__ void calo_find_clusters::calo_find_clusters_t::operator()(
dev_histo_ecal_cluster_et,
dev_histo_ecal_cluster_x,
dev_histo_ecal_cluster_y);
PrefixSum::prefix_sum<dev_ecal_neutral_cluster_offsets_t, host_total_sum_holder_t>(*this, arguments, context);
}
device/calo/clustering/src/CaloPreFilterClusters.cu
View file @ 9a8e05f4
......@@ -23,7 +23,7 @@ void calo_prefilter_clusters::calo_prefilter_clusters_t::set_arguments_size(
set_size<dev_num_prefiltered_clusters_t>(arguments, n_events);
set_size<dev_ecal_twocluster_offsets_t>(arguments, n_events + 1);
set_size<host_total_sum_holder_t>(arguments, 1);
set_size<dev_prefiltered_clusters_idx_t>(arguments, first<host_ecal_number_of_clusters_t>(arguments));
set_size<dev_prefiltered_clusters_idx_t>(arguments, first<host_ecal_number_of_neutral_clusters_t>(arguments));
}
void calo_prefilter_clusters::calo_prefilter_clusters_t::operator()(
......@@ -54,14 +54,25 @@ __global__ void calo_prefilter_clusters::calo_prefilter_clusters(
unsigned* prefiltered_clusters_idx = parameters.dev_prefiltered_clusters_idx + ecal_clusters_offset;
unsigned* num_prefiltered_clusters = parameters.dev_num_prefiltered_clusters + event_number;
__shared__ unsigned num_prefiltered_clusters_shared;
if (threadIdx.x == 0) {
num_prefiltered_clusters_shared = 0u;
}
__syncthreads();
for (unsigned i_cluster = threadIdx.x; i_cluster < ecal_num_clusters; i_cluster += blockDim.x) {
const auto particle = event_neutral_particles.particle(i_cluster);
const auto cluster = particle.cluster();
if (cluster.et > minEt_clusters && cluster.CaloNeutralE19 > minE19_clusters) {
const unsigned idx = atomicAdd(num_prefiltered_clusters, 1);
const unsigned idx = atomicAdd(&num_prefiltered_clusters_shared, 1);
prefiltered_clusters_idx[idx] = i_cluster;
}
}
__syncthreads();
if (threadIdx.x == 0) {
*num_prefiltered_clusters = num_prefiltered_clusters_shared;
}
}
__global__ void calo_prefilter_clusters::count_twoclusters(calo_prefilter_clusters::Parameters parameters)
......
device/calo/clustering/src/CaloSeedClusters.cu
View file @ 9a8e05f4
......@@ -22,14 +22,18 @@ __device__ void seed_clusters(
const int16_t min_adc,
unsigned* digit_is_seed)
{
__shared__ unsigned num_clusters_shared;
if (threadIdx.x == 0) num_clusters_shared = 0u;
__syncthreads();
// Loop over all CellIDs.
for (unsigned i = threadIdx.x; i < num_digits; i += blockDim.x) {
const auto digit = digits[i];
digit_is_seed[i] = 0;
if (digit.adc < min_adc || !digit.is_valid()) {
continue;
}
uint16_t* neighbors = &(geometry.neighbors[i * Calo::Constants::max_neighbours]);
const uint16_t* neighbors = &(geometry.neighbors[i * Calo::Constants::max_neighbours]);
bool is_max = true;
float energy = geometry.getE(i, digit.adc);
for (unsigned n = 0; n < Calo::Constants::max_neighbours; n++) {
......@@ -37,16 +41,21 @@ __device__ void seed_clusters(
if (neighbor_id == USHRT_MAX) {
continue;
}
auto const neighbor_digit = digits[neighbors[n]];
is_max = is_max && (digit.adc > neighbor_digit.adc || !neighbor_digit.is_valid());
if (neighbor_digit.is_valid()) energy += geometry.getE(neighbor_id, neighbor_digit.adc);
auto const neighbor_digit = digits[neighbor_id];
if (neighbor_digit.is_valid()) {
is_max = is_max && (digit.adc > neighbor_digit.adc);
energy += geometry.getE(neighbor_id, neighbor_digit.adc);
}
}
if (is_max) {
auto const id = atomicAdd(num_clusters.data(), 1);
clusters[id] = CaloSeedCluster(i, digits[i].adc, geometry.getX(i), geometry.getY(i), energy);
auto const id = atomicAdd(&num_clusters_shared, 1);
clusters[id] = CaloSeedCluster(i, digit.adc, geometry.getX(i), geometry.getY(i), energy);
digit_is_seed[i] = id;
}
}
__syncthreads();
if (threadIdx.x == 0) num_clusters[0] = num_clusters_shared;
}
__global__ void calo_seed_clusters::calo_seed_clusters(
......@@ -93,6 +102,7 @@ void calo_seed_clusters::calo_seed_clusters_t::operator()(
const Constants& constants,
Allen::Context const& context) const
{
Allen::memset_async<dev_ecal_digit_is_seed_t>(arguments, 0, context);
Allen::memset_async<dev_ecal_cluster_offsets_t>(arguments, 0, context);
// Find local maxima.
......
device/calo/clustering/src/MakeNeutralBasicParticles.cu
View file @ 9a8e05f4
......@@ -17,8 +17,8 @@ void make_neutral_basic_particles::make_neutral_particles_t::set_arguments_size(
const RuntimeOptions&,
const Constants&) const
{
auto n_clusters = first<host_number_of_clusters_t>(arguments);
set_size<dev_neutral_basic_particle_view_t>(arguments, n_clusters);
auto n_neutral_clusters = first<host_number_of_neutral_clusters_t>(arguments);
set_size<dev_neutral_basic_particle_view_t>(arguments, n_neutral_clusters);
set_size<dev_neutral_basic_particles_view_t>(arguments, first<host_number_of_events_t>(arguments));
set_size<dev_multi_event_neutral_particles_view_t>(arguments, 1);
set_size<dev_multi_event_container_neutral_particles_t>(arguments, 1);
......@@ -36,24 +36,39 @@ void make_neutral_basic_particles::make_neutral_particles_t::operator()(
void __global__ make_neutral_basic_particles::make_particles(make_neutral_basic_particles::Parameters parameters)
{
// WARNING: This algorithm creates "NeutralBasicParticle"s for all ECAL clusters,
// regardless of whether or not the cluster is matched to a charged track.
// WARNING: This algorithm creates "NeutralBasicParticle"s for ECAL clusters
// not matched to a charged track.
const unsigned number_of_events = parameters.dev_number_of_events[0];
const unsigned event_number = blockIdx.x;
const unsigned offset = parameters.dev_ecal_cluster_offsets[event_number];
const unsigned number_of_clusters = parameters.dev_ecal_cluster_offsets[event_number + 1] - offset;
const CaloCluster* event_clusters = parameters.dev_ecal_clusters + offset;
const unsigned offset_neutral = parameters.dev_ecal_neutral_cluster_offsets[event_number];
const unsigned number_of_neutral_clusters =
parameters.dev_ecal_neutral_cluster_offsets[event_number + 1] - offset_neutral;
for (unsigned i = threadIdx.x; i < number_of_clusters; i += blockDim.x) {
const auto* calo_cluster = event_clusters + i;
new (parameters.dev_neutral_basic_particle_view + offset + i)
Allen::Views::Physics::NeutralBasicParticle {calo_cluster};
if (number_of_neutral_clusters > 0) {
// Avoid events with no neutral clusters or (more importantly) events for which
// calo_find_clusters was masked and CaloClusters were not initialized
__shared__ unsigned i_neutral_particle;
if (threadIdx.x == 0) i_neutral_particle = 0;
__syncthreads();
for (unsigned i = threadIdx.x; i < number_of_clusters; i += blockDim.x) {
const auto* calo_cluster = event_clusters + i;
if (!calo_cluster->isTrackMatched) {
const unsigned id = atomicAdd(&i_neutral_particle, 1);
new (parameters.dev_neutral_basic_particle_view + offset_neutral + id)
Allen::Views::Physics::NeutralBasicParticle {calo_cluster};
}
}
}
if (threadIdx.x == 0) {
new (parameters.dev_neutral_basic_particles_view + event_number) Allen::Views::Physics::NeutralBasicParticles {
parameters.dev_neutral_basic_particle_view, parameters.dev_ecal_cluster_offsets, event_number};
parameters.dev_neutral_basic_particle_view, parameters.dev_ecal_neutral_cluster_offsets, event_number};
}
if (blockIdx.x == 0 && threadIdx.x == 0) {
......
device/calo/tools/include/BremRecovery.cuh
View file @ 9a8e05f4
......@@ -26,6 +26,7 @@ namespace brem_recovery {
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_velo_kalman_beamline_states;
// Calo digits
HOST_INPUT(host_ecal_number_of_digits_t, unsigned) host_ecal_number_of_digits;
DEVICE_INPUT(dev_ecal_digits_t, CaloDigit) dev_ecal_digits;
DEVICE_INPUT(dev_ecal_digits_offsets_t, unsigned) dev_ecal_digits_offsets;
// Output
......@@ -34,7 +35,7 @@ namespace brem_recovery {
DEVICE_OUTPUT(dev_brem_inECALacc_t, bool) dev_brem_inECALacc;
DEVICE_OUTPUT(dev_brem_ecal_digits_size_t, unsigned) dev_brem_ecal_digits_size;
DEVICE_OUTPUT(dev_brem_ecal_digits_t, std::array<unsigned, 4>) dev_brem_ecal_digits;
// Properties
DEVICE_OUTPUT(dev_ecal_digits_isBremMatched_t, bool) dev_ecal_digits_isBremMatched;
};
struct brem_recovery_t : public DeviceAlgorithm, Parameters {
......
device/vertex_fit/vertex_fitter/include/MFVertexFitter.cuh → device/calo/tools/include/FakeDigitMatching.cuh
View file @ 9a8e05f4
/*****************************************************************************\
* (c) Copyright 2018-2020 CERN for the benefit of the LHCb Collaboration *
* (c) Copyright 2024 CERN for the benefit of the LHCb Collaboration *
* *
* This software is distributed under the terms of the Apache License *
* version 2 (Apache-2.0), copied verbatim in the file "LICENSE". *
......@@ -10,29 +10,16 @@
\*****************************************************************************/
#pragma once
#include "VertexDefinitions.cuh"
#include "VertexFitDeviceFunctions.cuh"
#include "ParKalmanDefinitions.cuh"
#include "ParKalmanMath.cuh"
#include "AlgorithmTypes.cuh"
namespace MFVertexFit {
namespace fake_digit_matching {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_number_of_mf_svs_t, unsigned) host_number_of_mf_svs;
HOST_INPUT(host_selected_events_mf_t, unsigned) host_selected_events_mf;
DEVICE_INPUT(dev_kf_particles_t, Allen::Views::Physics::BasicParticles) dev_kf_particles;
DEVICE_INPUT(dev_mf_particles_t, Allen::Views::Physics::BasicParticles) dev_mf_particles;
DEVICE_INPUT(dev_mf_sv_offsets_t, unsigned) dev_mf_sv_offsets;
DEVICE_INPUT(dev_svs_kf_idx_t, unsigned) dev_svs_kf_idx;
DEVICE_INPUT(dev_svs_mf_idx_t, unsigned) dev_svs_mf_idx;
DEVICE_INPUT(dev_event_list_mf_t, unsigned) dev_event_list_mf;
DEVICE_OUTPUT(dev_mf_svs_t, VertexFit::TrackMVAVertex) dev_mf_svs;
HOST_INPUT(host_ecal_number_of_digits_t, unsigned) host_ecal_number_of_digits;
MASK_INPUT(dev_event_list_t) dev_event_list;
DEVICE_OUTPUT(dev_ecal_digits_isMatched_t, bool) dev_ecal_digits_isMatched;
};
__global__ void fit_mf_vertices(Parameters);
struct fit_mf_vertices_t : public DeviceAlgorithm, Parameters {
struct fake_digit_matching_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
void operator()(
......@@ -40,9 +27,5 @@ namespace MFVertexFit {
const RuntimeOptions&,
const Constants&,
const Allen::Context& context) const;
private:
Allen::Property<dim3> m_block_dim {this, "block_dim", {16, 16, 1}, "block dimensions"};
};
} // namespace MFVertexFit
\ No newline at end of file
} // namespace fake_digit_matching