diff --git a/CHANGELOG.md b/CHANGELOG.md
index 52793e6c50a..06e8a4184db 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -142,6 +142,7 @@ By @SupaMaggie70Incorporated in [#8206](https://github.com/gfx-rs/wgpu/pull/8206
 - 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
 
diff --git a/docs/api-specs/mesh_shading.md b/docs/api-specs/mesh_shading.md
index 4b28ec635e7..abb4fa49c1b 100644
--- a/docs/api-specs/mesh_shading.md
+++ b/docs/api-specs/mesh_shading.md
@@ -3,11 +3,10 @@
 🧪Experimental🧪
 
 `wgpu` supports an experimental version of mesh shading when `Features::EXPERIMENTAL_MESH_SHADER` is enabled.
-Currently `naga` has no support for parsing or writing mesh shaders.
-For this reason, all shaders must be created with `Device::create_shader_module_passthrough`.
+The status of the implementation is documented in [the mesh-shading issue](https://github.com/gfx-rs/wgpu/issues/7197).
 
 **Note**: The features documented here may have major bugs in them and are expected to be subject
-to breaking changes, suggestions for the API exposed by this should be posted on [the mesh-shading issue](https://github.com/gfx-rs/wgpu/issues/7197).
+to breaking changes. Suggestions for the API exposed by this should be posted on the issue above.
 
 ## Mesh shaders overview
 
@@ -83,10 +82,9 @@ An example of using mesh shaders to render a single triangle can be seen [here](
 ### Features
 * Using mesh shaders requires enabling `Features::EXPERIMENTAL_MESH_SHADER`.
 * Using mesh shaders with multiview requires enabling `Features::EXPERIMENTAL_MESH_SHADER_MULTIVIEW`.
-* Currently, only triangle rendering is tested
-* Line rendering is supported but untested
-* Point rendering is supported on vulkan. It is impossible on DirectX. Metal support hasn't been checked.
+* Using mesh shaders with point primitives requires enabling `Features::EXPERIMENTAL_MESH_SHADER_POINTS`.
 * Queries are unsupported
+* Primitive index support will be added once support lands in for them in general.
 
 ### Limits
 
@@ -97,16 +95,8 @@ An example of using mesh shaders to render a single triangle can be seen [here](
 * `max_mesh_multiview_count` - The maximum number of views used when multiview rendering with a mesh shader pipeline.
 * `max_mesh_output_layers` - the maximum number of output layers for a mesh shader pipeline.
 
-### Backend specific information
-* Only Vulkan is currently supported.
-* DirectX 12 doesn't support point rendering.
-* DirectX 12 support is planned.
-* Metal support is desired but not currently planned.
-
-
 ## Naga implementation
 
-
 ### Supported frontends
 * 🛠️ WGSL
 * ❌ SPIR-V
@@ -114,7 +104,7 @@ An example of using mesh shaders to render a single triangle can be seen [here](
 
 ### Supported backends
 * 🛠️ SPIR-V
-* ❌ HLSL
+* 🛠️ HLSL
 * ❌ MSL
 * 🚫 GLSL
 * 🚫 WGSL
@@ -128,9 +118,9 @@ An example of using mesh shaders to render a single triangle can be seen [here](
 
 The majority of changes relating to mesh shaders will be in WGSL and `naga`.
 
-Using any of these features in a `wgsl` program will require adding the `enable mesh_shading` directive to the top of a program.
+Using any of these features in a `wgsl` program will require adding the `enable wgpu_mesh_shader;` directive to the top of a program.
 
-Two new shader stages will be added to `WGSL`. Fragment shaders are also modified slightly. Both task shaders and mesh shaders are allowed to use any compute-specific functionality, such as subgroup operations.
+Two new shader stages will be added to `WGSL`. Fragment shaders are also modified slightly. Both task shaders and mesh shaders are allowed to use any compute-available functionality, including subgroup operations.
 
 ### Task shader
 
@@ -138,20 +128,24 @@ A function with the `@task` attribute is a **task shader entry point**. A mesh s
 
 A task shader entry point must have a `@workgroup_size` attribute, meeting the same requirements as one appearing on a compute shader entry point.
 
-A task shader entry point must also have a `@payload(G)` property, where `G` is the name of a global variable in the `task_payload` address space. Each task shader workgroup has its own instance of this variable, visible to all invocations in the workgroup. Whatever value the workgroup collectively stores in that global variable becomes the **task payload**, and is provided to all invocations in the mesh shader grid dispatched for the workgroup.
+A task shader entry point must also have a `@payload(G)` property, where `G` is the name of a global variable in the `task_payload` address space. Each task shader workgroup has its own instance of this variable, visible to all invocations in the workgroup. Whatever value the workgroup collectively stores in that global variable becomes the **task payload**, and is provided to all invocations in the mesh shader grid dispatched for the workgroup. A task payload variable must be at least 4 bytes in size.
 
 A task shader entry point must return a `vec3<u32>` value. The return value of each workgroup's first invocation (that is, the one whose `local_invocation_index` is `0`) is taken as the size of a **mesh shader grid** to dispatch, measured in workgroups. (If the task shader entry point returns `vec3(0, 0, 0)`, then no mesh shaders are dispatched.) Mesh shader grids are described in the next section.
 
 Each task shader workgroup dispatches an independent mesh shader grid: in mesh shader invocations, `@builtin` values like `workgroup_id` and `global_invocation_id` describe the position of the workgroup and invocation within that grid;
 and `@builtin(num_workgroups)` matches the task shader workgroup's return value. Mesh shaders dispatched for other task shader workgroups are not included in the count. If it is necessary for a mesh shader to know which task shader workgroup dispatched it, the task shader can include its own workgroup id in the task payload.
 
+Task shaders must return a value of type `vec3<u32>` decorated with `@builtin(mesh_task_size)`.
+
+Task shaders can use compute and subgroup builtin inputs, in addition to `view_index` and `draw_id`.
+
 ### Mesh shader
 
 A function with the `@mesh` attribute is a **mesh shader entry point**. Mesh shaders must not return anything.
 
 Like compute shaders, mesh shaders are invoked in a grid of workgroups, called a **mesh shader grid**. If the mesh shader pipeline has a task shader, then each task shader workgroup determines the size of a mesh shader grid to be dispatched, as described above. Otherwise, the three-component size passed to `draw_mesh_tasks`, or drawn from the indirect buffer for its indirect variants, specifies the size of the mesh shader grid directly, as the number of workgroups along each of the grid's three axes.
 
-If the mesh shader pipeline has a task shader entry point, then the pipeline's mesh shader entry point must also have a `@payload(G)` attribute, naming the same variable, and the sizes must match. Mesh shader invocations can read, but not write, this variable, which is initialized to whatever value was written to it by the task shader workgroup that dispatched this mesh shader grid.
+If the mesh shader pipeline has a task shader entry point, then the pipeline's mesh shader entry point must also have a `@payload(G)` attribute, and the sizes of the variables must match. Mesh shader invocations can read from, but not write to, this variable, which is initialized to whatever value was written to it by the task shader workgroup that dispatched this mesh shader grid.
 
 If the mesh shader pipeline does not have a task shader entry point, then the mesh shader entry point must not have any `@payload` attribute.
 
@@ -159,17 +153,19 @@ A mesh shader entry point must have the following attributes:
 
 - `@workgroup_size`: this has the same meaning as when it appears on a compute shader entry point.
 
-- `@vertex_output(V, NV)`: This indicates that the mesh shader workgroup will generate at most `NV` vertex values, each of type `V`.
-
-- `@primitive_output(P, NP)`: This indicates that the mesh shader workgroup will generate at most `NP` primitives, each of type `P`.
+- `@mesh(VAR)`: Here, `VAR` represents a workgroup variable storing the output information.
 
-Each mesh shader entry point invocation must call the `setMeshOutputs(numVertices: u32, numPrimitives: u32)` builtin function at least once. The values passed by each workgroup's first invocation (that is, the one whose `local_invocation_index` is `0`) determine how many vertices (values of type `V`) and primitives (values of type `P`) the workgroup must produce. The user can still write past these indices, but they won't be used in the output.
+All mesh shader outputs are per-workgroup, and taken from the workgroup variable specified above. The type must have exactly 4 fields:
+- A field decorated with `@builtin(vertex_count)`, with type `u32`: this field represents the number of vertices that will be drawn
+- A field decorated with `@builtin(primitive_count)`, with type `u32`: this field represents the number of primitives that will be drawn
+- A field decorated with `@builtin(vertices)`, typed as an array of `V`, where `V` is the vertex output type as specified below
+- A field decorated with `@builtin(primitives)`, typed as an array of `P`, where `P` is the primitive output type as specified below
 
-The `numVertices` and `numPrimitives` arguments must be no greater than `NV` and `NP` from the `@vertex_output` and `@primitive_output` attributes.
+For a vertex count `NV`, the first `NV` elements of the vertex array above are outputted. Therefore, `NV` must be less than or equal to the size of the vertex array. The same is true for primitives with `NP`.
 
-To produce vertex data, the workgroup as a whole must make `numVertices` calls to the `setVertex(i: u32, vertex: V)` builtin function. This establishes `vertex` as the value of the `i`'th vertex, where `i` is less than the maximum number of output vertices in the `@vertex_output` attribute. `V` is the type given in the `@vertex_output` attribute. `V` must meet the same requirements as a struct type returned by a `@vertex` entry point: all members must have either `@builtin` or `@location` attributes, there must be a `@builtin(position)`, and so on.
+The vertex output type `V` must meet the same requirements as a struct type returned by a `@vertex` entry point: all members must have either `@builtin` or `@location` attributes, there must be a `@builtin(position)`, and so on.
 
-To produce primitives, the workgroup as a whole must make `numPrimitives` calls to the `setPrimitive(i: u32, primitive: P)` builtin function. This establishes `primitive` as the value of the `i`'th primitive, where `i` is less than the maximum number of output primitives in the `@primitive_output` attribute. `P` is the type given in the `@primitive_output` attribute. `P` must be a struct type, every member of which either has a `@location` or `@builtin` attribute. The following `@builtin` attributes are allowed:
+The primitive output type `P` must be a struct type, every member of which either has a `@location` or `@builtin` attribute. All members decorated with `@location` must also be decorated with `@per_primitive`, as must the corresponding fragment input. The `@per_primitive` decoration may only be applied to members decorated with `@location`. The following `@builtin` attributes are allowed:
 
 - `triangle_indices`, `line_indices`, or `point_index`: The annotated member must be of type `vec3<u32>`, `vec2<u32>`, or `u32`.
 
@@ -179,15 +175,15 @@ To produce primitives, the workgroup as a whole must make `numPrimitives` calls
 
 - `cull_primitive`: The annotated member must be of type `bool`. If it is true, then the primitive is skipped during rendering.
 
-Every member of `P` with a `@location` attribute must either have a `@per_primitive` attribute, or be part of a struct type that appears in the primitive data as a struct member with the `@per_primitive` attribute.
-
 The `@location` attributes of `P` and `V` must not overlap, since they are merged to produce the user-defined inputs to the fragment shader.
 
-It is possible to write to the same vertex or primitive index repeatedly. Since the implicit arrays written by `setVertex` and `setPrimitive` are shared by the workgroup, data races on writes to the same index for a given type are undefined behavior.
+Mesh shaders may write to the `primitive_index` builtin. This is treated just like a field decorated with `@location`, so if the mesh shader outputs `primitive_index` the fragment shader must input it, and if the fragment shader inputs it, the mesh shader must write it (unlike vertex shader pipelines).
+
+Mesh shaders can use compute and mesh shader builtin inputs, in addition to `view_index`, and if no task shader is present, `draw_id`.
 
 ### Fragment shader
 
-Fragment shaders can access vertex output data as if it is from a vertex shader. They can also access primitive output data, provided the input is decorated with `@per_primitive`. The `@per_primitive` attribute can be applied to a value directly, such as `@per_primitive @location(1) value: vec4<f32>`, to a struct such as `@per_primitive primitive_input: PrimitiveInput` where `PrimitiveInput` is a struct containing fields decorated with `@location` and `@builtin`, or to members of a struct that are themselves decorated with `@location` or `@builtin`.
+Fragment shaders can access vertex output data as if it is from a vertex shader. They can also access primitive output data, provided the input is decorated with `@per_primitive`. The `@per_primitive` decoration may only be applied to inputs or struct members decorated with `@location`.
 
 The primitive state is part of the fragment input and must match the output of the mesh shader in the pipeline. Using `@per_primitive` also requires enabling the mesh shader extension. Additionally, the locations of vertex and primitive input cannot overlap.
 
@@ -196,75 +192,81 @@ The primitive state is part of the fragment input and must match the output of t
 The following is a full example of WGSL shaders that could be used to create a mesh shader pipeline, showing off many of the features.
 
 ```wgsl
-enable mesh_shading;
+enable wgpu_mesh_shader;
 
 const positions = array(
-	vec4(0.,1.,0.,1.),
-	vec4(-1.,-1.,0.,1.),
-	vec4(1.,-1.,0.,1.)
+    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.)
+    vec4(0., 1., 0., 1.),
+    vec4(0., 0., 1., 1.),
+    vec4(1., 0., 0., 1.)
 );
+
 struct TaskPayload {
-	colorMask: vec4<f32>,
-	visible: bool,
+    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>,
+    @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>,
+    @builtin(triangle_indices) indices: vec3<u32>,
+    @builtin(cull_primitive) cull: bool,
+    @per_primitive @location(1) colorMask: vec4<f32>,
 }
 struct PrimitiveInput {
-	@per_primitive @location(1) colorMask: vec4<f32>,
+    @per_primitive @location(1) colorMask: vec4<f32>,
 }
 
+var<task_payload> taskPayload: TaskPayload;
+var<workgroup> workgroupData: 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);
+    workgroupData = 1.0;
+    taskPayload.colorMask = vec4(1.0, 1.0, 0.0, 1.0);
+    taskPayload.visible = true;
+    return vec3(1, 1, 1);
 }
-@mesh
+
+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)
-@vertex_output(VertexOutput, 3) @primitive_output(PrimitiveOutput, 1)
 @workgroup_size(1)
 fn ms_main(@builtin(local_invocation_index) index: u32, @builtin(global_invocation_id) id: vec3<u32>) {
-	setMeshOutputs(3, 1);
-	workgroupData = 2.0;
-	var v: VertexOutput;
-
-	v.position = positions[0];
-	v.color = colors[0] * taskPayload.colorMask;
-	setVertex(0, v);
-
-	v.position = positions[1];
-	v.color = colors[1] * taskPayload.colorMask;
-	setVertex(1, v);
-
-	v.position = positions[2];
-	v.color = colors[2] * taskPayload.colorMask;
-	setVertex(2, v);
-
-	var p: PrimitiveOutput;
-	p.index = vec3<u32>(0, 1, 2);
-	p.cull = !taskPayload.visible;
-	p.colorMask = vec4<f32>(1.0, 0.0, 1.0, 1.0);
-	setPrimitive(0, p);
+    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].indices = 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;
+    return vertex.color * primitive.colorMask;
 }
 ```
diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs
index 1f1396eccff..826dad1c219 100644
--- a/naga/src/back/dot/mod.rs
+++ b/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"));
diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs
index 521e2dcade7..062734b049e 100644
--- a/naga/src/back/glsl/mod.rs
+++ b/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!()
         }
     }
diff --git a/naga/src/back/hlsl/conv.rs b/naga/src/back/hlsl/conv.rs
index b4d7af86ed6..6cd3679e817 100644
--- a/naga/src/back/hlsl/conv.rs
+++ b/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!(),
         })
     }
 }
diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs
index f13601bff91..e55472460e4 100644
--- a/naga/src/back/hlsl/writer.rs
+++ b/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();
diff --git a/naga/src/back/msl/mod.rs b/naga/src/back/msl/mod.rs
index f2df74d8c2d..2456cdbae8b 100644
--- a/naga/src/back/msl/mod.rs
+++ b/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}")?;
             }
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 484142630d2..ca7da02a930 100644
--- a/naga/src/back/msl/writer.rs
+++ b/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("");
diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs
index 109cc591e74..de643b82fab 100644
--- a/naga/src/back/pipeline_constants.rs
+++ b/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
diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs
index d0556acdc53..dd9a3811687 100644
--- a/naga/src/back/spv/block.rs
+++ b/naga/src/back/spv/block.rs
@@ -3655,7 +3655,6 @@ impl BlockContext<'_> {
                 } => {
                     self.write_subgroup_gather(mode, argument, result, &mut block)?;
                 }
