wgpu 29 update

Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "naga_oil"
-version = "0.21.0"
+version = "0.22.0"
 edition = "2021"
 license = "MIT OR Apache-2.0"
 description = "a crate for combining and manipulating shaders using naga IR"
@@ -18,7 +18,7 @@ prune = []
 allow_deprecated = []
 
 [dependencies]
-naga = { version = "28", features = ["wgsl-in", "wgsl-out", "termcolor"] }
+naga = { version = "29", features = ["wgsl-in", "wgsl-out", "termcolor"] }
 tracing = "0.1"
 regex = "1.8"
 thiserror = "2.0"
@@ -29,6 +29,6 @@ unicode-ident = "1"
 indexmap = "2"
 
 [dev-dependencies]
-wgpu = { version = "28", features = ["naga-ir"] }
+wgpu = { version = "29", features = ["naga-ir"] }
 futures-lite = "2"
 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::new_without_display_handle());
     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::new_without_display_handle());
     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::new_without_display_handle());
         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/atomics.txt

@@ -6,26 +6,26 @@ fn entry_pointX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX() -> f32 {
     atomicStore((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 1u);
     let _e3: u32 = atomicLoad((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX));
     y = _e3;
-    let _e7: u32 = atomicAdd((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 2u);
-    let _e8: u32 = y;
-    y = (_e8 + _e7);
-    let _e12: u32 = atomicSub((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 1u);
-    let _e13: u32 = y;
-    y = (_e13 + _e12);
-    let _e17: u32 = atomicMax((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 5u);
-    let _e18: u32 = y;
-    y = (_e18 + _e17);
-    let _e22: u32 = atomicMin((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 4u);
-    let _e23: u32 = y;
-    y = (_e23 + _e22);
+    let _e5: u32 = y;
+    let _e8: u32 = atomicAdd((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 2u);
+    y = (_e5 + _e8);
+    let _e10: u32 = y;
+    let _e13: u32 = atomicSub((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 1u);
+    y = (_e10 + _e13);
+    let _e15: u32 = y;
+    let _e18: u32 = atomicMax((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 5u);
+    y = (_e15 + _e18);
+    let _e20: u32 = y;
+    let _e23: u32 = atomicMin((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 4u);
+    y = (_e20 + _e23);
     let _e25: u32 = y;
-    let _e27: u32 = atomicExchange((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), _e25);
-    let _e28: u32 = y;
-    y = (_e28 + _e27);
+    let _e26: u32 = y;
+    let _e28: u32 = atomicExchange((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), _e26);
+    y = (_e25 + _e28);
     let _e33: _atomic_compare_exchange_result_Uint_4_ = atomicCompareExchangeWeak((&atomX_naga_oil_mod_XORSXG5C7NVXWI5LMMUX), 12u, 0u);
     if _e33.exchanged {
-        let _e36: u32 = y;
-        y = (_e36 + _e33.old_value);
+        let _e35: u32 = y;
+        y = (_e35 + _e33.old_value);
     }
     let _e38: u32 = y;
     return f32(_e38);

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/compose/tests/expected/use_shared_global.txt

@@ -1,8 +1,8 @@
 var<private> aX_naga_oil_mod_XNVXWIX: f32 = 0f;
 
 fn add() {
-    let _e2: f32 = aX_naga_oil_mod_XNVXWIX;
-    aX_naga_oil_mod_XNVXWIX = (_e2 + 1f);
+    let _e1: f32 = aX_naga_oil_mod_XNVXWIX;
+    aX_naga_oil_mod_XNVXWIX = (_e1 + 1f);
     return;
 }
 

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);
@@ -212,6 +222,7 @@ impl<'a> DerivedModule<'a> {
                 binding: gv.binding,
                 ty: self.import_type(&gv.ty),
                 init: gv.init.map(|c| self.import_global_expression(c)),
+                memory_decorations: gv.memory_decorations,
             };
 
             let span = self
@@ -319,7 +330,6 @@ impl<'a> DerivedModule<'a> {
             .iter()
             .map(|stmt| {
                 match stmt {
-                    // remap function calls
                     Statement::Call {
                         function,
                         arguments,
@@ -329,8 +339,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 +369,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 +519,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 +734,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 +768,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 +936,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 { .. } => (),
             }
         }
     }