naga 28 update

Cargo.toml

@@ -16,9 +16,11 @@ glsl = ["naga/glsl-in", "naga/glsl-out"]
 override_any = []
 prune = []
 allow_deprecated = []
+# forces composable modules to include "enable wgpu_ray_query"
+override_enable_ray_query = []
 
 [dependencies]
-naga = { version = "27", features = ["wgsl-in", "wgsl-out", "termcolor"] }
+naga = { version = "28", features = ["wgsl-in", "wgsl-out", "termcolor"] }
 tracing = "0.1"
 regex = "1.8"
 thiserror = "2.0"
@@ -29,6 +31,6 @@ unicode-ident = "1"
 indexmap = "2"
 
 [dev-dependencies]
-wgpu = { version = "27", features = ["naga-ir"] }
+wgpu = { version = "28", features = ["naga-ir"] }
 futures-lite = "2"
 tracing-subscriber = { version = "0.3", features = ["std", "fmt"] }

examples/pbr_compose_test.rs

@@ -140,11 +140,11 @@ 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 adapter = instance
-        .enumerate_adapters(wgpu::Backends::all())
-        .into_iter()
-        .next()
-        .unwrap();
+    let adapter =
+        futures_lite::future::block_on(instance.enumerate_adapters(wgpu::Backends::all()))
+            .into_iter()
+            .next()
+            .unwrap();
     let device =
         futures_lite::future::block_on(adapter.request_device(&wgpu::DeviceDescriptor::default()))
             .unwrap()
@@ -163,11 +163,11 @@ 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 adapter = instance
-        .enumerate_adapters(wgpu::Backends::all())
-        .into_iter()
-        .next()
-        .unwrap();
+    let adapter =
+        futures_lite::future::block_on(instance.enumerate_adapters(wgpu::Backends::all()))
+            .into_iter()
+            .next()
+            .unwrap();
     let device =
         futures_lite::future::block_on(adapter.request_device(&wgpu::DeviceDescriptor::default()))
             .unwrap()

src/compose/mod.rs

@@ -533,6 +533,8 @@ impl Composer {
                     early_depth_test: None,
                     workgroup_size: [0, 0, 0],
                     workgroup_size_overrides: None,
+                    mesh_info: None,
+                    task_payload: None,
                 };
 
                 naga_module.entry_points.push(ep);
@@ -1387,7 +1389,17 @@ impl Composer {
             true,
             &preprocessed_source,
             imports,
-            None,
+            if cfg!(feature = "override_enable_ray_query") {
+                Some(WgslDirectives {
+                    enables: vec![EnableDirective {
+                        extensions: vec!["wgpu_ray_query".to_string()],
+                        source_location: 0,
+                    }],
+                    ..Default::default()
+                })
+            } else {
+                None
+            },
         )
         .map_err(|err| err.into())
     }
@@ -1834,6 +1846,8 @@ impl Composer {
                 early_depth_test: ep.early_depth_test,
                 workgroup_size: ep.workgroup_size,
                 workgroup_size_overrides: ep.workgroup_size_overrides,
+                mesh_info: ep.mesh_info.clone(),
+                task_payload: ep.task_payload,
             });
         }
         let mut naga_module = naga::Module {

src/compose/test.rs

@@ -1298,6 +1298,7 @@ mod test {
         output_eq!(wgsl, "tests/expected/atomics.txt");
     }
 
+    #[cfg(feature = "override_enable_ray_query")]
     #[test]
     fn test_raycasts() {
         let mut composer =
@@ -1489,11 +1490,11 @@ mod test {
             .unwrap();
 
         let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::default());
-        let adapter = instance
-            .enumerate_adapters(wgpu::Backends::all())
-            .into_iter()
-            .next()
-            .unwrap();
+        let adapter =
+            futures_lite::future::block_on(instance.enumerate_adapters(wgpu::Backends::all()))
+                .into_iter()
+                .next()
+                .unwrap();
         let (device, queue) =
             futures_lite::future::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
                 required_features: Features::MAPPABLE_PRIMARY_BUFFERS,
@@ -1533,7 +1534,7 @@ mod test {
                 &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                     label: None,
                     bind_group_layouts: &[&layout],
-                    push_constant_ranges: &[],
+                    immediate_size: 0,
                 }),
             ),
             module: &shader_module,

