diff --git a/crates/gpu/src/ctx.rs b/crates/gpu/src/ctx.rs
index f3fa914a0..e84dbf08d 100644
--- a/crates/gpu/src/ctx.rs
+++ b/crates/gpu/src/ctx.rs
@@ -865,6 +865,17 @@ pub fn groups1(n: usize) -> [u32; 3] {
 
 #[cfg(test)]
 mod tests {
+    /// Every kernel module, as it is handed to the device (the generated constants, the bindings, `common.wgsl`, the
+    /// module), parses and validates without a GPU.
+    #[test]
+    fn every_kernel_module_validates_without_a_gpu() {
+        let consts = super::constants();
+        for m in super::MODULES {
+            let name = m.entries.first().copied().unwrap_or("?");
+            assert_eq!(crate::validate_wgsl(name, &super::module_source(m, &consts)), Ok(()));
+        }
+    }
+
     #[test]
     fn limits_override_parses() {
         assert_eq!(super::parse_limits("webgpu"), Some((128 << 20, 256 << 20)));
diff --git a/crates/gpu/src/faces.rs b/crates/gpu/src/faces.rs
index bee9929ef..919267786 100644
--- a/crates/gpu/src/faces.rs
+++ b/crates/gpu/src/faces.rs
@@ -817,6 +817,13 @@ static PER_CONV: AtomicBool = AtomicBool::new(false);
 mod tests {
     use super::*;
 
+    /// The face detector's kernels validate without a GPU.
+    #[test]
+    fn the_face_kernels_validate_without_a_gpu() {
+        assert_eq!(crate::validate_wgsl("faces_conv", include_str!("wgsl/faces_conv.wgsl")), Ok(()));
+        assert_eq!(crate::validate_wgsl("faces_ops", include_str!("wgsl/faces_ops.wgsl")), Ok(()));
+    }
+
     fn gpu() -> Option<Arc<FaceGpu>> {
         match accelerator() {
             Ok(a) => Some(a),
diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs
index e66a8fadd..619ab9596 100644
--- a/crates/gpu/src/lib.rs
+++ b/crates/gpu/src/lib.rs
@@ -572,10 +572,33 @@ pub fn try_render(
     }
 }
 
+/// Tests: parse and validate a WGSL module with naga, as wgpu does before it hands a shader to the driver. Needs no
+/// GPU, so kernels that don't build (a duplicate definition after a merge, a type error) fail every test run, not only
+/// the GPU ones.
+#[cfg(test)]
+pub(crate) fn validate_wgsl(name: &str, src: &str) -> Result<(), String> {
+    use wgpu::naga;
+    let module = naga::front::wgsl::parse_str(src).map_err(|e| format!("{name}: {}", e.emit_to_string(src)))?;
+    naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::all())
+        .validate(&module)
+        .map(|_| ())
+        .map_err(|e| format!("{name}: {}", e.emit_to_string(src)))
+}
+
 #[cfg(test)]
 mod tests {
     use lightcraft_pipeline::{OutputDepth, Overlay};
 
+    /// The helper refuses what a driver would: a function defined twice (what an automatic merge of two branches
+    /// touching the same kernel can leave).
+    #[test]
+    fn wgsl_validation_catches_a_duplicate_definition() {
+        let ok = "fn f() -> f32 { return 1.0; }\n@compute @workgroup_size(1) fn main() { _ = f(); }\n";
+        assert_eq!(super::validate_wgsl("ok", ok), Ok(()));
+        let twice = format!("fn f() -> f32 {{ return 2.0; }}\n{ok}");
+        assert!(super::validate_wgsl("twice", &twice).is_err());
+    }
+
     // Given a request the 8-bit kernel can't draw (high bit depth, a soft proof, Visualize HDR),
     // the GPU doesn't take it; the loupe and diagnostic overlays it can draw stay on the GPU
     #[test]
diff --git a/crates/gpu/src/nn.rs b/crates/gpu/src/nn.rs
index 0183d878e..235ba9be6 100644
--- a/crates/gpu/src/nn.rs
+++ b/crates/gpu/src/nn.rs
@@ -925,6 +925,25 @@ impl TileRunner for NetRunner {
 #[cfg(test)]
 mod tests {
     use super::*;
+
+    /// The denoise kernels validate without a GPU: the pooling kernel, and the convolution for every block size the
+    /// planner picks for these layer widths and channel counts.
+    #[test]
+    fn the_network_kernels_validate_without_a_gpu() {
+        assert_eq!(crate::validate_wgsl("nn_pool", include_str!("wgsl/nn_pool.wgsl")), Ok(()));
+        let mut sizes: Vec<(u32, u32, u32)> = Vec::new();
+        for n in [1usize, 3, 8, 16, 24, 32, 48, 64, 96, 128, 256, 512] {
+            for cin in [1usize, 3, 4, 8, 16, 32, 48, 64, 128, 256, 512] {
+                let b = blocks(n, cin);
+                if !sizes.contains(&b) {
+                    sizes.push(b);
+                }
+            }
+        }
+        for (bm, bn, bk) in sizes {
+            assert_eq!(crate::validate_wgsl(&format!("nn_conv {bm}x{bn}x{bk}"), &conv_source(bm, bn, bk)), Ok(()));
+        }
+    }
     use lightcraft_denoise::{onnx, reference, synthetic};
 
     fn net(tile: u64, ch: u64, depth: usize, seed: u64) -> Net {