Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

ECAL skip GPU unpacking of the rest of the block if a bad block is detected - 132x #42394

Merged
merged 3 commits into from
Jul 31, 2023
Merged

ECAL skip GPU unpacking of the rest of the block if a bad block is detected - 132x #42394

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
4795d11
Stop ECAL unpacking on GPU for the rest of the block if a bad block i…
thomreis Jul 18, 2023
cb32efe
Additional syncthreads and replace break with continue.
thomreis Jul 25, 2023
55fc14b
Put the syncthreads at the right place.
thomreis Jul 26, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
65 changes: 45 additions & 20 deletions EventFilter/EcalRawToDigi/plugins/UnpackGPU.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -272,32 +272,57 @@ namespace ecal {
// get the next channel coordinates
uint32_t nchannels = (block_length - 1) / 3;

bool bad_block = false;
__shared__ uint32_t ch_with_bad_block;
// 1 threads per channel in this block
for (uint32_t ich = 0; ich < nchannels; ich += NTHREADS) {
auto const i_to_access = ich + threadIdx.x;
// threads outside of the range -> leave the loop
if (i_to_access >= nchannels)
break;
if (i_to_access == 0) {
ch_with_bad_block = std::numeric_limits<uint32_t>::max();
}

// inc the channel's counter and get the pos where to store
auto const wdata = current_tower_block[1 + i_to_access * 3];
uint8_t const stripid = wdata & 0x7;
uint8_t const xtalid = (wdata >> 4) & 0x7;
// make sure the shared memory is initialised for all threads
__syncthreads();

// 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;
uint64_t wdata;
uint8_t stripid;
uint8_t xtalid;

// threads must be inside the range (no break here because of __syncthreads() afterwards)
if (i_to_access < nchannels && i_to_access < ch_with_bad_block) {
// inc the channel's counter and get the pos where to store
wdata = current_tower_block[1 + i_to_access * 3];
stripid = wdata & 0x7;
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) {
bad_block = true;
}
if (i_to_access > 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 = i_to_access - 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)) {
bad_block = true;
}
}
}

// 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
if (i_to_access >= nchannels || i_to_access >= ch_with_bad_block) {
continue;
}

ElectronicsIdGPU eid{fed2dcc(fed), ttid, stripid, xtalid};
Expand Down