src/compose/tests/raycast/top.wgsl

@@ -1,3 +1,5 @@
+enable wgpu_ray_query;
+
 #import test_module
 
 fn main() -> f32 {

src/derive.rs

@@ -908,6 +908,8 @@ impl<'a> DerivedModule<'a> {
                 workgroup_size: ep.workgroup_size,
                 function: self.localize_function(&ep.function),
                 workgroup_size_overrides: ep.workgroup_size_overrides,
+                mesh_info: ep.mesh_info.clone(),
+                task_payload: ep.task_payload,
             })
             .collect();
 

src/prune/mod.rs

@@ -83,6 +83,7 @@ impl FunctionReq {
             expressions,
             named_expressions,
             body,
+            diagnostic_filter_leaf: func.diagnostic_filter_leaf,
         }
     }
 
@@ -149,6 +150,7 @@ impl FunctionReq {
                 offset,
                 level,
                 depth_ref,
+                clamp_to_edge,
             } => Expression::ImageSample {
                 image: expr_map[image],
                 sampler: expr_map[sampler],
@@ -166,6 +168,7 @@ impl FunctionReq {
                     },
                 },
                 depth_ref: depth_ref.map(|e| expr_map[&e]),
+                clamp_to_edge: *clamp_to_edge,
             },
             Expression::ImageLoad {
                 image,
@@ -251,6 +254,7 @@ impl FunctionReq {
             Expression::Override(_) => expr.clone(),
             Expression::SubgroupBallotResult => expr.clone(),
             Expression::SubgroupOperationResult { .. } => expr.clone(),
+            Expression::RayQueryVertexPositions { query, committed } => todo!(),
         }
     }
 
@@ -402,7 +406,19 @@ impl FunctionReq {
                 }
             }
             (Statement::Kill, _) => Some(Statement::Kill),
-            (Statement::Barrier(b), _) => Some(Statement::Barrier(*b)),
+            // (Statement::Barrier(b), _) => Some(Statement::Barrier(*b)),
+            (Statement::ControlBarrier(b), _) => Some(Statement::ControlBarrier(*b)),
+            (Statement::MemoryBarrier(b), _) => Some(Statement::MemoryBarrier(*b)),
+            (
+                Statement::ImageAtomic {
+                    image,
+                    coordinate,
+                    array_index,
+                    fun,
+                    value,
+                },
+                _,
+            ) => todo!(),
             (Statement::Store { pointer, value }, StatementReq::Store(b)) => {
                 if !b {
                     return None;
@@ -1271,6 +1287,7 @@ impl<'a> Pruner<'a> {
                 offset: _offset,
                 level,
                 depth_ref,
+                clamp_to_edge,
             } => {
                 self.add_expression(function, func_req, context, *image, &PartReq::All);
                 self.add_expression(function, func_req, context, *sampler, &PartReq::All);
@@ -1384,6 +1401,7 @@ impl<'a> Pruner<'a> {
             Expression::SubgroupOperationResult { .. } => {
                 // nothing, handled by the statement
             }
+            Expression::RayQueryVertexPositions { query, committed } => todo!(),
         }
 
         func_req.exprs_required.insert(h_expr, part.clone());
@@ -1532,7 +1550,15 @@ impl<'a> Pruner<'a> {
                 Return(break_required)
             }
             Statement::Kill => Kill(),
-            Statement::Barrier(_) => Barrier(),
+            Statement::ControlBarrier(barrier) => Barrier(),
+            Statement::MemoryBarrier(barrier) => Barrier(),
+            Statement::ImageAtomic {
+                image,
+                coordinate,
+                array_index,
+                fun,
+                value,
+            } => todo!(),
             Statement::Store { pointer, value } => {
                 let var_ref = Self::resolve_var(function, *pointer, Vec::default());
                 let required = self.store_required(context, &var_ref);
@@ -1892,6 +1918,9 @@ impl<'a> Pruner<'a> {
                     early_depth_test: ep.early_depth_test,
                     workgroup_size: ep.workgroup_size,
                     function: mapped_f,
+                    mesh_info: None,
+                    task_payload: None,
+                    workgroup_size_overrides: None,
                 };
                 entry_points.push(new_ep);
             }