-                Statement::MeshFunction(_) => unreachable!(),
             }
         }
 
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index 1beb86577c8..ee1ea847739 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -2156,7 +2156,11 @@ impl Writer {
                     | Bi::CullPrimitive
                     | Bi::PointIndex
                     | Bi::LineIndices
-                    | Bi::TriangleIndices => unreachable!(),
+                    | Bi::TriangleIndices
+                    | Bi::VertexCount
+                    | Bi::PrimitiveCount
+                    | Bi::Vertices
+                    | Bi::Primitives => unreachable!(),
                 };
 
                 self.decorate(id, Decoration::BuiltIn, &[built_in as u32]);
diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs
index d1ebf62e6ee..daf32a7116f 100644
--- a/naga/src/back/wgsl/writer.rs
+++ b/naga/src/back/wgsl/writer.rs
@@ -856,7 +856,6 @@ impl<W: Write> Writer<W> {
                 }
             }
             Statement::RayQuery { .. } => unreachable!(),
-            Statement::MeshFunction(..) => unreachable!(),
             Statement::SubgroupBallot { result, predicate } => {
                 write!(self.out, "{level}")?;
                 let res_name = Baked(result).to_string();
diff --git a/naga/src/common/wgsl/to_wgsl.rs b/naga/src/common/wgsl/to_wgsl.rs
index 25847a5df7b..5e6178c049c 100644
--- a/naga/src/common/wgsl/to_wgsl.rs
+++ b/naga/src/common/wgsl/to_wgsl.rs
@@ -194,7 +194,11 @@ impl TryToWgsl for crate::BuiltIn {
             | Bi::TriangleIndices
             | Bi::LineIndices
             | Bi::MeshTaskSize
-            | Bi::PointIndex => return None,
+            | Bi::PointIndex
+            | Bi::VertexCount
+            | Bi::PrimitiveCount
+            | Bi::Vertices
+            | Bi::Primitives => return None,
         })
     }
 }
diff --git a/naga/src/compact/mod.rs b/naga/src/compact/mod.rs
index a7d3d463f11..1c9d4b55053 100644
--- a/naga/src/compact/mod.rs
+++ b/naga/src/compact/mod.rs
@@ -132,6 +132,46 @@ pub fn compact(module: &mut crate::Module, keep_unused: KeepUnused) {
                 }
             }
 
+            if let Some(task_payload) = e.task_payload {
+                module_tracer.global_variables_used.insert(task_payload);
+            }
+            if let Some(ref mesh_info) = e.mesh_info {
+                module_tracer
+                    .global_variables_used
+                    .insert(mesh_info.output_variable);
+                module_tracer
+                    .types_used
+                    .insert(mesh_info.vertex_output_type);
+                module_tracer
+                    .types_used
+                    .insert(mesh_info.primitive_output_type);
+                if let Some(max_vertices_override) = mesh_info.max_vertices_override {
+                    module_tracer
+                        .global_expressions_used
+                        .insert(max_vertices_override);
+                }
+                if let Some(max_primitives_override) = mesh_info.max_primitives_override {
+                    module_tracer
+                        .global_expressions_used
+                        .insert(max_primitives_override);
+                }
+            }
+            if e.stage == crate::ShaderStage::Task || e.stage == crate::ShaderStage::Mesh {
+                // u32 should always be there if the module is valid, as it is e.g. the type of some expressions
+                let u32_type = module
+                    .types
+                    .iter()
+                    .find_map(|tuple| {
+                        if tuple.1.inner == crate::TypeInner::Scalar(crate::Scalar::U32) {
+                            Some(tuple.0)
+                        } else {
+                            None
+                        }
+                    })
+                    .unwrap();
+                module_tracer.types_used.insert(u32_type);
+            }
+
             let mut used = module_tracer.as_function(&e.function);
             used.trace();
             FunctionMap::from(used)
@@ -221,45 +261,6 @@ pub fn compact(module: &mut crate::Module, keep_unused: KeepUnused) {
         }
     }
 
-    for entry in &module.entry_points {
-        if let Some(task_payload) = entry.task_payload {
-            module_tracer.global_variables_used.insert(task_payload);
-        }
-        if let Some(ref mesh_info) = entry.mesh_info {
-            module_tracer
-                .types_used
-                .insert(mesh_info.vertex_output_type);
-            module_tracer
-                .types_used
-                .insert(mesh_info.primitive_output_type);
-            if let Some(max_vertices_override) = mesh_info.max_vertices_override {
-                module_tracer
-                    .global_expressions_used
-                    .insert(max_vertices_override);
-            }
-            if let Some(max_primitives_override) = mesh_info.max_primitives_override {
-                module_tracer
-                    .global_expressions_used
-                    .insert(max_primitives_override);
-            }
-        }
-        if entry.stage == crate::ShaderStage::Task || entry.stage == crate::ShaderStage::Mesh {
-            // u32 should always be there if the module is valid, as it is e.g. the type of some expressions
-            let u32_type = module
-                .types
-                .iter()
-                .find_map(|tuple| {
-                    if tuple.1.inner == crate::TypeInner::Scalar(crate::Scalar::U32) {
-                        Some(tuple.0)
-                    } else {
-                        None
-                    }
-                })
-                .unwrap();
-            module_tracer.types_used.insert(u32_type);
-        }
-    }
-
     module_tracer.type_expression_tandem();
 
     // Now that we know what is used and what is never touched,
@@ -385,6 +386,7 @@ pub fn compact(module: &mut crate::Module, keep_unused: KeepUnused) {
             module_map.globals.adjust(task_payload);
         }
         if let Some(ref mut mesh_info) = entry.mesh_info {
+            module_map.globals.adjust(&mut mesh_info.output_variable);
             module_map.types.adjust(&mut mesh_info.vertex_output_type);
             module_map
                 .types
diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs
index b370501baca..39d6065f5f0 100644
--- a/naga/src/compact/statements.rs
+++ b/naga/src/compact/statements.rs
@@ -117,20 +117,6 @@ impl FunctionTracer<'_> {
                         self.expressions_used.insert(query);
                         self.trace_ray_query_function(fun);
                     }
