diff --git a/Cargo.toml b/Cargo.toml index 52f5e59..4bfc3b3 100644 --- a/Cargo.toml +++ b/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"] } diff --git a/examples/pbr_compose_test.rs b/examples/pbr_compose_test.rs index ba9037a..c161e20 100644 --- a/examples/pbr_compose_test.rs +++ b/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() diff --git a/src/compose/mod.rs b/src/compose/mod.rs index 7000d57..3e03e6e 100644 --- a/src/compose/mod.rs +++ b/src/compose/mod.rs @@ -275,6 +275,8 @@ pub struct ComposableModuleDefinition { modules: HashMap, // 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, - wgsl_directives: Option, + wgsl_directives: &WgslDirectives, ) -> Result { 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, - wgsl_directives: Option, ) -> Result { 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 { diff --git a/src/compose/test.rs b/src/compose/test.rs index 9ea3c65..2879309 100644 --- a/src/compose/test.rs +++ b/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, diff --git a/src/compose/tests/raycast/mod.wgsl b/src/compose/tests/raycast/mod.wgsl index 9cdac54..c7a0c7e 100644 --- a/src/compose/tests/raycast/mod.wgsl +++ b/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; diff --git a/src/compose/tests/raycast/top.wgsl b/src/compose/tests/raycast/top.wgsl index 1e4777f..e148964 100644 --- a/src/compose/tests/raycast/top.wgsl +++ b/src/compose/tests/raycast/top.wgsl @@ -1,3 +1,5 @@ +enable wgpu_ray_query; + #import test_module fn main() -> f32 { diff --git a/src/derive.rs b/src/derive.rs index 3a866b3..00e2316 100644 --- a/src/derive.rs +++ b/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(); diff --git a/src/prune/mod.rs b/src/prune/mod.rs index 9241675..fd6c3c1 100644 --- a/src/prune/mod.rs +++ b/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); }