From c105b0380f6a9a703990e7c7a934af74b7b829f3 Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Sat, 7 Mar 2026 14:51:26 -0800 Subject: [PATCH 1/8] naga-29 prep --- Cargo.toml | 4 ++ examples/pbr_compose_test.rs | 4 +- src/compose/mod.rs | 2 + src/compose/test.rs | 4 +- .../tests/expected/err_validation_1.txt | 3 +- .../tests/expected/err_validation_2.txt | 3 +- src/derive.rs | 44 +++++++++++++++---- src/redirect.rs | 4 +- 8 files changed, 52 insertions(+), 16 deletions(-) diff --git a/Cargo.toml b/Cargo.toml index 4bfc3b3..f58a4b8 100644 --- a/Cargo.toml +++ b/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" } diff --git a/examples/pbr_compose_test.rs b/examples/pbr_compose_test.rs index c161e20..e6e0975 100644 --- a/examples/pbr_compose_test.rs +++ b/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() diff --git a/src/compose/mod.rs b/src/compose/mod.rs index 3e03e6e..33022d3 100644 --- a/src/compose/mod.rs +++ b/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 { diff --git a/src/compose/test.rs b/src/compose/test.rs index 2879309..6e36969 100644 --- a/src/compose/test.rs +++ b/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, }), ), diff --git a/src/compose/tests/expected/err_validation_1.txt b/src/compose/tests/expected/err_validation_1.txt index c5067b6..38dfe0d 100644 --- a/src/compose/tests/expected/err_validation_1.txt +++ b/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]) diff --git a/src/compose/tests/expected/err_validation_2.txt b/src/compose/tests/expected/err_validation_2.txt index 040af39..69e3470 100644 --- a/src/compose/tests/expected/err_validation_2.txt +++ b/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]) diff --git a/src/derive.rs b/src/derive.rs index 00e2316..c3a5c90 100644 --- a/src/derive.rs +++ b/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.clone()) + } + 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(); diff --git a/src/redirect.rs b/src/redirect.rs index 96256b5..ec487f8 100644 --- a/src/redirect.rs +++ b/src/redirect.rs @@ -73,7 +73,9 @@ impl Redirector { | Statement::SubgroupBallot { .. } | Statement::SubgroupGather { .. } | Statement::SubgroupCollectiveOperation { .. } - | Statement::ImageAtomic { .. } => (), + | Statement::ImageAtomic { .. } + | Statement::RayPipelineFunction(..) + | Statement::CooperativeStore { .. } => (), } } } From dd20c3c6cb7a1c5431da3f54105fc9e20eb25f81 Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Sat, 7 Mar 2026 14:56:07 -0800 Subject: [PATCH 2/8] clippy --- src/derive.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/derive.rs b/src/derive.rs index c3a5c90..241eeb7 100644 --- a/src/derive.rs +++ b/src/derive.rs @@ -524,7 +524,7 @@ impl<'a> DerivedModule<'a> { | Statement::MemoryBarrier(_) | Statement::ControlBarrier(_) => stmt.clone(), Statement::RayPipelineFunction(ray_pipeline_function) => { - Statement::RayPipelineFunction(ray_pipeline_function.clone()) + Statement::RayPipelineFunction(*ray_pipeline_function) } Statement::CooperativeStore { target, data } => Statement::CooperativeStore { target: map_expr!(target), From ecb21871f2777d256bc1a896c8a03356d3096c77 Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Sat, 14 Mar 2026 20:16:15 -0700 Subject: [PATCH 3/8] coherent and volatile memory decorations --- src/derive.rs | 1 + 1 file changed, 1 insertion(+) diff --git a/src/derive.rs b/src/derive.rs index 241eeb7..e11c471 100644 --- a/src/derive.rs +++ b/src/derive.rs @@ -222,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 From 46f200eb828f809803b4623355f683aa2472d938 Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Sat, 14 Mar 2026 20:21:40 -0700 Subject: [PATCH 4/8] memory decorations and instance descriptor --- examples/pbr_compose_test.rs | 4 ++-- src/compose/test.rs | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/examples/pbr_compose_test.rs b/examples/pbr_compose_test.rs index e6e0975..e6b934a 100644 --- a/examples/pbr_compose_test.rs +++ b/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() diff --git a/src/compose/test.rs b/src/compose/test.rs index 6e36969..fa206e9 100644 --- a/src/compose/test.rs +++ b/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() From 32c55eef3a6bac5001d90ac20337b589b7984114 Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Mon, 16 Mar 2026 20:08:41 -0700 Subject: [PATCH 5/8] variable name changes --- src/compose/tests/expected/atomics.txt | 34 +++++++++---------- .../tests/expected/use_shared_global.txt | 4 +-- 2 files changed, 19 insertions(+), 19 deletions(-) diff --git a/src/compose/tests/expected/atomics.txt b/src/compose/tests/expected/atomics.txt index de7a9d2..c1a1ad1 100644 --- a/src/compose/tests/expected/atomics.txt +++ b/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); diff --git a/src/compose/tests/expected/use_shared_global.txt b/src/compose/tests/expected/use_shared_global.txt index c1f04e5..4ff5b5f 100644 --- a/src/compose/tests/expected/use_shared_global.txt +++ b/src/compose/tests/expected/use_shared_global.txt @@ -1,8 +1,8 @@ var 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; } From d116603149f3d010287617d448e8b1a1c93b7cfa Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Mon, 16 Mar 2026 20:10:20 -0700 Subject: [PATCH 6/8] clippy --- src/derive.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/derive.rs b/src/derive.rs index e11c471..06fc674 100644 --- a/src/derive.rs +++ b/src/derive.rs @@ -222,7 +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 + memory_decorations: gv.memory_decorations, }; let span = self From d4ebd9a3965430c61aba314dc1def7aa07ae262b Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Wed, 18 Mar 2026 18:07:43 -0700 Subject: [PATCH 7/8] bump to wgpu 29 and remove patches --- Cargo.toml | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/Cargo.toml b/Cargo.toml index f58a4b8..042969f 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -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,10 +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"] } - -[patch."crates-io"] -wgpu = { git = "https://github.com/gfx-rs/wgpu", branch = "trunk" } -naga = { git = "https://github.com/gfx-rs/wgpu", branch = "trunk" } From e43d2fd84848404de7e6ac4230de3f9c97f8c8c8 Mon Sep 17 00:00:00 2001 From: Christopher Biscardi Date: Wed, 18 Mar 2026 18:14:35 -0700 Subject: [PATCH 8/8] bump to v22 for naga_oil --- Cargo.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Cargo.toml b/Cargo.toml index 042969f..aa22519 100644 --- a/Cargo.toml +++ b/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"