diff --git a/configuration/python/AllenConf/HLT1_cosmics.py b/configuration/python/AllenConf/HLT1_cosmics.py
index 7711d046475cd2cd4865157740e89b1459165c16..1d4e0ec3369d95ea49408ef51ecaca9daf67e9d4 100644
--- a/configuration/python/AllenConf/HLT1_cosmics.py
+++ b/configuration/python/AllenConf/HLT1_cosmics.py
@@ -79,7 +79,8 @@ def setup_hlt1_node(enableRateValidator=True):
# Reconstruct objects needed as input for selection lines
decoded_calo = decode_calo()
- ecal_clusters = make_ecal_clusters(decoded_calo)
+ ecal_clusters = make_ecal_clusters(
+ decoded_calo, calo_find_clusters_name='calo_find_clusters_cosmics')
hlt1_config['reconstruction'] = {'ecal_clusters': ecal_clusters}
diff --git a/configuration/python/AllenConf/calo_reconstruction.py b/configuration/python/AllenConf/calo_reconstruction.py
index b0811c0155f2c4df62824ba45157b810e9a7d1cd..97b50026085133eb8946e1b03dc55605468e4d75 100755
--- a/configuration/python/AllenConf/calo_reconstruction.py
+++ b/configuration/python/AllenConf/calo_reconstruction.py
@@ -148,6 +148,7 @@ def make_track_matching(decoded_calo, velo_tracks, velo_states, long_tracks,
@configurable
def make_ecal_clusters(decoded_calo,
+ calo_find_clusters_name='calo_find_clusters',
seed_min_adc=50,
neighbour_min_adc=10,
min_et=400,
@@ -183,7 +184,7 @@ def make_ecal_clusters(decoded_calo,
calo_find_clusters = make_algorithm(
calo_find_clusters_t,
- name='calo_find_clusters_{hash}',
+ name=str(calo_find_clusters_name),
ecal_min_adc=neighbour_min_adc,
host_ecal_number_of_clusters_t=prefix_sum_ecal_num_clusters.
host_total_sum_holder_t,
@@ -273,8 +274,9 @@ def make_ecal_clusters(decoded_calo,
}
-def ecal_cluster_reco():
+def ecal_cluster_reco(calo_find_clusters_name='calo_find_clusters_reco'):
decoded_calo = decode_calo()
- ecal_clusters = make_ecal_clusters(decoded_calo)
+ ecal_clusters = make_ecal_clusters(
+ decoded_calo, calo_find_clusters_name=calo_find_clusters_name)
alg = ecal_clusters["dev_ecal_clusters"].producer
return alg
diff --git a/configuration/python/AllenConf/hlt1_reconstruction.py b/configuration/python/AllenConf/hlt1_reconstruction.py
index c0b4452ff8631f0ad3aa9e740f5052cd7fe6c5e4..41439ba13096dda6ca500ea53003fc60f71d786f 100644
--- a/configuration/python/AllenConf/hlt1_reconstruction.py
+++ b/configuration/python/AllenConf/hlt1_reconstruction.py
@@ -130,7 +130,8 @@ def hlt1_reconstruction(algorithm_name='',
if with_muon:
decoded_muon = decode_muon()
- muonID = is_muon(decoded_muon, long_tracks)
+ muonID = is_muon(
+ decoded_muon, long_tracks, is_muon_name=algorithm_name + 'is_muon')
# Replace long tracks with those containing muon hits.
long_tracks = muonID["long_tracks"]
else:
diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py
index ec2cfc2a3ab3777d0a5c807d70dbc7594d2403f8..e382e9dad8474c4bc1a97e5ab483955a1f74dfc1 100644
--- a/configuration/python/AllenConf/muon_reconstruction.py
+++ b/configuration/python/AllenConf/muon_reconstruction.py
@@ -115,7 +115,7 @@ def decode_muon(empty_banks=False):
}
-def is_muon(decoded_muon, long_tracks):
+def is_muon(decoded_muon, long_tracks, is_muon_name='is_muon'):
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"]
@@ -123,10 +123,11 @@ def is_muon(decoded_muon, long_tracks):
host_number_of_reconstructed_scifi_tracks = long_tracks[
"host_number_of_reconstructed_scifi_tracks"]
dev_scifi_states = long_tracks["dev_scifi_states"]
+ velo_kalman_filter = long_tracks["velo_kalman_filter"]
is_muon = make_algorithm(
is_muon_t,
- name='is_muon_{hash}',
+ name=str(is_muon_name),
host_number_of_events_t=host_number_of_events,
dev_number_of_events_t=dev_number_of_events,
host_number_of_reconstructed_scifi_tracks_t=
@@ -135,6 +136,8 @@ def is_muon(decoded_muon, long_tracks):
dev_long_tracks_view_t=long_tracks["dev_multi_event_long_tracks_view"],
dev_station_ocurrences_offset_t=decoded_muon[
"dev_station_ocurrences_offset"],
+ dev_velo_states_view_t=velo_kalman_filter[
+ "dev_velo_kalman_endvelo_states_view"],
dev_muon_hits_t=decoded_muon["dev_muon_hits"])
muon_hit_count_prefix_sum = make_algorithm(
@@ -206,6 +209,7 @@ def muon_id(algorithm_name=''):
from AllenConf.ut_reconstruction import decode_ut, make_ut_tracks
from AllenConf.scifi_reconstruction import decode_scifi, make_forward_tracks
+ if (algorithm_name != ''): algorithm_name = algorithm_name + '_'
decoded_velo = decode_velo()
velo_tracks = make_velo_tracks(decoded_velo)
decoded_ut = decode_ut()
@@ -216,9 +220,10 @@ def muon_id(algorithm_name=''):
ut_tracks,
velo_tracks["dev_accepted_velo_tracks"],
scifi_consolidate_tracks_name=algorithm_name +
- '_scifi_consolidate_tracks_muon_id')
+ 'scifi_consolidate_tracks_muon_id')
decoded_muon = decode_muon()
- muonID = is_muon(decoded_muon, long_tracks)
+ muonID = is_muon(
+ decoded_muon, long_tracks, is_muon_name=algorithm_name + 'is_muon')
alg = muonID["dev_is_muon"].producer
return alg
diff --git a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
index 45bdf874459fa2819ea169e7327bc30324c92fd9..886cd0724a63d970532f2af42cb69bd11cb85ed0 100644
--- a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
+++ b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
@@ -16,10 +16,6 @@ void pv_beamline_cleanup::pv_beamline_cleanup_t::init()
{
#ifndef ALLEN_STANDALONE
m_pvs = new Gaudi::Accumulators::AveragingCounter<> {this, "n_PVs"};
- histogram_n_pvs = new gaudi_monitoring::Lockable_Histogram<> {{this, "n_pvs_event", "n_pvs_event", {20, 0, 20}}, {}};
- histogram_pv_x = new gaudi_monitoring::Lockable_Histogram<> {{this, "pv_x", "pv_x", {100, -2.f, 2.f}}, {}};
- histogram_pv_y = new gaudi_monitoring::Lockable_Histogram<> {{this, "pv_y", "pv_y", {100, -2.f, 2.f}}, {}};
- histogram_pv_z = new gaudi_monitoring::Lockable_Histogram<> {{this, "pv_z", "pv_z", {100, -200.f, 200.f}}, {}};
histogram_n_smogpvs =
new gaudi_monitoring::Lockable_Histogram<> {{this, "n_smog2_PVs", "n_smog2_PVs", {10, -0.5f, 9.5f}}, {}};
@@ -29,6 +25,11 @@ void pv_beamline_cleanup::pv_beamline_cleanup_t::init()
"smogpv_z",
{property<nbins_histo_smogpvz_t>(), property<min_histo_smogpvz_t>(), property<max_histo_smogpvz_t>()}},
{}};
+ histogram_n_pvs =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "n_pvs_event", "n_pvs_event", {21, -0.5, 20.5}}, {}};
+ histogram_pv_x = new gaudi_monitoring::Lockable_Histogram<> {{this, "pv_x", "pv_x", {1000, -2.f, 2.f}}, {}};
+ histogram_pv_y = new gaudi_monitoring::Lockable_Histogram<> {{this, "pv_y", "pv_y", {1000, -2.f, 2.f}}, {}};
+ histogram_pv_z = new gaudi_monitoring::Lockable_Histogram<> {{this, "pv_z", "pv_z", {1000, -200.f, 200.f}}, {}};
#endif
}
@@ -48,14 +49,12 @@ void pv_beamline_cleanup::pv_beamline_cleanup_t::operator()(
const Allen::Context& context) const
{
auto dev_n_pvs_counter = make_device_buffer<unsigned>(arguments, 1u);
- auto dev_n_pvs_histo = make_device_buffer<unsigned>(arguments, 20u);
auto dev_n_smogpvs_histo = make_device_buffer<unsigned>(arguments, 10u);
-
- auto dev_pv_x_histo = make_device_buffer<unsigned>(arguments, 100u);
- auto dev_pv_y_histo = make_device_buffer<unsigned>(arguments, 100u);
- auto dev_pv_z_histo = make_device_buffer<unsigned>(arguments, 100u);
auto dev_smogpv_z_histo = make_device_buffer<unsigned>(arguments, property<nbins_histo_smogpvz_t>());
-
+ auto dev_n_pvs_histo = make_device_buffer<unsigned>(arguments, 21u);
+ auto dev_pv_x_histo = make_device_buffer<unsigned>(arguments, 1000u);
+ auto dev_pv_y_histo = make_device_buffer<unsigned>(arguments, 1000u);
+ auto dev_pv_z_histo = make_device_buffer<unsigned>(arguments, 1000u);
Allen::memset_async(dev_n_pvs_counter.data(), 0, dev_n_pvs_counter.size() * sizeof(unsigned), context);
Allen::memset_async(dev_n_pvs_histo.data(), 0, dev_n_pvs_histo.size() * sizeof(unsigned), context);
Allen::memset_async(dev_n_smogpvs_histo.data(), 0, dev_n_smogpvs_histo.size() * sizeof(unsigned), context);
@@ -81,8 +80,8 @@ void pv_beamline_cleanup::pv_beamline_cleanup_t::operator()(
arguments,
context,
std::tuple {std::tuple {dev_n_pvs_counter.get(), m_pvs},
- std::tuple {dev_n_pvs_histo.get(), histogram_n_pvs, 0, 20},
std::tuple {dev_n_smogpvs_histo.get(), histogram_n_smogpvs, -0.5f, 9.5f},
+ std::tuple {dev_n_pvs_histo.get(), histogram_n_pvs, -0.5, 20.5},
std::tuple {dev_pv_x_histo.get(), histogram_pv_x, -2, 2},
std::tuple {dev_pv_y_histo.get(), histogram_pv_y, -2, 2},
std::tuple {dev_pv_z_histo.get(), histogram_pv_z, -200, 200},
@@ -139,15 +138,15 @@ __global__ void pv_beamline_cleanup::pv_beamline_cleanup(
// monitoring
if (-2 < vertex1.position.x && vertex1.position.x < 2 && -200 < vertex1.position.z && vertex1.position.z < 200) {
- unsigned x_bin = std::floor(vertex1.position.x / 0.04f) + 50;
+ unsigned x_bin = std::floor(vertex1.position.x / 0.004f) + 500;
atomicAdd(&dev_pv_x_histo[x_bin], 1);
}
if (-2 < vertex1.position.y && vertex1.position.y < 2 && -200 < vertex1.position.z && vertex1.position.z < 200) {
- unsigned y_bin = std::floor(vertex1.position.y / 0.04f) + 50;
+ unsigned y_bin = std::floor(vertex1.position.y / 0.004f) + 500;
atomicAdd(&dev_pv_y_histo[y_bin], 1);
}
if (-200 < vertex1.position.z && vertex1.position.z < 200) {
- unsigned z_bin = std::floor(vertex1.position.z / 4) + 50;
+ unsigned z_bin = std::floor(vertex1.position.z / 0.4f) + 500;
atomicAdd(&dev_pv_z_histo[z_bin], 1);
}
if (parameters.min_histo_smogpvz < vertex1.position.z && vertex1.position.z < parameters.max_histo_smogpvz) {
@@ -163,8 +162,7 @@ __global__ void pv_beamline_cleanup::pv_beamline_cleanup(
__syncthreads();
parameters.dev_number_of_multi_final_vertices[event_number] = *tmp_number_vertices;
- if (*tmp_number_vertices < 20) atomicAdd(&dev_n_pvs_histo[*tmp_number_vertices], 1);
+ if (*tmp_number_vertices < 21) atomicAdd(&dev_n_pvs_histo[*tmp_number_vertices], 1);
if (*tmp_number_SMOG_vertices < 10) atomicAdd(&dev_n_smogpvs_histo[*tmp_number_SMOG_vertices], 1);
-
dev_n_pvs_counter[0] += *tmp_number_vertices;
}
diff --git a/device/SciFi/consolidate/include/ConsolidateSciFi.cuh b/device/SciFi/consolidate/include/ConsolidateSciFi.cuh
index 75163b6d12e3c1838c1284ef84c0cb8a6b0d6f51..3ef18bcacf19035d274e32dc575a93eadfb2b3f2 100644
--- a/device/SciFi/consolidate/include/ConsolidateSciFi.cuh
+++ b/device/SciFi/consolidate/include/ConsolidateSciFi.cuh
@@ -21,7 +21,6 @@
#include "CopyTrackParameters.cuh"
#ifndef ALLEN_STANDALONE
-#include <Gaudi/Accumulators.h>
#include "GaudiMonitoring.h"
#endif
@@ -183,9 +182,9 @@ namespace scifi_consolidate_tracks {
Property<histogram_long_track_forward_phi_min_t> m_histogramLongPhiMin {this, -4.f};
Property<histogram_long_track_forward_phi_max_t> m_histogramLongPhiMax {this, 4.f};
Property<histogram_long_track_forward_phi_nbins_t> m_histogramLongPhiNBins {this, 16u};
- Property<histogram_long_track_forward_nhits_min_t> m_histogramLongNhitsMin {this, 0.f};
- Property<histogram_long_track_forward_nhits_max_t> m_histogramLongNhitsMax {this, 50.f};
- Property<histogram_long_track_forward_nhits_nbins_t> m_histogramLongNhitsNBins {this, 50u};
+ Property<histogram_long_track_forward_nhits_min_t> m_histogramLongNhitsMin {this, -0.5f};
+ Property<histogram_long_track_forward_nhits_max_t> m_histogramLongNhitsMax {this, 50.5f};
+ Property<histogram_long_track_forward_nhits_nbins_t> m_histogramLongNhitsNBins {this, 51u};
#ifndef ALLEN_STANDALONE
private:
diff --git a/device/SciFi/consolidate/src/ConsolidateSciFi.cu b/device/SciFi/consolidate/src/ConsolidateSciFi.cu
index efc74b9334f3543cd6d6d9623a24bcbc342a3d12..fdae0e1d0991e374a9253aa5f51f758c1b8d4b19 100644
--- a/device/SciFi/consolidate/src/ConsolidateSciFi.cu
+++ b/device/SciFi/consolidate/src/ConsolidateSciFi.cu
@@ -155,7 +155,7 @@ void scifi_consolidate_tracks::scifi_consolidate_tracks_t::init()
#ifndef ALLEN_STANDALONE
m_long_tracks_forward = new Gaudi::Accumulators::Counter<>(this, "n_long_tracks_forward");
histogram_n_long_tracks_forward = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "n_long_tracks_forward_event", "n_long_tracks_forward_event", {80, 0, 200, {}, {}}}, {}};
+ {this, "n_long_tracks_forward_event", "n_long_tracks_forward_event", {201, -0.5, 200.5, {}, {}}}, {}};
histogram_long_track_forward_eta =
new gaudi_monitoring::Lockable_Histogram<> {{this,
"long_track_forward_eta",
@@ -200,7 +200,7 @@ void scifi_consolidate_tracks::scifi_consolidate_tracks_t::operator()(
make_device_buffer<unsigned>(arguments, property<histogram_long_track_forward_phi_nbins_t>());
auto dev_histogram_long_track_forward_nhits =
make_device_buffer<unsigned>(arguments, property<histogram_long_track_forward_nhits_nbins_t>());
- auto dev_histogram_n_long_tracks_forward = make_device_buffer<unsigned>(arguments, 80u);
+ auto dev_histogram_n_long_tracks_forward = make_device_buffer<unsigned>(arguments, 201u);
auto dev_n_long_tracks_forward_counter = make_device_buffer<unsigned>(arguments, 1u);
Allen::memset_async(
dev_histogram_long_track_forward_eta.data(),
@@ -254,7 +254,7 @@ void scifi_consolidate_tracks::scifi_consolidate_tracks_t::operator()(
histogram_long_track_forward_nhits,
property<histogram_long_track_forward_nhits_min_t>(),
property<histogram_long_track_forward_nhits_max_t>()},
- std::tuple {dev_histogram_n_long_tracks_forward.get(), histogram_n_long_tracks_forward, 0, 200},
+ std::tuple {dev_histogram_n_long_tracks_forward.get(), histogram_n_long_tracks_forward, -0.5f, 200.5f},
std::tuple {dev_n_long_tracks_forward_counter.get(), m_long_tracks_forward}});
#endif
}
@@ -395,10 +395,7 @@ __device__ void scifi_consolidate_tracks_impl(
auto used_scifi_hits = parameters.dev_used_scifi_hits.get();
auto accepted_velo_tracks = parameters.dev_accepted_and_unused_velo_tracks.get();
- if (number_of_tracks_event < 200) {
- unsigned bin = std::floor(number_of_tracks_event / 2.5);
- atomicAdd(&dev_histogram_n_long_tracks_forward[bin], 1);
- }
+ if (number_of_tracks_event < 201) atomicAdd(&dev_histogram_n_long_tracks_forward[number_of_tracks_event], 1);
dev_n_long_tracks_forward_counter[0] += number_of_tracks_event;
// Loop over tracks.
diff --git a/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh b/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
index fff2ea021ed11c6a08537c943b1e274ff0d8401c..df17c94a9646f6cf07cd8e8dda89bc1013e62b57 100644
--- a/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
+++ b/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
@@ -152,13 +152,13 @@ namespace seed_confirmTracks_consolidate {
Property<block_dim_t> m_block_dim {this, {{256, 1, 1}}};
Property<histogram_scifi_track_eta_min_t> m_histogramSciFiEtaMin {this, 0.f};
Property<histogram_scifi_track_eta_max_t> m_histogramSciFiEtaMax {this, 10.f};
- Property<histogram_scifi_track_eta_nbins_t> m_histogramSciFiEtaNBins {this, 40u};
+ Property<histogram_scifi_track_eta_nbins_t> m_histogramSciFiEtaNBins {this, 400u};
Property<histogram_scifi_track_phi_min_t> m_histogramSciFiPhiMin {this, -4.f};
Property<histogram_scifi_track_phi_max_t> m_histogramSciFiPhiMax {this, 4.f};
Property<histogram_scifi_track_phi_nbins_t> m_histogramSciFiPhiNBins {this, 16u};
- Property<histogram_scifi_track_nhits_min_t> m_histogramSciFiNhitsMin {this, 0.f};
- Property<histogram_scifi_track_nhits_max_t> m_histogramSciFiNhitsMax {this, 14.f};
- Property<histogram_scifi_track_nhits_nbins_t> m_histogramSciFiNhitsNBins {this, 14u};
+ Property<histogram_scifi_track_nhits_min_t> m_histogramSciFiNhitsMin {this, -0.5f};
+ Property<histogram_scifi_track_nhits_max_t> m_histogramSciFiNhitsMax {this, 14.5f};
+ Property<histogram_scifi_track_nhits_nbins_t> m_histogramSciFiNhitsNBins {this, 15u};
#ifndef ALLEN_STANDALONE
private:
Gaudi::Accumulators::Counter<>* m_seed_tracks;
diff --git a/device/SciFi/hybridseeding/consolidate/src/ConsolidateSciFiSeeding.cu b/device/SciFi/hybridseeding/consolidate/src/ConsolidateSciFiSeeding.cu
index 61ef2ffa591acd23c54d7d7fbafdb0c378e4e093..087c1776bbee6b823d4bca8fb741ad48eda13f81 100644
--- a/device/SciFi/hybridseeding/consolidate/src/ConsolidateSciFiSeeding.cu
+++ b/device/SciFi/hybridseeding/consolidate/src/ConsolidateSciFiSeeding.cu
@@ -72,7 +72,7 @@ void seed_confirmTracks_consolidate::seed_confirmTracks_consolidate_t::init()
#ifndef ALLEN_STANDALONE
m_seed_tracks = new Gaudi::Accumulators::Counter<>(this, "n_seed_tracks");
histogram_n_scifi_seeds = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "n_scifi_seeds_event", "n_scifi_seeds_event", {80, 0, 200, {}, {}}}, {}};
+ {this, "n_scifi_seeds_event", "n_scifi_seeds_event", {201, -0.5, 200.5, {}, {}}}, {}};
histogram_scifi_track_eta =
new gaudi_monitoring::Lockable_Histogram<> {{this,
"scifi_track_eta",
@@ -147,7 +147,7 @@ void seed_confirmTracks_consolidate::seed_confirmTracks_consolidate_t::operator(
make_device_buffer<unsigned>(arguments, property<histogram_scifi_track_phi_nbins_t>());
auto dev_histogram_scifi_track_nhits =
make_device_buffer<unsigned>(arguments, property<histogram_scifi_track_nhits_nbins_t>());
- auto dev_histogram_scifi_n_tracks = make_device_buffer<unsigned>(arguments, 80u);
+ auto dev_histogram_scifi_n_tracks = make_device_buffer<unsigned>(arguments, 201u);
auto dev_scifi_n_tracks_counter = make_device_buffer<unsigned>(arguments, 1u);
Allen::memset_async(
dev_histogram_scifi_track_eta.data(), 0, dev_histogram_scifi_track_eta.size() * sizeof(unsigned), context);
@@ -189,7 +189,7 @@ void seed_confirmTracks_consolidate::seed_confirmTracks_consolidate_t::operator(
histogram_scifi_track_nhits,
property<histogram_scifi_track_nhits_min_t>(),
property<histogram_scifi_track_nhits_max_t>()},
- std::tuple {dev_histogram_scifi_n_tracks.get(), histogram_n_scifi_seeds, 0, 200},
+ std::tuple {dev_histogram_scifi_n_tracks.get(), histogram_n_scifi_seeds, -0.5, 200.5},
std::tuple {dev_scifi_n_tracks_counter.get(), m_seed_tracks}});
#endif
}
@@ -237,9 +237,8 @@ __global__ void seed_confirmTracks_consolidate::seed_confirmTracks_consolidate(
// int * tracks_nY = parameters.dev_seeding_nY + parameters.dev_atomics_scifi[event_number];
auto used_scifi_hits = parameters.dev_used_scifi_hits.get();
- if (number_of_tracks_event < 200) {
- unsigned bin = std::floor(number_of_tracks_event / 2.5);
- atomicAdd(&dev_histogram_scifi_n_tracks[bin], 1);
+ if (number_of_tracks_event < 201) {
+ atomicAdd(&dev_histogram_scifi_n_tracks[number_of_tracks_event], 1);
}
dev_scifi_n_tracks_counter[0] += number_of_tracks_event;
diff --git a/device/calo/clustering/include/CaloFindClusters.cuh b/device/calo/clustering/include/CaloFindClusters.cuh
index 6b4f8df91f4c8771c533a47c331c598fc80557af..4208b9276e69e0cac767d85cea0027ff924a0647 100755
--- a/device/calo/clustering/include/CaloFindClusters.cuh
+++ b/device/calo/clustering/include/CaloFindClusters.cuh
@@ -17,6 +17,10 @@
#include "AlgorithmTypes.cuh"
#include <cfloat>
+#ifndef ALLEN_STANDALONE
+#include "GaudiMonitoring.h"
+#endif
+
namespace calo_find_clusters {
struct Parameters {
HOST_INPUT(host_ecal_number_of_clusters_t, unsigned) host_ecal_number_of_clusters;
@@ -33,11 +37,21 @@ namespace calo_find_clusters {
};
// Global function
- __global__ void calo_find_clusters(Parameters parameters, const char* raw_ecal_geometry, const int16_t min_adc);
+ __global__ void calo_find_clusters(
+ Parameters parameters,
+ const char* raw_ecal_geometry,
+ const int16_t min_adc,
+ gsl::span<unsigned> histogram_n_clusters,
+ gsl::span<unsigned> histogram_digit_e,
+ gsl::span<unsigned> histogram_cluster_e,
+ gsl::span<unsigned> histogram_cluster_et,
+ gsl::span<unsigned> histogram_cluster_x,
+ gsl::span<unsigned> histogram_cluster_y);
// Algorithm
struct calo_find_clusters_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters>, const RuntimeOptions&, const Constants&) const;
+ void init();
__host__ void operator()(
const ArgumentReferences<Parameters>& arguments,
@@ -49,5 +63,13 @@ namespace calo_find_clusters {
Property<block_dim_x_t> m_block_dim_x {this, 64};
Property<block_dim_y_t> m_block_dim_y {this, 16};
Property<ecal_min_adc_t> m_ecal_min_adc {this, 10};
+#ifndef ALLEN_STANDALONE
+ gaudi_monitoring::Lockable_Histogram<>* histogram_n_clusters;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_ecal_digit_e;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_ecal_cluster_e;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_ecal_cluster_et;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_ecal_cluster_x;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_ecal_cluster_y;
+#endif
};
} // namespace calo_find_clusters
diff --git a/device/calo/clustering/src/CaloFindClusters.cu b/device/calo/clustering/src/CaloFindClusters.cu
index e39a6c44dfd1dd50d11478e66e006ce45a0b93ea..32f57bbec4fe9c59dfe8a59ca01064bb5fc54211 100755
--- a/device/calo/clustering/src/CaloFindClusters.cu
+++ b/device/calo/clustering/src/CaloFindClusters.cu
@@ -13,6 +13,24 @@
INSTANTIATE_ALGORITHM(calo_find_clusters::calo_find_clusters_t)
+void calo_find_clusters::calo_find_clusters_t::init()
+{
+#ifndef ALLEN_STANDALONE
+ histogram_n_clusters =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "n_ecal_clusters", "NClusters", {401, -0.5, 400.5}}, {}};
+ histogram_ecal_digit_e =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "ecal_digit_e", "EcalDigitE", {1000, 0, 10000}}, {}};
+ histogram_ecal_cluster_e =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "ecal_cluster_e", "EcalClusterE", {5000, 0, 50000}}, {}};
+ histogram_ecal_cluster_et =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "ecal_cluster_et", "EcalClusterEt", {500, 0, 5000}}, {}};
+ histogram_ecal_cluster_x =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "ecal_cluster_x", "EcalClusterX", {800, -4000, 4000}}, {}};
+ histogram_ecal_cluster_y =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "ecal_cluster_y", "EcalClusterY", {800, -4000, 4000}}, {}};
+#endif
+}
+
__device__ void simple_clusters(
CaloDigit const* digits,
CaloSeedCluster const* seed_clusters,
@@ -20,7 +38,12 @@ __device__ void simple_clusters(
unsigned const num_clusters,
const CaloGeometry& calo,
const int16_t min_adc,
- float const* corrections)
+ float const* corrections,
+ gsl::span<unsigned> histogram_digit_e,
+ gsl::span<unsigned> histogram_cluster_e,
+ gsl::span<unsigned> histogram_cluster_et,
+ gsl::span<unsigned> histogram_cluster_x,
+ gsl::span<unsigned> histogram_cluster_y)
{
for (unsigned c = threadIdx.x; c < num_clusters; c += blockDim.x) {
auto const& seed_cluster = seed_clusters[c];
@@ -35,8 +58,14 @@ __device__ void simple_clusters(
}
auto const digit = digits[n_id];
if (digit.is_valid() && (digit.adc > min_adc)) {
- cluster.e += calo.getE(n_id, digit.adc);
+ const auto digit_e = calo.getE(n_id, digit.adc);
+ cluster.e += digit_e;
cluster.digits[n] = n_id;
+
+ if (digit_e < 10000) {
+ const unsigned bin = std::floor(digit_e / 10);
+ atomicAdd(&histogram_digit_e[bin], 1);
+ }
}
}
cluster.e -= corrections[c];
@@ -53,13 +82,37 @@ __device__ void simple_clusters(
}
cluster.CalcEt();
cluster.CaloNeutralE19 = calo.getE(seed_cluster.id, seed_cluster.adc) / cluster.e;
+
+ // Fill histograms
+ if (0 < cluster.e && cluster.e < 50000) {
+ const unsigned bin = std::floor(cluster.e / 10);
+ atomicAdd(&histogram_cluster_e[bin], 1);
+ }
+ if (0 < cluster.et && cluster.et < 5000) {
+ const unsigned bin = std::floor(cluster.et / 10);
+ atomicAdd(&histogram_cluster_et[bin], 1);
+ }
+ if (cluster.x < 4000 && cluster.x > -4000) {
+ const unsigned bin = std::floor(cluster.x / 10) + 400;
+ atomicAdd(&histogram_cluster_x[bin], 1);
+ }
+ if (cluster.y < 4000 && cluster.y > -4000) {
+ const unsigned bin = std::floor(cluster.y / 10) + 400;
+ atomicAdd(&histogram_cluster_y[bin], 1);
+ }
}
}
__global__ void calo_find_clusters::calo_find_clusters(
calo_find_clusters::Parameters parameters,
const char* raw_ecal_geometry,
- const int16_t min_adc)
+ const int16_t min_adc,
+ gsl::span<unsigned> histogram_n_clusters,
+ gsl::span<unsigned> histogram_digit_e,
+ gsl::span<unsigned> histogram_cluster_e,
+ gsl::span<unsigned> histogram_cluster_et,
+ gsl::span<unsigned> histogram_cluster_x,
+ gsl::span<unsigned> histogram_cluster_y)
{
// Get proper geometry.
auto ecal_geometry = CaloGeometry(raw_ecal_geometry);
@@ -70,6 +123,9 @@ __global__ void calo_find_clusters::calo_find_clusters(
unsigned const ecal_digits_offset = parameters.dev_ecal_digits_offsets[event_number];
unsigned const ecal_clusters_offset = parameters.dev_ecal_cluster_offsets[event_number];
unsigned const ecal_num_clusters = parameters.dev_ecal_cluster_offsets[event_number + 1] - ecal_clusters_offset;
+
+ if (ecal_num_clusters < 401) atomicAdd(&histogram_n_clusters[ecal_num_clusters], 1);
+
simple_clusters(
parameters.dev_ecal_digits + ecal_digits_offset,
parameters.dev_ecal_seed_clusters + Calo::Constants::ecal_max_index / 8 * event_number,
@@ -77,7 +133,12 @@ __global__ void calo_find_clusters::calo_find_clusters(
ecal_num_clusters,
ecal_geometry,
min_adc,
- parameters.dev_ecal_corrections + ecal_clusters_offset);
+ parameters.dev_ecal_corrections + ecal_clusters_offset,
+ histogram_digit_e,
+ histogram_cluster_e,
+ histogram_cluster_et,
+ histogram_cluster_x,
+ histogram_cluster_y);
}
void calo_find_clusters::calo_find_clusters_t::set_arguments_size(
@@ -94,8 +155,41 @@ __host__ void calo_find_clusters::calo_find_clusters_t::operator()(
const Constants& constants,
Allen::Context const& context) const
{
+ auto dev_histogram_n_clusters = make_device_buffer<unsigned>(arguments, 401u);
+ auto dev_histogram_digit_e = make_device_buffer<unsigned>(arguments, 1000u);
+ auto dev_histogram_cluster_e = make_device_buffer<unsigned>(arguments, 5000u);
+ auto dev_histogram_cluster_et = make_device_buffer<unsigned>(arguments, 500u);
+ auto dev_histogram_cluster_x = make_device_buffer<unsigned>(arguments, 800u);
+ auto dev_histogram_cluster_y = make_device_buffer<unsigned>(arguments, 800u);
+ Allen::memset_async(dev_histogram_n_clusters.data(), 0, dev_histogram_n_clusters.size() * sizeof(unsigned), context);
+ Allen::memset_async(dev_histogram_digit_e.data(), 0, dev_histogram_digit_e.size() * sizeof(unsigned), context);
+ Allen::memset_async(dev_histogram_cluster_e.data(), 0, dev_histogram_cluster_e.size() * sizeof(unsigned), context);
+ Allen::memset_async(dev_histogram_cluster_et.data(), 0, dev_histogram_cluster_et.size() * sizeof(unsigned), context);
+ Allen::memset_async(dev_histogram_cluster_x.data(), 0, dev_histogram_cluster_x.size() * sizeof(unsigned), context);
+ Allen::memset_async(dev_histogram_cluster_y.data(), 0, dev_histogram_cluster_y.size() * sizeof(unsigned), context);
+
// Find clusters.
global_function(calo_find_clusters)(
dim3(size<dev_event_list_t>(arguments)), dim3(property<block_dim_x_t>().get()), context)(
- arguments, constants.dev_ecal_geometry, property<ecal_min_adc_t>().get());
+ arguments,
+ constants.dev_ecal_geometry,
+ property<ecal_min_adc_t>().get(),
+ dev_histogram_n_clusters.get(),
+ dev_histogram_digit_e.get(),
+ dev_histogram_cluster_e.get(),
+ dev_histogram_cluster_et.get(),
+ dev_histogram_cluster_x.get(),
+ dev_histogram_cluster_y.get());
+
+#ifndef ALLEN_STANDALONE
+ gaudi_monitoring::fill(
+ arguments,
+ context,
+ std::tuple {std::tuple {dev_histogram_n_clusters.get(), histogram_n_clusters, -0.5, 400.5},
+ std::tuple {dev_histogram_digit_e.get(), histogram_ecal_digit_e, 0, 10000},
+ std::tuple {dev_histogram_cluster_e.get(), histogram_ecal_cluster_e, 0, 50000},
+ std::tuple {dev_histogram_cluster_et.get(), histogram_ecal_cluster_et, 0, 5000},
+ std::tuple {dev_histogram_cluster_x.get(), histogram_ecal_cluster_x, -4000.f, 4000.f},
+ std::tuple {dev_histogram_cluster_y.get(), histogram_ecal_cluster_y, -4000.f, 4000.f}});
+#endif
}
diff --git a/device/downstream/consolidate/src/DownstreamConsolidate.cu b/device/downstream/consolidate/src/DownstreamConsolidate.cu
index ccf30afeeeae0a80f3a71c4a3db6cce050812d32..044736d4cd4c2161773e8a38c01c4ae910617087 100644
--- a/device/downstream/consolidate/src/DownstreamConsolidate.cu
+++ b/device/downstream/consolidate/src/DownstreamConsolidate.cu
@@ -153,7 +153,7 @@ void downstream_consolidate::downstream_consolidate_t::init()
#ifndef ALLEN_STANDALONE
m_downstream_tracks = new Gaudi::Accumulators::Counter<>(this, "n_downstream_tracks");
histogram_n_downstream_tracks = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "n_downstream_tracks_event", "n_downstream_tracks_event", {80, 0, 200}}, {}};
+ {this, "n_downstream_tracks_event", "n_downstream_tracks_event", {201, -0.5, 200.5}}, {}};
histogram_downstream_track_eta =
new gaudi_monitoring::Lockable_Histogram<> {{this,
"downstream_track_eta",
@@ -200,7 +200,7 @@ void downstream_consolidate::downstream_consolidate_t::operator()(
make_device_buffer<unsigned>(arguments, property<histogram_downstream_track_phi_nbins_t>());
auto dev_histogram_downstream_track_nhits =
make_device_buffer<unsigned>(arguments, property<histogram_downstream_track_nhits_nbins_t>());
- auto dev_histogram_n_downstream_tracks = make_device_buffer<unsigned>(arguments, 80u);
+ auto dev_histogram_n_downstream_tracks = make_device_buffer<unsigned>(arguments, 201u);
auto dev_n_downstream_tracks_counter = make_device_buffer<unsigned>(arguments, 1u);
//
@@ -256,7 +256,7 @@ void downstream_consolidate::downstream_consolidate_t::operator()(
histogram_downstream_track_nhits,
property<histogram_downstream_track_nhits_min_t>(),
property<histogram_downstream_track_nhits_max_t>()},
- std::tuple {dev_histogram_n_downstream_tracks.get(), histogram_n_downstream_tracks, 0, 200},
+ std::tuple {dev_histogram_n_downstream_tracks.get(), histogram_n_downstream_tracks, -0.5f, 200.5f},
std::tuple {dev_n_downstream_tracks_counter.get(), m_downstream_tracks}});
#endif
}
@@ -301,10 +301,7 @@ __global__ void downstream_consolidate::downstream_consolidate(
//
// Monitoring fill
//
- if (downstream_tracks_size < 200) {
- unsigned bin = std::floor(downstream_tracks_size / 2.5);
- atomicAdd(dev_histogram_n_downstream_tracks.data() + bin, 1u);
- }
+ if (downstream_tracks_size < 201) atomicAdd(&dev_histogram_n_downstream_tracks[downstream_tracks_size], 1);
dev_n_downstream_tracks_counter[0] += downstream_tracks_size;
// Outputs
@@ -387,13 +384,13 @@ __device__ void downstream_consolidate::downstream_consolidate_t::monitor(
const unsigned int bin = static_cast<unsigned int>(
(eta - parameters.histogram_downstream_track_eta_min) * parameters.histogram_downstream_track_eta_nbins /
(parameters.histogram_downstream_track_eta_max - parameters.histogram_downstream_track_eta_min));
- atomicAdd(dev_histogram_downstream_track_eta.data() + bin, 1u);
+ atomicAdd(&dev_histogram_downstream_track_eta[bin], 1);
}
if (phi > parameters.histogram_downstream_track_phi_min && phi < parameters.histogram_downstream_track_phi_max) {
const unsigned int bin = static_cast<unsigned int>(
(phi - parameters.histogram_downstream_track_phi_min) * parameters.histogram_downstream_track_phi_nbins /
(parameters.histogram_downstream_track_phi_max - parameters.histogram_downstream_track_phi_min));
- atomicAdd(dev_histogram_downstream_track_phi.data() + bin, 1u);
+ atomicAdd(&dev_histogram_downstream_track_phi[bin], 1);
}
if (
nhits > parameters.histogram_downstream_track_nhits_min &&
@@ -401,6 +398,6 @@ __device__ void downstream_consolidate::downstream_consolidate_t::monitor(
const unsigned int bin = static_cast<unsigned int>(
(nhits - parameters.histogram_downstream_track_nhits_min) * parameters.histogram_downstream_track_nhits_nbins /
(parameters.histogram_downstream_track_nhits_max - parameters.histogram_downstream_track_nhits_min));
- atomicAdd(dev_histogram_downstream_track_nhits.data() + bin, 1u);
+ atomicAdd(&dev_histogram_downstream_track_nhits[bin], 1);
}
}
diff --git a/device/event_model/common/include/ParticleTypes.cuh b/device/event_model/common/include/ParticleTypes.cuh
index 772ce8f999e72556b40c1cb36de9cf4818754599..a5f2caff0942e8d4f790d7814cc5b881b44c0015 100644
--- a/device/event_model/common/include/ParticleTypes.cuh
+++ b/device/event_model/common/include/ParticleTypes.cuh
@@ -259,6 +259,16 @@ namespace Allen {
const float* qop) :
Track {velo_segment, ut_segment, scifi_segment, muon_segment, qop}
{}
+
+ __host__ __device__ float pt(Allen::Views::Physics::KalmanState velo_state) const
+ {
+ const auto qop = *m_qop;
+ const float tx = velo_state.tx();
+ const float ty = velo_state.ty();
+ const float slope2 = tx * tx + ty * ty;
+ const float pt = std::sqrt(slope2 / (1.0f + slope2)) / std::fabs(qop);
+ return pt;
+ }
};
struct LongTracks : ILHCbIDContainer<LongTracks> {
diff --git a/device/kalman/ParKalman/src/MakeLongTrackParticles.cu b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
index f64b55dfc91b3c5a089c7581cfd4b6732c8a916c..add01e2e21cd449e939eb8ecf8da756f153180a0 100644
--- a/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
+++ b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
@@ -17,11 +17,12 @@ void make_long_track_particles::make_long_track_particles_t::init()
{
#ifndef ALLEN_STANDALONE
histogram_n_trks = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "number_of_trks", "NTrks", {UT::Constants::max_num_tracks, 0, UT::Constants::max_num_tracks}}, {}};
- histogram_trk_eta = new gaudi_monitoring::Lockable_Histogram<> {{this, "trk_eta", "etaTrk", {100, 0, 5}}, {}};
- histogram_trk_phi = new gaudi_monitoring::Lockable_Histogram<> {{this, "trk_phi", "phiTrk", {100, -3.2, 3.2}}, {}};
+ {this, "number_of_trks", "NTrks", {UT::Constants::max_num_tracks + 1, -0.5, UT::Constants::max_num_tracks + .5}},
+ {}};
+ histogram_trk_eta = new gaudi_monitoring::Lockable_Histogram<> {{this, "trk_eta", "etaTrk", {1000, 0, 5}}, {}};
+ histogram_trk_phi = new gaudi_monitoring::Lockable_Histogram<> {{this, "trk_phi", "phiTrk", {1000, -3.2, 3.2}}, {}};
histogram_trk_pt =
- new gaudi_monitoring::Lockable_Histogram<> {{this, "trk_pt", "ptTrk", {100, 0, (unsigned) 1e4}}, {}};
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "trk_pt", "ptTrk", {1000, 0, (unsigned) 1e4}}, {}};
#endif
}
@@ -43,10 +44,10 @@ void make_long_track_particles::make_long_track_particles_t::operator()(
const Constants&,
const Allen::Context& context) const
{
- auto dev_histogram_n_trks = make_device_buffer<unsigned>(arguments, UT::Constants::max_num_tracks);
- auto dev_histogram_trk_eta = make_device_buffer<unsigned>(arguments, 100u);
- auto dev_histogram_trk_phi = make_device_buffer<unsigned>(arguments, 100u);
- auto dev_histogram_trk_pt = make_device_buffer<unsigned>(arguments, 100u);
+ auto dev_histogram_n_trks = make_device_buffer<unsigned>(arguments, UT::Constants::max_num_tracks + 1);
+ auto dev_histogram_trk_eta = make_device_buffer<unsigned>(arguments, 1000u);
+ auto dev_histogram_trk_phi = make_device_buffer<unsigned>(arguments, 1000u);
+ auto dev_histogram_trk_pt = make_device_buffer<unsigned>(arguments, 1000u);
Allen::memset_async(dev_histogram_n_trks.data(), 0, dev_histogram_n_trks.size() * sizeof(unsigned), context);
Allen::memset_async(dev_histogram_trk_eta.data(), 0, dev_histogram_trk_eta.size() * sizeof(unsigned), context);
Allen::memset_async(dev_histogram_trk_phi.data(), 0, dev_histogram_trk_phi.size() * sizeof(unsigned), context);
@@ -65,7 +66,7 @@ void make_long_track_particles::make_long_track_particles_t::operator()(
gaudi_monitoring::fill(
arguments,
context,
- std::tuple {std::tuple {dev_histogram_n_trks.get(), histogram_n_trks, 0, UT::Constants::max_num_tracks},
+ std::tuple {std::tuple {dev_histogram_n_trks.get(), histogram_n_trks, -0.5, UT::Constants::max_num_tracks + .5},
std::tuple {dev_histogram_trk_eta.get(), histogram_trk_eta, 0, 5},
std::tuple {dev_histogram_trk_phi.get(), histogram_trk_phi, -3.2f, 3.2f},
std::tuple {dev_histogram_trk_pt.get(), histogram_trk_pt, 0u, unsigned(1e4)}});
@@ -92,7 +93,7 @@ void __global__ make_long_track_particles::make_particles(
const unsigned number_of_tracks = event_long_tracks.size();
const auto pv_table = parameters.dev_kalman_pv_tables[event_number];
- if (number_of_tracks < UT::Constants::max_num_tracks) atomicAdd(&dev_histogram_n_trks[number_of_tracks], 1);
+ if (number_of_tracks < UT::Constants::max_num_tracks + 1) atomicAdd(&dev_histogram_n_trks[number_of_tracks], 1);
for (unsigned i = threadIdx.x; i < number_of_tracks; i += blockDim.x) {
const auto* long_track = &(event_long_tracks.track(i));
@@ -106,13 +107,19 @@ void __global__ make_long_track_particles::make_particles(
parameters.dev_lepton_id[offset + i]};
auto state = (parameters.dev_kalman_states_view + event_number)->state(i);
- const unsigned etabin = max(0u, min(99u, static_cast<unsigned>(state.eta() * 20)));
- atomicAdd(&dev_histogram_trk_eta[etabin], 1);
- const unsigned phibin =
- max(0u, min(99u, static_cast<unsigned>(std::atan2(state.ty(), state.tx()) * 15.625f + 50)));
- atomicAdd(&dev_histogram_trk_phi[phibin], 1);
- const unsigned ptbin = min(99u, static_cast<unsigned>(state.pt() * 0.01f));
- atomicAdd(&dev_histogram_trk_pt[ptbin], 1);
+ if (state.eta() > 0 && state.eta() < 5) {
+ unsigned bin = std::floor(state.eta() * 200);
+ atomicAdd(&dev_histogram_trk_eta[bin], 1);
+ }
+ float phi = std::atan2(state.ty(), state.tx());
+ if (phi > -3.2f && phi < 3.2f) {
+ unsigned bin = std::floor((phi + 3.2f) * 1000 / 6.4f);
+ atomicAdd(&dev_histogram_trk_phi[bin], 1);
+ }
+ if (state.pt() > 0 && state.pt() < float(1e4)) {
+ unsigned bin = std::floor(state.pt() * 0.1f);
+ atomicAdd(&dev_histogram_trk_pt[bin], 1);
+ }
}
}
diff --git a/device/muon/is_muon/include/IsMuon.cuh b/device/muon/is_muon/include/IsMuon.cuh
index 2519c2bf7dde7ab8a3f8ec8206bfaa9246574edb..107cccbcdfadd808ab6b4fa39e9f6f88ffbc8a84 100644
--- a/device/muon/is_muon/include/IsMuon.cuh
+++ b/device/muon/is_muon/include/IsMuon.cuh
@@ -16,6 +16,10 @@
#include "SciFiConsolidated.cuh"
#include "ParticleTypes.cuh"
+#ifndef ALLEN_STANDALONE
+#include "GaudiMonitoring.h"
+#endif
+
namespace is_muon {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
@@ -26,6 +30,7 @@ namespace is_muon {
DEVICE_INPUT(dev_scifi_states_t, MiniState) dev_scifi_states;
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_velo_states_view_t, Allen::Views::Physics::KalmanStates) dev_velo_states_view;
DEVICE_OUTPUT(dev_is_muon_t, bool) dev_is_muon;
DEVICE_OUTPUT(dev_muon_idxs_t, unsigned) dev_muon_idxs;
DEVICE_OUTPUT(dev_muon_hit_counts_t, unsigned) dev_muon_hit_counts;
@@ -33,12 +38,18 @@ namespace is_muon {
PROPERTY(block_dim_x_t, "block_dim_x", "block dimension X", unsigned) block_dim_x;
};
- __global__ void
- is_muon(Parameters, const Muon::Constants::FieldOfInterest* dev_muon_foi, const float* dev_muon_momentum_cuts);
+ __global__ void is_muon(
+ Parameters,
+ const Muon::Constants::FieldOfInterest* dev_muon_foi,
+ const float* dev_muon_momentum_cuts,
+ gsl::span<unsigned> dev_histogram_n_muons,
+ gsl::span<unsigned> dev_histogram_muon_n_stations,
+ gsl::span<unsigned> dev_histogram_muon_pt);
struct is_muon_t : public DeviceAlgorithm, Parameters {
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const;
+ void init();
void operator()(
const ArgumentReferences<Parameters>& arguments,
const RuntimeOptions& runtime_options,
@@ -47,5 +58,10 @@ namespace is_muon {
private:
Property<block_dim_x_t> m_block_dim_x {this, 128};
+#ifndef ALLEN_STANDALONE
+ gaudi_monitoring::Lockable_Histogram<>* histogram_n_muons;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_muon_n_stations;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_muon_pt;
+#endif
};
} // namespace is_muon
diff --git a/device/muon/is_muon/src/IsMuon.cu b/device/muon/is_muon/src/IsMuon.cu
index af0c20a16c223c006284bd63664e5dd0883df27b..6c4dc055743f597d175c4435b9aecae4e4896914 100644
--- a/device/muon/is_muon/src/IsMuon.cu
+++ b/device/muon/is_muon/src/IsMuon.cu
@@ -13,6 +13,17 @@
INSTANTIATE_ALGORITHM(is_muon::is_muon_t)
+void is_muon::is_muon_t::init()
+{
+#ifndef ALLEN_STANDALONE
+ histogram_n_muons = new gaudi_monitoring::Lockable_Histogram<> {{this, "n_muons", "# muons", {2, -0.5, 1.5}}, {}};
+ histogram_muon_n_stations =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "muon_n_stations", "# muon stations", {3, 1.5, 4.5}}, {}};
+ histogram_muon_pt =
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "muon_pt", "muon pt", {1000, 0, (unsigned) 1e4}}, {}};
+#endif
+}
+
void is_muon::is_muon_t::set_arguments_size(
ArgumentReferences<Parameters> arguments,
const RuntimeOptions&,
@@ -31,12 +42,34 @@ void is_muon::is_muon_t::operator()(
const Constants& constants,
const Allen::Context& context) const
{
+ auto dev_histogram_n_muons = make_device_buffer<unsigned>(arguments, 2u);
+ auto dev_histogram_muon_n_stations = make_device_buffer<unsigned>(arguments, 3u);
+ auto dev_histogram_muon_pt = make_device_buffer<unsigned>(arguments, 1000u);
+ Allen::memset_async(dev_histogram_n_muons.data(), 0, dev_histogram_n_muons.size() * sizeof(unsigned), context);
+ Allen::memset_async(
+ dev_histogram_muon_n_stations.data(), 0, dev_histogram_muon_n_stations.size() * sizeof(unsigned), context);
+ Allen::memset_async(dev_histogram_muon_pt.data(), 0, dev_histogram_muon_pt.size() * sizeof(unsigned), context);
+
Allen::memset_async<dev_is_muon_t>(arguments, 0, context);
Allen::memset_async<dev_lepton_id_t>(arguments, 0, context);
Allen::memset_async<dev_muon_hit_counts_t>(arguments, 0, context);
global_function(is_muon)(dim3(size<dev_event_list_t>(arguments)), dim3(property<block_dim_x_t>().get()), context)(
- arguments, constants.dev_muon_foi, constants.dev_muon_momentum_cuts);
+ arguments,
+ constants.dev_muon_foi,
+ constants.dev_muon_momentum_cuts,
+ dev_histogram_n_muons.get(),
+ dev_histogram_muon_n_stations.get(),
+ dev_histogram_muon_pt.get());
+
+#ifndef ALLEN_STANDALONE
+ gaudi_monitoring::fill(
+ arguments,
+ context,
+ std::tuple {std::tuple {dev_histogram_n_muons.get(), histogram_n_muons, -0.5, 1.5},
+ std::tuple {dev_histogram_muon_n_stations.get(), histogram_muon_n_stations, 1.5, 4.5},
+ std::tuple {dev_histogram_muon_pt.get(), histogram_muon_pt, 0, (unsigned) 1e4}});
+#endif
}
__device__ float elliptical_foi_window(const float a, const float b, const float c, const float momentum)
@@ -83,7 +116,10 @@ __device__ bool is_in_window(
__global__ void is_muon::is_muon(
is_muon::Parameters parameters,
const Muon::Constants::FieldOfInterest* dev_muon_foi,
- const float* dev_muon_momentum_cuts)
+ const float* dev_muon_momentum_cuts,
+ gsl::span<unsigned> dev_histogram_n_muons,
+ gsl::span<unsigned> dev_histogram_muon_n_stations,
+ gsl::span<unsigned> dev_histogram_muon_pt)
{
// Put foi parameters in shared memory
__shared__ int8_t shared_muon_foi_params_content[sizeof(Muon::Constants::FieldOfInterest)];
@@ -124,6 +160,7 @@ __global__ void is_muon::is_muon(
const auto& state = parameters.dev_scifi_states[event_offset + track_id];
if (momentum < dev_muon_momentum_cuts[0]) {
+ atomicAdd(&dev_histogram_n_muons[0], 1);
continue;
}
@@ -174,14 +211,83 @@ __global__ void is_muon::is_muon(
if (occupancies[0] != 0 && occupancies[1] != 0) {
if (momentum < dev_muon_momentum_cuts[1]) {
parameters.dev_is_muon[event_offset + track_id] = true;
+
+ // monitoring
+ atomicAdd(&dev_histogram_n_muons[1], 1);
+ unsigned n_stations = 0;
+ if (occupancies[2] != 0) {
+ ++n_stations;
+ }
+ if (occupancies[3] != 0) {
+ ++n_stations;
+ }
+ atomicAdd(&dev_histogram_muon_n_stations[n_stations], 1);
+ const auto long_track = long_tracks.track(track_id);
+ const auto velo_track = long_track.track_segment<Allen::Views::Physics::Track::segment::velo>();
+ const auto velo_track_index = velo_track.track_index();
+ const auto endvelo_states = parameters.dev_velo_states_view[event_number];
+ const auto velo_state = endvelo_states.state(velo_track_index);
+ const float pt = long_track.pt(velo_state);
+ if (pt < unsigned(1e4)) {
+ const unsigned bin = std::floor(pt / 10.f);
+ atomicAdd(&dev_histogram_muon_pt[bin], 1);
+ }
}
else if (momentum < dev_muon_momentum_cuts[2]) {
parameters.dev_is_muon[event_offset + track_id] = (occupancies[2] != 0) || (occupancies[3] != 0);
+
+ // monitoring
+ if ((occupancies[2] != 0) || (occupancies[3] != 0)) {
+ atomicAdd(&dev_histogram_n_muons[1], 1);
+ unsigned n_stations = 0;
+ if (occupancies[2] != 0) {
+ ++n_stations;
+ }
+ if (occupancies[3] != 0) {
+ ++n_stations;
+ }
+ atomicAdd(&dev_histogram_muon_n_stations[n_stations], 1);
+ const auto long_track = long_tracks.track(track_id);
+ const auto velo_track = long_track.track_segment<Allen::Views::Physics::Track::segment::velo>();
+ const auto velo_track_index = velo_track.track_index();
+ const auto endvelo_states = parameters.dev_velo_states_view[event_number];
+ const auto velo_state = endvelo_states.state(velo_track_index);
+ const float pt = long_track.pt(velo_state);
+ if (pt < unsigned(1e4)) {
+ const unsigned bin = std::floor(pt / 10.f);
+ atomicAdd(&dev_histogram_muon_pt[bin], 1);
+ }
+ }
+ else {
+ atomicAdd(&dev_histogram_n_muons[0], 1);
+ }
}
else {
parameters.dev_is_muon[event_offset + track_id] = (occupancies[2] != 0) && (occupancies[3] != 0);
+
+ // monitoring
+ if ((occupancies[2] != 0) && (occupancies[3] != 0)) {
+ atomicAdd(&dev_histogram_n_muons[1], 1);
+ atomicAdd(&dev_histogram_muon_n_stations[2], 1);
+ const auto long_track = long_tracks.track(track_id);
+ const auto velo_track = long_track.track_segment<Allen::Views::Physics::Track::segment::velo>();
+ const auto velo_track_index = velo_track.track_index();
+ const auto endvelo_states = parameters.dev_velo_states_view[event_number];
+ const auto velo_state = endvelo_states.state(velo_track_index);
+ const float pt = long_track.pt(velo_state);
+ if (pt < unsigned(1e4)) {
+ const unsigned bin = std::floor(pt / 10.f);
+ atomicAdd(&dev_histogram_muon_pt[bin], 1);
+ }
+ }
+ else {
+ atomicAdd(&dev_histogram_n_muons[0], 1);
+ }
}
parameters.dev_lepton_id[event_offset + track_id] = parameters.dev_is_muon[event_offset + track_id];
}
+ else {
+ atomicAdd(&dev_histogram_n_muons[0], 1);
+ }
}
}
diff --git a/device/selections/lines/calibration/include/D2KPiLine.cuh b/device/selections/lines/calibration/include/D2KPiLine.cuh
index 9d201bd350f14d0c1f72b80f0a490664acf681f0..34ab96da75e3d961e7a28842b24ede0747799d17 100644
--- a/device/selections/lines/calibration/include/D2KPiLine.cuh
+++ b/device/selections/lines/calibration/include/D2KPiLine.cuh
@@ -116,7 +116,7 @@ namespace d2kpi_line {
Property<histogram_d0_mass_min_t> m_histogramD0MassMin {this, 1765.f};
Property<histogram_d0_mass_max_t> m_histogramD0MassMax {this, 1965.f};
Property<histogram_d0_mass_nbins_t> m_histogramD0MassNBins {this, 100u};
- Property<histogram_d0_pt_min_t> m_histogramD0PtMin {this, 0.f};
+ Property<histogram_d0_pt_min_t> m_histogramD0PtMin {this, 800.f};
Property<histogram_d0_pt_max_t> m_histogramD0PtMax {this, 1e4};
Property<histogram_d0_pt_nbins_t> m_histogramD0PtNBins {this, 100u};
Property<enable_monitoring_t> m_enable_monitoring {this, false};
diff --git a/device/selections/lines/charm/include/D2KKLine.cuh b/device/selections/lines/charm/include/D2KKLine.cuh
index f4b0753f5e59cb463692700ff6011a66f7591c3d..c8e5e61d816663175b8c0dd5eeb6960416650150 100644
--- a/device/selections/lines/charm/include/D2KKLine.cuh
+++ b/device/selections/lines/charm/include/D2KKLine.cuh
@@ -125,7 +125,7 @@ namespace d2kk_line {
Property<histogram_d02kk_mass_min_t> m_histogramD0MassMin {this, 1765.f};
Property<histogram_d02kk_mass_max_t> m_histogramD0MassMax {this, 1965.f};
Property<histogram_d02kk_mass_nbins_t> m_histogramD0MassNBins {this, 100u};
- Property<histogram_d02kk_pt_min_t> m_histogramD0PtMin {this, 0.f};
+ Property<histogram_d02kk_pt_min_t> m_histogramD0PtMin {this, 800.f};
Property<histogram_d02kk_pt_max_t> m_histogramD0PtMax {this, 1e4};
Property<histogram_d02kk_pt_nbins_t> m_histogramD0PtNBins {this, 100u};
Property<enable_monitoring_t> m_enable_monitoring {this, false};
diff --git a/device/selections/lines/charm/include/D2PiPiLine.cuh b/device/selections/lines/charm/include/D2PiPiLine.cuh
index e5e58ee315094cd1e067d66634962d20aaf6126c..975734d8bb8f03cd852018ff79a2dbeedb73723e 100644
--- a/device/selections/lines/charm/include/D2PiPiLine.cuh
+++ b/device/selections/lines/charm/include/D2PiPiLine.cuh
@@ -133,7 +133,7 @@ namespace d2pipi_line {
Property<histogram_d02pipi_mass_min_t> m_histogramD0MassMin {this, 1765.f};
Property<histogram_d02pipi_mass_max_t> m_histogramD0MassMax {this, 1965.f};
Property<histogram_d02pipi_mass_nbins_t> m_histogramD0MassNBins {this, 100u};
- Property<histogram_d02pipi_pt_min_t> m_histogramD0PtMin {this, 0.f};
+ Property<histogram_d02pipi_pt_min_t> m_histogramD0PtMin {this, 800.f};
Property<histogram_d02pipi_pt_max_t> m_histogramD0PtMax {this, 1e4};
Property<histogram_d02pipi_pt_nbins_t> m_histogramD0PtNBins {this, 100u};
Property<enable_monitoring_t> m_enable_monitoring {this, false};
diff --git a/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh b/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
index cf8ecbd482c47d64923cfa6ce7a3dd1e8ba2491a..e599309f5e668532c1bf24b20bbb0b0cde90cefb 100644
--- a/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
+++ b/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
@@ -167,13 +167,13 @@ namespace matching_consolidate_tracks {
Property<block_dim_t> m_block_dim {this, {{256, 1, 1}}};
Property<histogram_long_track_matching_eta_min_t> m_histogramLongEtaMin {this, 0.f};
Property<histogram_long_track_matching_eta_max_t> m_histogramLongEtaMax {this, 10.f};
- Property<histogram_long_track_matching_eta_nbins_t> m_histogramLongEtaNBins {this, 40u};
+ Property<histogram_long_track_matching_eta_nbins_t> m_histogramLongEtaNBins {this, 400u};
Property<histogram_long_track_matching_phi_min_t> m_histogramLongPhiMin {this, -4.f};
Property<histogram_long_track_matching_phi_max_t> m_histogramLongPhiMax {this, 4.f};
Property<histogram_long_track_matching_phi_nbins_t> m_histogramLongPhiNBins {this, 16u};
- Property<histogram_long_track_matching_nhits_min_t> m_histogramLongNhitsMin {this, 0.f};
- Property<histogram_long_track_matching_nhits_max_t> m_histogramLongNhitsMax {this, 50.f};
- Property<histogram_long_track_matching_nhits_nbins_t> m_histogramLongNhitsNBins {this, 50u};
+ Property<histogram_long_track_matching_nhits_min_t> m_histogramLongNhitsMin {this, -0.5f};
+ Property<histogram_long_track_matching_nhits_max_t> m_histogramLongNhitsMax {this, 50.5f};
+ Property<histogram_long_track_matching_nhits_nbins_t> m_histogramLongNhitsNBins {this, 51u};
#ifndef ALLEN_STANDALONE
private:
diff --git a/device/track_matching/consolidate/src/ConsolidateMatchedTracks.cu b/device/track_matching/consolidate/src/ConsolidateMatchedTracks.cu
index 37119ec74416cda09c7a2aa13d5cf1135682d90a..2c092c676cf31b0443596903511b41abfe25237b 100644
--- a/device/track_matching/consolidate/src/ConsolidateMatchedTracks.cu
+++ b/device/track_matching/consolidate/src/ConsolidateMatchedTracks.cu
@@ -73,7 +73,7 @@ void matching_consolidate_tracks::matching_consolidate_tracks_t::init()
#ifndef ALLEN_STANDALONE
m_long_tracks_matching = new Gaudi::Accumulators::Counter<>(this, "n_long_tracks_matching");
histogram_n_long_tracks_matching = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "n_long_tracks_matching_event", "n_long_tracks_matching_event", {80, 0, 200}}, {}};
+ {this, "n_long_tracks_matching_event", "n_long_tracks_matching_event", {201, -0.5, 200.5}}, {}};
histogram_long_track_matching_eta =
new gaudi_monitoring::Lockable_Histogram<> {{this,
"long_track_matching_eta",
@@ -116,7 +116,7 @@ void matching_consolidate_tracks::matching_consolidate_tracks_t::operator()(
make_device_buffer<unsigned>(arguments, property<histogram_long_track_matching_phi_nbins_t>());
auto dev_histogram_long_track_matching_nhits =
make_device_buffer<unsigned>(arguments, property<histogram_long_track_matching_nhits_nbins_t>());
- auto dev_histogram_n_long_tracks_matching = make_device_buffer<unsigned>(arguments, 80u);
+ auto dev_histogram_n_long_tracks_matching = make_device_buffer<unsigned>(arguments, 201u);
auto dev_n_long_tracks_matching_counter = make_device_buffer<unsigned>(arguments, 1u);
Allen::memset_async(
dev_histogram_long_track_matching_eta.data(),
@@ -172,7 +172,7 @@ void matching_consolidate_tracks::matching_consolidate_tracks_t::operator()(
histogram_long_track_matching_nhits,
property<histogram_long_track_matching_nhits_min_t>(),
property<histogram_long_track_matching_nhits_max_t>()},
- std::tuple {dev_histogram_n_long_tracks_matching.get(), histogram_n_long_tracks_matching, 0, 200},
+ std::tuple {dev_histogram_n_long_tracks_matching.get(), histogram_n_long_tracks_matching, -0.5, 200.5},
std::tuple {dev_n_long_tracks_matching_counter.get(), m_long_tracks_matching}});
#endif
}
@@ -225,10 +225,7 @@ __global__ void matching_consolidate_tracks::matching_consolidate_tracks(
dev_histogram_long_track_matching_phi,
dev_histogram_long_track_matching_nhits);
- if (number_of_tracks_event < 200) {
- unsigned bin = std::floor(number_of_tracks_event / 2.5);
- atomicAdd(&dev_histogram_n_long_tracks_matching[bin], 1);
- }
+ if (number_of_tracks_event < 201) atomicAdd(&dev_histogram_n_long_tracks_matching[number_of_tracks_event], 1);
dev_n_long_tracks_matching_counter[0] += number_of_tracks_event;
#endif
}
diff --git a/device/velo/consolidate_tracks/src/VeloConsolidateTracks.cu b/device/velo/consolidate_tracks/src/VeloConsolidateTracks.cu
index c55981435e23acec131cfe76e12cb10173649e4a..468978a82f86887052a22747899e16b7697d5f67 100644
--- a/device/velo/consolidate_tracks/src/VeloConsolidateTracks.cu
+++ b/device/velo/consolidate_tracks/src/VeloConsolidateTracks.cu
@@ -77,7 +77,7 @@ void velo_consolidate_tracks::velo_consolidate_tracks_t::init()
m_velo_tracks = new Gaudi::Accumulators::Counter<>(this, "n_velo_tracks");
histogram_n_velo_tracks = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "n_velo_tracks_event", "n_velo_tracks_event", {100, 0, 500}}, {}};
+ {this, "n_velo_tracks_event", "n_velo_tracks_event", {501, 0, 500}}, {}};
#endif
}
@@ -94,7 +94,7 @@ void velo_consolidate_tracks::velo_consolidate_tracks_t::operator()(
Allen::memset_async<dev_velo_multi_event_tracks_view_t>(arguments, 0, context);
Allen::memset_async<dev_velo_tracks_view_t>(arguments, 0, context);
- auto dev_number_of_tracks_histo = make_device_buffer<unsigned>(arguments, 100u);
+ auto dev_number_of_tracks_histo = make_device_buffer<unsigned>(arguments, 501u);
auto dev_tracks_counter = make_device_buffer<unsigned>(arguments, 1u);
Allen::memset_async(
dev_number_of_tracks_histo.data(), 0, dev_number_of_tracks_histo.size() * sizeof(unsigned), context);
@@ -110,7 +110,7 @@ void velo_consolidate_tracks::velo_consolidate_tracks_t::operator()(
gaudi_monitoring::fill(
arguments,
context,
- std::tuple {std::tuple {dev_number_of_tracks_histo.get(), histogram_n_velo_tracks, 0, 500},
+ std::tuple {std::tuple {dev_number_of_tracks_histo.get(), histogram_n_velo_tracks, -0.5, 500.5},
std::tuple {dev_tracks_counter.get(), m_velo_tracks}});
#endif
}
@@ -147,10 +147,7 @@ __global__ void velo_consolidate_tracks::velo_consolidate_tracks(
const auto event_number_of_tracks_in_main_track_container =
event_total_number_of_tracks - event_number_of_three_hit_tracks_filtered;
- if (event_total_number_of_tracks < 500) {
- unsigned bin = std::floor(event_total_number_of_tracks / 5);
- atomicAdd(&dev_number_of_tracks_histo[bin], 1);
- }
+ if (event_total_number_of_tracks < 501) atomicAdd(&dev_number_of_tracks_histo[event_total_number_of_tracks], 1);
dev_tracks_counter[0] += event_total_number_of_tracks;
// Pointers to data within event
diff --git a/device/velo/simplified_kalman_filter/include/VeloKalmanFilter.cuh b/device/velo/simplified_kalman_filter/include/VeloKalmanFilter.cuh
index ba2057e605c663b64871b6eeaa6ea954e07384fc..6d323787adc372784cddf9aca5b7f9a10c79591b 100644
--- a/device/velo/simplified_kalman_filter/include/VeloKalmanFilter.cuh
+++ b/device/velo/simplified_kalman_filter/include/VeloKalmanFilter.cuh
@@ -296,9 +296,9 @@ namespace velo_kalman_filter {
Property<histogram_velo_track_phi_min_t> m_histogramVeloPhiMin {this, -4.f};
Property<histogram_velo_track_phi_max_t> m_histogramVeloPhiMax {this, 4.f};
Property<histogram_velo_track_phi_nbins_t> m_histogramVeloPhiNBins {this, 16u};
- Property<histogram_velo_track_nhits_min_t> m_histogramVeloNhitsMin {this, 0.f};
- Property<histogram_velo_track_nhits_max_t> m_histogramVeloNhitsMax {this, 50.f};
- Property<histogram_velo_track_nhits_nbins_t> m_histogramVeloNhitsNBins {this, 50u};
+ Property<histogram_velo_track_nhits_min_t> m_histogramVeloNhitsMin {this, -0.5f};
+ Property<histogram_velo_track_nhits_max_t> m_histogramVeloNhitsMax {this, 50.5f};
+ Property<histogram_velo_track_nhits_nbins_t> m_histogramVeloNhitsNBins {this, 51u};
#ifndef ALLEN_STANDALONE
private:
diff --git a/device/vertex_fit/vertex_fitter/src/VertexFitter.cu b/device/vertex_fit/vertex_fitter/src/VertexFitter.cu
index cf26ecc26d45fae7a9d309ce3db15f56a7c38d81..c815910a39f1527c1efdd2b754fa32660d404769 100644
--- a/device/vertex_fit/vertex_fitter/src/VertexFitter.cu
+++ b/device/vertex_fit/vertex_fitter/src/VertexFitter.cu
@@ -16,7 +16,7 @@ void VertexFit::fit_secondary_vertices_t::init()
{
#ifndef ALLEN_STANDALONE
histogram_nsvs = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "number_of_svs", "NSVs", {VertexFit::max_svs, 0, VertexFit::max_svs}}, {}};
+ {this, "number_of_svs", "NSVs", {VertexFit::max_svs + 1, -0.5, VertexFit::max_svs + .5}}, {}};
#endif
}
@@ -78,7 +78,7 @@ void VertexFit::fit_secondary_vertices_t::operator()(
const Constants&,
const Allen::Context& context) const
{
- auto dev_histogram_nsvs = make_device_buffer<unsigned>(arguments, VertexFit::max_svs);
+ auto dev_histogram_nsvs = make_device_buffer<unsigned>(arguments, VertexFit::max_svs + 1);
Allen::memset_async(dev_histogram_nsvs.data(), 0, dev_histogram_nsvs.size() * sizeof(unsigned), context);
Allen::memset_async<dev_two_track_composite_view_t>(arguments, 0, context);
@@ -91,7 +91,7 @@ void VertexFit::fit_secondary_vertices_t::operator()(
#ifndef ALLEN_STANDALONE
gaudi_monitoring::fill(
- arguments, context, std::tuple {dev_histogram_nsvs.get(), histogram_nsvs, 0u, VertexFit::max_svs});
+ arguments, context, std::tuple {dev_histogram_nsvs.get(), histogram_nsvs, -0.5, VertexFit::max_svs + .5});
#endif
}
@@ -136,7 +136,7 @@ __global__ void VertexFit::fit_secondary_vertices(
const unsigned* event_svs_trk2_idx = parameters.dev_svs_trk2_idx + idx_offset;
const float* event_poca = parameters.dev_sv_poca + 3 * idx_offset;
- if (n_svs < VertexFit::max_svs) atomicAdd(&dev_histogram_nsvs[n_svs], 1);
+ if (n_svs < VertexFit::max_svs + 1) atomicAdd(&dev_histogram_nsvs[n_svs], 1);
// Tracks.
const auto long_track_particles = parameters.dev_long_track_particles->container(event_number);
diff --git a/main/include/GaudiMonitoring.h b/main/include/GaudiMonitoring.h
index 0a9c5d868f54b942414fe1d7414c0f6f0c826291..302b345de5807cc08c904b9ca44e20a20d2ea958 100644
--- a/main/include/GaudiMonitoring.h
+++ b/main/include/GaudiMonitoring.h
@@ -76,7 +76,7 @@ namespace gaudi_monitoring {
histo->call([=](auto& locked_histo) {
for (unsigned bin = 0; bin < n_bins; bin++) {
if (data[bin] != 0) {
- float value = min + bin * bin_size +
+ float value = float(min) + bin * bin_size +
half_bin_size; // use middle of bin for increment to not rely on floating point calculation
locked_histo[value] += data[bin];
}