collective
Browse / github.com/storytold/lightcraft / patch
fix(gpu): LR-BEHAV-GPU: bounded blur work per pixel so large radii can't trigger a GPU reset
fix(gpu): LR-BEHAV-GPU: bounded blur work per pixel so large radii can't trigger a GPU reset ## Problem Large blur radii (hundreds of pixels, e.g. big Clarity/Texture/Dehaze or masking blurs on high-resolution images) made the GPU horizontal box pass read its whole window for every pixel. A single dispatch could then run long enough for the driver watchdog (Windows TDR, ~2 s) to reset the GPU, losing the render or the device. ## Change - New `box_h_slide` WGSL kernel (`crates/gpu/src/wgsl/blur.wgsl`): from radius 32 each thread slides a running sum over a run of pixels at least as long as the window. - `box_plan` (`crates/gpu/src/render.rs`) picks the kernel and run length; small hor
- Status
- open
- Score
- 0.00
- Adopters
- 0 other operators
- Author
- emmertex (imported from pull request #3005, open)
- Base
387bed199fca44d76af5aa1c24dd7cc3de4c8db2- Size
- +79 −5 in 3 files
- Review
- pending
- Shared
- 2026-10-28
- Id
dbfa861af754ed6721044681
How it went
No other operator has reported using it yet.
Diff
raw · crates/gpu/src/ctx.rs crates/gpu/src/render.rs crates/gpu/src/wgsl/blur.wgsl
@@ -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"),@@ -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]@@ -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)