From 8ec767962146e1c68f9c8a14e8bdd6e504e48cd4 Mon Sep 17 00:00:00 2001 From: Maarten Van Veghel Date: Sun, 7 Sep 2025 22:51:35 +0200 Subject: [PATCH 1/4] add gamma corrections to ecal clusters --- .../BinaryDumpers/src/DumpCaloGeometry.cpp | 13 ++++++++--- .../calo/clustering/src/CaloFindClusters.cu | 1 + .../event_model/calo/include/CaloGeometry.cuh | 23 +++++++++++++++---- 3 files changed, 30 insertions(+), 7 deletions(-) diff --git a/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp b/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp index 6db865594a7..7e5ad6881f3 100644 --- a/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp +++ b/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp @@ -54,7 +54,7 @@ namespace Dumpers { // start at this index. Using the num we can further index these 32 channels. // Wasted space: 32 * 16 bits per missing card code - const unsigned geom_version = det.nSourceIDs() == 0 ? 3 : 4; // version 3 for run2, version 4 for run3 + const unsigned geom_version = det.nSourceIDs() == 0 ? 3 : 5; // version 3 for run2, version 4 or higher for run3 const auto [cards_or_febs, feb_indices] = [&]() -> std::array, 2> { if (geom_version == 3) { // version 3 for run 2 @@ -66,7 +66,7 @@ namespace Dumpers { } return {cards, std::vector {}}; } - else { // version 4 for run 3 + else { // version 4 or higher for run 3 using MapType = std::map>; MapType map = det.getSourceIDsMap(); std::vector vec_febs(750, 0); @@ -159,6 +159,7 @@ namespace Dumpers { std::vector neighbors(indexSize * max_neighbors, USHRT_MAX); std::vector xy(indexSize * 2, 0.f); std::vector gain(indexSize, 0.f); + std::vector gamma(indexSize, 0.f); // Create neighbours per cellID. for (auto const& param : det.cellParams()) { auto const caloIndex = LHCb::Detector::Calo::Index {param.cellID()}; @@ -177,6 +178,7 @@ namespace Dumpers { xy[idx * 2] = param.x(); xy[idx * 2 + 1] = param.y(); gain[idx] = det.cellGain(param.cellID()); + gamma[idx] = det.getGamma(param.cellID()); } // Get toLocalMatrix @@ -243,13 +245,18 @@ namespace Dumpers { output.write(static_cast(digits_ranges.size())); output.write(digits_ranges); - if (geom_version == 4) { + if (geom_version >= 4) { output.write(static_cast(cards_or_febs.size())); output.write(cards_or_febs); output.write(static_cast(feb_indices.size())); output.write(feb_indices); } + if (geom_version == 5) { + output.write(static_cast(gamma.size())); + output.write(gamma); + } + data = output.buffer(); } }; diff --git a/device/calo/clustering/src/CaloFindClusters.cu b/device/calo/clustering/src/CaloFindClusters.cu index b124004fe98..6820376cfcc 100755 --- a/device/calo/clustering/src/CaloFindClusters.cu +++ b/device/calo/clustering/src/CaloFindClusters.cu @@ -57,6 +57,7 @@ __device__ void simple_clusters( } } cluster.e -= corrections[c]; + cluster.e += calo.getGamma(seed_cluster.id); for (uint16_t n = 0; n < Calo::Constants::max_neighbours; n++) { auto const n_id = cluster.digits[n]; diff --git a/device/event_model/calo/include/CaloGeometry.cuh b/device/event_model/calo/include/CaloGeometry.cuh index 5e2c615c2d6..49f9b46a595 100644 --- a/device/event_model/calo/include/CaloGeometry.cuh +++ b/device/event_model/calo/include/CaloGeometry.cuh @@ -35,6 +35,7 @@ struct CaloGeometry { int* vec_febs = nullptr; uint32_t vec_febIndices_size = 0; int* vec_febIndices = nullptr; + float* gamma = nullptr; __device__ __host__ CaloGeometry(const char* raw_geometry) { @@ -79,7 +80,7 @@ struct CaloGeometry { p += sizeof(uint32_t); // Skip digits_ranges_size digits_ranges = (uint32_t*) p; - if (geom_version == 4) { + if (geom_version >= 4) { p += sizeof(float) * digits_ranges_size; // Skip digits_ranges vec_febs_size = *((uint32_t*) p); p += sizeof(uint32_t); // Skip vec_febs_size @@ -88,19 +89,26 @@ struct CaloGeometry { vec_febIndices_size = *((uint32_t*) p); p += sizeof(uint32_t); // Skip vec_febIndices_size vec_febIndices = (int*) p; - // p += sizeof(int) * vec_febIndices_size; // Skip vec_febs + } + + if (geom_version == 5) { + p += sizeof(int) * vec_febIndices_size; // Skip vec_febs + // const uint32_t gamma_size = *((uint32_t*) p); + p += sizeof(uint32_t); // Skip gamma_size + gamma = (float*) p; + // p += sizeof(float) * gamma_size; // Skip gamma } } __device__ __host__ inline int getFEB(uint32_t source_id, int nFeb) const { - assert(geom_version == 4 && "getFEB is only available for decoding_version 4 or higher"); + assert(geom_version >= 4 && "getFEB is only available for decoding_version 4 or higher"); return vec_febs[3 * (source_id & 0x7ff) + nFeb]; } __device__ __host__ inline int getFEBindex(uint32_t source_id, int nFeb) const { - assert(geom_version == 4 && "getFEBIndex is only available for decoding_version 4 or higher"); + assert(geom_version >= 4 && "getFEBIndex is only available for decoding_version 4 or higher"); return vec_febIndices[3 * (source_id & 0x7ff) + nFeb]; } @@ -117,6 +125,13 @@ struct CaloGeometry { // Convert ADC to energy __device__ __host__ inline float getE(uint16_t cellid, int16_t adc) const { return gain[cellid] * (adc - pedestal); } + // Get the 'gamma' correction for clusters + __device__ __host__ inline float getGamma(uint16_t cellid) const + { + assert(geom_version >= 5 && "getGamma is only available for decoding_version 5 or higher"); + return gamma[cellid]; + } + // Intercept track with calo plane, where 0 is front, 1 is showermax, 2 is back __device__ __host__ inline float getZFromTrackToCaloplaneIntersection(MiniState state, int plane) const { -- GitLab From 6992f40842cc5fcb8770cebb3b32c47d08e52036 Mon Sep 17 00:00:00 2001 From: Maarten Van Veghel Date: Mon, 8 Sep 2025 10:15:04 +0200 Subject: [PATCH 2/4] set gamma to zero for early geometries --- device/event_model/calo/include/CaloGeometry.cuh | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/device/event_model/calo/include/CaloGeometry.cuh b/device/event_model/calo/include/CaloGeometry.cuh index 49f9b46a595..a23e9c2034b 100644 --- a/device/event_model/calo/include/CaloGeometry.cuh +++ b/device/event_model/calo/include/CaloGeometry.cuh @@ -126,11 +126,7 @@ struct CaloGeometry { __device__ __host__ inline float getE(uint16_t cellid, int16_t adc) const { return gain[cellid] * (adc - pedestal); } // Get the 'gamma' correction for clusters - __device__ __host__ inline float getGamma(uint16_t cellid) const - { - assert(geom_version >= 5 && "getGamma is only available for decoding_version 5 or higher"); - return gamma[cellid]; - } + __device__ __host__ inline float getGamma(uint16_t cellid) const { return gamma ? gamma[cellid] : 0.f; } // Intercept track with calo plane, where 0 is front, 1 is showermax, 2 is back __device__ __host__ inline float getZFromTrackToCaloplaneIntersection(MiniState state, int plane) const -- GitLab From 0c0e039a0071af6b25e6253a713ffb75f551cb95 Mon Sep 17 00:00:00 2001 From: Maarten Van Veghel Date: Mon, 8 Sep 2025 13:52:53 +0200 Subject: [PATCH 3/4] fix geometry version checking --- device/calo/decoding/src/CaloDecode.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/device/calo/decoding/src/CaloDecode.cu b/device/calo/decoding/src/CaloDecode.cu index e3f07b5f3af..f0e973f747c 100644 --- a/device/calo/decoding/src/CaloDecode.cu +++ b/device/calo/decoding/src/CaloDecode.cu @@ -278,7 +278,7 @@ void calo_decode::calo_decode_t::operator()( if (bank_version <= 3 && geom_version <= 3) { return true; } - else if ((bank_version == 4 || bank_version == 5) && geom_version == 4) { + else if ((bank_version == 4 || bank_version == 5) && (geom_version == 4 || geom_version == 5)) { return true; } return false; -- GitLab From e696b8614d53cb46756b18f4e698558204c6e913 Mon Sep 17 00:00:00 2001 From: Arthur Hennequin Date: Mon, 8 Sep 2025 19:01:48 +0200 Subject: [PATCH 4/4] Fix geometry version test when determining min and max card codes --- Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp b/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp index 7e5ad6881f3..7f2fb64b7fd 100644 --- a/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp +++ b/Dumpers/BinaryDumpers/src/DumpCaloGeometry.cpp @@ -108,7 +108,7 @@ namespace Dumpers { int max = 0; size_t max_channels = 0; for (int card : cards_or_febs) { - if (geom_version == 4 && card == 0) continue; + if (geom_version >= 4 && card == 0) continue; const auto curCode = geom_version == 3 ? det.cardCode(card) : det.getFEBindex(card); // get FEB numbers from TELL40Link map min = std::min(curCode, min); -- GitLab