Naga mesh shader SPIR-V writer

CHANGELOG.md

@@ -139,8 +139,8 @@ By @SupaMaggie70Incorporated in [#8206](https://github.com/gfx-rs/wgpu/pull/8206
 - Removed three features from `wgpu-hal` which did nothing useful: `"cargo-clippy"`, `"gpu-allocator"`, and `"rustc-hash"`. By @kpreid in [#8357](https://github.com/gfx-rs/wgpu/pull/8357).
 - `wgpu_types::PollError` now always implements the `Error` trait. By @kpreid in [#8384](https://github.com/gfx-rs/wgpu/pull/8384).
 - The texture subresources used by the color attachments of a render pass are no longer allowed to overlap when accessed via different texture views. By @andyleiserson in [#8402](https://github.com/gfx-rs/wgpu/pull/8402).
-- The `STORAGE_READ_ONLY` texture usage is now permitted to coexist with other read-only usages. By @andyleiserson in [#8490](https://github.com/gfx-rs/wgpu/pull/8490).
 - Validate that buffers are unmapped in `write_buffer` calls. By @ErichDonGubler in [#8454](https://github.com/gfx-rs/wgpu/pull/8454).
+- Add WGSL parsing for mesh shaders. By @inner-daemons in [#8370](https://github.com/gfx-rs/wgpu/pull/8370).
 
 #### DX12
 
@@ -149,6 +149,7 @@ By @SupaMaggie70Incorporated in [#8206](https://github.com/gfx-rs/wgpu/pull/8206
 #### Vulkan
 
 - Fixed a validation error regarding atomic memory semantics. By @atlv24 in [#8391](https://github.com/gfx-rs/wgpu/pull/8391).
+- Add mesh shader writer support, allowing WGSL shaders to be used on the vulkan backend. By @inner-daemons in [#8456](https://github.com/gfx-rs/wgpu/pull/8456).
 
 #### WebGPU
 

examples/features/src/lib.rs

@@ -49,6 +49,7 @@ fn all_tests() -> Vec<wgpu_test::GpuTestInitializer> {
  cube::TEST,
  cube::TEST_LINES,
  hello_synchronization::tests::SYNC,
+ mesh_shader::TEST,
  mipmap::TEST,
  mipmap::TEST_QUERY,
  msaa_line::TEST,

examples/features/src/mesh_shader/mod.rs

@@ -1,32 +1,9 @@
-use std::process::Stdio;
-
 // Same as in mesh shader tests
-fn compile_glsl(device: &wgpu::Device, shader_stage: &'static str) -> wgpu::ShaderModule {
- let cmd = std::process::Command::new("glslc")
- .args([
- &format!(
- "{}/src/mesh_shader/shader.{shader_stage}",
- env!("CARGO_MANIFEST_DIR")
- ),
- "-o",
- "-",
- "--target-env=vulkan1.2",
- "--target-spv=spv1.4",
- ])
- .stdin(Stdio::piped())
- .stdout(Stdio::piped())
- .spawn()
- .expect("Failed to call glslc");
- let output = cmd.wait_with_output().expect("Error waiting for glslc");
- assert!(output.status.success());
- unsafe {
- device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough {
- entry_point: "main".into(),
- label: None,
- spirv: Some(wgpu::util::make_spirv_raw(&output.stdout)),
- ..Default::default()
- })
- }
+fn compile_wgsl(device: &wgpu::Device) -> wgpu::ShaderModule {
+ device.create_shader_module(wgpu::ShaderModuleDescriptor {
+ label: None,
+ source: wgpu::ShaderSource::Wgsl(include_str!("shader.wgsl").into()),
+ })
 }
 fn compile_hlsl(device: &wgpu::Device, entry: &str, stage_str: &str) -> wgpu::ShaderModule {
  let out_path = format!(
@@ -71,21 +48,22 @@ impl crate::framework::Example for Example {
  device: &wgpu::Device,
  _queue: &wgpu::Queue,
  ) -> Self {
- let (ts, ms, fs) = if adapter.get_info().backend == wgpu::Backend::Vulkan {
- (
- compile_glsl(device, "task"),
- compile_glsl(device, "mesh"),
- compile_glsl(device, "frag"),
- )
- } else if adapter.get_info().backend == wgpu::Backend::Dx12 {
- (
- compile_hlsl(device, "Task", "as"),
- compile_hlsl(device, "Mesh", "ms"),
- compile_hlsl(device, "Frag", "ps"),
- )
- } else {
- panic!("Example can only run on vulkan or dx12");
- };
+ let (ts, ms, fs, ts_name, ms_name, fs_name) =
+ if adapter.get_info().backend == wgpu::Backend::Vulkan {
+ let s = compile_wgsl(device);
+ (s.clone(), s.clone(), s, "ts_main", "ms_main", "fs_main")
+ } else if adapter.get_info().backend == wgpu::Backend::Dx12 {
+ (
+ compile_hlsl(device, "Task", "as"),
+ compile_hlsl(device, "Mesh", "ms"),
+ compile_hlsl(device, "Frag", "ps"),
+ "main",
+ "main",
+ "main",
+ )
+ } else {
+ panic!("Example can only run on vulkan or dx12");
+ };
  let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
  label: None,
  bind_group_layouts: &[],
@@ -96,17 +74,17 @@ impl crate::framework::Example for Example {
  layout: Some(&pipeline_layout),
  task: Some(wgpu::TaskState {
  module: &ts,
- entry_point: Some("main"),
+ entry_point: Some(ts_name),
  compilation_options: Default::default(),
  }),
  mesh: wgpu::MeshState {
  module: &ms,
- entry_point: Some("main"),
+ entry_point: Some(ms_name),
  compilation_options: Default::default(),
  },
  fragment: Some(wgpu::FragmentState {
  module: &fs,
- entry_point: Some("main"),
+ entry_point: Some(fs_name),
  compilation_options: Default::default(),
  targets: &[Some(config.view_formats[0].into())],
  }),
@@ -179,3 +157,23 @@ impl crate::framework::Example for Example {
 pub fn main() {
  crate::framework::run::<Example>("mesh_shader");
 }
+
+#[cfg(test)]
+#[wgpu_test::gpu_test]
+pub static TEST: crate::framework::ExampleTestParams<Example> =
+ crate::framework::ExampleTestParams {
+ name: "mesh_shader",
+ // Generated on 1080ti on Vk/Windows
+ image_path: "/examples/features/src/mesh_shader/screenshot.png",
+ width: 1024,
+ height: 768,
+ optional_features: wgpu::Features::default(),
+ base_test_parameters: wgpu_test::TestParameters::default()
+ .features(
+ wgpu::Features::EXPERIMENTAL_MESH_SHADER
+ | wgpu::Features::EXPERIMENTAL_PASSTHROUGH_SHADERS,
+ )
+ .limits(wgpu::Limits::defaults().using_recommended_minimum_mesh_shader_values()),
+ comparisons: &[wgpu_test::ComparisonType::Mean(0.005)],
+ _phantom: std::marker::PhantomData::<Example>,
+ };

examples/features/src/mesh_shader/shader.wgsl

@@ -0,0 +1,74 @@
+enable mesh_shading;
+
+const positions = array(
+ vec4(0., 1., 0., 1.),
+ vec4(-1., -1., 0., 1.),
+ vec4(1., -1., 0., 1.)
+);
+const colors = array(
+ vec4(0., 1., 0., 1.),
+ vec4(0., 0., 1., 1.),
+ vec4(1., 0., 0., 1.)
+);
+struct TaskPayload {
+ colorMask: vec4<f32>,
+ visible: bool,
+}
+var<task_payload> taskPayload: TaskPayload;
+var<workgroup> workgroupData: f32;
+struct VertexOutput {
+ @builtin(position) position: vec4<f32>,
+ @location(0) color: vec4<f32>,
+}
+struct PrimitiveOutput {
+ @builtin(triangle_indices) index: vec3<u32>,
+ @builtin(cull_primitive) cull: bool,
+ @per_primitive @location(1) colorMask: vec4<f32>,
+}
+struct PrimitiveInput {
+ @per_primitive @location(1) colorMask: vec4<f32>,
+}
+
+@task
+@payload(taskPayload)
+@workgroup_size(1)
+fn ts_main() -> @builtin(mesh_task_size) vec3<u32> {
+ workgroupData = 1.0;
+ taskPayload.colorMask = vec4(1.0, 1.0, 0.0, 1.0);
+ taskPayload.visible = true;
+ return vec3(3, 1, 1);
+}
+
+struct MeshOutput {
+ @builtin(vertices) vertices: array<VertexOutput, 3>,
+ @builtin(primitives) primitives: array<PrimitiveOutput, 1>,
+ @builtin(vertex_count) vertex_count: u32,
+ @builtin(primitive_count) primitive_count: u32,
+}
+
+var<workgroup> mesh_output: MeshOutput;
+@mesh(mesh_output)
+@payload(taskPayload)
+@workgroup_size(1)
+fn ms_main(@builtin(local_invocation_index) index: u32, @builtin(global_invocation_id) id: vec3<u32>) {
+ mesh_output.vertex_count = 3;
+ mesh_output.primitive_count = 1;
+ workgroupData = 2.0;
+
+ mesh_output.vertices[0].position = positions[0];
+ mesh_output.vertices[0].color = colors[0] * taskPayload.colorMask;
+
+ mesh_output.vertices[1].position = positions[1];
+ mesh_output.vertices[1].color = colors[1] * taskPayload.colorMask;
+
+ mesh_output.vertices[2].position = positions[2];
+ mesh_output.vertices[2].color = colors[2] * taskPayload.colorMask;
+
+ mesh_output.primitives[0].index = vec3<u32>(0, 1, 2);
+ mesh_output.primitives[0].cull = !taskPayload.visible;
+ mesh_output.primitives[0].colorMask = vec4<f32>(1.0, 0.0, 1.0, 1.0);
+}
+@fragment
+fn fs_main(vertex: VertexOutput, primitive: PrimitiveInput) -> @location(0) vec4<f32> {
+ return vertex.color * primitive.colorMask;
+}

naga/src/back/dot/mod.rs

@@ -307,25 +307,6 @@ impl StatementGraph {
  crate::RayQueryFunction::Terminate => "RayQueryTerminate",
  }
  }
- S::MeshFunction(crate::MeshFunction::SetMeshOutputs {
- vertex_count,
- primitive_count,
- }) => {
- self.dependencies.push((id, vertex_count, "vertex_count"));
- self.dependencies
- .push((id, primitive_count, "primitive_count"));
- "SetMeshOutputs"
- }
- S::MeshFunction(crate::MeshFunction::SetVertex { index, value }) => {
- self.dependencies.push((id, index, "index"));
- self.dependencies.push((id, value, "value"));
- "SetVertex"
- }
- S::MeshFunction(crate::MeshFunction::SetPrimitive { index, value }) => {
- self.dependencies.push((id, index, "index"));
- self.dependencies.push((id, value, "value"));
- "SetPrimitive"
- }
  S::SubgroupBallot { result, predicate } => {
  if let Some(predicate) = predicate {
  self.dependencies.push((id, predicate, "predicate"));

naga/src/back/glsl/mod.rs

@@ -2675,11 +2675,6 @@ impl<'a, W: Write> Writer<'a, W> {
  self.write_image_atomic(ctx, image, coordinate, array_index, fun, value)?
  }
  Statement::RayQuery { .. } => unreachable!(),
- Statement::MeshFunction(
- crate::MeshFunction::SetMeshOutputs { .. }
- | crate::MeshFunction::SetVertex { .. }
- | crate::MeshFunction::SetPrimitive { .. },
- ) => unreachable!(),
  Statement::SubgroupBallot { result, predicate } => {
  write!(self.out, "{level}")?;
  let res_name = Baked(result).to_string();
@@ -5270,7 +5265,11 @@ const fn glsl_built_in(built_in: crate::BuiltIn, options: VaryingOptions) -> &'s
  | Bi::PointIndex
  | Bi::LineIndices
  | Bi::TriangleIndices
- | Bi::MeshTaskSize => {
+ | Bi::MeshTaskSize
+ | Bi::VertexCount
+ | Bi::PrimitiveCount
+ | Bi::Vertices
+ | Bi::Primitives => {
  unimplemented!()
  }
  }

naga/src/back/hlsl/conv.rs

@@ -187,7 +187,11 @@ impl crate::BuiltIn {
  }
  Self::CullPrimitive => "SV_CullPrimitive",
  Self::PointIndex | Self::LineIndices | Self::TriangleIndices => unimplemented!(),
- Self::MeshTaskSize => unreachable!(),
+ Self::MeshTaskSize
+ | Self::VertexCount
+ | Self::PrimitiveCount
+ | Self::Vertices
+ | Self::Primitives => unreachable!(),
  })
  }
 }

naga/src/back/hlsl/writer.rs

@@ -2608,19 +2608,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
  writeln!(self.out, ".Abort();")?;
  }
  },
- Statement::MeshFunction(crate::MeshFunction::SetMeshOutputs {
- vertex_count,
- primitive_count,
- }) => {
- write!(self.out, "{level}SetMeshOutputCounts(")?;
- self.write_expr(module, vertex_count, func_ctx)?;
- write!(self.out, ", ")?;
- self.write_expr(module, primitive_count, func_ctx)?;
- write!(self.out, ");")?;
- }
- Statement::MeshFunction(
- crate::MeshFunction::SetVertex { .. } | crate::MeshFunction::SetPrimitive { .. },
- ) => unimplemented!(),
  Statement::SubgroupBallot { result, predicate } => {
  write!(self.out, "{level}")?;
  let name = Baked(result).to_string();

naga/src/back/msl/mod.rs

@@ -714,7 +714,11 @@ impl ResolvedBinding {
  Bi::CullPrimitive => "primitive_culled",
  // TODO: figure out how to make this written as a function call
  Bi::PointIndex | Bi::LineIndices | Bi::TriangleIndices => unimplemented!(),
- Bi::MeshTaskSize => unreachable!(),
+ Bi::MeshTaskSize
+ | Bi::VertexCount
+ | Bi::PrimitiveCount
+ | Bi::Vertices
+ | Bi::Primitives => unreachable!(),
  };
  write!(out, "{name}")?;
  }

naga/src/back/msl/writer.rs

@@ -4063,14 +4063,6 @@ impl<W: Write> Writer<W> {
  }
  }
  }
- // TODO: write emitters for these
- crate::Statement::MeshFunction(crate::MeshFunction::SetMeshOutputs { .. }) => {
- unimplemented!()
- }
- crate::Statement::MeshFunction(
- crate::MeshFunction::SetVertex { .. }
- | crate::MeshFunction::SetPrimitive { .. },
- ) => unimplemented!(),
  crate::Statement::SubgroupBallot { result, predicate } => {
  write!(self.out, "{level}")?;
  let name = self.namer.call("");

naga/src/back/pipeline_constants.rs

@@ -860,26 +860,6 @@ fn adjust_stmt(new_pos: &HandleVec<Expression, Handle<Expression>>, stmt: &mut S
  crate::RayQueryFunction::Terminate => {}
  }
  }
- Statement::MeshFunction(crate::MeshFunction::SetMeshOutputs {
- ref mut vertex_count,
- ref mut primitive_count,
- }) => {
- adjust(vertex_count);
- adjust(primitive_count);
- }
- Statement::MeshFunction(
- crate::MeshFunction::SetVertex {
- ref mut index,
- ref mut value,
- }
- | crate::MeshFunction::SetPrimitive {
- ref mut index,
- ref mut value,
- },
- ) => {
- adjust(index);
- adjust(value);
- }
  Statement::Break
  | Statement::Continue
  | Statement::Kill