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<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