naga-29 prep

Cargo.toml

@@ -32,3 +32,7 @@ indexmap = "2"
 wgpu = { version = "28", features = ["naga-ir"] }
 futures-lite = "2"
 tracing-subscriber = { version = "0.3", features = ["std", "fmt"] }
+
+[patch."crates-io"]
+wgpu = { git = "https://github.com/gfx-rs/wgpu", branch = "trunk" }
+naga = { git = "https://github.com/gfx-rs/wgpu", branch = "trunk" }

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 =
         futures_lite::future::block_on(instance.enumerate_adapters(wgpu::Backends::all()))
             .into_iter()
@@ -162,7 +162,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 =
         futures_lite::future::block_on(instance.enumerate_adapters(wgpu::Backends::all()))
             .into_iter()

src/compose/mod.rs

@@ -537,6 +537,7 @@ impl Composer {
                     workgroup_size_overrides: None,
                     mesh_info: None,
                     task_payload: None,
+                    incoming_ray_payload: None,
                 };
 
                 naga_module.entry_points.push(ep);
@@ -1835,6 +1836,7 @@ impl Composer {
                 workgroup_size_overrides: ep.workgroup_size_overrides,
                 mesh_info: ep.mesh_info.clone(),
                 task_payload: ep.task_payload,
+                incoming_ray_payload: ep.incoming_ray_payload,
             });
         }
         let mut naga_module = naga::Module {

src/compose/test.rs

@@ -1488,7 +1488,7 @@ mod test {
             })
             .unwrap();
 
-        let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::default());
+        let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::default());
         let adapter =
             futures_lite::future::block_on(instance.enumerate_adapters(wgpu::Backends::all()))
                 .into_iter()
@@ -1532,7 +1532,7 @@ mod test {
             layout: Some(
                 &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                     label: None,
-                    bind_group_layouts: &[&layout],
+                    bind_group_layouts: &[Some(&layout)],
                     immediate_size: 0,
                 }),
             ),

src/compose/tests/expected/err_validation_1.txt

@@ -4,7 +4,8 @@ error: failed to build a valid final module: Function [1] 'func' is invalid
 7 │ ╭ fn func() -> f32 {
 8 │ │     return 1u;
   │ │            ^^ naga::ir::Expression [0]
-  │ ╰──────────────^ naga::ir::Function [1]
+9 │ │ }
+  │ ╰─^ naga::ir::Function [1]
   │  
   = The `return` expression Some([0]) does not match the declared return type Some([0])
 

src/compose/tests/expected/err_validation_2.txt

@@ -4,7 +4,8 @@ error: failed to build a valid final module: Function [0] 'valid_inc::func' is i
 7 │ ╭ fn func() -> f32 {
 8 │ │     return 1u;
   │ │            ^^ naga::ir::Expression [0]
-  │ ╰──────────────^ naga::ir::Function [0]
+9 │ │ }
+  │ ╰─^ naga::ir::Function [0]
   │  
   = The `return` expression Some([0]) does not match the declared return type Some([0])
 

src/derive.rs

@@ -131,7 +131,6 @@ impl<'a> DerivedModule<'a> {
                     | TypeInner::Atomic { .. }
                     | TypeInner::AccelerationStructure { .. }
                     | TypeInner::RayQuery { .. } => ty.inner.clone(),
-
                     TypeInner::Pointer { base, space } => TypeInner::Pointer {
                         base: self.import_type(base),
                         space: *space,
@@ -160,6 +159,17 @@ impl<'a> DerivedModule<'a> {
                         base: self.import_type(base),
                         size: *size,
                     },
+                    TypeInner::CooperativeMatrix {
+                        columns,
+                        rows,
+                        scalar,
+                        role,
+                    } => TypeInner::CooperativeMatrix {
+                        columns: *columns,
+                        rows: *rows,
+                        scalar: *scalar,
+                        role: *role,
+                    },
                 },
             };
             let span = self.shader.as_ref().unwrap().types.get_span(*h_type);
@@ -319,7 +329,6 @@ impl<'a> DerivedModule<'a> {
             .iter()
             .map(|stmt| {
                 match stmt {
-                    // remap function calls
                     Statement::Call {
                         function,
                         arguments,
@@ -329,8 +338,6 @@ impl<'a> DerivedModule<'a> {
                         arguments: arguments.iter().map(|expr| map_expr!(expr)).collect(),
                         result: result.as_ref().map(|result| map_expr!(result)),
                     },
-
-                    // recursively
                     Statement::Block(b) => Statement::Block(map_block!(b)),
                     Statement::If {
                         condition,
@@ -361,8 +368,6 @@ impl<'a> DerivedModule<'a> {
                         continuing: map_block!(continuing),
                         break_if: map_expr_opt!(break_if),
                     },
-
-                    // map expressions
                     Statement::Emit(exprs) => {
                         // iterate once to add expressions that should NOT be part of the emit statement
                         for expr in exprs.clone() {
@@ -513,12 +518,18 @@ impl<'a> DerivedModule<'a> {
                             value: map_expr!(value),
                         }
                     }
-                    // else just copy
                     Statement::Break
                     | Statement::Continue
                     | Statement::Kill
                     | Statement::MemoryBarrier(_)
                     | Statement::ControlBarrier(_) => stmt.clone(),
+                    Statement::RayPipelineFunction(ray_pipeline_function) => {
+                        Statement::RayPipelineFunction(*ray_pipeline_function)
+                    }
+                    Statement::CooperativeStore { target, data } => Statement::CooperativeStore {
+                        target: map_expr!(target),
+                        data: *data,
+                    },
                 }
             })
             .collect();
@@ -722,12 +733,10 @@ impl<'a> DerivedModule<'a> {
                 convert: *convert,
             },
             Expression::ArrayLength(expr) => Expression::ArrayLength(map_expr!(expr)),
-
             Expression::LocalVariable(_) | Expression::FunctionArgument(_) => {
                 is_external = true;
                 expr.clone()
             }
-
             Expression::AtomicResult { ty, comparison } => Expression::AtomicResult {
                 ty: self.import_type(ty),
                 comparison: *comparison,
@@ -758,6 +767,22 @@ impl<'a> DerivedModule<'a> {
             Expression::SubgroupOperationResult { ty } => Expression::SubgroupOperationResult {
                 ty: self.import_type(ty),
             },
+            Expression::CooperativeLoad {
+                columns,
+                rows,
+                role,
+                data,
+            } => Expression::CooperativeLoad {
+                columns: *columns,
+                rows: *rows,
+                role: *role,
+                data: *data,
+            },
+            Expression::CooperativeMultiplyAdd { a, b, c } => Expression::CooperativeMultiplyAdd {
+                a: map_expr!(a),
+                b: map_expr!(b),
+                c: map_expr!(c),
+            },
         };
 
         if !non_emitting_only || is_external {
@@ -910,6 +935,7 @@ impl<'a> DerivedModule<'a> {
                 workgroup_size_overrides: ep.workgroup_size_overrides,
                 mesh_info: ep.mesh_info.clone(),
                 task_payload: ep.task_payload,
+                incoming_ray_payload: ep.incoming_ray_payload,
             })
             .collect();
 

src/redirect.rs

@@ -73,7 +73,9 @@ impl Redirector {
                 | Statement::SubgroupBallot { .. }
                 | Statement::SubgroupGather { .. }
                 | Statement::SubgroupCollectiveOperation { .. }
-                | Statement::ImageAtomic { .. } => (),
+                | Statement::ImageAtomic { .. }
+                | Statement::RayPipelineFunction(..)
+                | Statement::CooperativeStore { .. } => (),
             }
         }
     }