From be2e3eea7c9a8aed7833148c4df94e7510cac285 Mon Sep 17 00:00:00 2001 From: Jean-Francois Marchand <jean-francois.marchand@cern.ch> Date: Thu, 25 Jan 2024 15:44:13 +0100 Subject: [PATCH 1/2] Changes from Daniel --- device/calo/decoding/include/CaloDecode.cuh | 1 - device/calo/decoding/src/CaloDecode.cu | 63 ++++++++------- .../event_model/calo/include/CaloRawEvent.cuh | 78 ------------------- 3 files changed, 33 insertions(+), 109 deletions(-) delete mode 100644 device/event_model/calo/include/CaloRawEvent.cuh diff --git a/device/calo/decoding/include/CaloDecode.cuh b/device/calo/decoding/include/CaloDecode.cuh index 099db76c8f4..eb1e93ab8e8 100644 --- a/device/calo/decoding/include/CaloDecode.cuh +++ b/device/calo/decoding/include/CaloDecode.cuh @@ -4,7 +4,6 @@ #pragma once -#include "CaloRawEvent.cuh" #include "CaloGeometry.cuh" #include "CaloDigit.cuh" #include "AlgorithmTypes.cuh" diff --git a/device/calo/decoding/src/CaloDecode.cu b/device/calo/decoding/src/CaloDecode.cu index 2c5f79b90b5..559ff0afab2 100644 --- a/device/calo/decoding/src/CaloDecode.cu +++ b/device/calo/decoding/src/CaloDecode.cu @@ -34,30 +34,31 @@ namespace { for (unsigned bank_number = threadIdx.x; bank_number < raw_event.number_of_raw_banks; bank_number += blockDim.x) { auto raw_bank = raw_event.raw_bank(bank_number); + auto raw_bank_data_u32 = reinterpret_cast<const uint32_t*>(raw_bank.data); if constexpr (decoding_version < 4) { // old decoding - while (raw_bank.data < raw_bank.end) { - uint32_t word = *raw_bank.data; + const auto raw_bank_end = reinterpret_cast<const uint32_t*>(raw_bank.data + raw_bank.size); + while (raw_bank_data_u32 < raw_bank_end) { + uint32_t word = *raw_bank_data_u32; uint16_t trig_size = word & 0x7F; uint16_t code = (word >> 14) & 0x1FF; // Skip header and trigger words - raw_bank.data += 1 + (trig_size + 3) / 4; + raw_bank_data_u32 += 1 + (trig_size + 3) / 4; // pattern bits - unsigned int pattern = *raw_bank.data; - // Loop over all cards in this front-env sub-bank. - uint32_t last_data = *(raw_bank.data + 1); - raw_bank.data += 2; + unsigned int pattern = *raw_bank_data_u32; + uint32_t last_data = *(raw_bank_data_u32 + 1); + raw_bank_data_u32 += 2; int16_t offset = 0; for (unsigned int bit_num = 0; 32 > bit_num; ++bit_num) { if (31 < offset) { offset -= 32; - last_data = *raw_bank.data; - raw_bank.data += 1; + last_data = *raw_bank_data_u32; + raw_bank_data_u32 += 1; } int adc; if (0 == (pattern & (1 << bit_num))) { //.. short coding @@ -70,8 +71,8 @@ namespace { if (28 == offset) adc &= 0xF; //== clean-up extra bits offset += 12; if (32 < offset) { //.. get the extra bits on next word - last_data = *raw_bank.data; - raw_bank.data += 1; + last_data = *raw_bank_data_u32; + raw_bank_data_u32 += 1; offset -= 32; int temp = (last_data << (12 - offset)) & 0xFFF; adc += temp; @@ -99,16 +100,17 @@ namespace { } auto raw_bank_fiberCheck = raw_event_fiberCheck.raw_bank(bank_number); + auto raw_bank_fiberCheck_data_u32 = reinterpret_cast<const uint32_t*>(raw_bank_fiberCheck.data); - auto get_data = [](uint32_t const* raw_data) { - auto d = *raw_data; + auto get_data = [](uint32_t const* raw_data_u32) { + auto d = *raw_data_u32; if constexpr (decoding_version == 4) { // big endian d = ((d >> 24) & 0x000000FF) | ((d >> 8) & 0x0000FF00) | ((d << 8) & 0x00FF0000) | ((d << 24) & 0xFF000000); } return d; }; - uint32_t pattern = *(raw_bank.data); + uint32_t pattern = *(raw_bank_data_u32); int offset = 0; uint32_t lastData = pattern; @@ -119,10 +121,10 @@ namespace { for (int ifeb = 0; ifeb < 3; ifeb++) { // First, remove 3 LLTs if (ifeb == 0) { - raw_bank.data += 3; - raw_bank_fiberCheck.data += 3; + raw_bank_data_u32 += 3; + raw_bank_fiberCheck_data_u32 += 3; } - lastData = get_data(raw_bank.data); + lastData = get_data(raw_bank_data_u32); int nADC = 0; bool isFiberOff = false; @@ -132,12 +134,12 @@ namespace { // ... and readout data for (unsigned int bitNum = 0; 32 > bitNum; bitNum++) { if (nADC % 8 == 0) { // Check fibers pattern, 1 fiber corresponds to 8 ADC (96b) - if (offset == 32) raw_bank_fiberCheck.data += 1; - uint32_t pattern1 = get_data(raw_bank_fiberCheck.data); - raw_bank_fiberCheck.data += 1; - uint32_t pattern2 = get_data(raw_bank_fiberCheck.data); - raw_bank_fiberCheck.data += 1; - uint32_t pattern3 = get_data(raw_bank_fiberCheck.data); + if (offset == 32) raw_bank_fiberCheck_data_u32 += 1; + uint32_t pattern1 = get_data(raw_bank_fiberCheck_data_u32); + raw_bank_fiberCheck_data_u32 += 1; + uint32_t pattern2 = get_data(raw_bank_fiberCheck_data_u32); + raw_bank_fiberCheck_data_u32 += 1; + uint32_t pattern3 = get_data(raw_bank_fiberCheck_data_u32); if (pattern1 == fibMask1 && pattern2 == fibMask2 && pattern3 == fibMask3) isFiberOff = true; else @@ -145,8 +147,8 @@ namespace { } if (31 < offset) { offset -= 32; - raw_bank.data += 1; - lastData = get_data(raw_bank.data); + raw_bank_data_u32 += 1; + lastData = get_data(raw_bank_data_u32); } int adc = 0; @@ -158,16 +160,16 @@ namespace { adc = ((lastData >> (20 - offset)) & 0xfff); if (28 == offset) { //.. get the extra bits on next word - raw_bank.data += 1; - lastData = get_data(raw_bank.data); + raw_bank_data_u32 += 1; + lastData = get_data(raw_bank_data_u32); int temp = (lastData >> (offset - 4)) & 0xFF; offset -= 32; adc = (adc << 8) + temp; } if (24 == offset) { //.. get the extra bits on next word - raw_bank.data += 1; - lastData = get_data(raw_bank.data); + raw_bank_data_u32 += 1; + lastData = get_data(raw_bank_data_u32); int temp = (lastData >> (offset + 4)) & 0xF; offset -= 32; adc = (adc << 4) + temp; @@ -235,7 +237,8 @@ __global__ void calo_decode_dispatch( auto ecal_geometry = CaloGeometry(raw_ecal_geometry); auto const ecal_digits_offset = parameters.dev_ecal_digits_offsets[event_number]; - decode<Calo::RawEvent<mep_layout>, decoding_version>( + + decode<Allen::RawEvent<mep_layout>, decoding_version>( parameters.dev_ecal_raw_input, parameters.dev_ecal_raw_input_offsets, parameters.dev_ecal_raw_input_sizes, diff --git a/device/event_model/calo/include/CaloRawEvent.cuh b/device/event_model/calo/include/CaloRawEvent.cuh deleted file mode 100644 index 71f4886e0d7..00000000000 --- a/device/event_model/calo/include/CaloRawEvent.cuh +++ /dev/null @@ -1,78 +0,0 @@ -/*****************************************************************************\ -* (c) Copyright 2021 CERN for the benefit of the LHCb Collaboration * -* * -* This software is distributed under the terms of the Apache License * -* version 2 (Apache-2.0), copied verbatim in the file "LICENSE". * -* * -* In applying this licence, CERN does not waive the privileges and immunities * -* granted to it by virtue of its status as an Intergovernmental Organization * -* or submit itself to any jurisdiction. * -\*****************************************************************************/ -#pragma once - -#include <type_traits> -#include "BackendCommon.h" -#include "MEPTools.h" - -namespace Calo { - - struct RawBank { - uint32_t source_id = 0; - uint32_t const* data = nullptr; - uint32_t const* end = nullptr; - uint8_t const type; - - // For Allen format - __device__ __host__ RawBank(const char* raw_bank, const uint16_t s, const uint8_t t) : - RawBank {*reinterpret_cast<uint32_t const*>(raw_bank), raw_bank + sizeof(uint32_t), s, t} - {} - - // For MEP format - __device__ __host__ RawBank(const uint32_t sid, const char* fragment, const uint16_t s, const uint8_t t) : - source_id {sid}, data {reinterpret_cast<uint32_t const*>(fragment)}, - end {reinterpret_cast<uint32_t const*>(fragment + s)}, type {t} - {} - }; - - template<bool mep_layout> - struct RawEvent { - - uint32_t number_of_raw_banks = 0; - const char* data = nullptr; - const uint32_t* offsets = nullptr; - typename std::conditional_t<mep_layout, uint32_t const, uint16_t const>* sizes = nullptr; - typename std::conditional_t<mep_layout, uint32_t const, uint8_t const>* types = nullptr; - const unsigned event = 0; - - // For Allen format - __device__ __host__ - RawEvent(char const* d, uint32_t const* o, uint32_t const* s, uint32_t const* t, unsigned const event_number) : - offsets {o}, - event {event_number} - { - if constexpr (mep_layout) { - data = d; - number_of_raw_banks = MEP::number_of_banks(o); - sizes = s; - types = t; - } - else { - data = d + offsets[event]; - number_of_raw_banks = reinterpret_cast<uint32_t const*>(data)[0]; - sizes = Allen::bank_sizes(s, event); - types = Allen::bank_types(t, event); - } - } - - __device__ __host__ RawBank raw_bank(unsigned const n) const - { - if constexpr (mep_layout) { - return MEP::raw_bank<RawBank>(data, offsets, sizes, types, event, n); - } - else { - uint32_t const* bank_offsets = reinterpret_cast<uint32_t const*>(data) + 1; - return RawBank {data + (number_of_raw_banks + 2) * sizeof(uint32_t) + bank_offsets[n], sizes[n], types[n]}; - } - } - }; -} // namespace Calo -- GitLab From 860fd46fbf2f0e605432e5f9bec33ecec35cd9e6 Mon Sep 17 00:00:00 2001 From: Jean-Francois Marchand <jean-francois.marchand@cern.ch> Date: Fri, 2 Feb 2024 12:42:22 +0100 Subject: [PATCH 2/2] fix formatting --- device/calo/decoding/src/CaloDecode.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/device/calo/decoding/src/CaloDecode.cu b/device/calo/decoding/src/CaloDecode.cu index 559ff0afab2..df0e16ebc84 100644 --- a/device/calo/decoding/src/CaloDecode.cu +++ b/device/calo/decoding/src/CaloDecode.cu @@ -237,7 +237,6 @@ __global__ void calo_decode_dispatch( auto ecal_geometry = CaloGeometry(raw_ecal_geometry); auto const ecal_digits_offset = parameters.dev_ecal_digits_offsets[event_number]; - decode<Allen::RawEvent<mep_layout>, decoding_version>( parameters.dev_ecal_raw_input, parameters.dev_ecal_raw_input_offsets, -- GitLab