From 9bf6a327712b50358f514dd5196c7c18a3746415 Mon Sep 17 00:00:00 2001
From: Kate Abigail Richardson <kaaricha@tdeb01.lbdaq.cern.ch>
Date: Mon, 17 Jul 2023 15:07:46 +0200
Subject: [PATCH 01/18] also update buffer sizes
---
.../python/AllenConf/muon_reconstruction.py | 3 +
.../PV/beamlinePV/src/pv_beamline_cleanup.cu | 20 ++-
.../consolidate/include/ConsolidateSciFi.cuh | 1 -
.../include/ConsolidateSciFiSeeding.cuh | 2 +-
.../src/ConsolidateSciFiSeeding.cu | 7 +-
.../clustering/include/CaloFindClusters.cuh | 17 ++-
.../calo/clustering/src/CaloFindClusters.cu | 93 ++++++++++++-
.../ParKalman/src/MakeLongTrackParticles.cu | 32 +++--
device/muon/is_muon/include/IsMuon.cuh | 13 +-
device/muon/is_muon/src/IsMuon.cu | 129 +++++++++++++++++-
.../Hlt1/include/GatherSelections.cuh | 3 +
.../selections/Hlt1/src/GatherSelections.cu | 15 ++
.../lines/calibration/include/D2KPiLine.cuh | 2 +-
.../lines/charm/include/D2KKLine.cuh | 2 +-
.../lines/charm/include/D2PiPiLine.cuh | 2 +-
.../include/ConsolidateMatchedTracks.cuh | 2 +-
16 files changed, 298 insertions(+), 45 deletions(-)
diff --git a/configuration/python/AllenConf/muon_reconstruction.py b/configuration/python/AllenConf/muon_reconstruction.py
index ec2cfc2a3ab..43b3bfbd6f4 100644
--- a/configuration/python/AllenConf/muon_reconstruction.py
+++ b/configuration/python/AllenConf/muon_reconstruction.py
@@ -123,6 +123,7 @@ 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,
@@ -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(
diff --git a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
index 45bdf874459..8cbdda7e402 100644
--- a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
+++ b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
@@ -17,9 +17,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 +26,9 @@ 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_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
}
@@ -50,12 +50,10 @@ void pv_beamline_cleanup::pv_beamline_cleanup_t::operator()(
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_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);
@@ -139,15 +137,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) {
diff --git a/device/SciFi/consolidate/include/ConsolidateSciFi.cuh b/device/SciFi/consolidate/include/ConsolidateSciFi.cuh
index 75163b6d12e..228beab0427 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
diff --git a/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh b/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
index fff2ea021ed..0dc464dcdba 100644
--- a/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
+++ b/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
@@ -152,7 +152,7 @@ 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};
diff --git a/device/SciFi/hybridseeding/consolidate/src/ConsolidateSciFiSeeding.cu b/device/SciFi/hybridseeding/consolidate/src/ConsolidateSciFiSeeding.cu
index 61ef2ffa591..43eb941d9c7 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", {200, 0, 200, {}, {}}}, {}};
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, 200u);
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);
@@ -238,8 +238,7 @@ __global__ void seed_confirmTracks_consolidate::seed_confirmTracks_consolidate(
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);
+ 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 6b4f8df91f4..ce32a2f4285 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,12 @@ 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 +54,15 @@ 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 e39a6c44dfd..4171a33ea1d 100755
--- a/device/calo/clustering/src/CaloFindClusters.cu
+++ b/device/calo/clustering/src/CaloFindClusters.cu
@@ -13,6 +13,19 @@
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", {400, 0, 400}}, {}};
+ 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 +33,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 +53,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);
+ ++histogram_digit_e[bin];
+ }
}
}
cluster.e -= corrections[c];
@@ -53,13 +77,37 @@ __device__ void simple_clusters(
}
cluster.CalcEt();
cluster.CaloNeutralE19 = calo.getE(seed_cluster.id, seed_cluster.adc) / cluster.e;
+
+ // Fill histograms
+ if (cluster.e < 50000) {
+ const unsigned bin = std::floor(cluster.e / 10);
+ ++histogram_cluster_e[bin];
+ }
+ if (cluster.et < 5000) {
+ const unsigned bin = std::floor(cluster.et / 10);
+ ++histogram_cluster_et[bin];
+ }
+ if (cluster.x < 4000 && cluster.x > -4000) {
+ const unsigned bin = std::floor(cluster.x / 10) + 400;
+ ++histogram_cluster_x[bin];
+ }
+ if (cluster.y < 4000 && cluster.y > -4000) {
+ const unsigned bin = std::floor(cluster.y / 10) + 400;
+ ++histogram_cluster_y[bin];
+ }
}
}
__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 +118,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 < 400) histogram_n_clusters[ecal_num_clusters]++;
+
simple_clusters(
parameters.dev_ecal_digits + ecal_digits_offset,
parameters.dev_ecal_seed_clusters + Calo::Constants::ecal_max_index / 8 * event_number,
@@ -77,7 +128,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 +150,35 @@ __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, 400u);
+ 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, 400},
+ 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/kalman/ParKalman/src/MakeLongTrackParticles.cu b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
index f64b55dfc91..e194dcdc09a 100644
--- a/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
+++ b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
@@ -18,10 +18,10 @@ 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}}, {}};
+ 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
}
@@ -44,9 +44,9 @@ void make_long_track_particles::make_long_track_particles_t::operator()(
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_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);
@@ -106,13 +106,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 2519c2bf7dd..ac809216bb8 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;
@@ -34,11 +39,12 @@ namespace is_muon {
};
__global__ void
- is_muon(Parameters, const Muon::Constants::FieldOfInterest* dev_muon_foi, const float* dev_muon_momentum_cuts);
+ 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 +53,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 af0c20a16c2..68e905095e9 100644
--- a/device/muon/is_muon/src/IsMuon.cu
+++ b/device/muon/is_muon/src/IsMuon.cu
@@ -13,6 +13,30 @@
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, 2}},
+ {}};
+ histogram_muon_n_stations =
+ new gaudi_monitoring::Lockable_Histogram<> {{this,
+ "muon_n_stations",
+ "# muon stations",
+ {3, 2, 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 +55,28 @@ 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, 2},
+ std::tuple {dev_histogram_muon_n_stations.get(), histogram_muon_n_stations, 2, 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 +123,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 +167,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]) {
+ ++dev_histogram_n_muons[0];
continue;
}
@@ -174,14 +218,91 @@ __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;
- }
- else if (momentum < dev_muon_momentum_cuts[2]) {
+
+ // monitoring
+ ++dev_histogram_n_muons[1];
+ unsigned n_stations = 0;
+ if (occupancies[2] != 0) {
+ ++n_stations;
+ }
+ if (occupancies[3] != 0) {
+ ++n_stations;
+ }
+ ++dev_histogram_muon_n_stations[n_stations];
+ 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 auto qop = long_track.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);
+ if (pt < unsigned(1e4)) {
+ const unsigned bin = std::floor(pt / 10.f);
+ ++dev_histogram_muon_pt[bin];
+ }
+ } 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)) {
+ ++dev_histogram_n_muons[1];
+ unsigned n_stations = 0;
+ if (occupancies[2] != 0) {
+ ++n_stations;
+ }
+ if (occupancies[3] != 0) {
+ ++n_stations;
+ }
+ ++dev_histogram_muon_n_stations[n_stations];
+ 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 auto qop = long_track.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);
+ if (pt < unsigned(1e4)) {
+ const unsigned bin = std::floor(pt / 10.f);
+ ++dev_histogram_muon_pt[bin];
+ }
+ } else {
+ ++dev_histogram_n_muons[0];
+ }
}
else {
parameters.dev_is_muon[event_offset + track_id] = (occupancies[2] != 0) && (occupancies[3] != 0);
+
+ // monitoring
+ if ((occupancies[2] != 0) && (occupancies[3] != 0)) {
+ ++dev_histogram_n_muons[1];
+ ++dev_histogram_muon_n_stations[2];
+ 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 auto qop = long_track.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);
+ if (pt < unsigned(1e4)) {
+ const unsigned bin = std::floor(pt / 10.f);
+ ++dev_histogram_muon_pt[bin];
+ }
+ } else {
+ ++dev_histogram_n_muons[0];
+ }
}
parameters.dev_lepton_id[event_offset + track_id] = parameters.dev_is_muon[event_offset + track_id];
+ } else {
+ ++dev_histogram_n_muons[0];
}
}
}
diff --git a/device/selections/Hlt1/include/GatherSelections.cuh b/device/selections/Hlt1/include/GatherSelections.cuh
index 2fa0fd5e0ea..50e893f6f6b 100644
--- a/device/selections/Hlt1/include/GatherSelections.cuh
+++ b/device/selections/Hlt1/include/GatherSelections.cuh
@@ -89,6 +89,9 @@ namespace gather_selections {
mutable std::vector<std::unique_ptr<Gaudi::Accumulators::Counter<>>> m_rate_counters;
gaudi_monitoring::Lockable_Histogram<>* histogram_line_passes;
gaudi_monitoring::Lockable_Histogram<>* histogram_line_rates;
+ gaudi_monitoring::Lockable_Histogram<>* histogram_line_rates_normalized;
+ bool rate_ref_line;
+ unsigned rate_ref_line_index;
#endif
};
} // namespace gather_selections
diff --git a/device/selections/Hlt1/src/GatherSelections.cu b/device/selections/Hlt1/src/GatherSelections.cu
index 0b4a292859d..4efbf24a890 100644
--- a/device/selections/Hlt1/src/GatherSelections.cu
+++ b/device/selections/Hlt1/src/GatherSelections.cu
@@ -156,7 +156,12 @@ void gather_selections::gather_selections_t::init()
std::istringstream is(line_names);
std::string line_name;
std::vector<std::string> line_labels;
+ rate_ref_line = false;
while (std::getline(is, line_name, ',')) {
+ if (line_name == "Hlt1ODIN1kHzLumi") {
+ rate_ref_line = true;
+ rate_ref_line_index = line_labels.size();
+ }
const std::string pass_counter_name {line_name + "Pass"};
const std::string rate_counter_name {line_name + "Rate"};
m_pass_counters.push_back(std::make_unique<Gaudi::Accumulators::Counter<>>(this, pass_counter_name));
@@ -169,6 +174,8 @@ void gather_selections::gather_selections_t::init()
{this, "line_passes", "line passes", {unsigned(n_lines), 0, n_lines, {}, line_labels}}, {}};
histogram_line_rates = new gaudi_monitoring::Lockable_Histogram<> {
{this, "line_rates", "line rates", {unsigned(n_lines), 0, n_lines, {}, line_labels}}, {}};
+ histogram_line_rates_normalized = new gaudi_monitoring::Lockable_Histogram<> {
+ {this, "line_rates_normalized", "line rates normalized", {unsigned(n_lines), 0, n_lines, {}, line_labels}}, {}};
#endif
}
@@ -335,19 +342,27 @@ void gather_selections::gather_selections_t::operator()(
// Monitoring
auto host_histo_line_passes = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
auto host_histo_line_rates = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
+ auto host_histo_line_rates_normalized = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
Allen::copy_async(host_histo_line_passes.get(), dev_histo_line_passes.get(), context, Allen::memcpyDeviceToHost);
Allen::copy_async(host_histo_line_rates.get(), dev_histo_line_rates.get(), context, Allen::memcpyDeviceToHost);
+ Allen::memset_async(host_histo_line_rates_normalized.data(), 0, host_histo_line_rates_normalized.size() * sizeof(unsigned), context);
Allen::synchronize(context);
+ float t = std::numeric_limits<float>::max();
+ if (rate_ref_line) t = host_histo_line_rates[rate_ref_line_index]/1000.f;
for (unsigned i = 0; i < first<host_number_of_active_lines_t>(arguments); i++) {
m_pass_counters[i]->buffer() += host_histo_line_passes[i];
m_rate_counters[i]->buffer() += host_histo_line_rates[i];
+ if (rate_ref_line) host_histo_line_rates_normalized[i] = host_histo_line_rates[i] / t;
}
gaudi_monitoring::details::fill_gaudi_histogram(
host_histo_line_passes.get(), histogram_line_passes, 0u, first<host_number_of_active_lines_t>(arguments));
gaudi_monitoring::details::fill_gaudi_histogram(
host_histo_line_rates.get(), histogram_line_rates, 0u, first<host_number_of_active_lines_t>(arguments));
+ gaudi_monitoring::details::fill_gaudi_histogram(
+ host_histo_line_rates_normalized.get(), histogram_line_rates_normalized
+ , 0u, first<host_number_of_active_lines_t>(arguments));
#endif
// Reduce output mask to its proper size
diff --git a/device/selections/lines/calibration/include/D2KPiLine.cuh b/device/selections/lines/calibration/include/D2KPiLine.cuh
index 9d201bd350f..34ab96da75e 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 f4b0753f5e5..c8e5e61d816 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 e5e58ee3150..975734d8bb8 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 cf8ecbd482c..53903e7b07a 100644
--- a/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
+++ b/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
@@ -167,7 +167,7 @@ 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};
--
GitLab
From 8ce84b7da076e2ca7ff3caedf4b49b3ab52be2ac Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Mon, 17 Jul 2023 09:50:04 -0400
Subject: [PATCH 02/18] center int histos
---
device/PV/beamlinePV/src/pv_beamline_cleanup.cu | 9 ++++-----
device/SciFi/consolidate/include/ConsolidateSciFi.cuh | 6 +++---
device/SciFi/consolidate/src/ConsolidateSciFi.cu | 11 ++++-------
.../consolidate/include/ConsolidateSciFiSeeding.cuh | 6 +++---
.../consolidate/src/ConsolidateSciFiSeeding.cu | 8 ++++----
device/kalman/ParKalman/src/MakeLongTrackParticles.cu | 8 ++++----
.../consolidate/include/ConsolidateMatchedTracks.cuh | 6 +++---
.../consolidate/src/ConsolidateMatchedTracks.cu | 11 ++++-------
.../consolidate_tracks/src/VeloConsolidateTracks.cu | 11 ++++-------
.../include/VeloKalmanFilter.cuh | 6 +++---
device/vertex_fit/vertex_fitter/src/VertexFitter.cu | 8 ++++----
11 files changed, 40 insertions(+), 50 deletions(-)
diff --git a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
index 8cbdda7e402..bbd027bb009 100644
--- a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
+++ b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
@@ -16,7 +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_n_smogpvs =
new gaudi_monitoring::Lockable_Histogram<> {{this, "n_smog2_PVs", "n_smog2_PVs", {10, -0.5f, 9.5f}}, {}};
@@ -26,6 +25,7 @@ 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}}, {}};
@@ -48,9 +48,9 @@ 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_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);
@@ -79,8 +79,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},
@@ -161,8 +161,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 228beab0427..3ef18bcacf1 100644
--- a/device/SciFi/consolidate/include/ConsolidateSciFi.cuh
+++ b/device/SciFi/consolidate/include/ConsolidateSciFi.cuh
@@ -182,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 efc74b9334f..fdae0e1d099 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 0dc464dcdba..df17c94a964 100644
--- a/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
+++ b/device/SciFi/hybridseeding/consolidate/include/ConsolidateSciFiSeeding.cuh
@@ -156,9 +156,9 @@ namespace seed_confirmTracks_consolidate {
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 43eb941d9c7..087c1776bbe 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", {200, 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, 200u);
+ 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,7 +237,7 @@ __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) {
+ 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/kalman/ParKalman/src/MakeLongTrackParticles.cu b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
index e194dcdc09a..d7f798eb9b3 100644
--- a/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
+++ b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
@@ -17,7 +17,7 @@ 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}}, {}};
+ {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 =
@@ -43,7 +43,7 @@ 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_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);
@@ -65,7 +65,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 +92,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));
diff --git a/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh b/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
index 53903e7b07a..e599309f5e6 100644
--- a/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
+++ b/device/track_matching/consolidate/include/ConsolidateMatchedTracks.cuh
@@ -171,9 +171,9 @@ namespace matching_consolidate_tracks {
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 37119ec7441..2c092c676cf 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 c55981435e2..468978a82f8 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 ba2057e605c..6d323787adc 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 cf26ecc26d4..e76c0299fc8 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);
--
GitLab
From c48dfe26e258e6e70ad3538fb85a54d570621d63 Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Mon, 17 Jul 2023 10:58:49 -0400
Subject: [PATCH 03/18] update new histos
---
configuration/python/AllenConf/HLT1_cosmics.py | 2 +-
configuration/python/AllenConf/calo_reconstruction.py | 7 ++++---
configuration/python/AllenConf/hlt1_reconstruction.py | 2 +-
configuration/python/AllenConf/muon_reconstruction.py | 9 +++++----
device/calo/clustering/src/CaloFindClusters.cu | 8 ++++----
device/muon/is_muon/src/IsMuon.cu | 8 ++++----
6 files changed, 19 insertions(+), 17 deletions(-)
diff --git a/configuration/python/AllenConf/HLT1_cosmics.py b/configuration/python/AllenConf/HLT1_cosmics.py
index 7711d046475..f4a0c8b1f8a 100644
--- a/configuration/python/AllenConf/HLT1_cosmics.py
+++ b/configuration/python/AllenConf/HLT1_cosmics.py
@@ -79,7 +79,7 @@ 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 b0811c0155f..58af863ed04 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,8 @@ 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 c0b4452ff86..0089d690804 100644
--- a/configuration/python/AllenConf/hlt1_reconstruction.py
+++ b/configuration/python/AllenConf/hlt1_reconstruction.py
@@ -130,7 +130,7 @@ 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 43b3bfbd6f4..a781c5c62d8 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"]
@@ -127,7 +127,7 @@ def is_muon(decoded_muon, long_tracks):
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=
@@ -209,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()
@@ -219,9 +220,9 @@ 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/calo/clustering/src/CaloFindClusters.cu b/device/calo/clustering/src/CaloFindClusters.cu
index 4171a33ea1d..0dfa6feaa26 100755
--- a/device/calo/clustering/src/CaloFindClusters.cu
+++ b/device/calo/clustering/src/CaloFindClusters.cu
@@ -17,7 +17,7 @@ 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", {400, 0, 400}}, {}};
+ {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}}, {}};
@@ -119,7 +119,7 @@ __global__ void calo_find_clusters::calo_find_clusters(
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 < 400) histogram_n_clusters[ecal_num_clusters]++;
+ if (ecal_num_clusters < 401) histogram_n_clusters[ecal_num_clusters]++;
simple_clusters(
parameters.dev_ecal_digits + ecal_digits_offset,
@@ -150,7 +150,7 @@ __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, 400u);
+ 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);
@@ -172,7 +172,7 @@ __host__ void calo_find_clusters::calo_find_clusters_t::operator()(
gaudi_monitoring::fill(
arguments,
context,
- std::tuple {std::tuple {dev_histogram_n_clusters.get(), histogram_n_clusters, 0, 400},
+ 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},
diff --git a/device/muon/is_muon/src/IsMuon.cu b/device/muon/is_muon/src/IsMuon.cu
index 68e905095e9..27989c82c1c 100644
--- a/device/muon/is_muon/src/IsMuon.cu
+++ b/device/muon/is_muon/src/IsMuon.cu
@@ -20,13 +20,13 @@ void is_muon::is_muon_t::init()
new gaudi_monitoring::Lockable_Histogram<> {{this,
"n_muons",
"# muons",
- {2, 0, 2}},
+ {2, -0.5, 1.5}},
{}};
histogram_muon_n_stations =
new gaudi_monitoring::Lockable_Histogram<> {{this,
"muon_n_stations",
"# muon stations",
- {3, 2, 5}},
+ {3, 1.5, 4.5}},
{}};
histogram_muon_pt =
new gaudi_monitoring::Lockable_Histogram<> {{this,
@@ -73,8 +73,8 @@ void is_muon::is_muon_t::operator()(
gaudi_monitoring::fill(
arguments,
context,
- std::tuple {std::tuple {dev_histogram_n_muons.get(), histogram_n_muons, 0, 2},
- std::tuple {dev_histogram_muon_n_stations.get(), histogram_muon_n_stations, 2, 5},
+ 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
}
--
GitLab
From e050f7997d794e0ded60775031f760b7ce522248 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Mon, 17 Jul 2023 15:20:28 +0000
Subject: [PATCH 04/18] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/31097686
---
.../python/AllenConf/HLT1_cosmics.py | 3 +-
.../python/AllenConf/calo_reconstruction.py | 3 +-
.../python/AllenConf/hlt1_reconstruction.py | 3 +-
.../python/AllenConf/muon_reconstruction.py | 5 ++-
.../clustering/include/CaloFindClusters.cuh | 13 ++++--
.../calo/clustering/src/CaloFindClusters.cu | 31 +++++++++-----
.../ParKalman/src/MakeLongTrackParticles.cu | 7 ++--
device/muon/is_muon/include/IsMuon.cuh | 9 +++-
device/muon/is_muon/src/IsMuon.cu | 41 +++++++++----------
.../selections/Hlt1/src/GatherSelections.cu | 14 ++++---
.../vertex_fitter/src/VertexFitter.cu | 6 +--
11 files changed, 82 insertions(+), 53 deletions(-)
diff --git a/configuration/python/AllenConf/HLT1_cosmics.py b/configuration/python/AllenConf/HLT1_cosmics.py
index f4a0c8b1f8a..1d4e0ec3369 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, calo_find_clusters_name='calo_find_clusters_cosmics')
+ 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 58af863ed04..97b50026085 100755
--- a/configuration/python/AllenConf/calo_reconstruction.py
+++ b/configuration/python/AllenConf/calo_reconstruction.py
@@ -276,6 +276,7 @@ def make_ecal_clusters(decoded_calo,
def ecal_cluster_reco(calo_find_clusters_name='calo_find_clusters_reco'):
decoded_calo = decode_calo()
- ecal_clusters = make_ecal_clusters(decoded_calo, calo_find_clusters_name=calo_find_clusters_name)
+ 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 0089d690804..41439ba1309 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, is_muon_name=algorithm_name+'is_muon')
+ 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 a781c5c62d8..e382e9dad84 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, is_muon_name = 'is_muon'):
+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"]
@@ -222,7 +222,8 @@ def muon_id(algorithm_name=''):
scifi_consolidate_tracks_name=algorithm_name +
'scifi_consolidate_tracks_muon_id')
decoded_muon = decode_muon()
- muonID = is_muon(decoded_muon, long_tracks, is_muon_name=algorithm_name + 'is_muon')
+ 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/calo/clustering/include/CaloFindClusters.cuh b/device/calo/clustering/include/CaloFindClusters.cuh
index ce32a2f4285..4208b9276e6 100755
--- a/device/calo/clustering/include/CaloFindClusters.cuh
+++ b/device/calo/clustering/include/CaloFindClusters.cuh
@@ -37,7 +37,16 @@ namespace calo_find_clusters {
};
// Global function
- __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);
+ __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 {
@@ -62,7 +71,5 @@ namespace calo_find_clusters {
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 0dfa6feaa26..0c974281bb2 100755
--- a/device/calo/clustering/src/CaloFindClusters.cu
+++ b/device/calo/clustering/src/CaloFindClusters.cu
@@ -16,13 +16,18 @@ 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}}, {}};
+ 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
}
@@ -166,7 +171,15 @@ __host__ void calo_find_clusters::calo_find_clusters_t::operator()(
// 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(), 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());
+ 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(
@@ -179,6 +192,4 @@ __host__ void calo_find_clusters::calo_find_clusters_t::operator()(
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/kalman/ParKalman/src/MakeLongTrackParticles.cu b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
index d7f798eb9b3..f5c0409f017 100644
--- a/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
+++ b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
@@ -17,7 +17,8 @@ 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+1, -0.5, UT::Constants::max_num_tracks+.5}}, {}};
+ {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 =
@@ -43,7 +44,7 @@ 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+1);
+ 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);
@@ -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.5, UT::Constants::max_num_tracks+.5},
+ 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)}});
diff --git a/device/muon/is_muon/include/IsMuon.cuh b/device/muon/is_muon/include/IsMuon.cuh
index ac809216bb8..107cccbcdfa 100644
--- a/device/muon/is_muon/include/IsMuon.cuh
+++ b/device/muon/is_muon/include/IsMuon.cuh
@@ -38,8 +38,13 @@ 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, gsl::span<unsigned> dev_histogram_n_muons, gsl::span<unsigned> dev_histogram_muon_n_stations, gsl::span<unsigned> dev_histogram_muon_pt);
+ __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;
diff --git a/device/muon/is_muon/src/IsMuon.cu b/device/muon/is_muon/src/IsMuon.cu
index 27989c82c1c..7c1b075db29 100644
--- a/device/muon/is_muon/src/IsMuon.cu
+++ b/device/muon/is_muon/src/IsMuon.cu
@@ -16,24 +16,11 @@ 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_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}},
- {}};
+ 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}},
- {}};
+ new gaudi_monitoring::Lockable_Histogram<> {{this, "muon_pt", "muon pt", {1000, 0, (unsigned) 1e4}}, {}};
#endif
}
@@ -59,7 +46,8 @@ void is_muon::is_muon_t::operator()(
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_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);
@@ -67,7 +55,12 @@ void is_muon::is_muon_t::operator()(
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, dev_histogram_n_muons.get(), dev_histogram_muon_n_stations.get(), dev_histogram_muon_pt.get());
+ 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(
@@ -243,7 +236,8 @@ __global__ void is_muon::is_muon(
const unsigned bin = std::floor(pt / 10.f);
++dev_histogram_muon_pt[bin];
}
- } else if (momentum < dev_muon_momentum_cuts[2]) {
+ }
+ else if (momentum < dev_muon_momentum_cuts[2]) {
parameters.dev_is_muon[event_offset + track_id] = (occupancies[2] != 0) || (occupancies[3] != 0);
// monitoring
@@ -271,7 +265,8 @@ __global__ void is_muon::is_muon(
const unsigned bin = std::floor(pt / 10.f);
++dev_histogram_muon_pt[bin];
}
- } else {
+ }
+ else {
++dev_histogram_n_muons[0];
}
}
@@ -296,12 +291,14 @@ __global__ void is_muon::is_muon(
const unsigned bin = std::floor(pt / 10.f);
++dev_histogram_muon_pt[bin];
}
- } else {
+ }
+ else {
++dev_histogram_n_muons[0];
}
}
parameters.dev_lepton_id[event_offset + track_id] = parameters.dev_is_muon[event_offset + track_id];
- } else {
+ }
+ else {
++dev_histogram_n_muons[0];
}
}
diff --git a/device/selections/Hlt1/src/GatherSelections.cu b/device/selections/Hlt1/src/GatherSelections.cu
index 4efbf24a890..71dc2525533 100644
--- a/device/selections/Hlt1/src/GatherSelections.cu
+++ b/device/selections/Hlt1/src/GatherSelections.cu
@@ -342,14 +342,16 @@ void gather_selections::gather_selections_t::operator()(
// Monitoring
auto host_histo_line_passes = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
auto host_histo_line_rates = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
- auto host_histo_line_rates_normalized = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
+ auto host_histo_line_rates_normalized =
+ make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
Allen::copy_async(host_histo_line_passes.get(), dev_histo_line_passes.get(), context, Allen::memcpyDeviceToHost);
Allen::copy_async(host_histo_line_rates.get(), dev_histo_line_rates.get(), context, Allen::memcpyDeviceToHost);
- Allen::memset_async(host_histo_line_rates_normalized.data(), 0, host_histo_line_rates_normalized.size() * sizeof(unsigned), context);
+ Allen::memset_async(
+ host_histo_line_rates_normalized.data(), 0, host_histo_line_rates_normalized.size() * sizeof(unsigned), context);
Allen::synchronize(context);
float t = std::numeric_limits<float>::max();
- if (rate_ref_line) t = host_histo_line_rates[rate_ref_line_index]/1000.f;
+ if (rate_ref_line) t = host_histo_line_rates[rate_ref_line_index] / 1000.f;
for (unsigned i = 0; i < first<host_number_of_active_lines_t>(arguments); i++) {
m_pass_counters[i]->buffer() += host_histo_line_passes[i];
m_rate_counters[i]->buffer() += host_histo_line_rates[i];
@@ -361,8 +363,10 @@ void gather_selections::gather_selections_t::operator()(
gaudi_monitoring::details::fill_gaudi_histogram(
host_histo_line_rates.get(), histogram_line_rates, 0u, first<host_number_of_active_lines_t>(arguments));
gaudi_monitoring::details::fill_gaudi_histogram(
- host_histo_line_rates_normalized.get(), histogram_line_rates_normalized
- , 0u, first<host_number_of_active_lines_t>(arguments));
+ host_histo_line_rates_normalized.get(),
+ histogram_line_rates_normalized,
+ 0u,
+ first<host_number_of_active_lines_t>(arguments));
#endif
// Reduce output mask to its proper size
diff --git a/device/vertex_fit/vertex_fitter/src/VertexFitter.cu b/device/vertex_fit/vertex_fitter/src/VertexFitter.cu
index e76c0299fc8..83832e3485d 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+1, -0.5, VertexFit::max_svs+.5}}, {}};
+ {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+1);
+ 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, -0.5, VertexFit::max_svs+.5});
+ arguments, context, std::tuple {dev_histogram_nsvs.get(), histogram_nsvs, -0.5, VertexFit::max_svs + .5});
#endif
}
--
GitLab
From f7e215fce478fe3a72de6972f51d5136e244bcb3 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Mon, 31 Jul 2023 10:28:09 +0000
Subject: [PATCH 05/18] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/31422514
---
device/PV/beamlinePV/src/pv_beamline_cleanup.cu | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
index bbd027bb009..886cd0724a6 100644
--- a/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
+++ b/device/PV/beamlinePV/src/pv_beamline_cleanup.cu
@@ -25,7 +25,8 @@ 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_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}}, {}};
--
GitLab
From c2da0655fca04a1a563bb1d6bd10a7437ecec428 Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Tue, 9 Jan 2024 15:54:09 -0500
Subject: [PATCH 06/18] move pt calculation to particletypes
---
device/event_model/common/include/ParticleTypes.cuh | 9 +++++++++
device/muon/is_muon/src/IsMuon.cu | 6 +-----
main/include/GaudiMonitoring.h | 2 +-
3 files changed, 11 insertions(+), 6 deletions(-)
diff --git a/device/event_model/common/include/ParticleTypes.cuh b/device/event_model/common/include/ParticleTypes.cuh
index 772ce8f999e..d0a16c73e98 100644
--- a/device/event_model/common/include/ParticleTypes.cuh
+++ b/device/event_model/common/include/ParticleTypes.cuh
@@ -80,6 +80,15 @@ namespace Allen {
{}
__host__ __device__ float qop() const { return *m_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;
+ }
+
enum struct segment { velo, ut, scifi, muon };
template<segment t>
diff --git a/device/muon/is_muon/src/IsMuon.cu b/device/muon/is_muon/src/IsMuon.cu
index 7c1b075db29..dc72cc20132 100644
--- a/device/muon/is_muon/src/IsMuon.cu
+++ b/device/muon/is_muon/src/IsMuon.cu
@@ -227,11 +227,7 @@ __global__ void is_muon::is_muon(
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 auto qop = long_track.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);
+ const float pt = long_track.pt(velo_state);
if (pt < unsigned(1e4)) {
const unsigned bin = std::floor(pt / 10.f);
++dev_histogram_muon_pt[bin];
diff --git a/main/include/GaudiMonitoring.h b/main/include/GaudiMonitoring.h
index 0a9c5d868f5..302b345de58 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];
}
--
GitLab
From 975d10f6f01fedbc2f489d8ebab88de89fe5a87e Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Tue, 9 Jan 2024 17:21:46 -0500
Subject: [PATCH 07/18] clean up
---
device/calo/clustering/src/CaloFindClusters.cu | 12 ++++++------
device/muon/is_muon/src/IsMuon.cu | 12 ++----------
2 files changed, 8 insertions(+), 16 deletions(-)
diff --git a/device/calo/clustering/src/CaloFindClusters.cu b/device/calo/clustering/src/CaloFindClusters.cu
index 0c974281bb2..fd85d142f5b 100755
--- a/device/calo/clustering/src/CaloFindClusters.cu
+++ b/device/calo/clustering/src/CaloFindClusters.cu
@@ -64,7 +64,7 @@ __device__ void simple_clusters(
if (digit_e < 10000) {
const unsigned bin = std::floor(digit_e / 10);
- ++histogram_digit_e[bin];
+ atomicAdd(&histogram_digit_e[bin], 1);
}
}
}
@@ -86,19 +86,19 @@ __device__ void simple_clusters(
// Fill histograms
if (cluster.e < 50000) {
const unsigned bin = std::floor(cluster.e / 10);
- ++histogram_cluster_e[bin];
+ atomicAdd(&histogram_cluster_e[bin], 1);
}
if (cluster.et < 5000) {
const unsigned bin = std::floor(cluster.et / 10);
- ++histogram_cluster_et[bin];
+ atomicAdd(&histogram_cluster_et[bin], 1);
}
if (cluster.x < 4000 && cluster.x > -4000) {
const unsigned bin = std::floor(cluster.x / 10) + 400;
- ++histogram_cluster_x[bin];
+ atomicAdd(&histogram_cluster_x[bin], 1);
}
if (cluster.y < 4000 && cluster.y > -4000) {
const unsigned bin = std::floor(cluster.y / 10) + 400;
- ++histogram_cluster_y[bin];
+ atomicAdd(&histogram_cluster_y[bin], 1);
}
}
}
@@ -124,7 +124,7 @@ __global__ void calo_find_clusters::calo_find_clusters(
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) histogram_n_clusters[ecal_num_clusters]++;
+ if (ecal_num_clusters < 401) atomicAdd(&histogram_n_clusters[ecal_num_clusters], 1);
simple_clusters(
parameters.dev_ecal_digits + ecal_digits_offset,
diff --git a/device/muon/is_muon/src/IsMuon.cu b/device/muon/is_muon/src/IsMuon.cu
index dc72cc20132..285990c0fc9 100644
--- a/device/muon/is_muon/src/IsMuon.cu
+++ b/device/muon/is_muon/src/IsMuon.cu
@@ -252,11 +252,7 @@ __global__ void is_muon::is_muon(
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 auto qop = long_track.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);
+ const float pt = long_track.pt(velo_state);
if (pt < unsigned(1e4)) {
const unsigned bin = std::floor(pt / 10.f);
++dev_histogram_muon_pt[bin];
@@ -278,11 +274,7 @@ __global__ void is_muon::is_muon(
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 auto qop = long_track.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);
+ const float pt = long_track.pt(velo_state);
if (pt < unsigned(1e4)) {
const unsigned bin = std::floor(pt / 10.f);
++dev_histogram_muon_pt[bin];
--
GitLab
From 4b2448dae9058124948ecb34a5be7fc25ffd7567 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Tue, 9 Jan 2024 22:22:30 +0000
Subject: [PATCH 08/18] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/35136853
---
device/event_model/common/include/ParticleTypes.cuh | 3 ++-
device/vertex_fit/vertex_fitter/src/VertexFitter.cu | 2 +-
2 files changed, 3 insertions(+), 2 deletions(-)
diff --git a/device/event_model/common/include/ParticleTypes.cuh b/device/event_model/common/include/ParticleTypes.cuh
index d0a16c73e98..859025adb15 100644
--- a/device/event_model/common/include/ParticleTypes.cuh
+++ b/device/event_model/common/include/ParticleTypes.cuh
@@ -80,7 +80,8 @@ namespace Allen {
{}
__host__ __device__ float qop() const { return *m_qop; }
- __host__ __device__ float pt(Allen::Views::Physics::KalmanState velo_state) const {
+ __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();
diff --git a/device/vertex_fit/vertex_fitter/src/VertexFitter.cu b/device/vertex_fit/vertex_fitter/src/VertexFitter.cu
index 83832e3485d..c815910a39f 100644
--- a/device/vertex_fit/vertex_fitter/src/VertexFitter.cu
+++ b/device/vertex_fit/vertex_fitter/src/VertexFitter.cu
@@ -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+1) 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);
--
GitLab
From 6f14350472b782073066bf8e692d917cf9081c65 Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Wed, 10 Jan 2024 12:53:17 -0500
Subject: [PATCH 09/18] move pt to long track
---
.../common/include/ParticleTypes.cuh | 20 +++++++++----------
1 file changed, 10 insertions(+), 10 deletions(-)
diff --git a/device/event_model/common/include/ParticleTypes.cuh b/device/event_model/common/include/ParticleTypes.cuh
index 859025adb15..a5f2caff094 100644
--- a/device/event_model/common/include/ParticleTypes.cuh
+++ b/device/event_model/common/include/ParticleTypes.cuh
@@ -80,16 +80,6 @@ namespace Allen {
{}
__host__ __device__ float qop() const { return *m_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;
- }
-
enum struct segment { velo, ut, scifi, muon };
template<segment t>
@@ -269,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> {
--
GitLab
From d5d56c1b827763bf0d541e716f760fb1a4fe0c02 Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Wed, 10 Jan 2024 13:16:28 -0500
Subject: [PATCH 10/18] use atomicadds in downstreamconsolidate
---
.../consolidate/src/DownstreamConsolidate.cu | 17 +++++++----------
1 file changed, 7 insertions(+), 10 deletions(-)
diff --git a/device/downstream/consolidate/src/DownstreamConsolidate.cu b/device/downstream/consolidate/src/DownstreamConsolidate.cu
index ccf30afeeea..e7ffba15944 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);
}
}
--
GitLab
From 04447f8648cded43d992ceb72e482f8538b8fa25 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Wed, 10 Jan 2024 18:17:04 +0000
Subject: [PATCH 11/18] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/35163486
---
.../consolidate/src/DownstreamConsolidate.cu | 474 +++++++++---------
1 file changed, 239 insertions(+), 235 deletions(-)
diff --git a/device/downstream/consolidate/src/DownstreamConsolidate.cu b/device/downstream/consolidate/src/DownstreamConsolidate.cu
index e7ffba15944..5309218e08c 100644
--- a/device/downstream/consolidate/src/DownstreamConsolidate.cu
+++ b/device/downstream/consolidate/src/DownstreamConsolidate.cu
@@ -152,252 +152,256 @@ 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<> {
+ histogram_n_downstream_tracks = new gaudi_monitoring::Lockable_Histogram<>
+ {
{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",
- "#eta",
- {property<histogram_downstream_track_eta_nbins_t>(),
- property<histogram_downstream_track_eta_min_t>(),
- property<histogram_downstream_track_eta_max_t>()}},
- {}};
- histogram_downstream_track_phi =
- new gaudi_monitoring::Lockable_Histogram<> {{this,
- "downstream_track_phi",
- "#phi",
- {property<histogram_downstream_track_phi_nbins_t>(),
- property<histogram_downstream_track_phi_min_t>(),
- property<histogram_downstream_track_phi_max_t>()}},
- {}};
- histogram_downstream_track_nhits =
- new gaudi_monitoring::Lockable_Histogram<> {{this,
- "downstream_track_nhits",
- "N. hits / track",
- {property<histogram_downstream_track_nhits_nbins_t>(),
- property<histogram_downstream_track_nhits_min_t>(),
- property<histogram_downstream_track_nhits_max_t>()}},
- {}};
+ histogram_downstream_track_eta =
+ new gaudi_monitoring::Lockable_Histogram<> {{this,
+ "downstream_track_eta",
+ "#eta",
+ {property<histogram_downstream_track_eta_nbins_t>(),
+ property<histogram_downstream_track_eta_min_t>(),
+ property<histogram_downstream_track_eta_max_t>()}},
+ {}};
+ histogram_downstream_track_phi =
+ new gaudi_monitoring::Lockable_Histogram<> {{this,
+ "downstream_track_phi",
+ "#phi",
+ {property<histogram_downstream_track_phi_nbins_t>(),
+ property<histogram_downstream_track_phi_min_t>(),
+ property<histogram_downstream_track_phi_max_t>()}},
+ {}};
+ histogram_downstream_track_nhits =
+ new gaudi_monitoring::Lockable_Histogram<> {{this,
+ "downstream_track_nhits",
+ "N. hits / track",
+ {property<histogram_downstream_track_nhits_nbins_t>(),
+ property<histogram_downstream_track_nhits_min_t>(),
+ property<histogram_downstream_track_nhits_max_t>()}},
+ {}};
#endif
-}
-
-void downstream_consolidate::downstream_consolidate_t::operator()(
- const ArgumentReferences<Parameters>& arguments,
- const RuntimeOptions&,
- const Constants& constants,
- const Allen::Context& context) const
-{
- // Initialize container to avoid invalid std::function destructor
- Allen::memset_async<dev_multi_event_downstream_tracks_view_t>(arguments, 0, context);
- Allen::memset_async<dev_downstream_tracks_view_t>(arguments, 0, context);
-
- //
- // Create monitoring container
- //
- auto dev_histogram_downstream_track_eta =
- make_device_buffer<unsigned>(arguments, property<histogram_downstream_track_eta_nbins_t>());
- auto dev_histogram_downstream_track_phi =
- 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, 201u);
- auto dev_n_downstream_tracks_counter = make_device_buffer<unsigned>(arguments, 1u);
+ }
- //
- // Initialize monitoring container
- //
- Allen::memset_async(
- dev_histogram_downstream_track_eta.data(),
- 0,
- dev_histogram_downstream_track_eta.size() * sizeof(unsigned),
- context);
- Allen::memset_async(
- dev_histogram_downstream_track_phi.data(),
- 0,
- dev_histogram_downstream_track_phi.size() * sizeof(unsigned),
- context);
- Allen::memset_async(
- dev_histogram_downstream_track_nhits.data(),
- 0,
- dev_histogram_downstream_track_nhits.size() * sizeof(unsigned),
- context);
- Allen::memset_async(
- dev_histogram_n_downstream_tracks.data(), 0, dev_histogram_n_downstream_tracks.size() * sizeof(unsigned), context);
- Allen::memset_async(
- dev_n_downstream_tracks_counter.data(), 0, dev_n_downstream_tracks_counter.size() * sizeof(unsigned), context);
-
- // Fill the consolidation memory
- global_function(downstream_consolidate)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
- arguments,
- constants.dev_unique_x_sector_layer_offsets.data(),
- dev_histogram_n_downstream_tracks.get(),
- dev_n_downstream_tracks_counter.get());
-
- // Create views
- global_function(downstream_create_tracks_view)(dim3(first<host_number_of_events_t>(arguments)), 256, context)(
- arguments,
- dev_histogram_downstream_track_eta.get(),
- dev_histogram_downstream_track_phi.get(),
- dev_histogram_downstream_track_nhits.get());
+ void downstream_consolidate::downstream_consolidate_t::operator()(
+ const ArgumentReferences<Parameters>& arguments,
+ const RuntimeOptions&,
+ const Constants& constants,
+ const Allen::Context& context) const
+ {
+ // Initialize container to avoid invalid std::function destructor
+ Allen::memset_async<dev_multi_event_downstream_tracks_view_t>(arguments, 0, context);
+ Allen::memset_async<dev_downstream_tracks_view_t>(arguments, 0, context);
+
+ //
+ // Create monitoring container
+ //
+ auto dev_histogram_downstream_track_eta =
+ make_device_buffer<unsigned>(arguments, property<histogram_downstream_track_eta_nbins_t>());
+ auto dev_histogram_downstream_track_phi =
+ 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, 201u);
+ auto dev_n_downstream_tracks_counter = make_device_buffer<unsigned>(arguments, 1u);
+
+ //
+ // Initialize monitoring container
+ //
+ Allen::memset_async(
+ dev_histogram_downstream_track_eta.data(),
+ 0,
+ dev_histogram_downstream_track_eta.size() * sizeof(unsigned),
+ context);
+ Allen::memset_async(
+ dev_histogram_downstream_track_phi.data(),
+ 0,
+ dev_histogram_downstream_track_phi.size() * sizeof(unsigned),
+ context);
+ Allen::memset_async(
+ dev_histogram_downstream_track_nhits.data(),
+ 0,
+ dev_histogram_downstream_track_nhits.size() * sizeof(unsigned),
+ context);
+ Allen::memset_async(
+ dev_histogram_n_downstream_tracks.data(),
+ 0,
+ dev_histogram_n_downstream_tracks.size() * sizeof(unsigned),
+ context);
+ Allen::memset_async(
+ dev_n_downstream_tracks_counter.data(), 0, dev_n_downstream_tracks_counter.size() * sizeof(unsigned), context);
+
+ // Fill the consolidation memory
+ global_function(downstream_consolidate)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
+ arguments,
+ constants.dev_unique_x_sector_layer_offsets.data(),
+ dev_histogram_n_downstream_tracks.get(),
+ dev_n_downstream_tracks_counter.get());
+
+ // Create views
+ global_function(downstream_create_tracks_view)(dim3(first<host_number_of_events_t>(arguments)), 256, context)(
+ arguments,
+ dev_histogram_downstream_track_eta.get(),
+ dev_histogram_downstream_track_phi.get(),
+ dev_histogram_downstream_track_nhits.get());
#ifndef ALLEN_STANDALONE
- gaudi_monitoring::fill(
- arguments,
- context,
- std::tuple {std::tuple {dev_histogram_downstream_track_eta.get(),
- histogram_downstream_track_eta,
- property<histogram_downstream_track_eta_min_t>(),
- property<histogram_downstream_track_eta_max_t>()},
- std::tuple {dev_histogram_downstream_track_phi.get(),
- histogram_downstream_track_phi,
- property<histogram_downstream_track_phi_min_t>(),
- property<histogram_downstream_track_phi_max_t>()},
- std::tuple {dev_histogram_downstream_track_nhits.get(),
- 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.5f, 200.5f},
- std::tuple {dev_n_downstream_tracks_counter.get(), m_downstream_tracks}});
+ gaudi_monitoring::fill(
+ arguments,
+ context,
+ std::tuple {std::tuple {dev_histogram_downstream_track_eta.get(),
+ histogram_downstream_track_eta,
+ property<histogram_downstream_track_eta_min_t>(),
+ property<histogram_downstream_track_eta_max_t>()},
+ std::tuple {dev_histogram_downstream_track_phi.get(),
+ histogram_downstream_track_phi,
+ property<histogram_downstream_track_phi_min_t>(),
+ property<histogram_downstream_track_phi_max_t>()},
+ std::tuple {dev_histogram_downstream_track_nhits.get(),
+ 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.5f, 200.5f},
+ std::tuple {dev_n_downstream_tracks_counter.get(), m_downstream_tracks}});
#endif
-}
-
-__global__ void downstream_consolidate::downstream_consolidate(
- downstream_consolidate::Parameters parameters,
- const unsigned* dev_unique_x_sector_layer_offsets,
- // Monitoring
- gsl::span<unsigned> dev_histogram_n_downstream_tracks,
- gsl::span<unsigned> dev_n_downstream_tracks_counter)
-{
- // Basic
- const unsigned event_number = parameters.dev_event_list[blockIdx.x];
- const unsigned number_of_events = parameters.dev_number_of_events[0];
-
- // Tracks
- const auto downstream_tracks_memory =
- parameters.dev_downstream_tracks + event_number * UT::DownstreamTracks::TotalMemorySize;
- UT::DownstreamTracks_Const downstream_tracks {downstream_tracks_memory};
-
- // Event offsets for tracks
- const auto downstream_tracks_offset = parameters.dev_offsets_downstream_tracks[event_number];
- const auto downstream_tracks_size =
- parameters.dev_offsets_downstream_tracks[event_number + 1] - downstream_tracks_offset;
-
- // Event offsets for hit ofsets
- const auto downstream_hit_number_offsets = parameters.dev_offsets_downstream_hit_numbers + downstream_tracks_offset;
-
- // Total numbers
- const auto downstream_total_number_of_tracks = parameters.dev_offsets_downstream_tracks[number_of_events];
- const auto downstream_total_number_of_hits =
- parameters.dev_offsets_downstream_hit_numbers[downstream_total_number_of_tracks];
-
- // UT hits
- const unsigned number_of_unique_x_sectors = dev_unique_x_sector_layer_offsets[UT::Constants::n_layers];
- const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * number_of_unique_x_sectors];
- const UT::HitOffsets ut_hit_offsets {
- parameters.dev_ut_hit_offsets, event_number, number_of_unique_x_sectors, dev_unique_x_sector_layer_offsets};
- const auto event_hit_offset = ut_hit_offsets.event_offset();
- UT::ConstHits ut_hits {parameters.dev_ut_hits, total_number_of_hits, event_hit_offset};
-
- //
- // Monitoring fill
- //
- 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
- const auto downstream_track_scifi_indices = parameters.dev_downstream_track_scifi_idx + downstream_tracks_offset;
- const auto downstream_track_qops = parameters.dev_downstream_track_qops + downstream_tracks_offset;
-
- UT::Consolidated::Hits downstream_track_hits(
- parameters.dev_downstream_track_hits, 0, downstream_total_number_of_hits);
-
- UT::Consolidated::Tracks output_tracks(
- parameters.dev_offsets_downstream_tracks,
- parameters.dev_offsets_downstream_hit_numbers,
- event_number,
- number_of_events);
-
- Velo::Consolidated::States output_states(
- parameters.dev_downstream_track_states, downstream_total_number_of_tracks, downstream_tracks_offset);
-
- // Fill states
- for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
-
- // First AOS part of state
- output_states.x(track_idx) = downstream_tracks.x(track_idx);
- output_states.y(track_idx) = downstream_tracks.y(track_idx);
- output_states.z(track_idx) = UT::Constants::zMidUT;
- output_states.tx(track_idx) = downstream_tracks.tx(track_idx);
- output_states.ty(track_idx) = downstream_tracks.ty(track_idx);
- output_states.qop(track_idx) = downstream_tracks.qop(track_idx);
-
- // Fill extra qop for downstream track
- downstream_track_qops[track_idx] = downstream_tracks.qop(track_idx);
-
- // Second AOS part of state
- output_states.c00(track_idx) = 0.f;
- output_states.c20(track_idx) = 0.f;
- output_states.c22(track_idx) = 0.f;
- output_states.c11(track_idx) = 0.f;
- output_states.c31(track_idx) = 0.f;
- output_states.c33(track_idx) = 0.f;
- output_states.chi2(track_idx) = downstream_tracks.chi2(track_idx);
- output_states.ndof(track_idx) = downstream_tracks.n_hits(track_idx) - 1u;
}
- // Scifi idx
- for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
- downstream_track_scifi_indices[track_idx] = downstream_tracks.scifi(track_idx);
- }
+ __global__ void downstream_consolidate::downstream_consolidate(
+ downstream_consolidate::Parameters parameters,
+ const unsigned* dev_unique_x_sector_layer_offsets,
+ // Monitoring
+ gsl::span<unsigned> dev_histogram_n_downstream_tracks,
+ gsl::span<unsigned> dev_n_downstream_tracks_counter)
+ {
+ // Basic
+ const unsigned event_number = parameters.dev_event_list[blockIdx.x];
+ const unsigned number_of_events = parameters.dev_number_of_events[0];
+
+ // Tracks
+ const auto downstream_tracks_memory =
+ parameters.dev_downstream_tracks + event_number * UT::DownstreamTracks::TotalMemorySize;
+ UT::DownstreamTracks_Const downstream_tracks {downstream_tracks_memory};
+
+ // Event offsets for tracks
+ const auto downstream_tracks_offset = parameters.dev_offsets_downstream_tracks[event_number];
+ const auto downstream_tracks_size =
+ parameters.dev_offsets_downstream_tracks[event_number + 1] - downstream_tracks_offset;
+
+ // Event offsets for hit ofsets
+ const auto downstream_hit_number_offsets = parameters.dev_offsets_downstream_hit_numbers + downstream_tracks_offset;
+
+ // Total numbers
+ const auto downstream_total_number_of_tracks = parameters.dev_offsets_downstream_tracks[number_of_events];
+ const auto downstream_total_number_of_hits =
+ parameters.dev_offsets_downstream_hit_numbers[downstream_total_number_of_tracks];
+
+ // UT hits
+ const unsigned number_of_unique_x_sectors = dev_unique_x_sector_layer_offsets[UT::Constants::n_layers];
+ const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * number_of_unique_x_sectors];
+ const UT::HitOffsets ut_hit_offsets {
+ parameters.dev_ut_hit_offsets, event_number, number_of_unique_x_sectors, dev_unique_x_sector_layer_offsets};
+ const auto event_hit_offset = ut_hit_offsets.event_offset();
+ UT::ConstHits ut_hits {parameters.dev_ut_hits, total_number_of_hits, event_hit_offset};
+
+ //
+ // Monitoring fill
+ //
+ 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
+ const auto downstream_track_scifi_indices = parameters.dev_downstream_track_scifi_idx + downstream_tracks_offset;
+ const auto downstream_track_qops = parameters.dev_downstream_track_qops + downstream_tracks_offset;
+
+ UT::Consolidated::Hits downstream_track_hits(
+ parameters.dev_downstream_track_hits, 0, downstream_total_number_of_hits);
+
+ UT::Consolidated::Tracks output_tracks(
+ parameters.dev_offsets_downstream_tracks,
+ parameters.dev_offsets_downstream_hit_numbers,
+ event_number,
+ number_of_events);
+
+ Velo::Consolidated::States output_states(
+ parameters.dev_downstream_track_states, downstream_total_number_of_tracks, downstream_tracks_offset);
+
+ // Fill states
+ for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
+
+ // First AOS part of state
+ output_states.x(track_idx) = downstream_tracks.x(track_idx);
+ output_states.y(track_idx) = downstream_tracks.y(track_idx);
+ output_states.z(track_idx) = UT::Constants::zMidUT;
+ output_states.tx(track_idx) = downstream_tracks.tx(track_idx);
+ output_states.ty(track_idx) = downstream_tracks.ty(track_idx);
+ output_states.qop(track_idx) = downstream_tracks.qop(track_idx);
+
+ // Fill extra qop for downstream track
+ downstream_track_qops[track_idx] = downstream_tracks.qop(track_idx);
+
+ // Second AOS part of state
+ output_states.c00(track_idx) = 0.f;
+ output_states.c20(track_idx) = 0.f;
+ output_states.c22(track_idx) = 0.f;
+ output_states.c11(track_idx) = 0.f;
+ output_states.c31(track_idx) = 0.f;
+ output_states.c33(track_idx) = 0.f;
+ output_states.chi2(track_idx) = downstream_tracks.chi2(track_idx);
+ output_states.ndof(track_idx) = downstream_tracks.n_hits(track_idx) - 1u;
+ }
- // Fill hits
- for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
- const auto nhits = downstream_tracks.n_hits(track_idx);
- const auto target_hit_offset = downstream_hit_number_offsets[track_idx];
- for (unsigned hit_idx = 0; hit_idx < nhits; hit_idx++) {
- downstream_track_hits.set(
- target_hit_offset + hit_idx, ut_hits.getHit(downstream_tracks.hits(track_idx, hit_idx)));
+ // Scifi idx
+ for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
+ downstream_track_scifi_indices[track_idx] = downstream_tracks.scifi(track_idx);
}
- }
- __syncthreads();
-}
+ // Fill hits
+ for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
+ const auto nhits = downstream_tracks.n_hits(track_idx);
+ const auto target_hit_offset = downstream_hit_number_offsets[track_idx];
+ for (unsigned hit_idx = 0; hit_idx < nhits; hit_idx++) {
+ downstream_track_hits.set(
+ target_hit_offset + hit_idx, ut_hits.getHit(downstream_tracks.hits(track_idx, hit_idx)));
+ }
+ }
-__device__ void downstream_consolidate::downstream_consolidate_t::monitor(
- const downstream_consolidate::Parameters& parameters,
- const Allen::Views::Physics::DownstreamTrack downstream_track,
- const Allen::Views::Physics::KalmanState downstream_state,
- gsl::span<unsigned> dev_histogram_downstream_track_eta,
- gsl::span<unsigned> dev_histogram_downstream_track_phi,
- gsl::span<unsigned> dev_histogram_downstream_track_nhits)
-{
- const auto nhits = downstream_track.number_of_hits();
- const auto tx = downstream_state.tx();
- const auto ty = downstream_state.ty();
- const auto slope2 = tx * tx + ty * ty;
- const auto rho = std::sqrt(slope2);
- const auto eta = eta_from_rho(rho);
- const auto phi = std::atan2(ty, tx);
-
- // Filling histograms
- if (eta > parameters.histogram_downstream_track_eta_min && eta < parameters.histogram_downstream_track_eta_max) {
- 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[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[bin], 1);
+ __syncthreads();
}
- if (
- nhits > parameters.histogram_downstream_track_nhits_min &&
- nhits < parameters.histogram_downstream_track_nhits_max) {
- 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[bin], 1);
+
+ __device__ void downstream_consolidate::downstream_consolidate_t::monitor(
+ const downstream_consolidate::Parameters& parameters,
+ const Allen::Views::Physics::DownstreamTrack downstream_track,
+ const Allen::Views::Physics::KalmanState downstream_state,
+ gsl::span<unsigned> dev_histogram_downstream_track_eta,
+ gsl::span<unsigned> dev_histogram_downstream_track_phi,
+ gsl::span<unsigned> dev_histogram_downstream_track_nhits)
+ {
+ const auto nhits = downstream_track.number_of_hits();
+ const auto tx = downstream_state.tx();
+ const auto ty = downstream_state.ty();
+ const auto slope2 = tx * tx + ty * ty;
+ const auto rho = std::sqrt(slope2);
+ const auto eta = eta_from_rho(rho);
+ const auto phi = std::atan2(ty, tx);
+
+ // Filling histograms
+ if (eta > parameters.histogram_downstream_track_eta_min && eta < parameters.histogram_downstream_track_eta_max) {
+ 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[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[bin], 1);
+ }
+ if (
+ nhits > parameters.histogram_downstream_track_nhits_min &&
+ nhits < parameters.histogram_downstream_track_nhits_max) {
+ 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[bin], 1);
+ }
}
-}
--
GitLab
From 2868cef9e415b7aa2a1e08cfc493b0f73d0f154e Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Thu, 11 Jan 2024 13:26:15 -0500
Subject: [PATCH 12/18] fix } issue
---
device/downstream/consolidate/src/DownstreamConsolidate.cu | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/device/downstream/consolidate/src/DownstreamConsolidate.cu b/device/downstream/consolidate/src/DownstreamConsolidate.cu
index 5309218e08c..70e66be2422 100644
--- a/device/downstream/consolidate/src/DownstreamConsolidate.cu
+++ b/device/downstream/consolidate/src/DownstreamConsolidate.cu
@@ -153,8 +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", {201, -0.5, 200.5}, {}};
+ {{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",
@@ -180,7 +179,7 @@ void downstream_consolidate::downstream_consolidate_t::init()
property<histogram_downstream_track_nhits_max_t>()}},
{}};
#endif
- }
+}
void downstream_consolidate::downstream_consolidate_t::operator()(
const ArgumentReferences<Parameters>& arguments,
--
GitLab
From 11b9a7b675cbf922e54be64bbf2c5d56d96980b6 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Thu, 11 Jan 2024 18:27:14 +0000
Subject: [PATCH 13/18] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/35191291
---
.../consolidate/src/DownstreamConsolidate.cu | 473 +++++++++---------
1 file changed, 235 insertions(+), 238 deletions(-)
diff --git a/device/downstream/consolidate/src/DownstreamConsolidate.cu b/device/downstream/consolidate/src/DownstreamConsolidate.cu
index 70e66be2422..044736d4cd4 100644
--- a/device/downstream/consolidate/src/DownstreamConsolidate.cu
+++ b/device/downstream/consolidate/src/DownstreamConsolidate.cu
@@ -152,255 +152,252 @@ 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", {201, -0.5, 200.5}}, {}};
- histogram_downstream_track_eta =
- new gaudi_monitoring::Lockable_Histogram<> {{this,
- "downstream_track_eta",
- "#eta",
- {property<histogram_downstream_track_eta_nbins_t>(),
- property<histogram_downstream_track_eta_min_t>(),
- property<histogram_downstream_track_eta_max_t>()}},
- {}};
- histogram_downstream_track_phi =
- new gaudi_monitoring::Lockable_Histogram<> {{this,
- "downstream_track_phi",
- "#phi",
- {property<histogram_downstream_track_phi_nbins_t>(),
- property<histogram_downstream_track_phi_min_t>(),
- property<histogram_downstream_track_phi_max_t>()}},
- {}};
- histogram_downstream_track_nhits =
- new gaudi_monitoring::Lockable_Histogram<> {{this,
- "downstream_track_nhits",
- "N. hits / track",
- {property<histogram_downstream_track_nhits_nbins_t>(),
- property<histogram_downstream_track_nhits_min_t>(),
- property<histogram_downstream_track_nhits_max_t>()}},
- {}};
+ histogram_n_downstream_tracks = new gaudi_monitoring::Lockable_Histogram<> {
+ {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",
+ "#eta",
+ {property<histogram_downstream_track_eta_nbins_t>(),
+ property<histogram_downstream_track_eta_min_t>(),
+ property<histogram_downstream_track_eta_max_t>()}},
+ {}};
+ histogram_downstream_track_phi =
+ new gaudi_monitoring::Lockable_Histogram<> {{this,
+ "downstream_track_phi",
+ "#phi",
+ {property<histogram_downstream_track_phi_nbins_t>(),
+ property<histogram_downstream_track_phi_min_t>(),
+ property<histogram_downstream_track_phi_max_t>()}},
+ {}};
+ histogram_downstream_track_nhits =
+ new gaudi_monitoring::Lockable_Histogram<> {{this,
+ "downstream_track_nhits",
+ "N. hits / track",
+ {property<histogram_downstream_track_nhits_nbins_t>(),
+ property<histogram_downstream_track_nhits_min_t>(),
+ property<histogram_downstream_track_nhits_max_t>()}},
+ {}};
#endif
}
- void downstream_consolidate::downstream_consolidate_t::operator()(
- const ArgumentReferences<Parameters>& arguments,
- const RuntimeOptions&,
- const Constants& constants,
- const Allen::Context& context) const
- {
- // Initialize container to avoid invalid std::function destructor
- Allen::memset_async<dev_multi_event_downstream_tracks_view_t>(arguments, 0, context);
- Allen::memset_async<dev_downstream_tracks_view_t>(arguments, 0, context);
-
- //
- // Create monitoring container
- //
- auto dev_histogram_downstream_track_eta =
- make_device_buffer<unsigned>(arguments, property<histogram_downstream_track_eta_nbins_t>());
- auto dev_histogram_downstream_track_phi =
- 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, 201u);
- auto dev_n_downstream_tracks_counter = make_device_buffer<unsigned>(arguments, 1u);
-
- //
- // Initialize monitoring container
- //
- Allen::memset_async(
- dev_histogram_downstream_track_eta.data(),
- 0,
- dev_histogram_downstream_track_eta.size() * sizeof(unsigned),
- context);
- Allen::memset_async(
- dev_histogram_downstream_track_phi.data(),
- 0,
- dev_histogram_downstream_track_phi.size() * sizeof(unsigned),
- context);
- Allen::memset_async(
- dev_histogram_downstream_track_nhits.data(),
- 0,
- dev_histogram_downstream_track_nhits.size() * sizeof(unsigned),
- context);
- Allen::memset_async(
- dev_histogram_n_downstream_tracks.data(),
- 0,
- dev_histogram_n_downstream_tracks.size() * sizeof(unsigned),
- context);
- Allen::memset_async(
- dev_n_downstream_tracks_counter.data(), 0, dev_n_downstream_tracks_counter.size() * sizeof(unsigned), context);
-
- // Fill the consolidation memory
- global_function(downstream_consolidate)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
- arguments,
- constants.dev_unique_x_sector_layer_offsets.data(),
- dev_histogram_n_downstream_tracks.get(),
- dev_n_downstream_tracks_counter.get());
-
- // Create views
- global_function(downstream_create_tracks_view)(dim3(first<host_number_of_events_t>(arguments)), 256, context)(
- arguments,
- dev_histogram_downstream_track_eta.get(),
- dev_histogram_downstream_track_phi.get(),
- dev_histogram_downstream_track_nhits.get());
+void downstream_consolidate::downstream_consolidate_t::operator()(
+ const ArgumentReferences<Parameters>& arguments,
+ const RuntimeOptions&,
+ const Constants& constants,
+ const Allen::Context& context) const
+{
+ // Initialize container to avoid invalid std::function destructor
+ Allen::memset_async<dev_multi_event_downstream_tracks_view_t>(arguments, 0, context);
+ Allen::memset_async<dev_downstream_tracks_view_t>(arguments, 0, context);
+
+ //
+ // Create monitoring container
+ //
+ auto dev_histogram_downstream_track_eta =
+ make_device_buffer<unsigned>(arguments, property<histogram_downstream_track_eta_nbins_t>());
+ auto dev_histogram_downstream_track_phi =
+ 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, 201u);
+ auto dev_n_downstream_tracks_counter = make_device_buffer<unsigned>(arguments, 1u);
+
+ //
+ // Initialize monitoring container
+ //
+ Allen::memset_async(
+ dev_histogram_downstream_track_eta.data(),
+ 0,
+ dev_histogram_downstream_track_eta.size() * sizeof(unsigned),
+ context);
+ Allen::memset_async(
+ dev_histogram_downstream_track_phi.data(),
+ 0,
+ dev_histogram_downstream_track_phi.size() * sizeof(unsigned),
+ context);
+ Allen::memset_async(
+ dev_histogram_downstream_track_nhits.data(),
+ 0,
+ dev_histogram_downstream_track_nhits.size() * sizeof(unsigned),
+ context);
+ Allen::memset_async(
+ dev_histogram_n_downstream_tracks.data(), 0, dev_histogram_n_downstream_tracks.size() * sizeof(unsigned), context);
+ Allen::memset_async(
+ dev_n_downstream_tracks_counter.data(), 0, dev_n_downstream_tracks_counter.size() * sizeof(unsigned), context);
+
+ // Fill the consolidation memory
+ global_function(downstream_consolidate)(dim3(size<dev_event_list_t>(arguments)), property<block_dim_t>(), context)(
+ arguments,
+ constants.dev_unique_x_sector_layer_offsets.data(),
+ dev_histogram_n_downstream_tracks.get(),
+ dev_n_downstream_tracks_counter.get());
+
+ // Create views
+ global_function(downstream_create_tracks_view)(dim3(first<host_number_of_events_t>(arguments)), 256, context)(
+ arguments,
+ dev_histogram_downstream_track_eta.get(),
+ dev_histogram_downstream_track_phi.get(),
+ dev_histogram_downstream_track_nhits.get());
#ifndef ALLEN_STANDALONE
- gaudi_monitoring::fill(
- arguments,
- context,
- std::tuple {std::tuple {dev_histogram_downstream_track_eta.get(),
- histogram_downstream_track_eta,
- property<histogram_downstream_track_eta_min_t>(),
- property<histogram_downstream_track_eta_max_t>()},
- std::tuple {dev_histogram_downstream_track_phi.get(),
- histogram_downstream_track_phi,
- property<histogram_downstream_track_phi_min_t>(),
- property<histogram_downstream_track_phi_max_t>()},
- std::tuple {dev_histogram_downstream_track_nhits.get(),
- 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.5f, 200.5f},
- std::tuple {dev_n_downstream_tracks_counter.get(), m_downstream_tracks}});
+ gaudi_monitoring::fill(
+ arguments,
+ context,
+ std::tuple {std::tuple {dev_histogram_downstream_track_eta.get(),
+ histogram_downstream_track_eta,
+ property<histogram_downstream_track_eta_min_t>(),
+ property<histogram_downstream_track_eta_max_t>()},
+ std::tuple {dev_histogram_downstream_track_phi.get(),
+ histogram_downstream_track_phi,
+ property<histogram_downstream_track_phi_min_t>(),
+ property<histogram_downstream_track_phi_max_t>()},
+ std::tuple {dev_histogram_downstream_track_nhits.get(),
+ 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.5f, 200.5f},
+ std::tuple {dev_n_downstream_tracks_counter.get(), m_downstream_tracks}});
#endif
- }
+}
- __global__ void downstream_consolidate::downstream_consolidate(
- downstream_consolidate::Parameters parameters,
- const unsigned* dev_unique_x_sector_layer_offsets,
- // Monitoring
- gsl::span<unsigned> dev_histogram_n_downstream_tracks,
- gsl::span<unsigned> dev_n_downstream_tracks_counter)
- {
- // Basic
- const unsigned event_number = parameters.dev_event_list[blockIdx.x];
- const unsigned number_of_events = parameters.dev_number_of_events[0];
-
- // Tracks
- const auto downstream_tracks_memory =
- parameters.dev_downstream_tracks + event_number * UT::DownstreamTracks::TotalMemorySize;
- UT::DownstreamTracks_Const downstream_tracks {downstream_tracks_memory};
-
- // Event offsets for tracks
- const auto downstream_tracks_offset = parameters.dev_offsets_downstream_tracks[event_number];
- const auto downstream_tracks_size =
- parameters.dev_offsets_downstream_tracks[event_number + 1] - downstream_tracks_offset;
-
- // Event offsets for hit ofsets
- const auto downstream_hit_number_offsets = parameters.dev_offsets_downstream_hit_numbers + downstream_tracks_offset;
-
- // Total numbers
- const auto downstream_total_number_of_tracks = parameters.dev_offsets_downstream_tracks[number_of_events];
- const auto downstream_total_number_of_hits =
- parameters.dev_offsets_downstream_hit_numbers[downstream_total_number_of_tracks];
-
- // UT hits
- const unsigned number_of_unique_x_sectors = dev_unique_x_sector_layer_offsets[UT::Constants::n_layers];
- const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * number_of_unique_x_sectors];
- const UT::HitOffsets ut_hit_offsets {
- parameters.dev_ut_hit_offsets, event_number, number_of_unique_x_sectors, dev_unique_x_sector_layer_offsets};
- const auto event_hit_offset = ut_hit_offsets.event_offset();
- UT::ConstHits ut_hits {parameters.dev_ut_hits, total_number_of_hits, event_hit_offset};
-
- //
- // Monitoring fill
- //
- 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
- const auto downstream_track_scifi_indices = parameters.dev_downstream_track_scifi_idx + downstream_tracks_offset;
- const auto downstream_track_qops = parameters.dev_downstream_track_qops + downstream_tracks_offset;
-
- UT::Consolidated::Hits downstream_track_hits(
- parameters.dev_downstream_track_hits, 0, downstream_total_number_of_hits);
-
- UT::Consolidated::Tracks output_tracks(
- parameters.dev_offsets_downstream_tracks,
- parameters.dev_offsets_downstream_hit_numbers,
- event_number,
- number_of_events);
-
- Velo::Consolidated::States output_states(
- parameters.dev_downstream_track_states, downstream_total_number_of_tracks, downstream_tracks_offset);
-
- // Fill states
- for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
-
- // First AOS part of state
- output_states.x(track_idx) = downstream_tracks.x(track_idx);
- output_states.y(track_idx) = downstream_tracks.y(track_idx);
- output_states.z(track_idx) = UT::Constants::zMidUT;
- output_states.tx(track_idx) = downstream_tracks.tx(track_idx);
- output_states.ty(track_idx) = downstream_tracks.ty(track_idx);
- output_states.qop(track_idx) = downstream_tracks.qop(track_idx);
-
- // Fill extra qop for downstream track
- downstream_track_qops[track_idx] = downstream_tracks.qop(track_idx);
-
- // Second AOS part of state
- output_states.c00(track_idx) = 0.f;
- output_states.c20(track_idx) = 0.f;
- output_states.c22(track_idx) = 0.f;
- output_states.c11(track_idx) = 0.f;
- output_states.c31(track_idx) = 0.f;
- output_states.c33(track_idx) = 0.f;
- output_states.chi2(track_idx) = downstream_tracks.chi2(track_idx);
- output_states.ndof(track_idx) = downstream_tracks.n_hits(track_idx) - 1u;
- }
+__global__ void downstream_consolidate::downstream_consolidate(
+ downstream_consolidate::Parameters parameters,
+ const unsigned* dev_unique_x_sector_layer_offsets,
+ // Monitoring
+ gsl::span<unsigned> dev_histogram_n_downstream_tracks,
+ gsl::span<unsigned> dev_n_downstream_tracks_counter)
+{
+ // Basic
+ const unsigned event_number = parameters.dev_event_list[blockIdx.x];
+ const unsigned number_of_events = parameters.dev_number_of_events[0];
- // Scifi idx
- for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
- downstream_track_scifi_indices[track_idx] = downstream_tracks.scifi(track_idx);
- }
+ // Tracks
+ const auto downstream_tracks_memory =
+ parameters.dev_downstream_tracks + event_number * UT::DownstreamTracks::TotalMemorySize;
+ UT::DownstreamTracks_Const downstream_tracks {downstream_tracks_memory};
- // Fill hits
- for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
- const auto nhits = downstream_tracks.n_hits(track_idx);
- const auto target_hit_offset = downstream_hit_number_offsets[track_idx];
- for (unsigned hit_idx = 0; hit_idx < nhits; hit_idx++) {
- downstream_track_hits.set(
- target_hit_offset + hit_idx, ut_hits.getHit(downstream_tracks.hits(track_idx, hit_idx)));
- }
- }
+ // Event offsets for tracks
+ const auto downstream_tracks_offset = parameters.dev_offsets_downstream_tracks[event_number];
+ const auto downstream_tracks_size =
+ parameters.dev_offsets_downstream_tracks[event_number + 1] - downstream_tracks_offset;
+
+ // Event offsets for hit ofsets
+ const auto downstream_hit_number_offsets = parameters.dev_offsets_downstream_hit_numbers + downstream_tracks_offset;
- __syncthreads();
+ // Total numbers
+ const auto downstream_total_number_of_tracks = parameters.dev_offsets_downstream_tracks[number_of_events];
+ const auto downstream_total_number_of_hits =
+ parameters.dev_offsets_downstream_hit_numbers[downstream_total_number_of_tracks];
+
+ // UT hits
+ const unsigned number_of_unique_x_sectors = dev_unique_x_sector_layer_offsets[UT::Constants::n_layers];
+ const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * number_of_unique_x_sectors];
+ const UT::HitOffsets ut_hit_offsets {
+ parameters.dev_ut_hit_offsets, event_number, number_of_unique_x_sectors, dev_unique_x_sector_layer_offsets};
+ const auto event_hit_offset = ut_hit_offsets.event_offset();
+ UT::ConstHits ut_hits {parameters.dev_ut_hits, total_number_of_hits, event_hit_offset};
+
+ //
+ // Monitoring fill
+ //
+ 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
+ const auto downstream_track_scifi_indices = parameters.dev_downstream_track_scifi_idx + downstream_tracks_offset;
+ const auto downstream_track_qops = parameters.dev_downstream_track_qops + downstream_tracks_offset;
+
+ UT::Consolidated::Hits downstream_track_hits(
+ parameters.dev_downstream_track_hits, 0, downstream_total_number_of_hits);
+
+ UT::Consolidated::Tracks output_tracks(
+ parameters.dev_offsets_downstream_tracks,
+ parameters.dev_offsets_downstream_hit_numbers,
+ event_number,
+ number_of_events);
+
+ Velo::Consolidated::States output_states(
+ parameters.dev_downstream_track_states, downstream_total_number_of_tracks, downstream_tracks_offset);
+
+ // Fill states
+ for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
+
+ // First AOS part of state
+ output_states.x(track_idx) = downstream_tracks.x(track_idx);
+ output_states.y(track_idx) = downstream_tracks.y(track_idx);
+ output_states.z(track_idx) = UT::Constants::zMidUT;
+ output_states.tx(track_idx) = downstream_tracks.tx(track_idx);
+ output_states.ty(track_idx) = downstream_tracks.ty(track_idx);
+ output_states.qop(track_idx) = downstream_tracks.qop(track_idx);
+
+ // Fill extra qop for downstream track
+ downstream_track_qops[track_idx] = downstream_tracks.qop(track_idx);
+
+ // Second AOS part of state
+ output_states.c00(track_idx) = 0.f;
+ output_states.c20(track_idx) = 0.f;
+ output_states.c22(track_idx) = 0.f;
+ output_states.c11(track_idx) = 0.f;
+ output_states.c31(track_idx) = 0.f;
+ output_states.c33(track_idx) = 0.f;
+ output_states.chi2(track_idx) = downstream_tracks.chi2(track_idx);
+ output_states.ndof(track_idx) = downstream_tracks.n_hits(track_idx) - 1u;
}
- __device__ void downstream_consolidate::downstream_consolidate_t::monitor(
- const downstream_consolidate::Parameters& parameters,
- const Allen::Views::Physics::DownstreamTrack downstream_track,
- const Allen::Views::Physics::KalmanState downstream_state,
- gsl::span<unsigned> dev_histogram_downstream_track_eta,
- gsl::span<unsigned> dev_histogram_downstream_track_phi,
- gsl::span<unsigned> dev_histogram_downstream_track_nhits)
- {
- const auto nhits = downstream_track.number_of_hits();
- const auto tx = downstream_state.tx();
- const auto ty = downstream_state.ty();
- const auto slope2 = tx * tx + ty * ty;
- const auto rho = std::sqrt(slope2);
- const auto eta = eta_from_rho(rho);
- const auto phi = std::atan2(ty, tx);
-
- // Filling histograms
- if (eta > parameters.histogram_downstream_track_eta_min && eta < parameters.histogram_downstream_track_eta_max) {
- 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[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[bin], 1);
- }
- if (
- nhits > parameters.histogram_downstream_track_nhits_min &&
- nhits < parameters.histogram_downstream_track_nhits_max) {
- 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[bin], 1);
+ // Scifi idx
+ for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
+ downstream_track_scifi_indices[track_idx] = downstream_tracks.scifi(track_idx);
+ }
+
+ // Fill hits
+ for (unsigned track_idx = threadIdx.x; track_idx < downstream_tracks_size; track_idx += blockDim.x) {
+ const auto nhits = downstream_tracks.n_hits(track_idx);
+ const auto target_hit_offset = downstream_hit_number_offsets[track_idx];
+ for (unsigned hit_idx = 0; hit_idx < nhits; hit_idx++) {
+ downstream_track_hits.set(
+ target_hit_offset + hit_idx, ut_hits.getHit(downstream_tracks.hits(track_idx, hit_idx)));
}
}
+
+ __syncthreads();
+}
+
+__device__ void downstream_consolidate::downstream_consolidate_t::monitor(
+ const downstream_consolidate::Parameters& parameters,
+ const Allen::Views::Physics::DownstreamTrack downstream_track,
+ const Allen::Views::Physics::KalmanState downstream_state,
+ gsl::span<unsigned> dev_histogram_downstream_track_eta,
+ gsl::span<unsigned> dev_histogram_downstream_track_phi,
+ gsl::span<unsigned> dev_histogram_downstream_track_nhits)
+{
+ const auto nhits = downstream_track.number_of_hits();
+ const auto tx = downstream_state.tx();
+ const auto ty = downstream_state.ty();
+ const auto slope2 = tx * tx + ty * ty;
+ const auto rho = std::sqrt(slope2);
+ const auto eta = eta_from_rho(rho);
+ const auto phi = std::atan2(ty, tx);
+
+ // Filling histograms
+ if (eta > parameters.histogram_downstream_track_eta_min && eta < parameters.histogram_downstream_track_eta_max) {
+ 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[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[bin], 1);
+ }
+ if (
+ nhits > parameters.histogram_downstream_track_nhits_min &&
+ nhits < parameters.histogram_downstream_track_nhits_max) {
+ 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[bin], 1);
+ }
+}
--
GitLab
From b21754f4b43accdb5b574eeeb64b6f9ebb97689c Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Tue, 16 Jan 2024 18:36:30 -0500
Subject: [PATCH 14/18] atomicadds in ismuon
---
device/muon/is_muon/src/IsMuon.cu | 26 +++++++++++++-------------
1 file changed, 13 insertions(+), 13 deletions(-)
diff --git a/device/muon/is_muon/src/IsMuon.cu b/device/muon/is_muon/src/IsMuon.cu
index 285990c0fc9..6c4dc055743 100644
--- a/device/muon/is_muon/src/IsMuon.cu
+++ b/device/muon/is_muon/src/IsMuon.cu
@@ -160,7 +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]) {
- ++dev_histogram_n_muons[0];
+ atomicAdd(&dev_histogram_n_muons[0], 1);
continue;
}
@@ -213,7 +213,7 @@ __global__ void is_muon::is_muon(
parameters.dev_is_muon[event_offset + track_id] = true;
// monitoring
- ++dev_histogram_n_muons[1];
+ atomicAdd(&dev_histogram_n_muons[1], 1);
unsigned n_stations = 0;
if (occupancies[2] != 0) {
++n_stations;
@@ -221,7 +221,7 @@ __global__ void is_muon::is_muon(
if (occupancies[3] != 0) {
++n_stations;
}
- ++dev_histogram_muon_n_stations[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();
@@ -230,7 +230,7 @@ __global__ void is_muon::is_muon(
const float pt = long_track.pt(velo_state);
if (pt < unsigned(1e4)) {
const unsigned bin = std::floor(pt / 10.f);
- ++dev_histogram_muon_pt[bin];
+ atomicAdd(&dev_histogram_muon_pt[bin], 1);
}
}
else if (momentum < dev_muon_momentum_cuts[2]) {
@@ -238,7 +238,7 @@ __global__ void is_muon::is_muon(
// monitoring
if ((occupancies[2] != 0) || (occupancies[3] != 0)) {
- ++dev_histogram_n_muons[1];
+ atomicAdd(&dev_histogram_n_muons[1], 1);
unsigned n_stations = 0;
if (occupancies[2] != 0) {
++n_stations;
@@ -246,7 +246,7 @@ __global__ void is_muon::is_muon(
if (occupancies[3] != 0) {
++n_stations;
}
- ++dev_histogram_muon_n_stations[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();
@@ -255,11 +255,11 @@ __global__ void is_muon::is_muon(
const float pt = long_track.pt(velo_state);
if (pt < unsigned(1e4)) {
const unsigned bin = std::floor(pt / 10.f);
- ++dev_histogram_muon_pt[bin];
+ atomicAdd(&dev_histogram_muon_pt[bin], 1);
}
}
else {
- ++dev_histogram_n_muons[0];
+ atomicAdd(&dev_histogram_n_muons[0], 1);
}
}
else {
@@ -267,8 +267,8 @@ __global__ void is_muon::is_muon(
// monitoring
if ((occupancies[2] != 0) && (occupancies[3] != 0)) {
- ++dev_histogram_n_muons[1];
- ++dev_histogram_muon_n_stations[2];
+ 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();
@@ -277,17 +277,17 @@ __global__ void is_muon::is_muon(
const float pt = long_track.pt(velo_state);
if (pt < unsigned(1e4)) {
const unsigned bin = std::floor(pt / 10.f);
- ++dev_histogram_muon_pt[bin];
+ atomicAdd(&dev_histogram_muon_pt[bin], 1);
}
}
else {
- ++dev_histogram_n_muons[0];
+ atomicAdd(&dev_histogram_n_muons[0], 1);
}
}
parameters.dev_lepton_id[event_offset + track_id] = parameters.dev_is_muon[event_offset + track_id];
}
else {
- ++dev_histogram_n_muons[0];
+ atomicAdd(&dev_histogram_n_muons[0], 1);
}
}
}
--
GitLab
From 2e1de4013ab822e73085c9c9e816648bdf513e17 Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Mon, 22 Jan 2024 12:19:02 -0500
Subject: [PATCH 15/18] remove rate histogram
---
.../Hlt1/include/GatherSelections.cuh | 3 ---
.../selections/Hlt1/src/GatherSelections.cu | 19 -------------------
2 files changed, 22 deletions(-)
diff --git a/device/selections/Hlt1/include/GatherSelections.cuh b/device/selections/Hlt1/include/GatherSelections.cuh
index 50e893f6f6b..2fa0fd5e0ea 100644
--- a/device/selections/Hlt1/include/GatherSelections.cuh
+++ b/device/selections/Hlt1/include/GatherSelections.cuh
@@ -89,9 +89,6 @@ namespace gather_selections {
mutable std::vector<std::unique_ptr<Gaudi::Accumulators::Counter<>>> m_rate_counters;
gaudi_monitoring::Lockable_Histogram<>* histogram_line_passes;
gaudi_monitoring::Lockable_Histogram<>* histogram_line_rates;
- gaudi_monitoring::Lockable_Histogram<>* histogram_line_rates_normalized;
- bool rate_ref_line;
- unsigned rate_ref_line_index;
#endif
};
} // namespace gather_selections
diff --git a/device/selections/Hlt1/src/GatherSelections.cu b/device/selections/Hlt1/src/GatherSelections.cu
index 71dc2525533..0b4a292859d 100644
--- a/device/selections/Hlt1/src/GatherSelections.cu
+++ b/device/selections/Hlt1/src/GatherSelections.cu
@@ -156,12 +156,7 @@ void gather_selections::gather_selections_t::init()
std::istringstream is(line_names);
std::string line_name;
std::vector<std::string> line_labels;
- rate_ref_line = false;
while (std::getline(is, line_name, ',')) {
- if (line_name == "Hlt1ODIN1kHzLumi") {
- rate_ref_line = true;
- rate_ref_line_index = line_labels.size();
- }
const std::string pass_counter_name {line_name + "Pass"};
const std::string rate_counter_name {line_name + "Rate"};
m_pass_counters.push_back(std::make_unique<Gaudi::Accumulators::Counter<>>(this, pass_counter_name));
@@ -174,8 +169,6 @@ void gather_selections::gather_selections_t::init()
{this, "line_passes", "line passes", {unsigned(n_lines), 0, n_lines, {}, line_labels}}, {}};
histogram_line_rates = new gaudi_monitoring::Lockable_Histogram<> {
{this, "line_rates", "line rates", {unsigned(n_lines), 0, n_lines, {}, line_labels}}, {}};
- histogram_line_rates_normalized = new gaudi_monitoring::Lockable_Histogram<> {
- {this, "line_rates_normalized", "line rates normalized", {unsigned(n_lines), 0, n_lines, {}, line_labels}}, {}};
#endif
}
@@ -342,31 +335,19 @@ void gather_selections::gather_selections_t::operator()(
// Monitoring
auto host_histo_line_passes = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
auto host_histo_line_rates = make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
- auto host_histo_line_rates_normalized =
- make_host_buffer<unsigned>(arguments, first<host_number_of_active_lines_t>(arguments));
Allen::copy_async(host_histo_line_passes.get(), dev_histo_line_passes.get(), context, Allen::memcpyDeviceToHost);
Allen::copy_async(host_histo_line_rates.get(), dev_histo_line_rates.get(), context, Allen::memcpyDeviceToHost);
- Allen::memset_async(
- host_histo_line_rates_normalized.data(), 0, host_histo_line_rates_normalized.size() * sizeof(unsigned), context);
Allen::synchronize(context);
- float t = std::numeric_limits<float>::max();
- if (rate_ref_line) t = host_histo_line_rates[rate_ref_line_index] / 1000.f;
for (unsigned i = 0; i < first<host_number_of_active_lines_t>(arguments); i++) {
m_pass_counters[i]->buffer() += host_histo_line_passes[i];
m_rate_counters[i]->buffer() += host_histo_line_rates[i];
- if (rate_ref_line) host_histo_line_rates_normalized[i] = host_histo_line_rates[i] / t;
}
gaudi_monitoring::details::fill_gaudi_histogram(
host_histo_line_passes.get(), histogram_line_passes, 0u, first<host_number_of_active_lines_t>(arguments));
gaudi_monitoring::details::fill_gaudi_histogram(
host_histo_line_rates.get(), histogram_line_rates, 0u, first<host_number_of_active_lines_t>(arguments));
- gaudi_monitoring::details::fill_gaudi_histogram(
- host_histo_line_rates_normalized.get(),
- histogram_line_rates_normalized,
- 0u,
- first<host_number_of_active_lines_t>(arguments));
#endif
// Reduce output mask to its proper size
--
GitLab
From d81d1850e302603bbd399e9348f710207266edab Mon Sep 17 00:00:00 2001
From: kaaricha <kate.abigail.richardson@cern.ch>
Date: Fri, 1 Mar 2024 20:47:09 -0500
Subject: [PATCH 16/18] fix failing test
---
device/calo/clustering/src/CaloFindClusters.cu | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/device/calo/clustering/src/CaloFindClusters.cu b/device/calo/clustering/src/CaloFindClusters.cu
index fd85d142f5b..2ceff91a96f 100755
--- a/device/calo/clustering/src/CaloFindClusters.cu
+++ b/device/calo/clustering/src/CaloFindClusters.cu
@@ -84,11 +84,11 @@ __device__ void simple_clusters(
cluster.CaloNeutralE19 = calo.getE(seed_cluster.id, seed_cluster.adc) / cluster.e;
// Fill histograms
- if (cluster.e < 50000) {
+ if (0 < cluster.e && cluster.e < 50000) {
const unsigned bin = std::floor(cluster.e / 10);
atomicAdd(&histogram_cluster_e[bin], 1);
}
- if (cluster.et < 5000) {
+ if (0< cluster.et && cluster.et < 5000) {
const unsigned bin = std::floor(cluster.et / 10);
atomicAdd(&histogram_cluster_et[bin], 1);
}
--
GitLab
From 483d058ab5a94e934693e430c7fe4724064920c7 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Sat, 2 Mar 2024 01:47:49 +0000
Subject: [PATCH 17/18] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/36645534
---
device/calo/clustering/src/CaloFindClusters.cu | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/device/calo/clustering/src/CaloFindClusters.cu b/device/calo/clustering/src/CaloFindClusters.cu
index 2ceff91a96f..32f57bbec4f 100755
--- a/device/calo/clustering/src/CaloFindClusters.cu
+++ b/device/calo/clustering/src/CaloFindClusters.cu
@@ -88,7 +88,7 @@ __device__ void simple_clusters(
const unsigned bin = std::floor(cluster.e / 10);
atomicAdd(&histogram_cluster_e[bin], 1);
}
- if (0< cluster.et && cluster.et < 5000) {
+ if (0 < cluster.et && cluster.et < 5000) {
const unsigned bin = std::floor(cluster.et / 10);
atomicAdd(&histogram_cluster_et[bin], 1);
}
--
GitLab
From d072be47e8e5a7cebe0ce6400cdf631c2703324e Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Sat, 2 Mar 2024 21:19:45 +0000
Subject: [PATCH 18/18] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/36655454
---
device/kalman/ParKalman/src/MakeLongTrackParticles.cu | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/device/kalman/ParKalman/src/MakeLongTrackParticles.cu b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
index f5c0409f017..add01e2e21c 100644
--- a/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
+++ b/device/kalman/ParKalman/src/MakeLongTrackParticles.cu
@@ -93,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+1) 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));
--
GitLab