Wgpu 24

Cargo.toml

@@ -17,7 +17,7 @@ prune = []
 allow_deprecated = []
 
 [dependencies]
-naga = { version = "23", features = ["wgsl-in", "wgsl-out"] }
+naga = { version = "24", features = ["wgsl-in", "wgsl-out"] }
 tracing = "0.1"
 regex = "1.8"
 regex-syntax = "0.8"
@@ -31,6 +31,6 @@ once_cell = "1.17.0"
 indexmap = "2"
 
 [dev-dependencies]
-wgpu = { version = "23", features = ["naga-ir"] }
+wgpu = { version = "24", features = ["naga-ir"] }
 futures-lite = "1"
 tracing-subscriber = { version = "0.3", features = ["std", "fmt"] }

examples/pbr_compose_test.rs

@@ -139,7 +139,7 @@ fn test_compose_final_module(n: usize, composer: &mut Composer) {
 
 // make shader module from string
 fn test_wgsl_string_compile(n: usize) {
-    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::default());
+    let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::default());
     let adapter = instance
         .enumerate_adapters(wgpu::Backends::all())
         .into_iter()
@@ -163,7 +163,7 @@ fn test_wgsl_string_compile(n: usize) {
 
 // make shader module from composed naga
 fn test_composer_compile(n: usize, composer: &mut Composer) {
-    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::default());
+    let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::default());
     let adapter = instance
         .enumerate_adapters(wgpu::Backends::all())
         .into_iter()

src/compose/mod.rs

@@ -521,13 +521,15 @@ impl Composer {
                     expressions: Default::default(),
                     named_expressions: Default::default(),
                     body: Default::default(),
+                    diagnostic_filter_leaf: Default::default(),
                 };
                 let ep = EntryPoint {
                     name: dummy_entry_point.clone(),
                     stage: naga::ShaderStage::Vertex,
                     function: func,
                     early_depth_test: None,
                     workgroup_size: [0, 0, 0],
+                    workgroup_size_overrides: None,
                 };
 
                 naga_module.entry_points.push(ep);
@@ -1043,6 +1045,7 @@ impl Composer {
                         expressions: Default::default(),
                         named_expressions: Default::default(),
                         body: Default::default(),
+                        diagnostic_filter_leaf: None,
                     };
 
                     // record owned function
@@ -1080,6 +1083,7 @@ impl Composer {
                     expressions: Default::default(),
                     named_expressions: Default::default(),
                     body: Default::default(),
+                    diagnostic_filter_leaf: None,
                 };
 
                 owned_functions.insert(ep.function.name.clone().unwrap(), (None, header_function));
@@ -1784,6 +1788,7 @@ impl Composer {
                 stage: stage.unwrap_or(ep.stage),
                 early_depth_test: ep.early_depth_test,
                 workgroup_size: ep.workgroup_size,
+                workgroup_size_overrides: ep.workgroup_size_overrides,
             });
         }
 

src/compose/test.rs

@@ -1263,6 +1263,47 @@ mod test {
         output_eq!(wgsl, "tests/expected/atomics.txt");
     }
 
+    #[test]
+    #[should_panic] // Diagnostic filters not yet supported
+    fn test_diagnostic_filters() {
+        let mut composer = Composer::default();
+
+        composer
+            .add_composable_module(ComposableModuleDescriptor {
+                source: include_str!("tests/diagnostic_filters/filters.wgsl"),
+                file_path: "tests/diagnostic_filters/filters.wgsl",
+                ..Default::default()
+            })
+            .unwrap();
+
+        // TODO enable this test when HLSL support is available
+        if cfg!(feature = "test_shader") && false {
+            assert_eq!(test_shader(&mut composer), 28.0);
+        }
+
+        let module = composer
+            .make_naga_module(NagaModuleDescriptor {
+                source: include_str!("tests/diagnostic_filters/top.wgsl"),
+                file_path: "tests/diagnostic_filters/top.wgsl",
+                ..Default::default()
+            })
+            .unwrap();
+
+        let info = composer.create_validator().validate(&module).unwrap();
+        let wgsl = naga::back::wgsl::write_string(
+            &module,
+            &info,
+            naga::back::wgsl::WriterFlags::EXPLICIT_TYPES,
+        )
+        .unwrap();
+
+        // let mut f = std::fs::File::create("atomics.txt").unwrap();
+        // f.write_all(wgsl.as_bytes()).unwrap();
+        // drop(f);
+
+        output_eq!(wgsl, "tests/expected/diagnostic_filters.txt");
+    }
+
     #[cfg(feature = "test_shader")]
     #[test]
     fn effective_defs() {
@@ -1354,7 +1395,7 @@ mod test {
             })
             .unwrap();
 
