mx/Jungfraujoch
2
1
Fork 0
Code Issues Pull Requests Actions 4 Packages Projects Releases 67 Wiki Activity

spot_finding: keep the GPU wave inside the image

Browse Source 
rowsPerWave is rounded up, so with 32 waves the last waves can start at
or past the last row: rmin was never clamped and only the drain loop
checked front against height. On any detector below about 1500 rows -
including the module-converted 500K and 1M geometries and the kernel's
own unit tests - the priming and steady-state loops read whole rows past
the end of the image buffer, and those garbage rows entered the sliding
background window of the bottom rows.

Blocks with no rows to write now return before the first __syncthreads
(rmin depends only on blockIdx.y, so the block leaves together and the
collective ops stay well formed), and both remaining reads are bounded by
height. Rows past the end keep the INT32_MIN sentinel, which the window
already treats as "not counted".

The raw read in the steady-state loop is left as it is: making it apply
the prev_out substitution that the other two read sites use would change
which pixels are found, which is a separate question from this fix.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
This commit is contained in:
leonarski_f
2026-07-31 14:43:49 +02:00
co-authored by Claude Opus 5
parent 4a8a8e69bf
commit d2d1d78545
1 changed files with 11 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
image_analysis/spot_finding/ImageSpotFinderGPU.cu
+11 -2
View File
@@ -114,6 +114,12 @@ __global__ void analyze_pixel(const int32_t *in, uint32_t *prev_out, uint32_t *o
const int32_t rmin = blockIdx.y * rowsPerWave; // lowest result row for this wave
const int32_t rmax = min(rmin + rowsPerWave, static_cast<int32_t>(params.height)); // past highest result row for this wave
// rowsPerWave is rounded up, so the last waves can start at or past the last row and have no
// rows to write. rmin depends only on blockIdx.y, so the whole block leaves together and the
// __syncthreads/__ballot_sync below stay collective.
if (rmin >= static_cast<int32_t>(params.height))
return;
const int32_t left = max(static_cast<int32_t>(threadIdx.x) - static_cast<int32_t>(ImageSpotFinder::NBX), 0); // leftmost column touched by this thread
const int32_t right = min(static_cast<int32_t>(threadIdx.x) + static_cast<int32_t>(ImageSpotFinder::NBX) + 1, static_cast<int32_t>(params.width)); // past rightmost column touched by this thread
@@ -139,7 +145,8 @@ __global__ void analyze_pixel(const int32_t *in, uint32_t *prev_out, uint32_t *o
// wave front up to rmin + nby + 1
do {
if (data_col) { // read at the front end of the wave
// Rows past the last one contribute nothing; shared_val keeps its INT32_MIN initialiser.
if (data_col && (front < static_cast<int32_t>(params.height))) { // read at the front end of the wave
const int32_t npixel = front * params.width + col;
const bool sat = ((prev_out[npixel / 32] & (1U << (npixel % 32))) != 0);
const int32_t val = sat ? INT32_MAX : in[npixel];
@@ -175,7 +182,9 @@ __global__ void analyze_pixel(const int32_t *in, uint32_t *prev_out, uint32_t *o
old = 0; // no effect value
cnt = 1; // bring count to normal
}
int32_t val = in[front * params.width + col];
int32_t val = INT32_MIN; // past the last row nothing enters the window
if (front < static_cast<int32_t>(params.height))
val = in[front * params.width + col];
shared_val[(front % window) * blockDim.x + threadIdx.x] = val;
if (val == INT32_MAX || val == INT32_MIN) {
val = 0; // no effect value