-                    St::MeshFunction(crate::MeshFunction::SetMeshOutputs {
-                        vertex_count,
-                        primitive_count,
-                    }) => {
-                        self.expressions_used.insert(vertex_count);
-                        self.expressions_used.insert(primitive_count);
-                    }
-                    St::MeshFunction(
-                        crate::MeshFunction::SetPrimitive { index, value }
-                        | crate::MeshFunction::SetVertex { index, value },
-                    ) => {
-                        self.expressions_used.insert(index);
-                        self.expressions_used.insert(value);
-                    }
                     St::SubgroupBallot { result, predicate } => {
                         if let Some(predicate) = predicate {
                             self.expressions_used.insert(predicate);
@@ -349,26 +335,6 @@ impl FunctionMap {
                         adjust(query);
                         self.adjust_ray_query_function(fun);
                     }
-                    St::MeshFunction(crate::MeshFunction::SetMeshOutputs {
-                        ref mut vertex_count,
-                        ref mut primitive_count,
-                    }) => {
-                        adjust(vertex_count);
-                        adjust(primitive_count);
-                    }
-                    St::MeshFunction(
-                        crate::MeshFunction::SetVertex {
-                            ref mut index,
-                            ref mut value,
-                        }
-                        | crate::MeshFunction::SetPrimitive {
-                            ref mut index,
-                            ref mut value,
-                        },
-                    ) => {
-                        adjust(index);
-                        adjust(value);
-                    }
                     St::SubgroupBallot {
                         ref mut result,
                         ref mut predicate,
diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs
index 2a3a971a8bf..ac9eaf8306f 100644
--- a/naga/src/front/spv/mod.rs
+++ b/naga/src/front/spv/mod.rs
@@ -4661,7 +4661,6 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                 | S::Atomic { .. }
                 | S::ImageAtomic { .. }
                 | S::RayQuery { .. }
-                | S::MeshFunction(..)
                 | S::SubgroupBallot { .. }
                 | S::SubgroupCollectiveOperation { .. }
                 | S::SubgroupGather { .. } => {}
diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs
index 17dab5cb0ea..0cd7e11c737 100644
--- a/naga/src/front/wgsl/error.rs
+++ b/naga/src/front/wgsl/error.rs
@@ -406,6 +406,9 @@ pub(crate) enum Error<'a> {
         accept_span: Span,
         accept_type: String,
     },
+    ExpectedGlobalVariable {
+        name_span: Span,
+    },
     StructMemberTooLarge {
         member_name_span: Span,
     },
@@ -1370,6 +1373,11 @@ impl<'a> Error<'a> {
                 ],
                 notes: vec![],
             },
+            Error::ExpectedGlobalVariable { name_span } => ParseError {
+                message: "expected global variable".to_string(),
+                labels: vec![(name_span, "variable used here".into())],
+                notes: vec![],
+            },
             Error::StructMemberTooLarge { member_name_span } => ParseError {
                 message: "struct member is too large".into(),
                 labels: vec![(member_name_span, "this member exceeds the maximum size".into())],
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
index 2066d7cf2c8..33a1de6d579 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/naga/src/front/wgsl/lower/mod.rs
@@ -1479,47 +1479,93 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
             .collect();
 
         if let Some(ref entry) = f.entry_point {
-            let workgroup_size_info = if let Some(workgroup_size) = entry.workgroup_size {
-                // TODO: replace with try_map once stabilized
-                let mut workgroup_size_out = [1; 3];
-                let mut workgroup_size_overrides_out = [None; 3];
-                for (i, size) in workgroup_size.into_iter().enumerate() {
-                    if let Some(size_expr) = size {
-                        match self.const_u32(size_expr, &mut ctx.as_const()) {
-                            Ok(value) => {
-                                workgroup_size_out[i] = value.0;
-                            }
-                            Err(err) => {
-                                if let Error::ConstantEvaluatorError(ref ty, _) = *err {
-                                    match **ty {
-                                        proc::ConstantEvaluatorError::OverrideExpr => {
-                                            workgroup_size_overrides_out[i] =
-                                                Some(self.workgroup_size_override(
-                                                    size_expr,
-                                                    &mut ctx.as_override(),
-                                                )?);
-                                        }
-                                        _ => {
-                                            return Err(err);
+            let (workgroup_size, workgroup_size_overrides) =
+                if let Some(workgroup_size) = entry.workgroup_size {
+                    // TODO: replace with try_map once stabilized
+                    let mut workgroup_size_out = [1; 3];
+                    let mut workgroup_size_overrides_out = [None; 3];
+                    for (i, size) in workgroup_size.into_iter().enumerate() {
+                        if let Some(size_expr) = size {
+                            match self.const_u32(size_expr, &mut ctx.as_const()) {
+                                Ok(value) => {
+                                    workgroup_size_out[i] = value.0;
+                                }
+                                Err(err) => {
+                                    if let Error::ConstantEvaluatorError(ref ty, _) = *err {
+                                        match **ty {
+                                            proc::ConstantEvaluatorError::OverrideExpr => {
+                                                workgroup_size_overrides_out[i] =
+                                                    Some(self.workgroup_size_override(
+                                                        size_expr,
+                                                        &mut ctx.as_override(),
+                                                    )?);
+                                            }
+                                            _ => {
+                                                return Err(err);
+                                            }
                                         }
+                                    } else {
+                                        return Err(err);
                                     }
-                                } else {
-                                    return Err(err);
                                 }
                             }
                         }
                     }
-                }
-                if workgroup_size_overrides_out.iter().all(|x| x.is_none()) {
-                    (workgroup_size_out, None)
+                    if workgroup_size_overrides_out.iter().all(|x| x.is_none()) {
+                        (workgroup_size_out, None)
+                    } else {
+                        (workgroup_size_out, Some(workgroup_size_overrides_out))
+                    }
                 } else {
-                    (workgroup_size_out, Some(workgroup_size_overrides_out))
+                    ([0; 3], None)
+                };
+
+            let mesh_info = if let Some((var_name, var_span)) = entry.mesh_output_variable {
+                let var = match ctx.globals.get(var_name) {
+                    Some(&LoweredGlobalDecl::Var(handle)) => handle,
+                    Some(_) => {
+                        return Err(Box::new(Error::ExpectedGlobalVariable {
+                            name_span: var_span,
+                        }))
+                    }
+                    None => return Err(Box::new(Error::UnknownIdent(var_span, var_name))),
+                };
+
+                let mut info = ctx.module.analyze_mesh_shader_info(var);
+                if let Some(h) = info.1[0] {
+                    info.0.max_vertices_override = Some(
+                        ctx.module
+                            .global_expressions
+                            .append(crate::Expression::Override(h), Span::UNDEFINED),
+                    );
+                }
+                if let Some(h) = info.1[1] {
+                    info.0.max_primitives_override = Some(
+                        ctx.module
+                            .global_expressions
+                            .append(crate::Expression::Override(h), Span::UNDEFINED),
+                    );
                 }
+
+                Some(info.0)
+            } else {
+                None
+            };
+
+            let task_payload = if let Some((var_name, var_span)) = entry.task_payload {
+                Some(match ctx.globals.get(var_name) {
+                    Some(&LoweredGlobalDecl::Var(handle)) => handle,
+                    Some(_) => {
+                        return Err(Box::new(Error::ExpectedGlobalVariable {
+                            name_span: var_span,
+                        }))
+                    }
+                    None => return Err(Box::new(Error::UnknownIdent(var_span, var_name))),
+                })
             } else {
-                ([0; 3], None)
+                None
             };
 
-            let (workgroup_size, workgroup_size_overrides) = workgroup_size_info;
             ctx.module.entry_points.push(ir::EntryPoint {
                 name: f.name.name.to_string(),
                 stage: entry.stage,
@@ -1527,8 +1573,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                 workgroup_size,
                 workgroup_size_overrides,
                 function,
-                mesh_info: None,
-                task_payload: None,
+                mesh_info,
+                task_payload,
             });
             Ok(LoweredGlobalDecl::EntryPoint(
                 ctx.module.entry_points.len() - 1,
@@ -4059,6 +4105,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                 interpolation,
                 sampling,
                 blend_src,
+                per_primitive,
             }) => {
                 let blend_src = if let Some(blend_src) = blend_src {
                     Some(self.const_u32(blend_src, &mut ctx.as_const())?.0)
@@ -4071,7 +4118,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                     interpolation,
                     sampling,
                     blend_src,
-                    per_primitive: false,
+                    per_primitive,
                 };
                 binding.apply_default_interpolation(&ctx.module.types[ty].inner);
                 Some(binding)
diff --git a/naga/src/front/wgsl/parse/ast.rs b/naga/src/front/wgsl/parse/ast.rs
index 345e9c4c486..04964e7ba5f 100644
--- a/naga/src/front/wgsl/parse/ast.rs
+++ b/naga/src/front/wgsl/parse/ast.rs
@@ -128,6 +128,8 @@ pub struct EntryPoint<'a> {
     pub stage: crate::ShaderStage,
     pub early_depth_test: Option<crate::EarlyDepthTest>,
     pub workgroup_size: Option<[Option<Handle<Expression<'a>>>; 3]>,
+    pub mesh_output_variable: Option<(&'a str, Span)>,
+    pub task_payload: Option<(&'a str, Span)>,
 }
 
 #[cfg(doc)]
@@ -152,6 +154,7 @@ pub enum Binding<'a> {
         interpolation: Option<crate::Interpolation>,
         sampling: Option<crate::Sampling>,
         blend_src: Option<Handle<Expression<'a>>>,
+        per_primitive: bool,
     },
 }
 
diff --git a/naga/src/front/wgsl/parse/conv.rs b/naga/src/front/wgsl/parse/conv.rs
index de07ba2e391..5431f0d2525 100644
--- a/naga/src/front/wgsl/parse/conv.rs
+++ b/naga/src/front/wgsl/parse/conv.rs
@@ -6,7 +6,11 @@ use crate::Span;
 
 use alloc::boxed::Box;
 
-pub fn map_address_space(word: &str, span: Span) -> Result<'_, crate::AddressSpace> {
+pub fn map_address_space<'a>(
+    word: &str,
+    span: Span,
+    enable_extensions: &EnableExtensions,
+) -> Result<'a, crate::AddressSpace> {
     match word {
         "private" => Ok(crate::AddressSpace::Private),
         "workgroup" => Ok(crate::AddressSpace::WorkGroup),
@@ -16,6 +20,16 @@ pub fn map_address_space(word: &str, span: Span) -> Result<'_, crate::AddressSpa
         }),
         "push_constant" => Ok(crate::AddressSpace::PushConstant),
         "function" => Ok(crate::AddressSpace::Function),
+        "task_payload" => {
+            if enable_extensions.contains(ImplementedEnableExtension::WgpuMeshShader) {
+                Ok(crate::AddressSpace::TaskPayload)
+            } else {
+                Err(Box::new(Error::EnableExtensionNotEnabled {
+                    span,
+                    kind: ImplementedEnableExtension::WgpuMeshShader.into(),
+                }))
+            }
+        }
         _ => Err(Box::new(Error::UnknownAddressSpace(span))),
     }
 }
@@ -50,6 +64,17 @@ pub fn map_built_in(
         "subgroup_id" => crate::BuiltIn::SubgroupId,
         "subgroup_size" => crate::BuiltIn::SubgroupSize,
         "subgroup_invocation_id" => crate::BuiltIn::SubgroupInvocationId,
+        // mesh
+        "cull_primitive" => crate::BuiltIn::CullPrimitive,
+        "point_index" => crate::BuiltIn::PointIndex,
+        "line_indices" => crate::BuiltIn::LineIndices,
+        "triangle_indices" => crate::BuiltIn::TriangleIndices,
+        "mesh_task_size" => crate::BuiltIn::MeshTaskSize,
+        // mesh global variable
+        "vertex_count" => crate::BuiltIn::VertexCount,
+        "vertices" => crate::BuiltIn::Vertices,
+        "primitive_count" => crate::BuiltIn::PrimitiveCount,
+        "primitives" => crate::BuiltIn::Primitives,
         _ => return Err(Box::new(Error::UnknownBuiltin(span))),
     };
     match built_in {
@@ -61,6 +86,21 @@ pub fn map_built_in(
                 }));
             }
         }
+        crate::BuiltIn::CullPrimitive
+        | crate::BuiltIn::PointIndex
+        | crate::BuiltIn::LineIndices
+        | crate::BuiltIn::TriangleIndices
+        | crate::BuiltIn::VertexCount
+        | crate::BuiltIn::Vertices
+        | crate::BuiltIn::PrimitiveCount
+        | crate::BuiltIn::Primitives => {
+            if !enable_extensions.contains(ImplementedEnableExtension::WgpuMeshShader) {
+                return Err(Box::new(Error::EnableExtensionNotEnabled {
+                    span,
+                    kind: ImplementedEnableExtension::WgpuMeshShader.into(),
+                }));
+            }
+        }
         _ => {}
     }
     Ok(built_in)
diff --git a/naga/src/front/wgsl/parse/directive/enable_extension.rs b/naga/src/front/wgsl/parse/directive/enable_extension.rs
index 38d6d6719ca..7d1e1b2df6e 100644
--- a/naga/src/front/wgsl/parse/directive/enable_extension.rs
+++ b/naga/src/front/wgsl/parse/directive/enable_extension.rs
@@ -10,6 +10,7 @@ use alloc::boxed::Box;
 /// Tracks the status of every enable-extension known to Naga.
 #[derive(Clone, Debug, Eq, PartialEq)]
 pub struct EnableExtensions {
+    wgpu_mesh_shader: bool,
     dual_source_blending: bool,
     /// Whether `enable f16;` was written earlier in the shader module.
     f16: bool,
@@ -19,6 +20,7 @@ pub struct EnableExtensions {
 impl EnableExtensions {
     pub(crate) const fn empty() -> Self {
         Self {
+            wgpu_mesh_shader: false,
             f16: false,
             dual_source_blending: false,
             clip_distances: false,
@@ -28,6 +30,7 @@ impl EnableExtensions {
     /// Add an enable-extension to the set requested by a module.
     pub(crate) fn add(&mut self, ext: ImplementedEnableExtension) {
         let field = match ext {
+            ImplementedEnableExtension::WgpuMeshShader => &mut self.wgpu_mesh_shader,
             ImplementedEnableExtension::DualSourceBlending => &mut self.dual_source_blending,
             ImplementedEnableExtension::F16 => &mut self.f16,
             ImplementedEnableExtension::ClipDistances => &mut self.clip_distances,
@@ -38,6 +41,7 @@ impl EnableExtensions {
     /// Query whether an enable-extension tracked here has been requested.
     pub(crate) const fn contains(&self, ext: ImplementedEnableExtension) -> bool {
         match ext {
+            ImplementedEnableExtension::WgpuMeshShader => self.wgpu_mesh_shader,
             ImplementedEnableExtension::DualSourceBlending => self.dual_source_blending,
             ImplementedEnableExtension::F16 => self.f16,
             ImplementedEnableExtension::ClipDistances => self.clip_distances,
@@ -70,6 +74,7 @@ impl EnableExtension {
     const F16: &'static str = "f16";
     const CLIP_DISTANCES: &'static str = "clip_distances";
     const DUAL_SOURCE_BLENDING: &'static str = "dual_source_blending";
+    const MESH_SHADER: &'static str = "wgpu_mesh_shader";
     const SUBGROUPS: &'static str = "subgroups";
     const PRIMITIVE_INDEX: &'static str = "primitive_index";
 
@@ -81,6 +86,7 @@ impl EnableExtension {
             Self::DUAL_SOURCE_BLENDING => {
                 Self::Implemented(ImplementedEnableExtension::DualSourceBlending)
             }
+            Self::MESH_SHADER => Self::Implemented(ImplementedEnableExtension::WgpuMeshShader),
             Self::SUBGROUPS => Self::Unimplemented(UnimplementedEnableExtension::Subgroups),
             Self::PRIMITIVE_INDEX => {
                 Self::Unimplemented(UnimplementedEnableExtension::PrimitiveIndex)
@@ -93,6 +99,7 @@ impl EnableExtension {
     pub const fn to_ident(self) -> &'static str {
         match self {
             Self::Implemented(kind) => match kind {
+                ImplementedEnableExtension::WgpuMeshShader => Self::MESH_SHADER,
                 ImplementedEnableExtension::DualSourceBlending => Self::DUAL_SOURCE_BLENDING,
                 ImplementedEnableExtension::F16 => Self::F16,
                 ImplementedEnableExtension::ClipDistances => Self::CLIP_DISTANCES,
@@ -126,6 +133,8 @@ pub enum ImplementedEnableExtension {
     ///
     /// [`enable clip_distances;`]: https://www.w3.org/TR/WGSL/#extension-clip_distances
     ClipDistances,
+    /// Enables the `wgpu_mesh_shader` extension, native only
+    WgpuMeshShader,
 }
 
 /// A variant of [`EnableExtension::Unimplemented`].
diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs
index c01ba4de30f..838a4dd1a88 100644
--- a/naga/src/front/wgsl/parse/mod.rs
+++ b/naga/src/front/wgsl/parse/mod.rs
@@ -178,6 +178,7 @@ struct BindingParser<'a> {
     sampling: ParsedAttribute<crate::Sampling>,
     invariant: ParsedAttribute<bool>,
     blend_src: ParsedAttribute<Handle<ast::Expression<'a>>>,
+    per_primitive: ParsedAttribute<()>,
 }
 
 impl<'a> BindingParser<'a> {
@@ -238,6 +239,18 @@ impl<'a> BindingParser<'a> {
                 lexer.skip(Token::Separator(','));
                 lexer.expect(Token::Paren(')'))?;
             }
+            "per_primitive" => {
+                if !lexer
+                    .enable_extensions
+                    .contains(ImplementedEnableExtension::WgpuMeshShader)
+                {
+                    return Err(Box::new(Error::EnableExtensionNotEnabled {
+                        span: name_span,
+                        kind: ImplementedEnableExtension::WgpuMeshShader.into(),
+                    }));
+                }
+                self.per_primitive.set((), name_span)?;
+            }
             _ => return Err(Box::new(Error::UnknownAttribute(name_span))),
         }
         Ok(())
@@ -251,9 +264,10 @@ impl<'a> BindingParser<'a> {
             self.sampling.value,
             self.invariant.value.unwrap_or_default(),
             self.blend_src.value,
+            self.per_primitive.value,
         ) {
-            (None, None, None, None, false, None) => Ok(None),
-            (Some(location), None, interpolation, sampling, false, blend_src) => {
+            (None, None, None, None, false, None, None) => Ok(None),
+            (Some(location), None, interpolation, sampling, false, blend_src, per_primitive) => {
                 // Before handing over the completed `Module`, we call
                 // `apply_default_interpolation` to ensure that the interpolation and
                 // sampling have been explicitly specified on all vertex shader output and fragment
@@ -263,17 +277,18 @@ impl<'a> BindingParser<'a> {
                     interpolation,
                     sampling,
                     blend_src,
+                    per_primitive: per_primitive.is_some(),
                 }))
             }
-            (None, Some(crate::BuiltIn::Position { .. }), None, None, invariant, None) => {
+            (None, Some(crate::BuiltIn::Position { .. }), None, None, invariant, None, None) => {
                 Ok(Some(ast::Binding::BuiltIn(crate::BuiltIn::Position {
                     invariant,
                 })))
             }
-            (None, Some(built_in), None, None, false, None) => {
+            (None, Some(built_in), None, None, false, None, None) => {
                 Ok(Some(ast::Binding::BuiltIn(built_in)))
             }
-            (_, _, _, _, _, _) => Err(Box::new(Error::InconsistentBinding(span))),
+            (_, _, _, _, _, _, _) => Err(Box::new(Error::InconsistentBinding(span))),
         }
     }
 }
@@ -1318,7 +1333,7 @@ impl Parser {
                     };
                     crate::AddressSpace::Storage { access }
                 }
-                _ => conv::map_address_space(class_str, span)?,
+                _ => conv::map_address_space(class_str, span, &lexer.enable_extensions)?,
             };
             lexer.expect(Token::Paren('>'))?;
         }
@@ -1691,7 +1706,7 @@ impl Parser {
             "ptr" => {
                 lexer.expect_generic_paren('<')?;
                 let (ident, span) = lexer.next_ident_with_span()?;
-                let mut space = conv::map_address_space(ident, span)?;
+                let mut space = conv::map_address_space(ident, span, &lexer.enable_extensions)?;
                 lexer.expect(Token::Separator(','))?;
                 let base = self.type_decl(lexer, ctx)?;
                 if let crate::AddressSpace::Storage { ref mut access } = space {
@@ -2790,12 +2805,14 @@ impl Parser {
         // read attributes
         let mut binding = None;
         let mut stage = ParsedAttribute::default();
-        let mut compute_span = Span::new(0, 0);
+        let mut compute_like_span = Span::new(0, 0);
         let mut workgroup_size = ParsedAttribute::default();
         let mut early_depth_test = ParsedAttribute::default();
         let (mut bind_index, mut bind_group) =
             (ParsedAttribute::default(), ParsedAttribute::default());
         let mut id = ParsedAttribute::default();
+        let mut payload = ParsedAttribute::default();
+        let mut mesh_output = ParsedAttribute::default();
 
         let mut must_use: ParsedAttribute<Span> = ParsedAttribute::default();
 
@@ -2854,7 +2871,51 @@ impl Parser {
                 }
                 "compute" => {
                     stage.set(ShaderStage::Compute, name_span)?;
-                    compute_span = name_span;
+                    compute_like_span = name_span;
+                }
+                "task" => {
+                    if !lexer
+                        .enable_extensions
+                        .contains(ImplementedEnableExtension::WgpuMeshShader)
+                    {
+                        return Err(Box::new(Error::EnableExtensionNotEnabled {
+                            span: name_span,
+                            kind: ImplementedEnableExtension::WgpuMeshShader.into(),
+                        }));
+                    }
+                    stage.set(ShaderStage::Task, name_span)?;
+                    compute_like_span = name_span;
+                }
+                "mesh" => {
+                    if !lexer
+                        .enable_extensions
+                        .contains(ImplementedEnableExtension::WgpuMeshShader)
+                    {
+                        return Err(Box::new(Error::EnableExtensionNotEnabled {
+                            span: name_span,
+                            kind: ImplementedEnableExtension::WgpuMeshShader.into(),
+                        }));
+                    }
+                    stage.set(ShaderStage::Mesh, name_span)?;
+                    compute_like_span = name_span;
+
+                    lexer.expect(Token::Paren('('))?;
+                    mesh_output.set(lexer.next_ident_with_span()?, name_span)?;
+                    lexer.expect(Token::Paren(')'))?;
+                }
+                "payload" => {
+                    if !lexer
+                        .enable_extensions
+                        .contains(ImplementedEnableExtension::WgpuMeshShader)
+                    {
+                        return Err(Box::new(Error::EnableExtensionNotEnabled {
+                            span: name_span,
+                            kind: ImplementedEnableExtension::WgpuMeshShader.into(),
+                        }));
+                    }
+                    lexer.expect(Token::Paren('('))?;
+                    payload.set(lexer.next_ident_with_span()?, name_span)?;
+                    lexer.expect(Token::Paren(')'))?;
                 }
                 "workgroup_size" => {
                     lexer.expect(Token::Paren('('))?;
@@ -3020,13 +3081,16 @@ impl Parser {
                 )?;
                 Some(ast::GlobalDeclKind::Fn(ast::Function {
                     entry_point: if let Some(stage) = stage.value {
-                        if stage == ShaderStage::Compute && workgroup_size.value.is_none() {
-                            return Err(Box::new(Error::MissingWorkgroupSize(compute_span)));
+                        if stage.compute_like() && workgroup_size.value.is_none() {
+                            return Err(Box::new(Error::MissingWorkgroupSize(compute_like_span)));
                         }
+
                         Some(ast::EntryPoint {
                             stage,
                             early_depth_test: early_depth_test.value,
                             workgroup_size: workgroup_size.value,
+                            mesh_output_variable: mesh_output.value,
+                            task_payload: payload.value,
                         })
                     } else {
                         None
diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs
index 4093d823b4b..5aa411b72e5 100644
--- a/naga/src/ir/mod.rs
+++ b/naga/src/ir/mod.rs
@@ -409,7 +409,7 @@ pub enum BuiltIn {
     PointCoord,
     /// Read in fragment shaders
     FrontFacing,
-    /// Read in fragment shaders, in the future may written in mesh shaders
+    /// Read in fragment shaders, written in mesh shaders
     PrimitiveIndex,
     /// Read in fragment shaders
     Barycentric,
@@ -450,6 +450,15 @@ pub enum BuiltIn {
     LineIndices,
     /// Written in mesh shaders
     TriangleIndices,
+
+    /// Written to a workgroup variable in mesh shaders
+    VertexCount,
+    /// Written to a workgroup variable in mesh shaders
+    Vertices,
+    /// Written to a workgroup variable in mesh shaders
+    PrimitiveCount,
+    /// Written to a workgroup variable in mesh shaders
+    Primitives,
 }
 
 /// Number of bytes per scalar.
@@ -2211,8 +2220,6 @@ pub enum Statement {
         /// The specific operation we're performing on `query`.
         fun: RayQueryFunction,
     },
-    /// A mesh shader intrinsic.
-    MeshFunction(MeshFunction),
     /// Calculate a bitmask using a boolean from each active thread in the subgroup
     SubgroupBallot {
         /// The [`SubgroupBallotResult`] expression representing this load's result.
@@ -2569,7 +2576,7 @@ pub struct DocComments {
 }
 
 /// The output topology for a mesh shader. Note that mesh shaders don't allow things like triangle-strips.
-#[derive(Debug, Clone, Copy)]
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
 #[cfg_attr(feature = "serialize", derive(Serialize))]
 #[cfg_attr(feature = "deserialize", derive(Deserialize))]
 #[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
@@ -2583,7 +2590,7 @@ pub enum MeshOutputTopology {
 }
 
 /// Information specific to mesh shader entry points.
-#[derive(Debug, Clone)]
+#[derive(Debug, Clone, PartialEq, Eq)]
 #[cfg_attr(feature = "serialize", derive(Serialize))]
 #[cfg_attr(feature = "deserialize", derive(Deserialize))]
 #[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
@@ -2603,29 +2610,8 @@ pub struct MeshStageInfo {
     pub vertex_output_type: Handle<Type>,
     /// The type used by primitive outputs, i.e. what is passed to `setPrimitive`.
     pub primitive_output_type: Handle<Type>,
-}
-
-/// Mesh shader intrinsics
-#[derive(Debug, Clone, Copy)]
-#[cfg_attr(feature = "serialize", derive(Serialize))]
-#[cfg_attr(feature = "deserialize", derive(Deserialize))]
-#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
-pub enum MeshFunction {
-    /// Sets the number of vertices and primitives that will be outputted.
-    SetMeshOutputs {
-        vertex_count: Handle<Expression>,
-        primitive_count: Handle<Expression>,
-    },
-    /// Sets the output vertex at a given index.
-    SetVertex {
-        index: Handle<Expression>,
-        value: Handle<Expression>,
-    },
-    /// Sets the output primitive at a given index.
-    SetPrimitive {
-        index: Handle<Expression>,
-        value: Handle<Expression>,
-    },
+    /// The global variable holding the outputted vertices, primitives, and counts
+    pub output_variable: Handle<GlobalVariable>,
 }
 
 /// Shader module.
diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs
index eca63ee4fb5..64da0a9661e 100644
--- a/naga/src/proc/mod.rs
+++ b/naga/src/proc/mod.rs
@@ -27,6 +27,8 @@ use thiserror::Error;
 pub use type_methods::min_max_float_representable_by;
 pub use typifier::{compare_types, ResolveContext, ResolveError, TypeResolution};
 
+use crate::non_max_u32::NonMaxU32;
+
 impl From<super::StorageFormat> for super::Scalar {
     fn from(format: super::StorageFormat) -> Self {
         use super::{ScalarKind as Sk, StorageFormat as Sf};
@@ -653,3 +655,178 @@ fn test_matrix_size() {
         48,
     );
 }
+
+impl crate::Module {
+    /// Extracts mesh shader info from a mesh output global variable. Used in frontends
+    /// and by validators. This only validates the output variable itself, and not the
+    /// vertex and primitive output types.
+    ///
+    /// The output contains the extracted mesh stage info, with overrides unset,
+    /// and then the overrides separately. This is because the overrides should be
+    /// treated as expressions elsewhere, but that requires mutably modifying the
+    /// module and the expressions should only be created at parse time, not validation
+    /// time.
+    #[allow(clippy::type_complexity)]
+    pub fn analyze_mesh_shader_info(
+        &self,
+        gv: crate::Handle<crate::GlobalVariable>,
+    ) -> (
+        crate::MeshStageInfo,
+        [Option<crate::Handle<crate::Override>>; 2],
+        Option<crate::WithSpan<crate::valid::EntryPointError>>,
+    ) {
+        use crate::span::AddSpan;
+        use crate::valid::EntryPointError;
+        #[derive(Default)]
+        struct OutError {
+            pub inner: Option<EntryPointError>,
+        }
+        impl OutError {
+            pub fn set(&mut self, err: EntryPointError) {
+                if self.inner.is_none() {
+                    self.inner = Some(err);
+                }
+            }
+        }
+
+        // Used to temporarily initialize stuff
+        let null_type = crate::Handle::new(NonMaxU32::new(0).unwrap());
+        let mut output = crate::MeshStageInfo {
+            topology: crate::MeshOutputTopology::Triangles,
+            max_vertices: 0,
+            max_vertices_override: None,
+            max_primitives: 0,
+            max_primitives_override: None,
+            vertex_output_type: null_type,
+            primitive_output_type: null_type,
+            output_variable: gv,
+        };
+        // Stores the error to output, if any.
+        let mut error = OutError::default();
+        let r#type = &self.types[self.global_variables[gv].ty].inner;
+
+        let mut topology = output.topology;
+        // Max, max override, type
+        let mut vertex_info = (0, None, null_type);
+        let mut primitive_info = (0, None, null_type);
+
+        match r#type {
+            &crate::TypeInner::Struct { ref members, .. } => {
+                let mut builtins = crate::FastHashSet::default();
+                for member in members {
+                    match member.binding {
+                        Some(crate::Binding::BuiltIn(crate::BuiltIn::VertexCount)) => {
+                            // Must have type u32
+                            if self.types[member.ty].inner.scalar() != Some(crate::Scalar::U32) {
+                                error.set(EntryPointError::BadMeshOutputVariableField);
+                            }
+                            // Each builtin should only occur once
+                            if builtins.contains(&crate::BuiltIn::VertexCount) {
+                                error.set(EntryPointError::BadMeshOutputVariableType);
+                            }
+                            builtins.insert(crate::BuiltIn::VertexCount);
+                        }
+                        Some(crate::Binding::BuiltIn(crate::BuiltIn::PrimitiveCount)) => {
+                            // Must have type u32
+                            if self.types[member.ty].inner.scalar() != Some(crate::Scalar::U32) {
+                                error.set(EntryPointError::BadMeshOutputVariableField);
+                            }
+                            // Each builtin should only occur once
+                            if builtins.contains(&crate::BuiltIn::PrimitiveCount) {
+                                error.set(EntryPointError::BadMeshOutputVariableType);
+                            }
+                            builtins.insert(crate::BuiltIn::PrimitiveCount);
+                        }
+                        Some(crate::Binding::BuiltIn(
+                            crate::BuiltIn::Vertices | crate::BuiltIn::Primitives,
+                        )) => {
+                            let ty = &self.types[member.ty].inner;
+                            // Analyze the array type to determine size and vertex/primitive type
+                            let (a, b, c) = match ty {
+                                &crate::TypeInner::Array { base, size, .. } => {
+                                    let ty = base;
+                                    let (max, max_override) = match size {
+                                        crate::ArraySize::Constant(a) => (a.get(), None),
+                                        crate::ArraySize::Pending(o) => (0, Some(o)),
+                                        crate::ArraySize::Dynamic => {
+                                            error.set(EntryPointError::BadMeshOutputVariableField);
+                                            (0, None)
+                                        }
+                                    };
+                                    (max, max_override, ty)
+                                }
+                                _ => {
+                                    error.set(EntryPointError::BadMeshOutputVariableField);
+                                    (0, None, null_type)
+                                }
+                            };
+                            if matches!(
+                                member.binding,
+                                Some(crate::Binding::BuiltIn(crate::BuiltIn::Primitives))
+                            ) {
+                                // Primitives require special analysis to determine topology
+                                primitive_info = (a, b, c);
+                                match self.types[c].inner {
+                                    crate::TypeInner::Struct { ref members, .. } => {
+                                        for member in members {
+                                            match member.binding {
+                                                Some(crate::Binding::BuiltIn(
+                                                    crate::BuiltIn::PointIndex,
+                                                )) => {
+                                                    topology = crate::MeshOutputTopology::Points;
+                                                }
+                                                Some(crate::Binding::BuiltIn(
+                                                    crate::BuiltIn::LineIndices,
+                                                )) => {
+                                                    topology = crate::MeshOutputTopology::Lines;
+                                                }
+                                                Some(crate::Binding::BuiltIn(
+                                                    crate::BuiltIn::TriangleIndices,
+                                                )) => {
+                                                    topology = crate::MeshOutputTopology::Triangles;
+                                                }
+                                                _ => (),
+                                            }
+                                        }
+                                    }
+                                    _ => (),
+                                }
+                                // Each builtin should only occur once
+                                if builtins.contains(&crate::BuiltIn::Primitives) {
+                                    error.set(EntryPointError::BadMeshOutputVariableType);
+                                }
+                                builtins.insert(crate::BuiltIn::Primitives);
+                            } else {
+                                vertex_info = (a, b, c);
+                                // Each builtin should only occur once
+                                if builtins.contains(&crate::BuiltIn::Vertices) {
+                                    error.set(EntryPointError::BadMeshOutputVariableType);
+                                }
+                                builtins.insert(crate::BuiltIn::Vertices);
+                            }
+                        }
+                        _ => error.set(EntryPointError::BadMeshOutputVariableType),
+                    }
+                }
+                output = crate::MeshStageInfo {
+                    topology,
+                    max_vertices: vertex_info.0,
+                    max_vertices_override: None,
+                    vertex_output_type: vertex_info.2,
+                    max_primitives: primitive_info.0,
+                    max_primitives_override: None,
+                    primitive_output_type: primitive_info.2,
+                    ..output
+                }
+            }
+            _ => error.set(EntryPointError::BadMeshOutputVariableType),
+        }
+        (
+            output,
+            [vertex_info.1, primitive_info.1],
+            error
+                .inner
+                .map(|a| a.with_span_handle(self.global_variables[gv].ty, &self.types)),
+        )
+    }
+}
diff --git a/naga/src/proc/terminator.rs b/naga/src/proc/terminator.rs
index f76d4c06a3b..b29ccb054a3 100644
--- a/naga/src/proc/terminator.rs
+++ b/naga/src/proc/terminator.rs
@@ -36,7 +36,6 @@ pub fn ensure_block_returns(block: &mut crate::Block) {
             | S::ImageStore { .. }
             | S::Call { .. }
             | S::RayQuery { .. }
-            | S::MeshFunction(..)
             | S::Atomic { .. }
             | S::ImageAtomic { .. }
             | S::WorkGroupUniformLoad { .. }
diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs
index 6ef2ca0988d..e01a7b0b735 100644
--- a/naga/src/valid/analyzer.rs
+++ b/naga/src/valid/analyzer.rs
@@ -85,25 +85,6 @@ struct FunctionUniformity {
     exit: ExitFlags,
 }
 
-/// Mesh shader related characteristics of a function.
-#[derive(Debug, Clone, Default)]
-#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
-#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
-#[cfg_attr(test, derive(PartialEq))]
-pub struct FunctionMeshShaderInfo {
-    /// The type of value this function passes to [`SetVertex`], and the
-    /// expression that first established it.
-    ///
-    /// [`SetVertex`]: crate::ir::MeshFunction::SetVertex
-    pub vertex_type: Option<(Handle<crate::Type>, Handle<crate::Expression>)>,
-
-    /// The type of value this function passes to [`SetPrimitive`], and the
-    /// expression that first established it.
-    ///
-    /// [`SetPrimitive`]: crate::ir::MeshFunction::SetPrimitive
-    pub primitive_type: Option<(Handle<crate::Type>, Handle<crate::Expression>)>,
-}
-
 impl ops::BitOr for FunctionUniformity {
     type Output = Self;
     fn bitor(self, other: Self) -> Self {
@@ -321,9 +302,6 @@ pub struct FunctionInfo {
     /// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
     /// validation.
     diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
-
-    /// Mesh shader info for this function and its callees.
-    pub mesh_shader_info: FunctionMeshShaderInfo,
 }
 
 impl FunctionInfo {
@@ -520,9 +498,6 @@ impl FunctionInfo {
             *mine |= *other;
         }
 
-        // Inherit mesh output types from our callees.
-        self.try_update_mesh_info(&callee.mesh_shader_info)?;
-
         Ok(FunctionUniformity {
             result: callee.uniformity.clone(),
             exit: if callee.may_kill {
@@ -1155,36 +1130,6 @@ impl FunctionInfo {
                     }
                     FunctionUniformity::new()
                 }
-                S::MeshFunction(func) => {
-                    self.available_stages |= ShaderStages::MESH;
-                    match &func {
-                        // TODO: double check all of this uniformity stuff. I frankly don't fully understand all of it.
-                        &crate::MeshFunction::SetMeshOutputs {
-                            vertex_count,
-                            primitive_count,
-                        } => {
-                            let _ = self.add_ref(vertex_count);
-                            let _ = self.add_ref(primitive_count);
-                            FunctionUniformity::new()
-                        }
-                        &crate::MeshFunction::SetVertex { index, value }
-                        | &crate::MeshFunction::SetPrimitive { index, value } => {
-                            let _ = self.add_ref(index);
-                            let _ = self.add_ref(value);
-                            let ty = self.expressions[value.index()].ty.handle().ok_or(
-                                FunctionError::InvalidMeshShaderOutputType(value).with_span(),
-                            )?;
-
-                            if matches!(func, crate::MeshFunction::SetVertex { .. }) {
-                                self.try_update_mesh_vertex_type(ty, value)?;
-                            } else {
-                                self.try_update_mesh_primitive_type(ty, value)?;
-                            };
-
-                            FunctionUniformity::new()
-                        }
-                    }
-                }
                 S::SubgroupBallot {
                     result: _,
                     predicate,
@@ -1230,72 +1175,6 @@ impl FunctionInfo {
         }
         Ok(combined_uniformity)
     }
-
-    /// Note the type of value passed to [`SetVertex`].
-    ///
-    /// Record that this function passed a value of type `ty` as the second
-    /// argument to the [`SetVertex`] builtin function. All calls to
-    /// `SetVertex` must pass the same type, and this must match the
-    /// function's [`vertex_output_type`].
-    ///
-    /// [`SetVertex`]: crate::ir::MeshFunction::SetVertex
-    /// [`vertex_output_type`]: crate::ir::MeshStageInfo::vertex_output_type
-    fn try_update_mesh_vertex_type(
-        &mut self,
-        ty: Handle<crate::Type>,
-        value: Handle<crate::Expression>,
-    ) -> Result<(), WithSpan<FunctionError>> {
-        if let &Some(ref existing) = &self.mesh_shader_info.vertex_type {
-            if existing.0 != ty {
-                return Err(
-                    FunctionError::ConflictingMeshOutputTypes(existing.1, value).with_span()
-                );
-            }
-        } else {
-            self.mesh_shader_info.vertex_type = Some((ty, value));
-        }
-        Ok(())
-    }
-
-    /// Note the type of value passed to [`SetPrimitive`].
-    ///
-    /// Record that this function passed a value of type `ty` as the second
-    /// argument to the [`SetPrimitive`] builtin function. All calls to
-    /// `SetPrimitive` must pass the same type, and this must match the
-    /// function's [`primitive_output_type`].
-    ///
-    /// [`SetPrimitive`]: crate::ir::MeshFunction::SetPrimitive
-    /// [`primitive_output_type`]: crate::ir::MeshStageInfo::primitive_output_type
-    fn try_update_mesh_primitive_type(
-        &mut self,
-        ty: Handle<crate::Type>,
-        value: Handle<crate::Expression>,
-    ) -> Result<(), WithSpan<FunctionError>> {
-        if let &Some(ref existing) = &self.mesh_shader_info.primitive_type {
-            if existing.0 != ty {
-                return Err(
-                    FunctionError::ConflictingMeshOutputTypes(existing.1, value).with_span()
-                );
-            }
-        } else {
-            self.mesh_shader_info.primitive_type = Some((ty, value));
-        }
-        Ok(())
-    }
-
-    /// Update this function's mesh shader info, given that it calls `callee`.
-    fn try_update_mesh_info(
-        &mut self,
-        callee: &FunctionMeshShaderInfo,
-    ) -> Result<(), WithSpan<FunctionError>> {
-        if let &Some(ref other_vertex) = &callee.vertex_type {
-            self.try_update_mesh_vertex_type(other_vertex.0, other_vertex.1)?;
-        }
-        if let &Some(ref other_primitive) = &callee.primitive_type {
-            self.try_update_mesh_primitive_type(other_primitive.0, other_primitive.1)?;
-        }
-        Ok(())
-    }
 }
 
 impl ModuleInfo {
@@ -1331,7 +1210,6 @@ impl ModuleInfo {
             sampling: crate::FastHashSet::default(),
             dual_source_blending: false,
             diagnostic_filter_leaf: fun.diagnostic_filter_leaf,
-            mesh_shader_info: FunctionMeshShaderInfo::default(),
         };
         let resolve_context =
             ResolveContext::with_locals(module, &fun.local_variables, &fun.arguments);
@@ -1465,7 +1343,6 @@ fn uniform_control_flow() {
         sampling: crate::FastHashSet::default(),
         dual_source_blending: false,
         diagnostic_filter_leaf: None,
-        mesh_shader_info: FunctionMeshShaderInfo::default(),
     };
     let resolve_context = ResolveContext {
         constants: &Arena::new(),
diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs
index 0216c6ef7f6..abf6bc430a6 100644
--- a/naga/src/valid/function.rs
+++ b/naga/src/valid/function.rs
@@ -1547,41 +1547,6 @@ impl super::Validator {
                         crate::RayQueryFunction::Terminate => {}
                     }
                 }
-                S::MeshFunction(func) => {
-                    let ensure_u32 =
-                        |expr: Handle<crate::Expression>| -> Result<(), WithSpan<FunctionError>> {
-                            let u32_ty = TypeResolution::Value(Ti::Scalar(crate::Scalar::U32));
-                            let ty = context
-                                .resolve_type_impl(expr, &self.valid_expression_set)
-                                .map_err_inner(|source| {
-                                    FunctionError::Expression {
-                                        source,
-                                        handle: expr,
-                                    }
-                                    .with_span_handle(expr, context.expressions)
-                                })?;
-                            if !context.compare_types(&u32_ty, ty) {
-                                return Err(FunctionError::InvalidMeshFunctionCall(expr)
-                                    .with_span_handle(expr, context.expressions));
-                            }
-                            Ok(())
-                        };
-                    match func {
-                        crate::MeshFunction::SetMeshOutputs {
-                            vertex_count,
-                            primitive_count,
-                        } => {
-                            ensure_u32(vertex_count)?;
-                            ensure_u32(primitive_count)?;
-                        }
-                        crate::MeshFunction::SetVertex { index, value: _ }
-                        | crate::MeshFunction::SetPrimitive { index, value: _ } => {
-                            ensure_u32(index)?;
-                            // Value is validated elsewhere (since the value type isn't known ahead of time but must match for all calls
-                            // in a function or the function's called functions)
-                        }
-                    }
-                }
                 S::SubgroupBallot { result, predicate } => {
                     stages &= self.subgroup_stages;
                     if !self.capabilities.contains(super::Capabilities::SUBGROUP) {
diff --git a/naga/src/valid/handles.rs b/naga/src/valid/handles.rs
index adb9f355c11..5b7fb3fab75 100644
--- a/naga/src/valid/handles.rs
+++ b/naga/src/valid/handles.rs
@@ -237,6 +237,7 @@ impl super::Validator {
                 Self::validate_global_variable_handle(task_payload, global_variables)?;
             }
             if let Some(ref mesh_info) = entry_point.mesh_info {
+                Self::validate_global_variable_handle(mesh_info.output_variable, global_variables)?;
                 validate_type(mesh_info.vertex_output_type)?;
                 validate_type(mesh_info.primitive_output_type)?;
                 for ov in mesh_info
@@ -815,22 +816,6 @@ impl super::Validator {
                 }
                 Ok(())
             }
-            crate::Statement::MeshFunction(func) => match func {
-                crate::MeshFunction::SetMeshOutputs {
-                    vertex_count,
-                    primitive_count,
-                } => {
-                    validate_expr(vertex_count)?;
-                    validate_expr(primitive_count)?;
-                    Ok(())
-                }
-                crate::MeshFunction::SetVertex { index, value }
-                | crate::MeshFunction::SetPrimitive { index, value } => {
-                    validate_expr(index)?;
-                    validate_expr(value)?;
-                    Ok(())
-                }
-            },
             crate::Statement::SubgroupBallot { result, predicate } => {
                 validate_expr_opt(predicate)?;
                 validate_expr(result)?;
diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs
index 779667c04ec..a18b3738e55 100644
--- a/naga/src/valid/interface.rs
+++ b/naga/src/valid/interface.rs
@@ -98,8 +98,6 @@ pub enum VaryingError {
     InvalidPerPrimitive,
     #[error("Non-builtin members of a mesh primitive output struct must be decorated with `@per_primitive`")]
     MissingPerPrimitive,
-    #[error("The `MESH_SHADER` capability must be enabled to use per-primitive fragment inputs.")]
-    PerPrimitiveNotAllowed,
 }
 
 #[derive(Clone, Debug, thiserror::Error)]
@@ -131,6 +129,9 @@ pub enum EntryPointError {
     InvalidIntegerInterpolation { location: u32 },
     #[error(transparent)]
     Function(#[from] FunctionError),
+    #[error("Capability {0:?} is not supported")]
+    UnsupportedCapability(Capabilities),
+
     #[error("mesh shader entry point missing mesh shader attributes")]
     ExpectedMeshShaderAttributes,
     #[error("Non mesh shader entry point cannot have mesh shader attributes")]
@@ -141,24 +142,28 @@ pub enum EntryPointError {
     TaskPayloadWrongAddressSpace,
     #[error("For a task payload to be used, it must be declared with @payload")]
     WrongTaskPayloadUsed,
-    #[error("A function can only set vertex and primitive types that correspond to the mesh shader attributes")]
-    WrongMeshOutputType,
-    #[error("Only mesh shader entry points can write to mesh output vertices and primitives")]
-    UnexpectedMeshShaderOutput,
-    #[error("Mesh shader entry point cannot have a return type")]
-    UnexpectedMeshShaderEntryResult,
     #[error("Task shader entry point must return @builtin(mesh_task_size) vec3<u32>")]
     WrongTaskShaderEntryResult,
-    #[error("Mesh output type must be a user-defined struct.")]
-    InvalidMeshOutputType,
-    #[error("Mesh primitive outputs must have exactly one of `@builtin(triangle_indices)`, `@builtin(line_indices)`, or `@builtin(point_index)`")]
-    InvalidMeshPrimitiveOutputType,
     #[error("Task shaders must declare a task payload output")]
     ExpectedTaskPayload,
     #[error(
-        "The `MESH_SHADER` capability must be enabled to compile mesh shaders and task shaders."
+        "Mesh shader output variable must be a struct with fields that are all allowed builtins"
     )]
-    MeshShaderCapabilityDisabled,
+    BadMeshOutputVariableType,
+    #[error("Mesh shader output variable fields must have types that are in accordance with the mesh shader spec")]
+    BadMeshOutputVariableField,
+    #[error("Mesh shader entry point cannot have a return type")]
+    UnexpectedMeshShaderEntryResult,
+    #[error(
+        "Mesh output type must be a user-defined struct with fields in alignment with the mesh shader spec"
+    )]
+    InvalidMeshOutputType,
+    #[error("Mesh primitive outputs must have exactly one of `@builtin(triangle_indices)`, `@builtin(line_indices)`, or `@builtin(point_index)`")]
+    InvalidMeshPrimitiveOutputType,
+    #[error("Mesh output global variable must live in the workgroup address space")]
+    WrongMeshOutputAddressSpace,
+    #[error("Task payload must be at least 4 bytes, but is {0} bytes")]
+    TaskPayloadTooSmall(u32),
 }
 
 fn storage_usage(access: crate::StorageAccess) -> GlobalUse {
@@ -312,7 +317,10 @@ impl VaryingContext<'_> {
                         *ty_inner == Ti::Scalar(crate::Scalar::BOOL),
                     ),
                     Bi::PrimitiveIndex => (
-                        self.stage == St::Fragment && !self.output,
+                        (self.stage == St::Fragment && !self.output)
+                            || (self.stage == St::Mesh
+                                && self.output
+                                && self.mesh_output_type == MeshOutputType::PrimitiveOutput),
                         *ty_inner == Ti::Scalar(crate::Scalar::U32),
                     ),
                     Bi::Barycentric => (
@@ -390,7 +398,29 @@ impl VaryingContext<'_> {
                                 scalar: crate::Scalar::U32,
                             },
                     ),
+                    // Validated elsewhere, shouldn't be here
+                    Bi::VertexCount | Bi::PrimitiveCount | Bi::Vertices | Bi::Primitives => {
+                        (false, true)
+                    }
                 };
+                match built_in {
+                    Bi::CullPrimitive
+                    | Bi::PointIndex
+                    | Bi::LineIndices
+                    | Bi::TriangleIndices
+                    | Bi::MeshTaskSize
+                    | Bi::VertexCount
+                    | Bi::PrimitiveCount
+                    | Bi::Vertices
+                    | Bi::Primitives => {
+                        if !self.capabilities.contains(Capabilities::MESH_SHADER) {
+                            return Err(VaryingError::UnsupportedCapability(
+                                Capabilities::MESH_SHADER,
+                            ));
+                        }
+                    }
+                    _ => (),
+                }
 
                 if !visible {
                     return Err(VaryingError::InvalidBuiltInStage(built_in));
@@ -408,7 +438,9 @@ impl VaryingContext<'_> {
                 per_primitive,
             } => {
                 if per_primitive && !self.capabilities.contains(Capabilities::MESH_SHADER) {
-                    return Err(VaryingError::PerPrimitiveNotAllowed);
+                    return Err(VaryingError::UnsupportedCapability(
+                        Capabilities::MESH_SHADER,
+                    ));
                 }
                 // Only IO-shareable types may be stored in locations.
                 if !self.type_info[ty.index()]
@@ -836,7 +868,9 @@ impl super::Validator {
             crate::ShaderStage::Task | crate::ShaderStage::Mesh
         ) && !self.capabilities.contains(Capabilities::MESH_SHADER)
         {
-            return Err(EntryPointError::MeshShaderCapabilityDisabled.with_span());
+            return Err(
+                EntryPointError::UnsupportedCapability(Capabilities::MESH_SHADER).with_span(),
+            );
         }
         if ep.early_depth_test.is_some() {
             let required = Capabilities::EARLY_DEPTH_TEST;
@@ -868,6 +902,7 @@ impl super::Validator {
             (crate::ShaderStage::Mesh, &None) => {
                 return Err(EntryPointError::ExpectedMeshShaderAttributes.with_span());
             }
+            (crate::ShaderStage::Mesh, &Some(..)) => {}
             (_, &Some(_)) => {
                 return Err(EntryPointError::UnexpectedMeshShaderAttributes.with_span());
             }
@@ -901,6 +936,9 @@ impl super::Validator {
                     }
                     info.insert_global_use(GlobalUse::READ, handle);
                 }
+                if let Some(ref mesh_info) = ep.mesh_info {
+                    info.insert_global_use(GlobalUse::READ, mesh_info.output_variable);
+                }
             }
 
             // Other stages must not have a payload.
@@ -1020,6 +1058,11 @@ impl super::Validator {
                     return Err(EntryPointError::WrongTaskPayloadUsed
                         .with_span_handle(var_handle, &module.global_variables));
                 }
+                let size = module.types[var.ty].inner.size(module.to_ctx());
+                if size < 4 {
+                    return Err(EntryPointError::TaskPayloadTooSmall(size)
+                        .with_span_handle(var_handle, &module.global_variables));
+                }
             }
 
             let allowed_usage = match var.space {
@@ -1074,21 +1117,48 @@ impl super::Validator {
         // If this is a `Mesh` entry point, check its vertex and primitive output types.
         // We verified previously that only mesh shaders can have `mesh_info`.
         if let &Some(ref mesh_info) = &ep.mesh_info {
-            // Mesh shaders don't return any value. All their results are supplied through
-            // [`SetVertex`] and [`SetPrimitive`] calls.
-            if let Some((used_vertex_type, _)) = info.mesh_shader_info.vertex_type {
-                if used_vertex_type != mesh_info.vertex_output_type {
-                    return Err(EntryPointError::WrongMeshOutputType
-                        .with_span_handle(mesh_info.vertex_output_type, &module.types));
+            if module.global_variables[mesh_info.output_variable].space
+                != crate::AddressSpace::WorkGroup
+            {
+                return Err(EntryPointError::WrongMeshOutputAddressSpace.with_span());
+            }
+
+            let mut implied = module.analyze_mesh_shader_info(mesh_info.output_variable);
+            if let Some(e) = implied.2 {
+                return Err(e);
+            }
+
+            if let Some(e) = mesh_info.max_vertices_override {
+                if let crate::Expression::Override(o) = module.global_expressions[e] {
+                    if implied.1[0] != Some(o) {
+                        return Err(EntryPointError::BadMeshOutputVariableType.with_span());
+                    }
                 }
             }
-            if let Some((used_primitive_type, _)) = info.mesh_shader_info.primitive_type {
-                if used_primitive_type != mesh_info.primitive_output_type {
-                    return Err(EntryPointError::WrongMeshOutputType
-                        .with_span_handle(mesh_info.primitive_output_type, &module.types));
+            if let Some(e) = mesh_info.max_primitives_override {
+                if let crate::Expression::Override(o) = module.global_expressions[e] {
+                    if implied.1[1] != Some(o) {
+                        return Err(EntryPointError::BadMeshOutputVariableType.with_span());
+                    }
                 }
             }
 
+            implied.0.max_vertices_override = mesh_info.max_vertices_override;
+            implied.0.max_primitives_override = mesh_info.max_primitives_override;
+            if implied.0 != *mesh_info {
+                return Err(EntryPointError::BadMeshOutputVariableType.with_span());
+            }
+            if mesh_info.topology == crate::MeshOutputTopology::Points
+                && !self
+                    .capabilities
+                    .contains(Capabilities::MESH_SHADER_POINT_TOPOLOGY)
+            {
+                return Err(EntryPointError::UnsupportedCapability(
+                    Capabilities::MESH_SHADER_POINT_TOPOLOGY,
+                )
+                .with_span());
+            }
+
             self.validate_mesh_output_type(
                 ep,
                 module,
@@ -1101,14 +1171,6 @@ impl super::Validator {
                 mesh_info.primitive_output_type,
                 MeshOutputType::PrimitiveOutput,
             )?;
-        } else {
-            // This is not a `Mesh` entry point, so ensure that it never tries to produce
-            // vertices or primitives.
-            if info.mesh_shader_info.vertex_type.is_some()
-                || info.mesh_shader_info.primitive_type.is_some()
-            {
-                return Err(EntryPointError::UnexpectedMeshShaderOutput.with_span());
-            }
         }
 
         Ok(info)
diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs
index a5ec5affceb..0ec5c5184f2 100644
--- a/naga/src/valid/mod.rs
+++ b/naga/src/valid/mod.rs
@@ -190,6 +190,8 @@ bitflags::bitflags! {
         const SHADER_BARYCENTRICS = 1 << 29;
         /// Support for task shaders, mesh shaders, and per-primitive fragment inputs
         const MESH_SHADER = 1 << 30;
+        /// Support for mesh shaders which output points.
+        const MESH_SHADER_POINT_TOPOLOGY = 1 << 30;
     }
 }
 
diff --git a/naga/tests/in/wgsl/mesh-shader-empty.toml b/naga/tests/in/wgsl/mesh-shader-empty.toml
new file mode 100644
index 00000000000..8500399f936
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader-empty.toml
@@ -0,0 +1,2 @@
+god_mode = true
+targets = "IR | ANALYSIS"
diff --git a/naga/tests/in/wgsl/mesh-shader-empty.wgsl b/naga/tests/in/wgsl/mesh-shader-empty.wgsl
new file mode 100644
index 00000000000..98a6bf8448b
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader-empty.wgsl
@@ -0,0 +1,37 @@
+// An empty WGSL shader to check that task payload/mesh output
+// are still properly written without being used in the shader
+
+enable wgpu_mesh_shader;
+
+struct TaskPayload {
+    dummy: u32,
+}
+struct VertexOutput {
+    @builtin(position) position: vec4<f32>,
+}
+struct PrimitiveOutput {
+    @builtin(triangle_indices) indices: vec3<u32>,
+}
+
+var<task_payload> taskPayload: TaskPayload;
+
+@task
+@payload(taskPayload)
+@workgroup_size(1)
+fn ts_main() -> @builtin(mesh_task_size) vec3<u32> {
+    return vec3(1, 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() {}
diff --git a/naga/tests/in/wgsl/mesh-shader-lines.toml b/naga/tests/in/wgsl/mesh-shader-lines.toml
new file mode 100644
index 00000000000..8500399f936
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader-lines.toml
@@ -0,0 +1,2 @@
+god_mode = true
+targets = "IR | ANALYSIS"
diff --git a/naga/tests/in/wgsl/mesh-shader-lines.wgsl b/naga/tests/in/wgsl/mesh-shader-lines.wgsl
new file mode 100644
index 00000000000..c475ff10619
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader-lines.wgsl
@@ -0,0 +1,37 @@
+// An empty WGSL shader to check that task payload/mesh output
+// are still properly written without being used in the shader
+
+enable wgpu_mesh_shader;
+
+struct TaskPayload {
+    dummy: u32,
+}
+struct VertexOutput {
+    @builtin(position) position: vec4<f32>,
+}
+struct PrimitiveOutput {
+    @builtin(line_indices) indices: vec2<u32>,
+}
+
+var<task_payload> taskPayload: TaskPayload;
+
+@task
+@payload(taskPayload)
+@workgroup_size(1)
+fn ts_main() -> @builtin(mesh_task_size) vec3<u32> {
+    return vec3(1, 1, 1);
+}
+
+struct MeshOutput {
+    @builtin(vertices) vertices: array<VertexOutput, 2>,
+    @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() {}
diff --git a/naga/tests/in/wgsl/mesh-shader-points.toml b/naga/tests/in/wgsl/mesh-shader-points.toml
new file mode 100644
index 00000000000..8500399f936
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader-points.toml
@@ -0,0 +1,2 @@
+god_mode = true
+targets = "IR | ANALYSIS"
diff --git a/naga/tests/in/wgsl/mesh-shader-points.wgsl b/naga/tests/in/wgsl/mesh-shader-points.wgsl
new file mode 100644
index 00000000000..84516ee8f2a
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader-points.wgsl
@@ -0,0 +1,37 @@
+// An empty WGSL shader to check that task payload/mesh output
+// are still properly written without being used in the shader
+
+enable wgpu_mesh_shader;
+
+struct TaskPayload {
+    dummy: u32,
+}
+struct VertexOutput {
+    @builtin(position) position: vec4<f32>,
+}
+struct PrimitiveOutput {
+    @builtin(point_index) indices: u32,
+}
+
+var<task_payload> taskPayload: TaskPayload;
+
+@task
+@payload(taskPayload)
+@workgroup_size(1)
+fn ts_main() -> @builtin(mesh_task_size) vec3<u32> {
+    return vec3(1, 1, 1);
+}
+
+struct MeshOutput {
+    @builtin(vertices) vertices: array<VertexOutput, 1>,
+    @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() {}
diff --git a/naga/tests/in/wgsl/mesh-shader.toml b/naga/tests/in/wgsl/mesh-shader.toml
new file mode 100644
index 00000000000..8500399f936
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader.toml
@@ -0,0 +1,2 @@
+god_mode = true
+targets = "IR | ANALYSIS"
diff --git a/naga/tests/in/wgsl/mesh-shader.wgsl b/naga/tests/in/wgsl/mesh-shader.wgsl
new file mode 100644
index 00000000000..4ed3a18cfdb
--- /dev/null
+++ b/naga/tests/in/wgsl/mesh-shader.wgsl
@@ -0,0 +1,79 @@
+// Main mesh shader test file. Tests most features.
+
+enable wgpu_mesh_shader;
+
+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,
+}
+struct VertexOutput {
+    @builtin(position) position: vec4<f32>,
+    @location(0) color: vec4<f32>,
+}
+struct PrimitiveOutput {
+    @builtin(triangle_indices) indices: vec3<u32>,
+    @builtin(cull_primitive) cull: bool,
+    @per_primitive @location(1) colorMask: vec4<f32>,
+}
+struct PrimitiveInput {
+    @per_primitive @location(1) colorMask: vec4<f32>,
+}
+
+var<task_payload> taskPayload: TaskPayload;
+var<workgroup> workgroupData: 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(1, 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].indices = 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;
+}
diff --git a/naga/tests/out/analysis/spv-shadow.info.ron b/naga/tests/out/analysis/spv-shadow.info.ron
index b08a28438ed..381f841d5d9 100644
--- a/naga/tests/out/analysis/spv-shadow.info.ron
+++ b/naga/tests/out/analysis/spv-shadow.info.ron
@@ -413,10 +413,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -1595,10 +1591,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     entry_points: [
@@ -1693,10 +1685,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     const_expression_types: [
diff --git a/naga/tests/out/analysis/wgsl-access.info.ron b/naga/tests/out/analysis/wgsl-access.info.ron
index d297b09a404..c22cd768f2e 100644
--- a/naga/tests/out/analysis/wgsl-access.info.ron
+++ b/naga/tests/out/analysis/wgsl-access.info.ron
@@ -1197,10 +1197,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -2527,10 +2523,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -2571,10 +2563,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -2624,10 +2612,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -2671,10 +2655,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -2769,10 +2749,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -2894,10 +2870,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -2950,10 +2922,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3009,10 +2977,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3065,10 +3029,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3124,10 +3084,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3192,10 +3148,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3269,10 +3221,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3349,10 +3297,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3453,10 +3397,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -3653,10 +3593,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     entry_points: [
@@ -4354,10 +4290,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -4810,10 +4742,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -4884,10 +4812,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     const_expression_types: [
diff --git a/naga/tests/out/analysis/wgsl-collatz.info.ron b/naga/tests/out/analysis/wgsl-collatz.info.ron
index 2796f544510..219e016f8d7 100644
--- a/naga/tests/out/analysis/wgsl-collatz.info.ron
+++ b/naga/tests/out/analysis/wgsl-collatz.info.ron
@@ -275,10 +275,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     entry_points: [
@@ -434,10 +430,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     const_expression_types: [],
diff --git a/naga/tests/out/analysis/wgsl-mesh-shader-empty.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader-empty.info.ron
new file mode 100644
index 00000000000..9f9feac9c09
--- /dev/null
+++ b/naga/tests/out/analysis/wgsl-mesh-shader-empty.info.ron
@@ -0,0 +1,99 @@
+(
+    type_flags: [
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+    ],
+    functions: [],
+    entry_points: [
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ | WRITE"),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(4),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ"),
+                ("READ"),
+            ],
+            expressions: [],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+    ],
+    const_expression_types: [],
+)
\ No newline at end of file
diff --git a/naga/tests/out/analysis/wgsl-mesh-shader-lines.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader-lines.info.ron
new file mode 100644
index 00000000000..b4c8790508d
--- /dev/null
+++ b/naga/tests/out/analysis/wgsl-mesh-shader-lines.info.ron
@@ -0,0 +1,100 @@
+(
+    type_flags: [
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+    ],
+    functions: [],
+    entry_points: [
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ | WRITE"),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(6),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ"),
+                ("READ"),
+            ],
+            expressions: [],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+    ],
+    const_expression_types: [],
+)
\ No newline at end of file
diff --git a/naga/tests/out/analysis/wgsl-mesh-shader-points.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader-points.info.ron
new file mode 100644
index 00000000000..fa62867576b
--- /dev/null
+++ b/naga/tests/out/analysis/wgsl-mesh-shader-points.info.ron
@@ -0,0 +1,99 @@
+(
+    type_flags: [
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+    ],
+    functions: [],
+    entry_points: [
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ | WRITE"),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(5),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ"),
+                ("READ"),
+            ],
+            expressions: [],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+    ],
+    const_expression_types: [],
+)
\ No newline at end of file
diff --git a/naga/tests/out/analysis/wgsl-mesh-shader.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader.info.ron
new file mode 100644
index 00000000000..eacd33ad0f1
--- /dev/null
+++ b/naga/tests/out/analysis/wgsl-mesh-shader.info.ron
@@ -0,0 +1,1469 @@
+(
+    type_flags: [
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | IO_SHAREABLE | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | HOST_SHAREABLE | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+        ("DATA | SIZED | COPY | CREATION_RESOLVED | ARGUMENT | CONSTRUCTIBLE"),
+    ],
+    functions: [],
+    entry_points: [
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ | WRITE"),
+                ("WRITE"),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(1),
+                    ty: Value(Pointer(
+                        base: 0,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 3,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 3,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 2,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Bool,
+                        width: 1,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(6),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: None,
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                ("READ"),
+                ("WRITE"),
+                ("READ | WRITE"),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 0,
+                    assignable_global: None,
+                    ty: Handle(5),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(1),
+                        requirements: (""),
+                    ),
+                    ref_count: 0,
+                    assignable_global: None,
+                    ty: Handle(6),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 5,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 5,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(1),
+                    ty: Value(Pointer(
+                        base: 0,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 9,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 4,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 9,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 4,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 3,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 9,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 4,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 9,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 4,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 3,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 9,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 4,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 9,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 4,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 3,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 10,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 7,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 6,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Uint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(6),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 10,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 7,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 2,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 3,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(0),
+                    ty: Value(Pointer(
+                        base: 2,
+                        space: TaskPayload,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 11,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 10,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 7,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: Some(2),
+                    ty: Value(Pointer(
+                        base: 1,
+                        space: WorkGroup,
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Float,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
+            uniformity: (
+                non_uniform_result: Some(0),
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                (""),
+                (""),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(4),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(1),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(8),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(1),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(1),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+            diagnostic_filter_leaf: None,
+        ),
+    ],
+    const_expression_types: [],
+)
\ No newline at end of file
diff --git a/naga/tests/out/analysis/wgsl-overrides.info.ron b/naga/tests/out/analysis/wgsl-overrides.info.ron
index a76c9c89c9b..92e99112e53 100644
--- a/naga/tests/out/analysis/wgsl-overrides.info.ron
+++ b/naga/tests/out/analysis/wgsl-overrides.info.ron
@@ -201,10 +201,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     const_expression_types: [
diff --git a/naga/tests/out/analysis/wgsl-storage-textures.info.ron b/naga/tests/out/analysis/wgsl-storage-textures.info.ron
index 35b5a7e320c..8bb298a6450 100644
--- a/naga/tests/out/analysis/wgsl-storage-textures.info.ron
+++ b/naga/tests/out/analysis/wgsl-storage-textures.info.ron
@@ -184,10 +184,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
@@ -400,10 +396,6 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
-            mesh_shader_info: (
-                vertex_type: None,
-                primitive_type: None,
-            ),
         ),
     ],
     const_expression_types: [],
diff --git a/naga/tests/out/ir/wgsl-mesh-shader-empty.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader-empty.compact.ron
new file mode 100644
index 00000000000..8bdaa72da69
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader-empty.compact.ron
@@ -0,0 +1,234 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("dummy"),
+                        ty: 0,
+                        binding: None,
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 2,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 4,
+                        binding: Some(BuiltIn(TriangleIndices)),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 3,
+                size: Constant(3),
+                stride: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 5,
+                size: Constant(1),
+                stride: 16,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 6,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 7,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 48,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 64,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 68,
+                    ),
+                ],
+                span: 80,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 1,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 8,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 4,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 4,
+                        components: [
+                            0,
+                            1,
+                            2,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Return(
+                        value: Some(3),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Triangles,
+                max_vertices: 3,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 3,
+                primitive_output_type: 5,
+                output_variable: 1,
+            )),
+            task_payload: Some(0),
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/naga/tests/out/ir/wgsl-mesh-shader-empty.ron b/naga/tests/out/ir/wgsl-mesh-shader-empty.ron
new file mode 100644
index 00000000000..8bdaa72da69
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader-empty.ron
@@ -0,0 +1,234 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("dummy"),
+                        ty: 0,
+                        binding: None,
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 2,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 4,
+                        binding: Some(BuiltIn(TriangleIndices)),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 3,
+                size: Constant(3),
+                stride: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 5,
+                size: Constant(1),
+                stride: 16,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 6,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 7,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 48,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 64,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 68,
+                    ),
+                ],
+                span: 80,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 1,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 8,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 4,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 4,
+                        components: [
+                            0,
+                            1,
+                            2,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Return(
+                        value: Some(3),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Triangles,
+                max_vertices: 3,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 3,
+                primitive_output_type: 5,
+                output_variable: 1,
+            )),
+            task_payload: Some(0),
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/naga/tests/out/ir/wgsl-mesh-shader-lines.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader-lines.compact.ron
new file mode 100644
index 00000000000..b298f296985
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader-lines.compact.ron
@@ -0,0 +1,244 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("dummy"),
+                        ty: 0,
+                        binding: None,
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 2,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Bi,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 4,
+                        binding: Some(BuiltIn(LineIndices)),
+                        offset: 0,
+                    ),
+                ],
+                span: 8,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 3,
+                size: Constant(2),
+                stride: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 5,
+                size: Constant(1),
+                stride: 8,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 7,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 8,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 32,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 40,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 44,
+                    ),
+                ],
+                span: 48,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 1,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 9,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 6,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 6,
+                        components: [
+                            0,
+                            1,
+                            2,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Return(
+                        value: Some(3),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Lines,
+                max_vertices: 2,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 3,
+                primitive_output_type: 5,
+                output_variable: 1,
+            )),
+            task_payload: Some(0),
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/naga/tests/out/ir/wgsl-mesh-shader-lines.ron b/naga/tests/out/ir/wgsl-mesh-shader-lines.ron
new file mode 100644
index 00000000000..b298f296985
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader-lines.ron
@@ -0,0 +1,244 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("dummy"),
+                        ty: 0,
+                        binding: None,
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 2,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Bi,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 4,
+                        binding: Some(BuiltIn(LineIndices)),
+                        offset: 0,
+                    ),
+                ],
+                span: 8,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 3,
+                size: Constant(2),
+                stride: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 5,
+                size: Constant(1),
+                stride: 8,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 7,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 8,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 32,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 40,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 44,
+                    ),
+                ],
+                span: 48,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 1,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 9,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 6,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 6,
+                        components: [
+                            0,
+                            1,
+                            2,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Return(
+                        value: Some(3),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Lines,
+                max_vertices: 2,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 3,
+                primitive_output_type: 5,
+                output_variable: 1,
+            )),
+            task_payload: Some(0),
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/naga/tests/out/ir/wgsl-mesh-shader-points.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader-points.compact.ron
new file mode 100644
index 00000000000..558a88e28d8
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader-points.compact.ron
@@ -0,0 +1,234 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("dummy"),
+                        ty: 0,
+                        binding: None,
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 2,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PointIndex)),
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 3,
+                size: Constant(1),
+                stride: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 4,
+                size: Constant(1),
+                stride: 4,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 6,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 7,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 16,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 20,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 24,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 1,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 8,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 5,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 5,
+                        components: [
+                            0,
+                            1,
+                            2,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Return(
+                        value: Some(3),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Points,
+                max_vertices: 1,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 3,
+                primitive_output_type: 4,
+                output_variable: 1,
+            )),
+            task_payload: Some(0),
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/naga/tests/out/ir/wgsl-mesh-shader-points.ron b/naga/tests/out/ir/wgsl-mesh-shader-points.ron
new file mode 100644
index 00000000000..558a88e28d8
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader-points.ron
@@ -0,0 +1,234 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("dummy"),
+                        ty: 0,
+                        binding: None,
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 2,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PointIndex)),
+                        offset: 0,
+                    ),
+                ],
+                span: 4,
+            ),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 3,
+                size: Constant(1),
+                stride: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 4,
+                size: Constant(1),
+                stride: 4,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 6,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 7,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 16,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 20,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 0,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 24,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 1,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 8,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 5,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 5,
+                        components: [
+                            0,
+                            1,
+                            2,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Return(
+                        value: Some(3),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [],
+                result: None,
+                local_variables: [],
+                expressions: [],
+                named_expressions: {},
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Points,
+                max_vertices: 1,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 3,
+                primitive_output_type: 4,
+                output_variable: 1,
+            )),
+            task_payload: Some(0),
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/naga/tests/out/ir/wgsl-mesh-shader.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader.compact.ron
new file mode 100644
index 00000000000..9cb3c6479c3
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader.compact.ron
@@ -0,0 +1,980 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Float,
+                width: 4,
+            )),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: None,
+            inner: Scalar((
+                kind: Bool,
+                width: 1,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("colorMask"),
+                        ty: 1,
+                        binding: None,
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("visible"),
+                        ty: 2,
+                        binding: None,
+                        offset: 16,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 1,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("color"),
+                        ty: 1,
+                        binding: Some(Location(
+                            location: 0,
+                            interpolation: Some(Perspective),
+                            sampling: Some(Center),
+                            blend_src: None,
+                            per_primitive: false,
+                        )),
+                        offset: 16,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 6,
+                        binding: Some(BuiltIn(TriangleIndices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("cull"),
+                        ty: 2,
+                        binding: Some(BuiltIn(CullPrimitive)),
+                        offset: 12,
+                    ),
+                    (
+                        name: Some("colorMask"),
+                        ty: 1,
+                        binding: Some(Location(
+                            location: 1,
+                            interpolation: Some(Perspective),
+                            sampling: Some(Center),
+                            blend_src: None,
+                            per_primitive: true,
+                        )),
+                        offset: 16,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+        (
+            name: Some("PrimitiveInput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("colorMask"),
+                        ty: 1,
+                        binding: Some(Location(
+                            location: 1,
+                            interpolation: Some(Perspective),
+                            sampling: Some(Center),
+                            blend_src: None,
+                            per_primitive: true,
+                        )),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 4,
+                size: Constant(3),
+                stride: 32,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 7,
+                size: Constant(1),
+                stride: 32,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 9,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 10,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 96,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 5,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 128,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 5,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 132,
+                    ),
+                ],
+                span: 144,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 3,
+            init: None,
+        ),
+        (
+            name: Some("workgroupData"),
+            space: WorkGroup,
+            binding: None,
+            ty: 0,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 11,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 6,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    GlobalVariable(1),
+                    Literal(F32(1.0)),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 2,
+                        index: 0,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            4,
+                            5,
+                            6,
+                            7,
+                        ],
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 9,
+                        index: 1,
+                    ),
+                    Literal(Bool(true)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 6,
+                        components: [
+                            12,
+                            13,
+                            14,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Store(
+                        pointer: 0,
+                        value: 1,
+                    ),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 8,
+                        end: 9,
+                    )),
+                    Store(
+                        pointer: 3,
+                        value: 8,
+                    ),
+                    Emit((
+                        start: 10,
+                        end: 11,
+                    )),
+                    Store(
+                        pointer: 10,
+                        value: 11,
+                    ),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 15,
+                        end: 16,
+                    )),
+                    Return(
+                        value: Some(15),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [
+                    (
+                        name: Some("index"),
+                        ty: 5,
+                        binding: Some(BuiltIn(LocalInvocationIndex)),
+                    ),
+                    (
+                        name: Some("id"),
+                        ty: 6,
+                        binding: Some(BuiltIn(GlobalInvocationId)),
+                    ),
+                ],
+                result: None,
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    FunctionArgument(1),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 2,
+                        index: 2,
+                    ),
+                    Literal(U32(3)),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 5,
+                        index: 3,
+                    ),
+                    Literal(U32(1)),
+                    GlobalVariable(1),
+                    Literal(F32(2.0)),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 10,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 11,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 12,
+                        index: 0,
+                    ),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            14,
+                            15,
+                            16,
+                            17,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 19,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 20,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 21,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 23,
+                        index: 0,
+                    ),
+                    Load(
+                        pointer: 24,
+                    ),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            26,
+                            27,
+                            28,
+                            29,
+                        ],
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 30,
+                        right: 25,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 32,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 33,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 34,
+                        index: 0,
+                    ),
+                    Literal(F32(-1.0)),
+                    Literal(F32(-1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            36,
+                            37,
+                            38,
+                            39,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 41,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 42,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 43,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 45,
+                        index: 0,
+                    ),
+                    Load(
+                        pointer: 46,
+                    ),
+                    Literal(F32(0.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            48,
+                            49,
+                            50,
+                            51,
+                        ],
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 52,
+                        right: 47,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 54,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 55,
+                        index: 2,
+                    ),
+                    AccessIndex(
+                        base: 56,
+                        index: 0,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(-1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            58,
+                            59,
+                            60,
+                            61,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 63,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 64,
+                        index: 2,
+                    ),
+                    AccessIndex(
+                        base: 65,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 67,
+                        index: 0,
+                    ),
+                    Load(
+                        pointer: 68,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            70,
+                            71,
+                            72,
+                            73,
+                        ],
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 74,
+                        right: 69,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 76,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 77,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 78,
+                        index: 0,
+                    ),
+                    Literal(U32(0)),
+                    Literal(U32(1)),
+                    Literal(U32(2)),
+                    Compose(
+                        ty: 6,
+                        components: [
+                            80,
+                            81,
+                            82,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 84,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 85,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 86,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 88,
+                        index: 1,
+                    ),
+                    Load(
+                        pointer: 89,
+                    ),
+                    Unary(
+                        op: LogicalNot,
+                        expr: 90,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 92,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 93,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 94,
+                        index: 2,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            96,
+                            97,
+                            98,
+                            99,
+                        ],
+                    ),
+                ],
+                named_expressions: {
+                    0: "index",
+                    1: "id",
+                },
+                body: [
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Store(
+                        pointer: 3,
+                        value: 4,
+                    ),
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Store(
+                        pointer: 6,
+                        value: 7,
+                    ),
+                    Store(
+                        pointer: 8,
+                        value: 9,
+                    ),
+                    Emit((
+                        start: 11,
+                        end: 12,
+                    )),
+                    Emit((
+                        start: 12,
+                        end: 14,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 18,
+                        end: 19,
+                    )),
+                    Store(
+                        pointer: 13,
+                        value: 18,
+                    ),
+                    Emit((
+                        start: 20,
+                        end: 21,
+                    )),
+                    Emit((
+                        start: 21,
+                        end: 23,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 24,
+                        end: 26,
+                    )),
+                    Emit((
+                        start: 30,
+                        end: 32,
+                    )),
+                    Store(
+                        pointer: 22,
+                        value: 31,
+                    ),
+                    Emit((
+                        start: 33,
+                        end: 34,
+                    )),
+                    Emit((
+                        start: 34,
+                        end: 36,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 40,
+                        end: 41,
+                    )),
+                    Store(
+                        pointer: 35,
+                        value: 40,
+                    ),
+                    Emit((
+                        start: 42,
+                        end: 43,
+                    )),
+                    Emit((
+                        start: 43,
+                        end: 45,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 46,
+                        end: 48,
+                    )),
+                    Emit((
+                        start: 52,
+                        end: 54,
+                    )),
+                    Store(
+                        pointer: 44,
+                        value: 53,
+                    ),
+                    Emit((
+                        start: 55,
+                        end: 56,
+                    )),
+                    Emit((
+                        start: 56,
+                        end: 58,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 62,
+                        end: 63,
+                    )),
+                    Store(
+                        pointer: 57,
+                        value: 62,
+                    ),
+                    Emit((
+                        start: 64,
+                        end: 65,
+                    )),
+                    Emit((
+                        start: 65,
+                        end: 67,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 68,
+                        end: 70,
+                    )),
+                    Emit((
+                        start: 74,
+                        end: 76,
+                    )),
+                    Store(
+                        pointer: 66,
+                        value: 75,
+                    ),
+                    Emit((
+                        start: 77,
+                        end: 78,
+                    )),
+                    Emit((
+                        start: 78,
+                        end: 80,
+                    )),
+                    Emit((
+                        start: 83,
+                        end: 84,
+                    )),
+                    Store(
+                        pointer: 79,
+                        value: 83,
+                    ),
+                    Emit((
+                        start: 85,
+                        end: 86,
+                    )),
+                    Emit((
+                        start: 86,
+                        end: 88,
+                    )),
+                    Emit((
+                        start: 89,
+                        end: 92,
+                    )),
+                    Store(
+                        pointer: 87,
+                        value: 91,
+                    ),
+                    Emit((
+                        start: 93,
+                        end: 94,
+                    )),
+                    Emit((
+                        start: 94,
+                        end: 96,
+                    )),
+                    Emit((
+                        start: 100,
+                        end: 101,
+                    )),
+                    Store(
+                        pointer: 95,
+                        value: 100,
+                    ),
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Triangles,
+                max_vertices: 3,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 4,
+                primitive_output_type: 7,
+                output_variable: 2,
+            )),
+            task_payload: Some(0),
+        ),
+        (
+            name: "fs_main",
+            stage: Fragment,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("fs_main"),
+                arguments: [
+                    (
+                        name: Some("vertex"),
+                        ty: 4,
+                        binding: None,
+                    ),
+                    (
+                        name: Some("primitive"),
+                        ty: 8,
+                        binding: None,
+                    ),
+                ],
+                result: Some((
+                    ty: 1,
+                    binding: Some(Location(
+                        location: 0,
+                        interpolation: Some(Perspective),
+                        sampling: Some(Center),
+                        blend_src: None,
+                        per_primitive: false,
+                    )),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    FunctionArgument(1),
+                    AccessIndex(
+                        base: 0,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 1,
+                        index: 0,
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 2,
+                        right: 3,
+                    ),
+                ],
+                named_expressions: {
+                    0: "vertex",
+                    1: "primitive",
+                },
+                body: [
+                    Emit((
+                        start: 2,
+                        end: 5,
+                    )),
+                    Return(
+                        value: Some(4),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: None,
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/naga/tests/out/ir/wgsl-mesh-shader.ron b/naga/tests/out/ir/wgsl-mesh-shader.ron
new file mode 100644
index 00000000000..9cb3c6479c3
--- /dev/null
+++ b/naga/tests/out/ir/wgsl-mesh-shader.ron
@@ -0,0 +1,980 @@
+(
+    types: [
+        (
+            name: None,
+            inner: Scalar((
+                kind: Float,
+                width: 4,
+            )),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Quad,
+                scalar: (
+                    kind: Float,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: None,
+            inner: Scalar((
+                kind: Bool,
+                width: 1,
+            )),
+        ),
+        (
+            name: Some("TaskPayload"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("colorMask"),
+                        ty: 1,
+                        binding: None,
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("visible"),
+                        ty: 2,
+                        binding: None,
+                        offset: 16,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+        (
+            name: Some("VertexOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("position"),
+                        ty: 1,
+                        binding: Some(BuiltIn(Position(
+                            invariant: false,
+                        ))),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("color"),
+                        ty: 1,
+                        binding: Some(Location(
+                            location: 0,
+                            interpolation: Some(Perspective),
+                            sampling: Some(Center),
+                            blend_src: None,
+                            per_primitive: false,
+                        )),
+                        offset: 16,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+        (
+            name: None,
+            inner: Scalar((
+                kind: Uint,
+                width: 4,
+            )),
+        ),
+        (
+            name: None,
+            inner: Vector(
+                size: Tri,
+                scalar: (
+                    kind: Uint,
+                    width: 4,
+                ),
+            ),
+        ),
+        (
+            name: Some("PrimitiveOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("indices"),
+                        ty: 6,
+                        binding: Some(BuiltIn(TriangleIndices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("cull"),
+                        ty: 2,
+                        binding: Some(BuiltIn(CullPrimitive)),
+                        offset: 12,
+                    ),
+                    (
+                        name: Some("colorMask"),
+                        ty: 1,
+                        binding: Some(Location(
+                            location: 1,
+                            interpolation: Some(Perspective),
+                            sampling: Some(Center),
+                            blend_src: None,
+                            per_primitive: true,
+                        )),
+                        offset: 16,
+                    ),
+                ],
+                span: 32,
+            ),
+        ),
+        (
+            name: Some("PrimitiveInput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("colorMask"),
+                        ty: 1,
+                        binding: Some(Location(
+                            location: 1,
+                            interpolation: Some(Perspective),
+                            sampling: Some(Center),
+                            blend_src: None,
+                            per_primitive: true,
+                        )),
+                        offset: 0,
+                    ),
+                ],
+                span: 16,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 4,
+                size: Constant(3),
+                stride: 32,
+            ),
+        ),
+        (
+            name: None,
+            inner: Array(
+                base: 7,
+                size: Constant(1),
+                stride: 32,
+            ),
+        ),
+        (
+            name: Some("MeshOutput"),
+            inner: Struct(
+                members: [
+                    (
+                        name: Some("vertices"),
+                        ty: 9,
+                        binding: Some(BuiltIn(Vertices)),
+                        offset: 0,
+                    ),
+                    (
+                        name: Some("primitives"),
+                        ty: 10,
+                        binding: Some(BuiltIn(Primitives)),
+                        offset: 96,
+                    ),
+                    (
+                        name: Some("vertex_count"),
+                        ty: 5,
+                        binding: Some(BuiltIn(VertexCount)),
+                        offset: 128,
+                    ),
+                    (
+                        name: Some("primitive_count"),
+                        ty: 5,
+                        binding: Some(BuiltIn(PrimitiveCount)),
+                        offset: 132,
+                    ),
+                ],
+                span: 144,
+            ),
+        ),
+    ],
+    special_types: (
+        ray_desc: None,
+        ray_intersection: None,
+        ray_vertex_return: None,
+        external_texture_params: None,
+        external_texture_transfer_function: None,
+        predeclared_types: {},
+    ),
+    constants: [],
+    overrides: [],
+    global_variables: [
+        (
+            name: Some("taskPayload"),
+            space: TaskPayload,
+            binding: None,
+            ty: 3,
+            init: None,
+        ),
+        (
+            name: Some("workgroupData"),
+            space: WorkGroup,
+            binding: None,
+            ty: 0,
+            init: None,
+        ),
+        (
+            name: Some("mesh_output"),
+            space: WorkGroup,
+            binding: None,
+            ty: 11,
+            init: None,
+        ),
+    ],
+    global_expressions: [],
+    functions: [],
+    entry_points: [
+        (
+            name: "ts_main",
+            stage: Task,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ts_main"),
+                arguments: [],
+                result: Some((
+                    ty: 6,
+                    binding: Some(BuiltIn(MeshTaskSize)),
+                )),
+                local_variables: [],
+                expressions: [
+                    GlobalVariable(1),
+                    Literal(F32(1.0)),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 2,
+                        index: 0,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            4,
+                            5,
+                            6,
+                            7,
+                        ],
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 9,
+                        index: 1,
+                    ),
+                    Literal(Bool(true)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Literal(U32(1)),
+                    Compose(
+                        ty: 6,
+                        components: [
+                            12,
+                            13,
+                            14,
+                        ],
+                    ),
+                ],
+                named_expressions: {},
+                body: [
+                    Store(
+                        pointer: 0,
+                        value: 1,
+                    ),
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 8,
+                        end: 9,
+                    )),
+                    Store(
+                        pointer: 3,
+                        value: 8,
+                    ),
+                    Emit((
+                        start: 10,
+                        end: 11,
+                    )),
+                    Store(
+                        pointer: 10,
+                        value: 11,
+                    ),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 15,
+                        end: 16,
+                    )),
+                    Return(
+                        value: Some(15),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: Some(0),
+        ),
+        (
+            name: "ms_main",
+            stage: Mesh,
+            early_depth_test: None,
+            workgroup_size: (1, 1, 1),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("ms_main"),
+                arguments: [
+                    (
+                        name: Some("index"),
+                        ty: 5,
+                        binding: Some(BuiltIn(LocalInvocationIndex)),
+                    ),
+                    (
+                        name: Some("id"),
+                        ty: 6,
+                        binding: Some(BuiltIn(GlobalInvocationId)),
+                    ),
+                ],
+                result: None,
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    FunctionArgument(1),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 2,
+                        index: 2,
+                    ),
+                    Literal(U32(3)),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 5,
+                        index: 3,
+                    ),
+                    Literal(U32(1)),
+                    GlobalVariable(1),
+                    Literal(F32(2.0)),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 10,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 11,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 12,
+                        index: 0,
+                    ),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            14,
+                            15,
+                            16,
+                            17,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 19,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 20,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 21,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 23,
+                        index: 0,
+                    ),
+                    Load(
+                        pointer: 24,
+                    ),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            26,
+                            27,
+                            28,
+                            29,
+                        ],
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 30,
+                        right: 25,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 32,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 33,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 34,
+                        index: 0,
+                    ),
+                    Literal(F32(-1.0)),
+                    Literal(F32(-1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            36,
+                            37,
+                            38,
+                            39,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 41,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 42,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 43,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 45,
+                        index: 0,
+                    ),
+                    Load(
+                        pointer: 46,
+                    ),
+                    Literal(F32(0.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            48,
+                            49,
+                            50,
+                            51,
+                        ],
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 52,
+                        right: 47,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 54,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 55,
+                        index: 2,
+                    ),
+                    AccessIndex(
+                        base: 56,
+                        index: 0,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(-1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            58,
+                            59,
+                            60,
+                            61,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 63,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 64,
+                        index: 2,
+                    ),
+                    AccessIndex(
+                        base: 65,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 67,
+                        index: 0,
+                    ),
+                    Load(
+                        pointer: 68,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            70,
+                            71,
+                            72,
+                            73,
+                        ],
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 74,
+                        right: 69,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 76,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 77,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 78,
+                        index: 0,
+                    ),
+                    Literal(U32(0)),
+                    Literal(U32(1)),
+                    Literal(U32(2)),
+                    Compose(
+                        ty: 6,
+                        components: [
+                            80,
+                            81,
+                            82,
+                        ],
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 84,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 85,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 86,
+                        index: 1,
+                    ),
+                    GlobalVariable(0),
+                    AccessIndex(
+                        base: 88,
+                        index: 1,
+                    ),
+                    Load(
+                        pointer: 89,
+                    ),
+                    Unary(
+                        op: LogicalNot,
+                        expr: 90,
+                    ),
+                    GlobalVariable(2),
+                    AccessIndex(
+                        base: 92,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 93,
+                        index: 0,
+                    ),
+                    AccessIndex(
+                        base: 94,
+                        index: 2,
+                    ),
+                    Literal(F32(1.0)),
+                    Literal(F32(0.0)),
+                    Literal(F32(1.0)),
+                    Literal(F32(1.0)),
+                    Compose(
+                        ty: 1,
+                        components: [
+                            96,
+                            97,
+                            98,
+                            99,
+                        ],
+                    ),
+                ],
+                named_expressions: {
+                    0: "index",
+                    1: "id",
+                },
+                body: [
+                    Emit((
+                        start: 3,
+                        end: 4,
+                    )),
+                    Store(
+                        pointer: 3,
+                        value: 4,
+                    ),
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Store(
+                        pointer: 6,
+                        value: 7,
+                    ),
+                    Store(
+                        pointer: 8,
+                        value: 9,
+                    ),
+                    Emit((
+                        start: 11,
+                        end: 12,
+                    )),
+                    Emit((
+                        start: 12,
+                        end: 14,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 18,
+                        end: 19,
+                    )),
+                    Store(
+                        pointer: 13,
+                        value: 18,
+                    ),
+                    Emit((
+                        start: 20,
+                        end: 21,
+                    )),
+                    Emit((
+                        start: 21,
+                        end: 23,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 24,
+                        end: 26,
+                    )),
+                    Emit((
+                        start: 30,
+                        end: 32,
+                    )),
+                    Store(
+                        pointer: 22,
+                        value: 31,
+                    ),
+                    Emit((
+                        start: 33,
+                        end: 34,
+                    )),
+                    Emit((
+                        start: 34,
+                        end: 36,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 40,
+                        end: 41,
+                    )),
+                    Store(
+                        pointer: 35,
+                        value: 40,
+                    ),
+                    Emit((
+                        start: 42,
+                        end: 43,
+                    )),
+                    Emit((
+                        start: 43,
+                        end: 45,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 46,
+                        end: 48,
+                    )),
+                    Emit((
+                        start: 52,
+                        end: 54,
+                    )),
+                    Store(
+                        pointer: 44,
+                        value: 53,
+                    ),
+                    Emit((
+                        start: 55,
+                        end: 56,
+                    )),
+                    Emit((
+                        start: 56,
+                        end: 58,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 62,
+                        end: 63,
+                    )),
+                    Store(
+                        pointer: 57,
+                        value: 62,
+                    ),
+                    Emit((
+                        start: 64,
+                        end: 65,
+                    )),
+                    Emit((
+                        start: 65,
+                        end: 67,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 0,
+                        end: 0,
+                    )),
+                    Emit((
+                        start: 68,
+                        end: 70,
+                    )),
+                    Emit((
+                        start: 74,
+                        end: 76,
+                    )),
+                    Store(
+                        pointer: 66,
+                        value: 75,
+                    ),
+                    Emit((
+                        start: 77,
+                        end: 78,
+                    )),
+                    Emit((
+                        start: 78,
+                        end: 80,
+                    )),
+                    Emit((
+                        start: 83,
+                        end: 84,
+                    )),
+                    Store(
+                        pointer: 79,
+                        value: 83,
+                    ),
+                    Emit((
+                        start: 85,
+                        end: 86,
+                    )),
+                    Emit((
+                        start: 86,
+                        end: 88,
+                    )),
+                    Emit((
+                        start: 89,
+                        end: 92,
+                    )),
+                    Store(
+                        pointer: 87,
+                        value: 91,
+                    ),
+                    Emit((
+                        start: 93,
+                        end: 94,
+                    )),
+                    Emit((
+                        start: 94,
+                        end: 96,
+                    )),
+                    Emit((
+                        start: 100,
+                        end: 101,
+                    )),
+                    Store(
+                        pointer: 95,
+                        value: 100,
+                    ),
+                    Return(
+                        value: None,
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: Some((
+                topology: Triangles,
+                max_vertices: 3,
+                max_vertices_override: None,
+                max_primitives: 1,
+                max_primitives_override: None,
+                vertex_output_type: 4,
+                primitive_output_type: 7,
+                output_variable: 2,
+            )),
+            task_payload: Some(0),
+        ),
+        (
+            name: "fs_main",
+            stage: Fragment,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            workgroup_size_overrides: None,
+            function: (
+                name: Some("fs_main"),
+                arguments: [
+                    (
+                        name: Some("vertex"),
+                        ty: 4,
+                        binding: None,
+                    ),
+                    (
+                        name: Some("primitive"),
+                        ty: 8,
+                        binding: None,
+                    ),
+                ],
+                result: Some((
+                    ty: 1,
+                    binding: Some(Location(
+                        location: 0,
+                        interpolation: Some(Perspective),
+                        sampling: Some(Center),
+                        blend_src: None,
+                        per_primitive: false,
+                    )),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    FunctionArgument(1),
+                    AccessIndex(
+                        base: 0,
+                        index: 1,
+                    ),
+                    AccessIndex(
+                        base: 1,
+                        index: 0,
+                    ),
+                    Binary(
+                        op: Multiply,
+                        left: 2,
+                        right: 3,
+                    ),
+                ],
+                named_expressions: {
+                    0: "vertex",
+                    1: "primitive",
+                },
+                body: [
+                    Emit((
+                        start: 2,
+                        end: 5,
+                    )),
+                    Return(
+                        value: Some(4),
+                    ),
+                ],
+                diagnostic_filter_leaf: None,
+            ),
+            mesh_info: None,
+            task_payload: None,
+        ),
+    ],
+    diagnostic_filters: [],
+    diagnostic_filter_leaf: None,
+    doc_comments: None,
+)
\ No newline at end of file
diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs
index ec5203f291b..8683b37dd0e 100644
--- a/wgpu-core/src/device/mod.rs
+++ b/wgpu-core/src/device/mod.rs
@@ -514,6 +514,14 @@ pub fn create_validator(
         Caps::SHADER_BARYCENTRICS,
         features.intersects(wgt::Features::SHADER_BARYCENTRICS),
     );
+    caps.set(
+        Caps::MESH_SHADER,
+        features.intersects(wgt::Features::EXPERIMENTAL_MESH_SHADER),
+    );
+    caps.set(
+        Caps::MESH_SHADER_POINT_TOPOLOGY,
+        features.intersects(wgt::Features::EXPERIMENTAL_MESH_SHADER_POINTS),
+    );
 
     naga::valid::Validator::new(flags, caps)
 }
diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs
index dbd61694fde..f890336cff4 100644
--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@@ -917,6 +917,10 @@ impl PhysicalDeviceFeatures {
             F::EXPERIMENTAL_MESH_SHADER,
             caps.supports_extension(ext::mesh_shader::NAME),
         );
+        features.set(
+            F::EXPERIMENTAL_MESH_SHADER_POINTS,
+            caps.supports_extension(ext::mesh_shader::NAME),
+        );
         if let Some(ref mesh_shader) = self.mesh_shader {
             features.set(
                 F::EXPERIMENTAL_MESH_SHADER_MULTIVIEW,
diff --git a/wgpu-types/src/features.rs b/wgpu-types/src/features.rs
index 234d01d8544..f3a0a69e8a2 100644
--- a/wgpu-types/src/features.rs
+++ b/wgpu-types/src/features.rs
@@ -1255,6 +1255,16 @@ bitflags_array! {
         ///
         /// While metal supports this in theory, the behavior of `view_index` differs from vulkan and dx12 so the feature isn't exposed.
         const SELECTIVE_MULTIVIEW = 1 << 54;
+
+        /// Enables the use of point-primitive outputs from mesh shaders. Making use of this feature also requires enabling
+        /// `Features::EXPERIMENTAL_MESH_SHADER`.
+        ///
+        /// Supported platforms
+        /// - Vulkan
+        /// - Metal
+        ///
+        /// This is a native only feature.
+        const EXPERIMENTAL_MESH_SHADER_POINTS = 1 << 55;
     }
 
     /// Features that are not guaranteed to be supported.
@@ -1532,6 +1542,7 @@ impl Features {
         Self::from_bits_truncate(FeatureBits([
             FeaturesWGPU::EXPERIMENTAL_MESH_SHADER.bits()
                 | FeaturesWGPU::EXPERIMENTAL_MESH_SHADER_MULTIVIEW.bits()
+                | FeaturesWGPU::EXPERIMENTAL_MESH_SHADER_POINTS.bits()
                 | FeaturesWGPU::EXPERIMENTAL_RAY_QUERY.bits()
                 | FeaturesWGPU::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN.bits()
                 | FeaturesWGPU::EXPERIMENTAL_PASSTHROUGH_SHADERS.bits(),