From 873bfa74dd30037e2f89de206ee276a84218cbb6 Mon Sep 17 00:00:00 2001 From: Thomas Date: Thu, 15 Jun 2023 16:36:15 +0200 Subject: [PATCH] Add integrity checks for strip and xtal ids to ECAL GPU unpacker. --- EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu b/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu index 14255b6823b93..b80831e104f84 100644 --- a/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu +++ b/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu @@ -283,6 +283,23 @@ namespace ecal { auto const wdata = current_tower_block[1 + i_to_access * 3]; uint8_t const stripid = wdata & 0x7; uint8_t const xtalid = (wdata >> 4) & 0x7; + + // check if the stripid and xtalid are in the allowed range and if not skip the rest of the block + if (stripid < ElectronicsIdGPU::MIN_STRIPID || stripid > ElectronicsIdGPU::MAX_STRIPID || + xtalid < ElectronicsIdGPU::MIN_XTALID || xtalid > ElectronicsIdGPU::MAX_XTALID) { + break; + } + if (ich > 0 || threadIdx.x > 0) { + // check if the stripid has increased or that the xtalid has increased from the previous data word. If not something is wrong and the rest of the block is skipped. + auto const prev_i_to_access = ich + threadIdx.x - 1; + auto const prevwdata = current_tower_block[1 + prev_i_to_access * 3]; + uint8_t const laststripid = prevwdata & 0x7; + uint8_t const lastxtalid = (prevwdata >> 4) & 0x7; + if ((stripid == laststripid && xtalid <= lastxtalid) || (stripid < laststripid)) { + break; + } + } + ElectronicsIdGPU eid{fed2dcc(fed), ttid, stripid, xtalid}; auto const didraw = isBarrel ? compute_ebdetid(eid) : eid2did[eid.linearIndex()]; // FIXME: what kind of channels are these guys