From 6c012bc7ff3b29889351285277f8f597b68c9b2b Mon Sep 17 00:00:00 2001 From: Thomas Date: Wed, 26 Jul 2023 10:37:42 +0200 Subject: [PATCH] Put the syncthreads at the right place. --- EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu b/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu index 78314f7dda1f5..2ba5c508c588b 100644 --- a/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu +++ b/EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu @@ -281,6 +281,9 @@ namespace ecal { ch_with_bad_block = std::numeric_limits::max(); } + // make sure the shared memory is initialised for all threads + __syncthreads(); + uint64_t wdata; uint8_t stripid; uint8_t xtalid; @@ -309,13 +312,12 @@ namespace ecal { } } - __syncthreads(); - // check if this thread has the lowest bad block if (bad_block && i_to_access < ch_with_bad_block) { atomicMin(&ch_with_bad_block, i_to_access); } + // make sure that all threads that have to have set the ch_with_bad_block shared memory __syncthreads(); // threads outside of the range or bad block detected in this thread or one working on a lower block -> stop this loop iteration here