-        let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::default());
+        let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::default());
         let adapter = instance
             .enumerate_adapters(wgpu::Backends::all())
             .into_iter()

src/compose/tests/diagnostic_filters/filters.wgsl

@@ -0,0 +1,12 @@
+#define_import_path filters
+
+diagnostic(warning, derivative_uniformity);
+
+fn diagnostic_test(s : sampler, tex : texture_2d<f32>, ro_buffer : array<f32, 4>) -> vec4f {
+    if ro_buffer[0] == 0 {
+        // Emits a derivative uniformity error during validation.
+        return textureSample(tex, s, vec2(0.,0.));
+    }
+
+      return vec4f(0.);
+}

src/compose/tests/diagnostic_filters/top.wgsl

@@ -0,0 +1,10 @@
+#import filters
+
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(2) var tex : texture_2d<f32>;
+@group(1) @binding(0) var<storage, read> ro_buffer : array<f32, 4>;
+
+@fragment
+fn main(@builtin(position) p : vec4f) -> @location(0) vec4f {
+    return filters::diagnostic_test();
+}

src/compose/tests/expected/diagnostic_filters.txt

@@ -0,0 +1,20 @@
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(2) var tex : texture_2d<f32>;
+@group(1) @binding(0) var<storage, read> ro_buffer : array<f32, 4>;
+
+@fragment
+fn main(@builtin(position) p : vec4f) -> @location(0) vec4f {
+    return filters::diagnostic_test();
+}
+
+diagnostic(warning, derivative_uniformity);
+
+fn diagnostic_test() -> vec4f {
+    diagnostic(off, derivative_uniformity);
+    if ro_buffer[0] == 0 {
+        // Emits a derivative uniformity error during validation.
+        return textureSample(tex, s, vec2(0.,0.));
+    }
+
+      return vec4f(0.);
+}

src/compose/tests/expected/err_parse.txt

@@ -1,8 +1,8 @@
-error: no definition in scope for identifier: 'zdd'
+error: no definition in scope for identifier: `zdd`
    ┌─ tests/error_test/wgsl_parse_err.wgsl:16:12
    │
 16 │     return zdd; 
    │            ^^^ unknown identifier
    │
-   = no definition in scope for identifier: 'zdd'
+   = no definition in scope for identifier: `zdd`
 

src/derive.rs

@@ -443,6 +443,29 @@ impl<'a> DerivedModule<'a> {
                         argument: map_expr!(argument),
                         result: map_expr!(result),
                     },
+                    Statement::ImageAtomic {
+                        image,
+                        coordinate,
+                        array_index,
+                        fun,
+                        value,
+                    } => {
+                        let fun = match fun {
+                            AtomicFunction::Exchange {
+                                compare: Some(compare_expr),
+                            } => AtomicFunction::Exchange {
+                                compare: Some(map_expr!(compare_expr)),
+                            },
+                            fun => *fun,
+                        };
+                        Statement::ImageAtomic {
+                            image: map_expr!(image),
+                            coordinate: map_expr!(coordinate),
+                            array_index: map_expr_opt!(array_index),
+                            fun,
+                            value: map_expr!(value),
+                        }
+                    }
                     // else just copy
                     Statement::Break
                     | Statement::Continue
@@ -769,6 +792,7 @@ impl<'a> DerivedModule<'a> {
             expressions: Rc::try_unwrap(expressions).unwrap().into_inner(),
             named_expressions,
             body,
+            diagnostic_filter_leaf: None,
         }
     }
 
@@ -819,6 +843,7 @@ impl<'a> DerivedModule<'a> {
                 early_depth_test: ep.early_depth_test,
                 workgroup_size: ep.workgroup_size,
                 function: self.localize_function(&ep.function),
+                workgroup_size_overrides: ep.workgroup_size_overrides,
             })
             .collect();
 
@@ -842,6 +867,8 @@ impl From<DerivedModule<'_>> for naga::Module {
             special_types: Default::default(),
             entry_points: Default::default(),
             overrides: derived.pipeline_overrides,
+            diagnostic_filters: Default::default(),
+            diagnostic_filter_leaf: None,
         }
     }
 }

src/redirect.rs

@@ -71,7 +71,8 @@ impl Redirector {
                 | Statement::RayQuery { .. }
                 | Statement::SubgroupBallot { .. }
                 | Statement::SubgroupGather { .. }
-                | Statement::SubgroupCollectiveOperation { .. } => (),
+                | Statement::SubgroupCollectiveOperation { .. }
+                | Statement::ImageAtomic { .. } => (),
             }
         }
     }