Skip to content

Commit

Permalink
Put the syncthreads at the right place.
Browse files Browse the repository at this point in the history
  • Loading branch information
thomreis committed Jul 28, 2023
1 parent 45ce573 commit 6c012bc
Showing 1 changed file with 4 additions and 2 deletions.
6 changes: 4 additions & 2 deletions EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu
Original file line number Diff line number Diff line change
Expand Up @@ -281,6 +281,9 @@ namespace ecal {
ch_with_bad_block = std::numeric_limits<uint32_t>::max();
}

// make sure the shared memory is initialised for all threads
__syncthreads();

uint64_t wdata;
uint8_t stripid;
uint8_t xtalid;
Expand Down Expand Up @@ -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
Expand Down

0 comments on commit 6c012bc

Please sign in to comment.