From 7c9b6d88e52e2bfe6e08ebf33f170d66a7d4cbc5 Mon Sep 17 00:00:00 2001
From: OneSlips <alurkoff@yahoo.com>
Date: Sun, 25 Oct 2026 19:41:17 +1000
Subject: [PATCH] GPU-TESTS: every WGSL kernel parses and validates in the test
 run, without a GPU

The kernels were only compiled when a device was created, so a module that doesn't build (a
function defined twice after an automatic merge, a type error) passed CI on machines without a
GPU and failed on users' machines. naga (re-exported by wgpu) now parses and validates every
module as it is handed to the device: the finish, blur, resize, map, mask and geometry modules
with their generated constants and bindings, the denoise convolution for every block size the
planner picks, the pooling kernel and the face detector's two. A test checks that a duplicate
definition is caught.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
---
 crates/gpu/src/ctx.rs   | 11 +++++++++++
 crates/gpu/src/faces.rs |  7 +++++++
 crates/gpu/src/lib.rs   | 23 +++++++++++++++++++++++
 crates/gpu/src/nn.rs    | 19 +++++++++++++++++++
 4 files changed, 60 insertions(+)

diff --git a/crates/gpu/src/ctx.rs b/crates/gpu/src/ctx.rs
index 33e02fbe4..a6f702c12 100644
--- a/crates/gpu/src/ctx.rs
+++ b/crates/gpu/src/ctx.rs
@@ -763,6 +763,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 452e2dd67..ff8c8dbd8 100644
--- a/crates/gpu/src/lib.rs
+++ b/crates/gpu/src/lib.rs
@@ -550,10 +550,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 55420598b..c1db6d2ae 100644
--- a/crates/gpu/src/nn.rs
+++ b/crates/gpu/src/nn.rs
@@ -881,6 +881,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 {