From d2d1d78545c34939fdc16b025ee6046a74155a38 Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Fri, 31 Jul 2026 14:43:49 +0200 Subject: [PATCH] spot_finding: keep the GPU wave inside the image 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) --- image_analysis/spot_finding/ImageSpotFinderGPU.cu | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/image_analysis/spot_finding/ImageSpotFinderGPU.cu b/image_analysis/spot_finding/ImageSpotFinderGPU.cu index 2a1afc76..b535d8e1 100644 --- a/image_analysis/spot_finding/ImageSpotFinderGPU.cu +++ b/image_analysis/spot_finding/ImageSpotFinderGPU.cu @@ -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(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(params.height)) + return; + const int32_t left = max(static_cast(threadIdx.x) - static_cast(ImageSpotFinder::NBX), 0); // leftmost column touched by this thread const int32_t right = min(static_cast(threadIdx.x) + static_cast(ImageSpotFinder::NBX) + 1, static_cast(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(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(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