naga 28 update (#132)

This is a working naga 28 update. I noticed some tests haven't passed
(specifically cargo test --all-features) since before 0.14, so this PR
doesn't attempt to make them pass.

## enumerate_adaptors

`instance.enumerate_adapters` is async now (and available on webgpu):
gfx-rs/wgpu#8230 . more details in the wgpu
release notes.

## ControlBarrier & MemoryBarrier

Barrier was split in two to support MemoryBarriers:
gfx-rs/wgpu#7630

From the PR, it seems like falling back to ControlBarrier is fine so
that's what I did.

## Ray Query enable

ray queries require `enable wgpu_ray_query;`:
gfx-rs/wgpu#8545

This doesn't currently seem to make it through, and the relevant test
fails.

## ImageAtomic

Image atomics were added in gfx-rs/wgpu#6706

## Mesh Shaders

Mesh shaders are a major feature of wgpu 28, ~~but I've set their fields
to None here in the interest of doing an upgrade and not a feature add
at the same time~~

https://github.com/gfx-rs/wgpu/releases/tag/v28.0.0

update: I found some time and built a wgpu mesh/task shader demo and
used that to validate some of the mesh shader functionality. I've used
this to successfully compile a task shader with naga-oil and run it, but
there's still something missing from the mesh shader module output here.

---------

Co-authored-by: robtfm <50659922+robtfm@users.noreply.github.com>

Author: ChristopherBiscardi

Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "naga_oil"
-version = "0.20.0"
+version = "0.21.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 = "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 +29,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

@@ -275,6 +275,8 @@ pub struct ComposableModuleDefinition {
     modules: HashMap<ModuleKey, ComposableModule>,
     // used in spans when this module is included
     module_index: usize,
+    // any directives used in this module
+    wgsl_directives: WgslDirectives,
 }
 
 impl ComposableModuleDefinition {
@@ -533,6 +535,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);
@@ -583,17 +587,12 @@ impl Composer {
         language: ShaderLanguage,
         imports: &[ImportDefinition],
         shader_defs: &HashMap<String, ShaderDefValue>,
-        wgsl_directives: Option<WgslDirectives>,
+        wgsl_directives: &WgslDirectives,
     ) -> Result<IrBuildResult, ComposerError> {
         debug!("creating IR for {} with defs: {:?}", name, shader_defs);
 
-        let mut wgsl_string = String::new();
-        if let Some(wgsl_directives) = &wgsl_directives {
-            trace!("adding WGSL directives for {}", name);
-            wgsl_string = wgsl_directives.to_wgsl_string();
-        }
         let mut module_string = match language {
-            ShaderLanguage::Wgsl => wgsl_string,
+            ShaderLanguage::Wgsl => wgsl_directives.to_wgsl_string(),
             #[cfg(feature = "glsl")]
             ShaderLanguage::Glsl => String::from("#version 450\n"),
         };
@@ -851,7 +850,6 @@ impl Composer {
         demote_entrypoints: bool,
         source: &str,
         imports: Vec<ImportDefWithOffset>,
-        wgsl_directives: Option<WgslDirectives>,
     ) -> Result<ComposableModule, ComposerError> {
         let mut imports: Vec<_> = imports
             .into_iter()
@@ -985,7 +983,7 @@ impl Composer {
             module_definition.language,
             &imports,
             shader_defs,
-            wgsl_directives,
+            &module_definition.wgsl_directives,
         )?;
 
         // from here on errors need to be reported using the modified source with start_offset
@@ -1387,7 +1385,6 @@ impl Composer {
             true,
             &preprocessed_source,
             imports,
-            None,
         )
         .map_err(|err| err.into())
     }
@@ -1532,6 +1529,7 @@ impl Composer {
             mut imports,
             mut effective_defs,
             cleaned_source,
+            wgsl_directives,
             ..
         } = self
             .preprocessor
@@ -1617,6 +1615,7 @@ impl Composer {
             shader_defs,
             module_index,
             modules: Default::default(),
+            wgsl_directives,
         };
 
         // invalidate dependent modules if this module already exists
@@ -1763,6 +1762,7 @@ impl Composer {
             all_imports: Default::default(),
             shader_defs: Default::default(),
             modules: Default::default(),
+            wgsl_directives,
         };
 
         let PreprocessOutput {
@@ -1789,7 +1789,6 @@ impl Composer {
                 false,
                 &preprocessed_source,
                 imports,
-                Some(wgsl_directives),
             )
             .map_err(|e| ComposerError {
                 inner: e.inner,
@@ -1834,6 +1833,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

@@ -1489,11 +1489,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 +1533,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/mod.wgsl

@@ -1,3 +1,5 @@
+enable wgpu_ray_query;
+
 #define_import_path test_module
 
 @group(0) @binding(0) var tlas: acceleration_structure;

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,18 @@ impl FunctionReq {
                 }
             }
             (Statement::Kill, _) => Some(Statement::Kill),
-            (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 +1286,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 +1400,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 +1549,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 +1917,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);
             }