diff --git a/crates/gpu/src/ctx.rs b/crates/gpu/src/ctx.rs
--- a/crates/gpu/src/ctx.rs
+++ b/crates/gpu/src/ctx.rs
@@ -32,7 +32,7 @@ const MODULES: &[Module] = &[
     Module {
         src: include_str!("wgsl/blur.wgsl"),
         bindings: &[("src", false, "f32"), ("dst", true, "f32")],
-        entries: &["box_h", "box_v", "median_h", "median_v"],
+        entries: &["box_h", "box_h_slide", "box_v", "median_h", "median_v"],
     },
     Module {
         src: include_str!("wgsl/resize.wgsl"),
diff --git a/crates/gpu/src/render.rs b/crates/gpu/src/render.rs
--- a/crates/gpu/src/render.rs
+++ b/crates/gpu/src/render.rs
@@ -236,10 +236,42 @@ const CHUNK: usize = 16;
 /// One box pass of radius `r` (clamped edges) over a `w × h` image of `nc` interleaved channels,
 /// `src → dst`: horizontal or vertical.
 fn box_pass(cx: &mut Cx<'_>, horizontal: bool, src: &Buf, dst: &Buf, (w, h): (usize, usize), nc: usize, r: usize) {
-    let chunk = (2 * r + 1).clamp(CHUNK, 128);
-    let p = [w as u32, h as u32, nc as u32, r as u32, chunk as u32];
-    let (k, groups) = if horizontal { ("box_h", groups2(w, h, [64, 4])) } else { ("box_v", groups2(w, h.div_ceil(chunk), [64, 4])) };
-    cx.run(k, &p, &[Some(src), Some(dst)], groups);
+    let plan = box_plan(horizontal, r);
+    let p = [w as u32, h as u32, nc as u32, r as u32, plan.chunk as u32];
+    let groups = match plan.kernel {
+        "box_h" => groups2(w, h, [64, 4]),
+        "box_h_slide" => groups2(w.div_ceil(plan.chunk), h, [64, 4]),
+        _ => groups2(w, h.div_ceil(plan.chunk), [64, 4]),
+    };
+    cx.run(plan.kernel, &p, &[Some(src), Some(dst)], groups);
+}
+
+/// Radius from which the horizontal box pass slides a running sum ([`box_plan`]).
+const BOX_H_SLIDE_FROM: usize = 32;
+
+/// How one box pass runs: its kernel, the run of pixels each sliding thread covers, and the source
+/// reads it costs per output pixel.
+#[derive(Clone, Copy, Debug)]
+struct BoxPlan {
+    kernel: &'static str,
+    chunk: usize,
+    #[cfg_attr(not(test), allow(dead_code))]
+    reads_per_px: f32,
+}
+
+/// The box pass for radius `r`. Work per pixel stays bounded whatever the radius (a few reads), so a
+/// large blur can't make one dispatch run long enough for a driver watchdog (Windows TDR, ~2 s) to reset
+/// the GPU: small horizontal radii read their window per pixel (cache friendly), larger ones and every
+/// vertical pass slide a running sum over a run at least as long as the window.
+fn box_plan(horizontal: bool, r: usize) -> BoxPlan {
+    let window = r.saturating_mul(2).saturating_add(1);
+    if horizontal && r < BOX_H_SLIDE_FROM {
+        return BoxPlan { kernel: "box_h", chunk: 1, reads_per_px: window as f32 };
+    }
+    let chunk = window.max(CHUNK);
+    // priming (window reads) plus two reads per pixel of the run
+    let reads_per_px = (window as f32 + 2.0 * chunk as f32) / chunk as f32;
+    BoxPlan { kernel: if horizontal { "box_h_slide" } else { "box_v" }, chunk, reads_per_px }
 }
 
 /// Gaussian blur of a `w × h` image of `nc` interleaved channels: three box passes each way
@@ -1161,6 +1193,19 @@ mod tests {
         assert!(covered > 0 && covered < w * h, "{covered} px covered");
     }
 
+    /// A box pass costs a bounded number of reads per pixel at any radius (large blurs once made a
+    /// single dispatch long enough to trip a driver watchdog).
+    #[test]
+    fn box_pass_work_per_pixel_is_bounded_for_any_radius() {
+        for horizontal in [true, false] {
+            for r in [0usize, 1, 5, 31, 32, 64, 120, 1000, 100_000] {
+                let p = box_plan(horizontal, r);
+                assert!(p.reads_per_px <= 2.0 * BOX_H_SLIDE_FROM as f32 + 1.0, "{horizontal} r={r}: {p:?}");
+                assert!(p.chunk >= 1);
+            }
+        }
+    }
+
     /// The blur against `lightcraft_raster::blur::gaussian`: widths that do and don't fill the last
     /// workgroup, windows wider than the image, and 1, 2 (as two planes) and 3 channels.
     #[test]
diff --git a/crates/gpu/src/wgsl/blur.wgsl b/crates/gpu/src/wgsl/blur.wgsl
--- a/crates/gpu/src/wgsl/blur.wgsl
+++ b/crates/gpu/src/wgsl/blur.wgsl
@@ -48,6 +48,35 @@ fn box_h(@builtin(global_invocation_id) g: vec3<u32>) {
     st(row + x, nc, acc * (1.0 / f32(2 * r + 1)));
 }
 
+// Horizontal for large radii: one thread per CH consecutive pixels of a row, sliding a running sum
+// (box_h's per-pixel window costs 2r + 1 reads a pixel, which for radii in the hundreds made one
+// dispatch long enough for a driver watchdog to reset the GPU). P: w, h, nc, r, ch.
+@compute @workgroup_size(64, 4)
+fn box_h_slide(@builtin(global_invocation_id) g: vec3<u32>) {
+    let w = pu(0u);
+    let h = pu(1u);
+    let nc = pu(2u);
+    let r = i32(pu(3u));
+    let ch = pu(4u);
+    let x0 = g.x * ch;
+    let y = g.y;
+    if (x0 >= w || y >= h) {
+        return;
+    }
+    let last = i32(w) - 1;
+    let row = y * w;
+    let inv = 1.0 / f32(2 * r + 1);
+    let x1 = min(x0 + ch, w);
+    var acc = vec3<f32>(0.0);
+    for (var k = -r; k <= r; k++) {
+        acc += ld(row + u32(clamp(i32(x0) + k, 0, last)), nc);
+    }
+    for (var x = x0; x < x1; x++) {
+        st(row + x, nc, acc * inv);
+        acc = acc + ld(row + u32(min(i32(x) + r + 1, last)), nc) - ld(row + u32(max(i32(x) - r, 0)), nc);
+    }
+}
+
 // Vertical: one thread per CH consecutive rows of a column (the CPU re-primes every BAND rows the
 // same way).
 @compute @workgroup_size(64, 4)