diff --git a/docs/api-specs/mesh_shading.md b/docs/api-specs/mesh_shading.md
index 46d7948d55e..df0a5149f9f 100644
--- a/docs/api-specs/mesh_shading.md
+++ b/docs/api-specs/mesh_shading.md
@@ -2,13 +2,37 @@
 
 🧪Experimental🧪
 
-`wgpu` supports an experimental version of mesh shading. Currently `naga` has no support for mesh shaders beyond recognizing the additional shader stages.
+`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`.
 
 **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).
 
-***This is not*** a thorough explanation of mesh shading and how it works. Those wishing to understand mesh shading more broadly should look elsewhere first.
+## Mesh shaders overview
+
+### What are mesh shaders
+Mesh shaders are a new kind of rasterization pipeline intended to address some of the shortfalls with the vertex shader pipeline. The core idea of mesh shaders is that the GPU decides how to render the many small parts of a scene instead of the CPU issuing a draw call for every small part or issuing an inefficient monolithic draw call for a large part of the scene.
+
+Mesh shaders are specifically designed to be used with **meshlet rendering**, a technique where every object is split into many subobjects called meshlets that are each rendered with their own parameters. With the standard vertex pipeline, each draw call specifies an exact number of primitives to render and the same parameters for all vertex shaders on an entire object (or even multiple objects). This doesn't leave room for different LODs for different parts of an object, for example a closer part having more detail, nor does it allow culling smaller sections (or primitives) of objects. With mesh shaders, each task workgroup might get assigned to a single object. It can then analyze the different meshlets(sections) of that object, determine which are visible and should actually be rendered, and for those meshlets determine what LOD to use based on the distance from the camera. It can then dispatch a mesh workgroup for each meshlet, with each mesh workgroup then reading the data for that specific LOD of its meshlet, determining which and how many vertices and primitives to output, determining which remaining primitives need to be culled, and passing the resulting primitives to the rasterizer.
+
+Mesh shaders are most effective in scenes with many polygons. They can allow skipping processing of entire groups of primitives that are facing away from the camera or otherwise occluded, which reduces the number of primitives that need to be processed by more than half in most cases, and they can reduce the number of primitives that need to be processed for more distant objects. Scenes that are not bottlenecked by geometry (perhaps instead by fragment processing or post processing) will not see much benefit from using them.
+
+Mesh shaders were first shown off in [NVIDIA's asteroids demo](https://www.youtube.com/watch?v=CRfZYJ_sk5E). Now, they form the basis for [Unreal Engine's Nanite](https://www.unrealengine.com/en-US/blog/unreal-engine-5-is-now-available-in-preview#Nanite).
+
+### Mesh shader pipeline
+A mesh shader pipeline is just like a standard render pipeline, except that the vertex shader stage is replaced by a mesh shader stage (and optionally a task shader stage). This functions as follows:
+
+* If there is a task shader stage, task shader workgroups are invoked first, with the number of workgroups determined by the draw call. Each task shader workgroup outputs a workgroup size and a task payload. A dispatch group of mesh shaders with the given workgroup size is then invoked with the task payload as a parameter.
+* Otherwise, a single dispatch group of mesh shaders with workgroup size given by the draw call is invoked.
+* Each mesh shader dispatch group functions exactly as a compute dispatch group, except that it has special outputs and may take a task payload as input. Mesh dispatch groups invoked by different task shader workgroups cannot interact.
+* Each workgroup within the mesh shader dispatch group can output vertices and primitives
+  * It determines how many vertices and primitives to write and then sets those vertices and primitives.
+  * Primitives have an indices field which determines the indices of the vertices of that primitive. The indices are based on the output of that mesh shader workgroup only; there is no sharing of vertices across workgroups (no vertex or index buffer equivalents).
+  * Primitives can then be culled by setting the appropriate builtin
+  * Each vertex output functions exactly as the output from a vertex shader would
+  * There can also be per-primitive outputs passed to fragment shaders; these are not interpolated or based on the vertices of the primitive in any way.
+* Once all of the primitives are written, those that weren't culled are are rasterized. From this point forward, the only difference from a standard render pipeline is that there may be some per-primitive inputs passed to fragment shaders.
 
 ## `wgpu` API
 
@@ -79,32 +103,36 @@ This shader stage can be selected by marking a function with `@task`. Task shade
 
 The output of this determines how many workgroups of mesh shaders will be dispatched. Once dispatched, global id variables will be local to the task shader workgroup dispatch, and mesh shaders won't know the position of their dispatch among all mesh shader dispatches unless this is passed through the payload. The output may be zero to skip dispatching any mesh shader workgroups for the task shader workgroup.
 
-If task shaders are marked with `@payload(someVar)`, where `someVar` is global variable declared like `var<workgroup> someVar: <type>`, task shaders may write to `someVar`. This payload is passed to the mesh shader workgroup that is invoked. The mesh shader can skip declaring `@payload` to ignore this input.
+Task shaders must be marked with `@payload(someVar)`, where `someVar` is global variable declared like `var<task_payload> someVar: <type>`. Task shaders may use `someVar` as if it is a read-write workgroup storage variable. This payload is passed to the mesh shader workgroup that is invoked.
 
 ### Mesh shader
 This shader stage can be selected by marking a function with `@mesh`. Mesh shaders must not return anything.
 
-Mesh shaders can be marked with `@payload(someVar)` similar to task shaders. Unlike task shaders, mesh shaders cannot write to this workgroup memory. Declaring `@payload` in a pipeline with no task shader, in a pipeline with a task shader that doesn't declare `@payload`, or in a task shader with an `@payload` that is statically sized and smaller than the mesh shader payload is illegal.
+Mesh shaders can be marked with `@payload(someVar)` similar to task shaders. Unlike task shaders, this is optional, and mesh shaders cannot write to this memory. Declaring `@payload` in a pipeline with no task shader or in a task shader with an `@payload` that is statically sized and differently than the mesh shader payload is illegal. The `@payload` attribute can only be ignored in pipelines that don't have a task shader.
 
-Mesh shaders must be marked with `@vertex_output(OutputType, numOutputs)`, where `numOutputs` is the maximum number of vertices to be output by a mesh shader, and `OutputType` is the data associated with vertices, similar to a standard vertex shader output.
+Mesh shaders must be marked with `@vertex_output(OutputType, numOutputs)`, where `numOutputs` is the maximum number of vertices to be output by a mesh shader, and `OutputType` is the data associated with vertices, similar to a standard vertex shader output, and must be a struct.
 
 Mesh shaders must also be marked with `@primitive_output(OutputType, numOutputs)`, which is similar to `@vertex_output` except it describes the primitive outputs.
 
 ### Mesh shader outputs
 
-Primitive outputs from mesh shaders have some additional builtins they can set. These include `@builtin(cull_primitive)`, which must be a boolean value. If this is set to true, then the primitive is skipped during rendering.
+Vertex outputs from mesh shaders function identically to outputs of vertex shaders, and as such must have a field with `@builtin(position)`.
+
+Primitive outputs from mesh shaders have some additional builtins they can set. These include `@builtin(cull_primitive)`, which must be a boolean value. If this is set to true, then the primitive is skipped during rendering. All non-builtin primitive outputs must be decorated with `@per_primitive`.
 
 Mesh shader primitive outputs must also specify exactly one of `@builtin(triangle_indices)`, `@builtin(line_indices)`, or `@builtin(point_index)`. This determines the output topology of the mesh shader, and must match the output topology of the pipeline descriptor the mesh shader is used with. These must be of type `vec3<u32>`, `vec2<u32>`, and `u32` respectively. When setting this, each of the indices must be less than the number of vertices declared in `setMeshOutputs`.
 
 Additionally, the `@location` attributes from the vertex and primitive outputs can't overlap.
 
-Before setting any vertices or indices, or exiting, the mesh shader must call `setMeshOutputs(numVertices: u32, numIndices: u32)`, which declares the number of vertices and indices that will be written to. These must be less than the corresponding maximums set in `@vertex_output` and `@primitive_output`. The mesh shader must then write to exactly these numbers of vertices and primitives.
+Before exiting, the mesh shader must call `setMeshOutputs(numVertices: u32, numIndices: u32)`, which declares the number of vertices and indices that will be written to. These must be less than the corresponding maximums set in `@vertex_output` and `@primitive_output`. The mesh shader must then write to exactly this range of vertices and primitives. A varying member with `@per_primitive` cannot be used in function interfaces except as a primitive output for mesh shaders or as input for fragment shaders.
 
 The mesh shader can write to vertices using the `setVertex(idx: u32, vertex: VertexOutput)` where `VertexOutput` is replaced with the vertex type declared in `@vertex_output`, and `idx` is the index of the vertex to write. Similarly, the mesh shader can write to vertices using `setPrimitive(idx: u32, primitive: PrimitiveOutput)`. These can be written to multiple times, however unsynchronized writes are undefined behavior. The primitives and indices are shared across the entire mesh shader workgroup.
 
 ### Fragment shader
 
-Fragment shaders may now be passed the primitive info from a mesh shader the same was as they are passed vertex inputs, for example `fn fs_main(vertex: VertexOutput, primitive: PrimitiveOutput)`. The primitive state is part of the fragment input and must match the output of the mesh shader in the pipeline.
+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`.
+
+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.
 
 ### Full example
 
@@ -114,9 +142,9 @@ The following is a full example of WGSL shaders that could be used to create a m
 enable mesh_shading;
 
 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.),
@@ -127,7 +155,7 @@ struct TaskPayload {
 	colorMask: vec4<f32>,
 	visible: bool,
 }
-var<workgroup> taskPayload: TaskPayload;
+var<task_payload> taskPayload: TaskPayload;
 var<workgroup> workgroupData: f32;
 struct VertexOutput {
 	@builtin(position) position: vec4<f32>,
@@ -136,14 +164,12 @@ struct VertexOutput {
 struct PrimitiveOutput {
 	@builtin(triangle_indices) index: vec3<u32>,
 	@builtin(cull_primitive) cull: bool,
-	@location(1) colorMask: vec4<f32>,
+	@per_primitive @location(1) colorMask: vec4<f32>,
 }
 struct PrimitiveInput {
-	@location(1) colorMask: vec4<f32>,
+	@per_primitive @location(1) colorMask: vec4<f32>,
 }
-fn test_function(input: u32) {
 
-}
 @task
 @payload(taskPayload)
 @workgroup_size(1)
@@ -162,8 +188,6 @@ fn ms_main(@builtin(local_invocation_index) index: u32, @builtin(global_invocati
 	workgroupData = 2.0;
 	var v: VertexOutput;
 
-	test_function(1);
-
 	v.position = positions[0];
 	v.color = colors[0] * taskPayload.colorMask;
 	setVertex(0, v);
diff --git a/naga-cli/src/bin/naga.rs b/naga-cli/src/bin/naga.rs
index bbdc63fad5b..0a46d3e8ea6 100644
--- a/naga-cli/src/bin/naga.rs
+++ b/naga-cli/src/bin/naga.rs
@@ -64,6 +64,12 @@ struct Args {
     #[argh(option)]
     shader_model: Option<ShaderModelArg>,
 
+    /// the SPIR-V version to use if targeting SPIR-V
+    ///
+    /// For example, 1.0, 1.4, etc
+    #[argh(option)]
+    spirv_version: Option<SpirvVersionArg>,
+
     /// the shader stage, for example 'frag', 'vert', or 'compute'.
     /// if the shader stage is unspecified it will be derived from
     /// the file extension.
@@ -189,6 +195,22 @@ impl FromStr for ShaderModelArg {
     }
 }
 
+#[derive(Debug, Clone)]
+struct SpirvVersionArg(u8, u8);
+
+impl FromStr for SpirvVersionArg {
+    type Err = String;
+
+    fn from_str(s: &str) -> Result<Self, Self::Err> {
+        let dot = s
+            .find(".")
+            .ok_or_else(|| "Missing dot separator".to_owned())?;
+        let major = s[..dot].parse::<u8>().map_err(|e| e.to_string())?;
+        let minor = s[dot + 1..].parse::<u8>().map_err(|e| e.to_string())?;
+        Ok(Self(major, minor))
+    }
+}
+
 /// Newtype so we can implement [`FromStr`] for `ShaderSource`.
 #[derive(Debug, Clone, Copy)]
 struct ShaderStage(naga::ShaderStage);
@@ -465,6 +487,9 @@ fn run() -> anyhow::Result<()> {
     if let Some(ref version) = args.metal_version {
         params.msl.lang_version = version.0;
     }
+    if let Some(ref version) = args.spirv_version {
+        params.spv_out.lang_version = (version.0, version.1);
+    }
     params.keep_coordinate_space = args.keep_coordinate_space;
 
     params.dot.cfg_only = args.dot_cfg_only;
diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs
index 826dad1c219..1f1396eccff 100644
--- a/naga/src/back/dot/mod.rs
+++ b/naga/src/back/dot/mod.rs
@@ -307,6 +307,25 @@ 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/features.rs b/naga/src/back/glsl/features.rs
index a6dfe4e3100..b884f08ac39 100644
--- a/naga/src/back/glsl/features.rs
+++ b/naga/src/back/glsl/features.rs
@@ -610,6 +610,7 @@ impl<W> Writer<'_, W> {
                     interpolation,
                     sampling,
                     blend_src,
+                    per_primitive: _,
                 } => {
                     if interpolation == Some(Interpolation::Linear) {
                         self.features.request(Features::NOPERSPECTIVE_QUALIFIER);
diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs
index e78af74c844..1af18528944 100644
--- a/naga/src/back/glsl/mod.rs
+++ b/naga/src/back/glsl/mod.rs
@@ -139,7 +139,8 @@ impl crate::AddressSpace {
             | crate::AddressSpace::Uniform
             | crate::AddressSpace::Storage { .. }
             | crate::AddressSpace::Handle
-            | crate::AddressSpace::PushConstant => false,
+            | crate::AddressSpace::PushConstant
+            | crate::AddressSpace::TaskPayload => false,
         }
     }
 }
@@ -1300,6 +1301,9 @@ impl<'a, W: Write> Writer<'a, W> {
             crate::AddressSpace::Storage { .. } => {
                 self.write_interface_block(handle, global)?;
             }
+            crate::AddressSpace::TaskPayload => {
+                self.write_interface_block(handle, global)?;
+            }
             // A global variable in the `Function` address space is a
             // contradiction in terms.
             crate::AddressSpace::Function => unreachable!(),
@@ -1614,6 +1618,7 @@ impl<'a, W: Write> Writer<'a, W> {
                 interpolation,
                 sampling,
                 blend_src,
+                per_primitive: _,
             } => (location, interpolation, sampling, blend_src),
             crate::Binding::BuiltIn(built_in) => {
                 match built_in {
@@ -1732,6 +1737,7 @@ impl<'a, W: Write> Writer<'a, W> {
                 interpolation: None,
                 sampling: None,
                 blend_src,
+                per_primitive: false,
             },
             stage: self.entry_point.stage,
             options: VaryingOptions::from_writer_options(self.options, output),
@@ -2669,6 +2675,11 @@ 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();
@@ -5247,6 +5258,15 @@ const fn glsl_built_in(built_in: crate::BuiltIn, options: VaryingOptions) -> &'s
         Bi::SubgroupId => "gl_SubgroupID",
         Bi::SubgroupSize => "gl_SubgroupSize",
         Bi::SubgroupInvocationId => "gl_SubgroupInvocationID",
+        // mesh
+        // TODO: figure out how to map these to glsl things as glsl treats them as arrays
+        Bi::CullPrimitive
+        | Bi::PointIndex
+        | Bi::LineIndices
+        | Bi::TriangleIndices
+        | Bi::MeshTaskSize => {
+            unimplemented!()
+        }
     }
 }
 
@@ -5262,6 +5282,7 @@ const fn glsl_storage_qualifier(space: crate::AddressSpace) -> Option<&'static s
         As::Handle => Some("uniform"),
         As::WorkGroup => Some("shared"),
         As::PushConstant => Some("uniform"),
+        As::TaskPayload => unreachable!(),
     }
 }
 
diff --git a/naga/src/back/hlsl/conv.rs b/naga/src/back/hlsl/conv.rs
index ed40cbe5102..d6ccc5ec6e4 100644
--- a/naga/src/back/hlsl/conv.rs
+++ b/naga/src/back/hlsl/conv.rs
@@ -183,6 +183,9 @@ impl crate::BuiltIn {
             Self::PointSize | Self::ViewIndex | Self::PointCoord | Self::DrawID => {
                 return Err(Error::Custom(format!("Unsupported builtin {self:?}")))
             }
+            Self::CullPrimitive => "SV_CullPrimitive",
+            Self::PointIndex | Self::LineIndices | Self::TriangleIndices => unimplemented!(),
+            Self::MeshTaskSize => unreachable!(),
         })
     }
 }
diff --git a/naga/src/back/hlsl/mod.rs b/naga/src/back/hlsl/mod.rs
index 8df06cf1323..f357c02bb3f 100644
--- a/naga/src/back/hlsl/mod.rs
+++ b/naga/src/back/hlsl/mod.rs
@@ -283,7 +283,8 @@ impl crate::ShaderStage {
             Self::Vertex => "vs",
             Self::Fragment => "ps",
             Self::Compute => "cs",
-            Self::Task | Self::Mesh => unreachable!(),
+            Self::Task => "ts",
+            Self::Mesh => "ms",
         }
     }
 }
diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs
index 357b8597521..9401766448f 100644
--- a/naga/src/back/hlsl/writer.rs
+++ b/naga/src/back/hlsl/writer.rs
@@ -507,7 +507,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
 
             self.write_wrapped_functions(module, &ctx)?;
 
-            if ep.stage == ShaderStage::Compute {
+            if ep.stage.compute_like() {
                 // HLSL is calling workgroup size "num threads"
                 let num_threads = ep.workgroup_size;
                 writeln!(
@@ -967,6 +967,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 self.write_type(module, global.ty)?;
                 ""
             }
+            crate::AddressSpace::TaskPayload => unimplemented!(),
             crate::AddressSpace::Uniform => {
                 // constant buffer declarations are expected to be inlined, e.g.
                 // `cbuffer foo: register(b0) { field1: type1; }`
@@ -2599,6 +2600,19 @@ 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();
@@ -3076,7 +3090,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                                 crate::AddressSpace::Function
                                 | crate::AddressSpace::Private
                                 | crate::AddressSpace::WorkGroup
-                                | crate::AddressSpace::PushConstant,
+                                | crate::AddressSpace::PushConstant
+                                | crate::AddressSpace::TaskPayload,
                             )
                             | None => true,
                             Some(crate::AddressSpace::Uniform) => {
diff --git a/naga/src/back/msl/mod.rs b/naga/src/back/msl/mod.rs
index 7e9180e0c25..edba50f63b8 100644
--- a/naga/src/back/msl/mod.rs
+++ b/naga/src/back/msl/mod.rs
@@ -531,6 +531,7 @@ impl Options {
                 interpolation,
                 sampling,
                 blend_src,
+                per_primitive: _,
             } => match mode {
                 LocationMode::VertexInput => Ok(ResolvedBinding::Attribute(location)),
                 LocationMode::FragmentOutput => {
@@ -688,6 +689,10 @@ impl ResolvedBinding {
                     Bi::CullDistance | Bi::ViewIndex | Bi::DrawID => {
                         return Err(Error::UnsupportedBuiltIn(built_in))
                     }
+                    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!(),
                 };
                 write!(out, "{name}")?;
             }
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 9c0ac10d160..e9f846f737d 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -594,7 +594,8 @@ impl crate::AddressSpace {
             | Self::Private
             | Self::WorkGroup
             | Self::PushConstant
-            | Self::Handle => true,
+            | Self::Handle
+            | Self::TaskPayload => true,
             Self::Function => false,
         }
     }
@@ -607,6 +608,7 @@ impl crate::AddressSpace {
             // may end up with "const" even if the binding is read-write,
             // and that should be OK.
             Self::Storage { .. } => true,
+            Self::TaskPayload => unimplemented!(),
             // These should always be read-write.
             Self::Private | Self::WorkGroup => false,
             // These translate to `constant` address space, no need for qualifiers.
@@ -623,6 +625,7 @@ impl crate::AddressSpace {
             Self::Storage { .. } => Some("device"),
             Self::Private | Self::Function => Some("thread"),
             Self::WorkGroup => Some("threadgroup"),
+            Self::TaskPayload => Some("object_data"),
         }
     }
 }
@@ -4060,6 +4063,14 @@ 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("");
@@ -6615,7 +6626,7 @@ template <typename A>
                     LocationMode::Uniform,
                     false,
                 ),
-                crate::ShaderStage::Task | crate::ShaderStage::Mesh => unreachable!(),
+                crate::ShaderStage::Task | crate::ShaderStage::Mesh => unimplemented!(),
             };
 
             // Should this entry point be modified to do vertex pulling?
@@ -6682,6 +6693,9 @@ template <typename A>
                                 break;
                             }
                         }
+                        crate::AddressSpace::TaskPayload => {
+                            unimplemented!()
+                        }
                         crate::AddressSpace::Function
                         | crate::AddressSpace::Private
                         | crate::AddressSpace::WorkGroup => {}
@@ -7694,7 +7708,7 @@ mod workgroup_mem_init {
             fun_info: &valid::FunctionInfo,
         ) -> bool {
             options.zero_initialize_workgroup_memory
-                && ep.stage == crate::ShaderStage::Compute
+                && ep.stage.compute_like()
                 && module.global_variables.iter().any(|(handle, var)| {
                     !fun_info[handle].is_empty() && var.space == crate::AddressSpace::WorkGroup
                 })
diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs
index d2b3ed70eda..c009082a3c9 100644
--- a/naga/src/back/pipeline_constants.rs
+++ b/naga/src/back/pipeline_constants.rs
@@ -39,6 +39,8 @@ pub enum PipelineConstantError {
     ValidationError(#[from] WithSpan<ValidationError>),
     #[error("workgroup_size override isn't strictly positive")]
     NegativeWorkgroupSize,
+    #[error("max vertices or max primitives is negative")]
+    NegativeMeshOutputMax,
 }
 
 /// Compact `module` and replace all overrides with constants.
@@ -243,6 +245,7 @@ pub fn process_overrides<'a>(
     for ep in entry_points.iter_mut() {
         process_function(&mut module, &override_map, &mut layouter, &mut ep.function)?;
         process_workgroup_size_override(&mut module, &adjusted_global_expressions, ep)?;
+        process_mesh_shader_overrides(&mut module, &adjusted_global_expressions, ep)?;
     }
     module.entry_points = entry_points;
     module.overrides = overrides;
@@ -296,6 +299,28 @@ fn process_workgroup_size_override(
     Ok(())
 }
 
+fn process_mesh_shader_overrides(
+    module: &mut Module,
+    adjusted_global_expressions: &HandleVec<Expression, Handle<Expression>>,
+    ep: &mut crate::EntryPoint,
+) -> Result<(), PipelineConstantError> {
+    if let Some(ref mut mesh_info) = ep.mesh_info {
+        if let Some(r#override) = mesh_info.max_vertices_override {
+            mesh_info.max_vertices = module
+                .to_ctx()
+                .eval_expr_to_u32(adjusted_global_expressions[r#override])
+                .map_err(|_| PipelineConstantError::NegativeWorkgroupSize)?;
+        }
+        if let Some(r#override) = mesh_info.max_primitives_override {
+            mesh_info.max_primitives = module
+                .to_ctx()
+                .eval_expr_to_u32(adjusted_global_expressions[r#override])
+                .map_err(|_| PipelineConstantError::NegativeWorkgroupSize)?;
+        }
+    }
+    Ok(())
+}
+
 /// Add a [`Constant`] to `module` for the override `old_h`.
 ///
 /// Add the new `Constant` to `override_map` and `adjusted_constant_initializers`.
@@ -835,6 +860,26 @@ 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 7758d86c414..258d6869c99 100644
--- a/naga/src/back/spv/block.rs
+++ b/naga/src/back/spv/block.rs
@@ -3654,6 +3654,7 @@ impl BlockContext<'_> {
                 } => {
                     self.write_subgroup_gather(mode, argument, result, &mut block)?;
                 }
+                Statement::MeshFunction(_) => unreachable!(),
             }
         }
 
diff --git a/naga/src/back/spv/helpers.rs b/naga/src/back/spv/helpers.rs
index 84e130efaa3..f6d26794e70 100644
--- a/naga/src/back/spv/helpers.rs
+++ b/naga/src/back/spv/helpers.rs
@@ -54,6 +54,7 @@ pub(super) const fn map_storage_class(space: crate::AddressSpace) -> spirv::Stor
         crate::AddressSpace::Uniform => spirv::StorageClass::Uniform,
         crate::AddressSpace::WorkGroup => spirv::StorageClass::Workgroup,
         crate::AddressSpace::PushConstant => spirv::StorageClass::PushConstant,
+        crate::AddressSpace::TaskPayload => unreachable!(),
     }
 }
 
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index 48d13f03c25..d313ab2a39d 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -1952,6 +1952,7 @@ impl Writer {
                 interpolation,
                 sampling,
                 blend_src,
+                per_primitive: _,
             } => {
                 self.decorate(id, Decoration::Location, &[location]);
 
@@ -2101,6 +2102,11 @@ impl Writer {
                         )?;
                         BuiltIn::SubgroupLocalInvocationId
                     }
+                    Bi::MeshTaskSize
+                    | Bi::CullPrimitive
+                    | Bi::PointIndex
+                    | Bi::LineIndices
+                    | Bi::TriangleIndices => 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 8982242daca..245bc40dd5d 100644
--- a/naga/src/back/wgsl/writer.rs
+++ b/naga/src/back/wgsl/writer.rs
@@ -207,7 +207,7 @@ impl<W: Write> Writer<W> {
                     Attribute::Stage(ShaderStage::Compute),
                     Attribute::WorkGroupSize(ep.workgroup_size),
                 ],
-                ShaderStage::Task | ShaderStage::Mesh => unreachable!(),
+                ShaderStage::Mesh | ShaderStage::Task => unreachable!(),
             };
 
             self.write_attributes(&attributes)?;
@@ -856,6 +856,7 @@ 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();
@@ -1822,6 +1823,7 @@ fn map_binding_to_attribute(binding: &crate::Binding) -> Vec<Attribute> {
             interpolation,
             sampling,
             blend_src: None,
+            per_primitive: _,
         } => vec![
             Attribute::Location(location),
             Attribute::Interpolate(interpolation, sampling),
@@ -1831,6 +1833,7 @@ fn map_binding_to_attribute(binding: &crate::Binding) -> Vec<Attribute> {
             interpolation,
             sampling,
             blend_src: Some(blend_src),
+            per_primitive: _,
         } => vec![
             Attribute::Location(location),
             Attribute::BlendSrc(blend_src),
diff --git a/naga/src/common/wgsl/to_wgsl.rs b/naga/src/common/wgsl/to_wgsl.rs
index 035c4eafb32..dc891aa5a3f 100644
--- a/naga/src/common/wgsl/to_wgsl.rs
+++ b/naga/src/common/wgsl/to_wgsl.rs
@@ -188,7 +188,12 @@ impl TryToWgsl for crate::BuiltIn {
             | Bi::PointSize
             | Bi::DrawID
             | Bi::PointCoord
-            | Bi::WorkGroupSize => return None,
+            | Bi::WorkGroupSize
+            | Bi::CullPrimitive
+            | Bi::TriangleIndices
+            | Bi::LineIndices
+            | Bi::MeshTaskSize
+            | Bi::PointIndex => return None,
         })
     }
 }
@@ -352,6 +357,7 @@ pub const fn address_space_str(
             As::WorkGroup => "workgroup",
             As::Handle => return (None, None),
             As::Function => "function",
+            As::TaskPayload => return (None, None),
         }),
         None,
     )
diff --git a/naga/src/compact/mod.rs b/naga/src/compact/mod.rs
index d059ba21e4f..a7d3d463f11 100644
--- a/naga/src/compact/mod.rs
+++ b/naga/src/compact/mod.rs
@@ -221,6 +221,45 @@ 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,
@@ -342,6 +381,23 @@ pub fn compact(module: &mut crate::Module, keep_unused: KeepUnused) {
             &module_map,
             &mut reused_named_expressions,
         );
+        if let Some(ref mut task_payload) = entry.task_payload {
+            module_map.globals.adjust(task_payload);
+        }
+        if let Some(ref mut mesh_info) = entry.mesh_info {
+            module_map.types.adjust(&mut mesh_info.vertex_output_type);
+            module_map
+                .types
+                .adjust(&mut mesh_info.primitive_output_type);
+            if let Some(ref mut max_vertices_override) = mesh_info.max_vertices_override {
+                module_map.global_expressions.adjust(max_vertices_override);
+            }
+            if let Some(ref mut max_primitives_override) = mesh_info.max_primitives_override {
+                module_map
+                    .global_expressions
+                    .adjust(max_primitives_override);
+            }
+        }
     }
 }
 
diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs
index 39d6065f5f0..b370501baca 100644
--- a/naga/src/compact/statements.rs
+++ b/naga/src/compact/statements.rs
@@ -117,6 +117,20 @@ 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);
@@ -335,6 +349,26 @@ 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/glsl/functions.rs b/naga/src/front/glsl/functions.rs
index 7de7364cd40..ba096a82b3b 100644
--- a/naga/src/front/glsl/functions.rs
+++ b/naga/src/front/glsl/functions.rs
@@ -1377,6 +1377,8 @@ impl Frontend {
                 result: ty.map(|ty| FunctionResult { ty, binding: None }),
                 ..Default::default()
             },
+            mesh_info: None,
+            task_payload: None,
         });
 
         Ok(())
@@ -1446,6 +1448,7 @@ impl Context<'_> {
                         interpolation,
                         sampling: None,
                         blend_src: None,
+                        per_primitive: false,
                     };
                     location += 1;
 
@@ -1482,6 +1485,7 @@ impl Context<'_> {
                                 interpolation,
                                 sampling: None,
                                 blend_src: None,
+                                per_primitive: false,
                             };
                             location += 1;
                             binding
diff --git a/naga/src/front/glsl/mod.rs b/naga/src/front/glsl/mod.rs
index 876add46a1c..e5eda6b3ad9 100644
--- a/naga/src/front/glsl/mod.rs
+++ b/naga/src/front/glsl/mod.rs
@@ -107,7 +107,7 @@ impl ShaderMetadata {
         self.version = 0;
         self.profile = Profile::Core;
         self.stage = stage;
-        self.workgroup_size = [u32::from(stage == ShaderStage::Compute); 3];
+        self.workgroup_size = [u32::from(stage.compute_like()); 3];
         self.early_fragment_tests = false;
         self.extensions.clear();
     }
diff --git a/naga/src/front/glsl/variables.rs b/naga/src/front/glsl/variables.rs
index ef98143b769..98871bd2f81 100644
--- a/naga/src/front/glsl/variables.rs
+++ b/naga/src/front/glsl/variables.rs
@@ -465,6 +465,7 @@ impl Frontend {
                         interpolation,
                         sampling,
                         blend_src,
+                        per_primitive: false,
                     },
                     handle,
                     storage,
diff --git a/naga/src/front/interpolator.rs b/naga/src/front/interpolator.rs
index e23cae0e7c2..126e860426c 100644
--- a/naga/src/front/interpolator.rs
+++ b/naga/src/front/interpolator.rs
@@ -44,6 +44,7 @@ impl crate::Binding {
             interpolation: ref mut interpolation @ None,
             ref mut sampling,
             blend_src: _,
+            per_primitive: _,
         } = *self
         {
             match ty.scalar_kind() {
diff --git a/naga/src/front/spv/function.rs b/naga/src/front/spv/function.rs
index 67cbf05f04f..48b23e7c4c4 100644
--- a/naga/src/front/spv/function.rs
+++ b/naga/src/front/spv/function.rs
@@ -596,6 +596,8 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
             workgroup_size: ep.workgroup_size,
             workgroup_size_overrides: None,
             function,
+            mesh_info: None,
+            task_payload: None,
         });
 
         Ok(())
diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs
index 960437ece58..396318f14dc 100644
--- a/naga/src/front/spv/mod.rs
+++ b/naga/src/front/spv/mod.rs
@@ -263,6 +263,7 @@ impl Decoration {
                 interpolation,
                 sampling,
                 blend_src: None,
+                per_primitive: false,
             }),
             _ => Err(Error::MissingDecoration(spirv::Decoration::Location)),
         }
@@ -4613,6 +4614,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                 | S::Atomic { .. }
                 | S::ImageAtomic { .. }
                 | S::RayQuery { .. }
+                | S::MeshFunction(..)
                 | S::SubgroupBallot { .. }
                 | S::SubgroupCollectiveOperation { .. }
                 | S::SubgroupGather { .. } => {}
@@ -4894,6 +4896,8 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                 spirv::ExecutionModel::Vertex => crate::ShaderStage::Vertex,
                 spirv::ExecutionModel::Fragment => crate::ShaderStage::Fragment,
                 spirv::ExecutionModel::GLCompute => crate::ShaderStage::Compute,
+                spirv::ExecutionModel::TaskEXT => crate::ShaderStage::Task,
+                spirv::ExecutionModel::MeshEXT => crate::ShaderStage::Mesh,
                 _ => return Err(Error::UnsupportedExecutionModel(exec_model as u32)),
             },
             name,
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
index e90d7eab0a8..2066d7cf2c8 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/naga/src/front/wgsl/lower/mod.rs
@@ -1527,6 +1527,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                 workgroup_size,
                 workgroup_size_overrides,
                 function,
+                mesh_info: None,
+                task_payload: None,
             });
             Ok(LoweredGlobalDecl::EntryPoint(
                 ctx.module.entry_points.len() - 1,
@@ -4069,6 +4071,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                     interpolation,
                     sampling,
                     blend_src,
+                    per_primitive: false,
                 };
                 binding.apply_default_interpolation(&ctx.module.types[ty].inner);
                 Some(binding)
diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs
index 257445952b8..2856872db27 100644
--- a/naga/src/ir/mod.rs
+++ b/naga/src/ir/mod.rs
@@ -325,10 +325,21 @@ pub enum ShaderStage {
     Vertex,
     Fragment,
     Compute,
+    // Mesh shader stages
     Task,
     Mesh,
 }
 
+impl ShaderStage {
+    // TODO: make more things respect this
+    pub const fn compute_like(self) -> bool {
+        match self {
+            Self::Vertex | Self::Fragment => false,
+            Self::Compute | Self::Task | Self::Mesh => true,
+        }
+    }
+}
+
 /// Addressing space of variables.
 #[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
 #[cfg_attr(feature = "serialize", derive(Serialize))]
@@ -363,6 +374,8 @@ pub enum AddressSpace {
     ///
     /// [`SHADER_FLOAT16`]: crate::valid::Capabilities::SHADER_FLOAT16
     PushConstant,
+    /// Task shader to mesh shader payload
+    TaskPayload,
 }
 
 /// Built-in inputs and outputs.
@@ -373,7 +386,7 @@ pub enum AddressSpace {
 pub enum BuiltIn {
     Position { invariant: bool },
     ViewIndex,
-    // vertex
+    // vertex (and often mesh)
     BaseInstance,
     BaseVertex,
     ClipDistance,
@@ -386,10 +399,10 @@ pub enum BuiltIn {
     FragDepth,
     PointCoord,
     FrontFacing,
-    PrimitiveIndex,
+    PrimitiveIndex, // Also for mesh output
     SampleIndex,
     SampleMask,
-    // compute
+    // compute (and task/mesh)
     GlobalInvocationId,
     LocalInvocationId,
     LocalInvocationIndex,
@@ -401,6 +414,12 @@ pub enum BuiltIn {
     SubgroupId,
     SubgroupSize,
     SubgroupInvocationId,
+    // mesh
+    MeshTaskSize,
+    CullPrimitive,
+    PointIndex,
+    LineIndices,
+    TriangleIndices,
 }
 
 /// Number of bytes per scalar.
@@ -966,6 +985,14 @@ pub enum Binding {
         /// Optional `blend_src` index used for dual source blending.
         /// See <https://www.w3.org/TR/WGSL/#attribute-blend_src>
         blend_src: Option<u32>,
+        /// Whether the binding is a per-primitive binding for use with mesh shaders.
+        /// This is required to match for mesh and fragment shader stages.
+        /// This is merely an extra attribute on a binding. You still may not have
+        /// a per-vertex and per-primitive input with the same location.
+        ///
+        /// Per primitive values are not interpolated at all and are not dependent on the vertices
+        /// or pixel location. For example, it may be used to store a non-interpolated normal vector.
+        per_primitive: bool,
     },
 }
 
@@ -2141,6 +2168,8 @@ 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.
@@ -2314,6 +2343,10 @@ pub struct EntryPoint {
     pub workgroup_size_overrides: Option<[Option<Handle<Expression>>; 3]>,
     /// The entrance function.
     pub function: Function,
+    /// The information relating to a mesh shader
+    pub mesh_info: Option<MeshStageInfo>,
+    /// The unique global variable used as a task payload from task shader to mesh shader
+    pub task_payload: Option<Handle<GlobalVariable>>,
 }
 
 /// Return types predeclared for the frexp, modf, and atomicCompareExchangeWeak built-in functions.
@@ -2578,3 +2611,48 @@ pub struct Module {
     /// Doc comments.
     pub doc_comments: Option<Box<DocComments>>,
 }
+
+#[derive(Debug, Clone, Copy)]
+#[cfg_attr(feature = "serialize", derive(Serialize))]
+#[cfg_attr(feature = "deserialize", derive(Deserialize))]
+#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
+pub enum MeshOutputTopology {
+    Points,
+    Lines,
+    Triangles,
+}
+
+#[derive(Debug, Clone)]
+#[cfg_attr(feature = "serialize", derive(Serialize))]
+#[cfg_attr(feature = "deserialize", derive(Deserialize))]
+#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
+#[allow(dead_code)]
+pub struct MeshStageInfo {
+    pub topology: MeshOutputTopology,
+    pub max_vertices: u32,
+    pub max_vertices_override: Option<Handle<Expression>>,
+    pub max_primitives: u32,
+    pub max_primitives_override: Option<Handle<Expression>>,
+    pub vertex_output_type: Handle<Type>,
+    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 {
+    SetMeshOutputs {
+        vertex_count: Handle<Expression>,
+        primitive_count: Handle<Expression>,
+    },
+    SetVertex {
+        index: Handle<Expression>,
+        value: Handle<Expression>,
+    },
+    SetPrimitive {
+        index: Handle<Expression>,
+        value: Handle<Expression>,
+    },
+}
diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs
index 413e49c1eed..434c6e3f724 100644
--- a/naga/src/proc/mod.rs
+++ b/naga/src/proc/mod.rs
@@ -177,6 +177,9 @@ impl super::AddressSpace {
             crate::AddressSpace::Storage { access } => access,
             crate::AddressSpace::Handle => Sa::LOAD,
             crate::AddressSpace::PushConstant => Sa::LOAD,
+            // TaskPayload isn't always writable, but this is checked for elsewhere,
+            // when not using multiple payloads and matching the entry payload is checked.
+            crate::AddressSpace::TaskPayload => Sa::LOAD | Sa::STORE,
         }
     }
 }
diff --git a/naga/src/proc/terminator.rs b/naga/src/proc/terminator.rs
index b29ccb054a3..f76d4c06a3b 100644
--- a/naga/src/proc/terminator.rs
+++ b/naga/src/proc/terminator.rs
@@ -36,6 +36,7 @@ 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 95ae40dcdb4..101ea046487 100644
--- a/naga/src/valid/analyzer.rs
+++ b/naga/src/valid/analyzer.rs
@@ -85,6 +85,16 @@ 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 {
+    pub vertex_type: Option<(Handle<crate::Type>, Handle<crate::Expression>)>,
+    pub primitive_type: Option<(Handle<crate::Type>, Handle<crate::Expression>)>,
+}
+
 impl ops::BitOr for FunctionUniformity {
     type Output = Self;
     fn bitor(self, other: Self) -> Self {
@@ -302,6 +312,8 @@ pub struct FunctionInfo {
     /// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
     /// validation.
     diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
+
+    pub mesh_shader_info: FunctionMeshShaderInfo,
 }
 
 impl FunctionInfo {
@@ -372,6 +384,14 @@ impl FunctionInfo {
         info.uniformity.non_uniform_result
     }
 
+    pub fn insert_global_use(
+        &mut self,
+        global_use: GlobalUse,
+        global: Handle<crate::GlobalVariable>,
+    ) {
+        self.global_uses[global.index()] |= global_use;
+    }
+
     /// Record a use of `expr` for its value.
     ///
     /// This is used for almost all expression references. Anything
@@ -482,6 +502,8 @@ impl FunctionInfo {
             *mine |= *other;
         }
 
+        self.try_update_mesh_info(&callee.mesh_shader_info)?;
+
         Ok(FunctionUniformity {
             result: callee.uniformity.clone(),
             exit: if callee.may_kill {
@@ -635,7 +657,8 @@ impl FunctionInfo {
                     // local data is non-uniform
                     As::Function | As::Private => false,
                     // workgroup memory is exclusively accessed by the group
-                    As::WorkGroup => true,
+                    // task payload memory is very similar to workgroup memory
+                    As::WorkGroup | As::TaskPayload => true,
                     // uniform data
                     As::Uniform | As::PushConstant => true,
                     // storage data is only uniform when read-only
@@ -1113,6 +1136,34 @@ impl FunctionInfo {
                     }
                     FunctionUniformity::new()
                 }
+                S::MeshFunction(func) => 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.clone().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,
@@ -1158,6 +1209,53 @@ impl FunctionInfo {
         }
         Ok(combined_uniformity)
     }
+
+    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(())
+    }
+
+    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(())
+    }
+
+    fn try_update_mesh_info(
+        &mut self,
+        other: &FunctionMeshShaderInfo,
+    ) -> Result<(), WithSpan<FunctionError>> {
+        if let &Some(ref other_vertex) = &other.vertex_type {
+            self.try_update_mesh_vertex_type(other_vertex.0, other_vertex.1)?;
+        }
+        if let &Some(ref other_primitive) = &other.vertex_type {
+            self.try_update_mesh_primitive_type(other_primitive.0, other_primitive.1)?;
+        }
+        Ok(())
+    }
 }
 
 impl ModuleInfo {
@@ -1193,6 +1291,7 @@ 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);
@@ -1326,6 +1425,7 @@ 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 dc19e191764..0ae2ffdb54f 100644
--- a/naga/src/valid/function.rs
+++ b/naga/src/valid/function.rs
@@ -217,6 +217,14 @@ pub enum FunctionError {
     EmitResult(Handle<crate::Expression>),
     #[error("Expression not visited by the appropriate statement")]
     UnvisitedExpression(Handle<crate::Expression>),
+    #[error("Expression {0:?} should be u32, but isn't")]
+    InvalidMeshFunctionCall(Handle<crate::Expression>),
+    #[error("Mesh output types differ from {0:?} to {1:?}")]
+    ConflictingMeshOutputTypes(Handle<crate::Expression>, Handle<crate::Expression>),
+    #[error("Task payload variables differ from {0:?} to {1:?}")]
+    ConflictingTaskPayloadVariables(Handle<crate::Expression>, Handle<crate::Expression>),
+    #[error("Mesh shader output at {0:?} is not a user-defined struct")]
+    InvalidMeshShaderOutputType(Handle<crate::Expression>),
 }
 
 bitflags::bitflags! {
@@ -1539,6 +1547,40 @@ 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)?;
+                            // TODO: ensure it is correct for the value
+                        }
+                    }
+                }
                 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 e8a69013434..a0153e9398c 100644
--- a/naga/src/valid/handles.rs
+++ b/naga/src/valid/handles.rs
@@ -801,6 +801,22 @@ 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 7c8cc903139..1fed0fda529 100644
--- a/naga/src/valid/interface.rs
+++ b/naga/src/valid/interface.rs
@@ -43,6 +43,8 @@ pub enum GlobalVariableError {
     StorageAddressSpaceWriteOnlyNotSupported,
     #[error("Type is not valid for use as a push constant")]
     InvalidPushConstantType(#[source] PushConstantError),
+    #[error("Task payload must not be zero-sized")]
+    ZeroSizedTaskPayload,
 }
 
 #[derive(Clone, Debug, thiserror::Error)]
@@ -92,6 +94,10 @@ pub enum VaryingError {
     },
     #[error("Workgroup size is multi dimensional, `@builtin(subgroup_id)` and `@builtin(subgroup_invocation_id)` are not supported.")]
     InvalidMultiDimensionalSubgroupBuiltIn,
+    #[error("The `@per_primitive` attribute can only be used in fragment shader inputs or mesh shader primitive outputs")]
+    InvalidPerPrimitive,
+    #[error("Non-builtin members of a mesh primitive output struct must be decorated with `@per_primitive`")]
+    MissingPerPrimitive,
 }
 
 #[derive(Clone, Debug, thiserror::Error)]
@@ -123,6 +129,28 @@ pub enum EntryPointError {
     InvalidIntegerInterpolation { location: u32 },
     #[error(transparent)]
     Function(#[from] FunctionError),
+    #[error("Non mesh shader entry point cannot have mesh shader attributes")]
+    UnexpectedMeshShaderAttributes,
+    #[error("Non mesh/task shader entry point cannot have task payload attribute")]
+    UnexpectedTaskPayload,
+    #[error("Task payload must be declared with `var<task_payload>`")]
+    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,
 }
 
 fn storage_usage(access: crate::StorageAccess) -> GlobalUse {
@@ -139,6 +167,13 @@ fn storage_usage(access: crate::StorageAccess) -> GlobalUse {
     storage_usage
 }
 
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+enum MeshOutputType {
+    None,
+    VertexOutput,
+    PrimitiveOutput,
+}
+
 struct VaryingContext<'a> {
     stage: crate::ShaderStage,
     output: bool,
@@ -149,6 +184,7 @@ struct VaryingContext<'a> {
     built_ins: &'a mut crate::FastHashSet<crate::BuiltIn>,
     capabilities: Capabilities,
     flags: super::ValidationFlags,
+    mesh_output_type: MeshOutputType,
 }
 
 impl VaryingContext<'_> {
@@ -236,10 +272,9 @@ impl VaryingContext<'_> {
                     ),
                     Bi::Position { .. } => (
                         match self.stage {
-                            St::Vertex => self.output,
+                            St::Vertex | St::Mesh => self.output,
                             St::Fragment => !self.output,
-                            St::Compute => false,
-                            St::Task | St::Mesh => unreachable!(),
+                            St::Compute | St::Task => false,
                         },
                         *ty_inner
                             == Ti::Vector {
@@ -276,7 +311,7 @@ impl VaryingContext<'_> {
                         *ty_inner == Ti::Scalar(crate::Scalar::U32),
                     ),
                     Bi::LocalInvocationIndex => (
-                        self.stage == St::Compute && !self.output,
+                        self.stage.compute_like() && !self.output,
                         *ty_inner == Ti::Scalar(crate::Scalar::U32),
                     ),
                     Bi::GlobalInvocationId
@@ -284,7 +319,7 @@ impl VaryingContext<'_> {
                     | Bi::WorkGroupId
                     | Bi::WorkGroupSize
                     | Bi::NumWorkGroups => (
-                        self.stage == St::Compute && !self.output,
+                        self.stage.compute_like() && !self.output,
                         *ty_inner
                             == Ti::Vector {
                                 size: Vs::Tri,
@@ -292,17 +327,48 @@ impl VaryingContext<'_> {
                             },
                     ),
                     Bi::NumSubgroups | Bi::SubgroupId => (
-                        self.stage == St::Compute && !self.output,
+                        self.stage.compute_like() && !self.output,
                         *ty_inner == Ti::Scalar(crate::Scalar::U32),
                     ),
                     Bi::SubgroupSize | Bi::SubgroupInvocationId => (
                         match self.stage {
-                            St::Compute | St::Fragment => !self.output,
+                            St::Compute | St::Fragment | St::Task | St::Mesh => !self.output,
                             St::Vertex => false,
-                            St::Task | St::Mesh => unreachable!(),
                         },
                         *ty_inner == Ti::Scalar(crate::Scalar::U32),
                     ),
+                    Bi::CullPrimitive => (
+                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
+                        *ty_inner == Ti::Scalar(crate::Scalar::BOOL),
+                    ),
+                    Bi::PointIndex => (
+                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
+                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
+                    ),
+                    Bi::LineIndices => (
+                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
+                        *ty_inner
+                            == Ti::Vector {
+                                size: Vs::Bi,
+                                scalar: crate::Scalar::U32,
+                            },
+                    ),
+                    Bi::TriangleIndices => (
+                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
+                        *ty_inner
+                            == Ti::Vector {
+                                size: Vs::Tri,
+                                scalar: crate::Scalar::U32,
+                            },
+                    ),
+                    Bi::MeshTaskSize => (
+                        self.stage == St::Task && self.output,
+                        *ty_inner
+                            == Ti::Vector {
+                                size: Vs::Tri,
+                                scalar: crate::Scalar::U32,
+                            },
+                    ),
                 };
 
                 if !visible {
@@ -318,6 +384,7 @@ impl VaryingContext<'_> {
                 interpolation,
                 sampling,
                 blend_src,
+                per_primitive,
             } => {
                 // Only IO-shareable types may be stored in locations.
                 if !self.type_info[ty.index()]
@@ -326,6 +393,14 @@ impl VaryingContext<'_> {
                 {
                     return Err(VaryingError::NotIOShareableType(ty));
                 }
+                if !per_primitive && self.mesh_output_type == MeshOutputType::PrimitiveOutput {
+                    return Err(VaryingError::MissingPerPrimitive);
+                } else if per_primitive
+                    && ((self.stage != crate::ShaderStage::Fragment || self.output)
+                        && self.mesh_output_type != MeshOutputType::PrimitiveOutput)
+                {
+                    return Err(VaryingError::InvalidPerPrimitive);
+                }
 
                 if let Some(blend_src) = blend_src {
                     // `blend_src` is only valid if dual source blending was explicitly enabled,
@@ -390,11 +465,12 @@ impl VaryingContext<'_> {
                     }
                 }
 
+                // TODO: update this to reflect the fact that per-primitive outputs aren't interpolated for fragment and mesh stages
                 let needs_interpolation = match self.stage {
                     crate::ShaderStage::Vertex => self.output,
                     crate::ShaderStage::Fragment => !self.output,
-                    crate::ShaderStage::Compute => false,
-                    crate::ShaderStage::Task | crate::ShaderStage::Mesh => unreachable!(),
+                    crate::ShaderStage::Compute | crate::ShaderStage::Task => false,
+                    crate::ShaderStage::Mesh => self.output,
                 };
 
                 // It doesn't make sense to specify a sampling when `interpolation` is `Flat`, but
@@ -595,7 +671,9 @@ impl super::Validator {
                 TypeFlags::CONSTRUCTIBLE | TypeFlags::CREATION_RESOLVED,
                 false,
             ),
-            crate::AddressSpace::WorkGroup => (TypeFlags::DATA | TypeFlags::SIZED, false),
+            crate::AddressSpace::WorkGroup | crate::AddressSpace::TaskPayload => {
+                (TypeFlags::DATA | TypeFlags::SIZED, false)
+            }
             crate::AddressSpace::PushConstant => {
                 if !self.capabilities.contains(Capabilities::PUSH_CONSTANT) {
                     return Err(GlobalVariableError::UnsupportedCapability(
@@ -628,6 +706,14 @@ impl super::Validator {
             }
         }
 
+        if var.space == crate::AddressSpace::TaskPayload {
+            let ty = &gctx.types[var.ty].inner;
+            // HLSL doesn't allow zero sized payloads.
+            if ty.try_size(gctx) == Some(0) {
+                return Err(GlobalVariableError::ZeroSizedTaskPayload);
+            }
+        }
+
         if let Some(init) = var.init {
             match var.space {
                 crate::AddressSpace::Private | crate::AddressSpace::Function => {}
@@ -671,7 +757,7 @@ impl super::Validator {
             }
         }
 
-        if ep.stage == crate::ShaderStage::Compute {
+        if ep.stage.compute_like() {
             if ep
                 .workgroup_size
                 .iter()
@@ -683,10 +769,30 @@ impl super::Validator {
             return Err(EntryPointError::UnexpectedWorkgroupSize.with_span());
         }
 
+        if ep.stage != crate::ShaderStage::Mesh && ep.mesh_info.is_some() {
+            return Err(EntryPointError::UnexpectedMeshShaderAttributes.with_span());
+        }
+
         let mut info = self
             .validate_function(&ep.function, module, mod_info, true)
             .map_err(WithSpan::into_other)?;
 
+        if let Some(handle) = ep.task_payload {
+            if ep.stage != crate::ShaderStage::Task && ep.stage != crate::ShaderStage::Mesh {
+                return Err(EntryPointError::UnexpectedTaskPayload.with_span());
+            }
+            if module.global_variables[handle].space != crate::AddressSpace::TaskPayload {
+                return Err(EntryPointError::TaskPayloadWrongAddressSpace.with_span());
+            }
+            // Make sure that this is always present in the outputted shader
+            let uses = if ep.stage == crate::ShaderStage::Mesh {
+                GlobalUse::READ
+            } else {
+                GlobalUse::READ | GlobalUse::WRITE
+            };
+            info.insert_global_use(uses, handle);
+        }
+
         {
             use super::ShaderStages;
 
@@ -694,7 +800,8 @@ impl super::Validator {
                 crate::ShaderStage::Vertex => ShaderStages::VERTEX,
                 crate::ShaderStage::Fragment => ShaderStages::FRAGMENT,
                 crate::ShaderStage::Compute => ShaderStages::COMPUTE,
-                crate::ShaderStage::Task | crate::ShaderStage::Mesh => unreachable!(),
+                crate::ShaderStage::Mesh => ShaderStages::MESH,
+                crate::ShaderStage::Task => ShaderStages::TASK,
             };
 
             if !info.available_stages.contains(stage_bit) {
@@ -716,6 +823,7 @@ impl super::Validator {
                 built_ins: &mut argument_built_ins,
                 capabilities: self.capabilities,
                 flags: self.flags,
+                mesh_output_type: MeshOutputType::None,
             };
             ctx.validate(ep, fa.ty, fa.binding.as_ref())
                 .map_err_inner(|e| EntryPointError::Argument(index as u32, e).with_span())?;
@@ -734,6 +842,7 @@ impl super::Validator {
                 built_ins: &mut result_built_ins,
                 capabilities: self.capabilities,
                 flags: self.flags,
+                mesh_output_type: MeshOutputType::None,
             };
             ctx.validate(ep, fr.ty, fr.binding.as_ref())
                 .map_err_inner(|e| EntryPointError::Result(e).with_span())?;
@@ -742,11 +851,26 @@ impl super::Validator {
             {
                 return Err(EntryPointError::MissingVertexOutputPosition.with_span());
             }
+            if ep.stage == crate::ShaderStage::Mesh
+                && (!result_built_ins.is_empty() || !self.location_mask.is_empty())
+            {
+                return Err(EntryPointError::UnexpectedMeshShaderEntryResult.with_span());
+            }
+            // Cannot have any other built-ins or @location outputs as those are per-vertex or per-primitive
+            if ep.stage == crate::ShaderStage::Task
+                && (!result_built_ins.contains(&crate::BuiltIn::MeshTaskSize)
+                    || result_built_ins.len() != 1
+                    || !self.location_mask.is_empty())
+            {
+                return Err(EntryPointError::WrongTaskShaderEntryResult.with_span());
+            }
             if !self.blend_src_mask.is_empty() {
                 info.dual_source_blending = true;
             }
         } else if ep.stage == crate::ShaderStage::Vertex {
             return Err(EntryPointError::MissingVertexOutputPosition.with_span());
+        } else if ep.stage == crate::ShaderStage::Task {
+            return Err(EntryPointError::WrongTaskShaderEntryResult.with_span());
         }
 
         {
@@ -764,6 +888,15 @@ impl super::Validator {
             }
         }
 
+        if let Some(task_payload) = ep.task_payload {
+            if module.global_variables[task_payload].space != crate::AddressSpace::TaskPayload {
+                return Err(EntryPointError::TaskPayloadWrongAddressSpace
+                    .with_span_handle(task_payload, &module.global_variables));
+            }
+        } else if ep.stage == crate::ShaderStage::Task {
+            return Err(EntryPointError::ExpectedTaskPayload.with_span());
+        }
+
         self.ep_resource_bindings.clear();
         for (var_handle, var) in module.global_variables.iter() {
             let usage = info[var_handle];
@@ -771,6 +904,13 @@ impl super::Validator {
                 continue;
             }
 
+            if var.space == crate::AddressSpace::TaskPayload {
+                if ep.task_payload != Some(var_handle) {
+                    return Err(EntryPointError::WrongTaskPayloadUsed
+                        .with_span_handle(var_handle, &module.global_variables));
+                }
+            }
+
             let allowed_usage = match var.space {
                 crate::AddressSpace::Function => unreachable!(),
                 crate::AddressSpace::Uniform => GlobalUse::READ | GlobalUse::QUERY,
@@ -792,6 +932,15 @@ impl super::Validator {
                 crate::AddressSpace::Private | crate::AddressSpace::WorkGroup => {
                     GlobalUse::READ | GlobalUse::WRITE | GlobalUse::QUERY
                 }
+                crate::AddressSpace::TaskPayload => {
+                    GlobalUse::READ
+                        | GlobalUse::QUERY
+                        | if ep.stage == crate::ShaderStage::Task {
+                            GlobalUse::WRITE
+                        } else {
+                            GlobalUse::empty()
+                        }
+                }
                 crate::AddressSpace::PushConstant => GlobalUse::READ,
             };
             if !allowed_usage.contains(usage) {
@@ -811,6 +960,77 @@ impl super::Validator {
             }
         }
 
+        if let &Some(ref mesh_info) = &ep.mesh_info {
+            // Technically it is allowed to not output anything
+            // TODO: check that only the allowed builtins are used here
+            if let Some(used_vertex_type) = info.mesh_shader_info.vertex_type {
+                if used_vertex_type.0 != mesh_info.vertex_output_type {
+                    return Err(EntryPointError::WrongMeshOutputType
+                        .with_span_handle(mesh_info.vertex_output_type, &module.types));
+                }
+            }
+            if let Some(used_primitive_type) = info.mesh_shader_info.primitive_type {
+                if used_primitive_type.0 != mesh_info.primitive_output_type {
+                    return Err(EntryPointError::WrongMeshOutputType
+                        .with_span_handle(mesh_info.primitive_output_type, &module.types));
+                }
+            }
+
+            for (ty, mesh_output_type) in [
+                (mesh_info.vertex_output_type, MeshOutputType::VertexOutput),
+                (
+                    mesh_info.primitive_output_type,
+                    MeshOutputType::PrimitiveOutput,
+                ),
+            ] {
+                if !matches!(module.types[ty].inner, crate::TypeInner::Struct { .. }) {
+                    return Err(
+                        EntryPointError::InvalidMeshOutputType.with_span_handle(ty, &module.types)
+                    );
+                }
+                let mut result_built_ins = crate::FastHashSet::default();
+                let mut ctx = VaryingContext {
+                    stage: ep.stage,
+                    output: true,
+                    types: &module.types,
+                    type_info: &self.types,
+                    location_mask: &mut self.location_mask,
+                    blend_src_mask: &mut self.blend_src_mask,
+                    built_ins: &mut result_built_ins,
+                    capabilities: self.capabilities,
+                    flags: self.flags,
+                    mesh_output_type,
+                };
+                ctx.validate(ep, ty, None)
+                    .map_err_inner(|e| EntryPointError::Result(e).with_span())?;
+                if mesh_output_type == MeshOutputType::PrimitiveOutput {
+                    let mut num_indices_builtins = 0;
+                    if result_built_ins.contains(&crate::BuiltIn::PointIndex) {
+                        num_indices_builtins += 1;
+                    }
+                    if result_built_ins.contains(&crate::BuiltIn::LineIndices) {
+                        num_indices_builtins += 1;
+                    }
+                    if result_built_ins.contains(&crate::BuiltIn::TriangleIndices) {
+                        num_indices_builtins += 1;
+                    }
+                    if num_indices_builtins != 1 {
+                        return Err(EntryPointError::InvalidMeshPrimitiveOutputType
+                            .with_span_handle(ty, &module.types));
+                    }
+                } else if mesh_output_type == MeshOutputType::VertexOutput
+                    && !result_built_ins.contains(&crate::BuiltIn::Position { invariant: false })
+                {
+                    return Err(EntryPointError::MissingVertexOutputPosition
+                        .with_span_handle(ty, &module.types));
+                }
+            }
+        } else 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 fe45d3bfb07..babea985244 100644
--- a/naga/src/valid/mod.rs
+++ b/naga/src/valid/mod.rs
@@ -240,6 +240,8 @@ bitflags::bitflags! {
         const VERTEX = 0x1;
         const FRAGMENT = 0x2;
         const COMPUTE = 0x4;
+        const MESH = 0x8;
+        const TASK = 0x10;
     }
 }
 
diff --git a/naga/src/valid/type.rs b/naga/src/valid/type.rs
index e8b83ff08f3..aa0633e1852 100644
--- a/naga/src/valid/type.rs
+++ b/naga/src/valid/type.rs
@@ -220,9 +220,12 @@ const fn ptr_space_argument_flag(space: crate::AddressSpace) -> TypeFlags {
     use crate::AddressSpace as As;
     match space {
         As::Function | As::Private => TypeFlags::ARGUMENT,
-        As::Uniform | As::Storage { .. } | As::Handle | As::PushConstant | As::WorkGroup => {
-            TypeFlags::empty()
-        }
+        As::Uniform
+        | As::Storage { .. }
+        | As::Handle
+        | As::PushConstant
+        | As::WorkGroup
+        | As::TaskPayload => TypeFlags::empty(),
     }
 }
 
diff --git a/naga/tests/out/analysis/spv-shadow.info.ron b/naga/tests/out/analysis/spv-shadow.info.ron
index 6ddda61f5c6..b08a28438ed 100644
--- a/naga/tests/out/analysis/spv-shadow.info.ron
+++ b/naga/tests/out/analysis/spv-shadow.info.ron
@@ -18,7 +18,7 @@
     functions: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(1),
                 requirements: (""),
@@ -413,10 +413,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(1),
                 requirements: (""),
@@ -1591,12 +1595,16 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
+            mesh_shader_info: (
+                vertex_type: None,
+                primitive_type: None,
+            ),
         ),
     ],
     entry_points: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(1),
                 requirements: (""),
@@ -1685,6 +1693,10 @@
             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 319f62bdf13..d297b09a404 100644
--- a/naga/tests/out/analysis/wgsl-access.info.ron
+++ b/naga/tests/out/analysis/wgsl-access.info.ron
@@ -42,7 +42,7 @@
     functions: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -1197,10 +1197,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -2523,10 +2527,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(0),
                 requirements: (""),
@@ -2563,10 +2571,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(0),
                 requirements: (""),
@@ -2612,10 +2624,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -2655,10 +2671,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -2749,10 +2769,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -2870,10 +2894,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(0),
                 requirements: (""),
@@ -2922,10 +2950,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -2977,10 +3009,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(0),
                 requirements: (""),
@@ -3029,10 +3065,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -3084,10 +3124,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(0),
                 requirements: (""),
@@ -3148,10 +3192,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(2),
                 requirements: (""),
@@ -3221,10 +3269,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(2),
                 requirements: (""),
@@ -3297,10 +3349,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -3397,10 +3453,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(1),
                 requirements: (""),
@@ -3593,12 +3653,16 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
+            mesh_shader_info: (
+                vertex_type: None,
+                primitive_type: None,
+            ),
         ),
     ],
     entry_points: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(0),
                 requirements: (""),
@@ -4290,10 +4354,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -4742,10 +4810,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(0),
                 requirements: (""),
@@ -4812,6 +4884,10 @@
             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 7ec5799d758..2796f544510 100644
--- a/naga/tests/out/analysis/wgsl-collatz.info.ron
+++ b/naga/tests/out/analysis/wgsl-collatz.info.ron
@@ -8,7 +8,7 @@
     functions: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(3),
                 requirements: (""),
@@ -275,12 +275,16 @@
             sampling: [],
             dual_source_blending: false,
             diagnostic_filter_leaf: None,
+            mesh_shader_info: (
+                vertex_type: None,
+                primitive_type: None,
+            ),
         ),
     ],
     entry_points: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: Some(3),
                 requirements: (""),
@@ -430,6 +434,10 @@
             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-overrides.info.ron b/naga/tests/out/analysis/wgsl-overrides.info.ron
index 0e0ae318042..a76c9c89c9b 100644
--- a/naga/tests/out/analysis/wgsl-overrides.info.ron
+++ b/naga/tests/out/analysis/wgsl-overrides.info.ron
@@ -8,7 +8,7 @@
     entry_points: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -201,6 +201,10 @@
             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 fbbf7206c33..35b5a7e320c 100644
--- a/naga/tests/out/analysis/wgsl-storage-textures.info.ron
+++ b/naga/tests/out/analysis/wgsl-storage-textures.info.ron
@@ -11,7 +11,7 @@
     entry_points: [
         (
             flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -184,10 +184,14 @@
             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"),
-            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"),
             uniformity: (
                 non_uniform_result: None,
                 requirements: (""),
@@ -396,6 +400,10 @@
             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/spv-fetch_depth.compact.ron b/naga/tests/out/ir/spv-fetch_depth.compact.ron
index 1fbee2deb35..98f4426c3eb 100644
--- a/naga/tests/out/ir/spv-fetch_depth.compact.ron
+++ b/naga/tests/out/ir/spv-fetch_depth.compact.ron
@@ -196,6 +196,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/spv-fetch_depth.ron b/naga/tests/out/ir/spv-fetch_depth.ron
index 186f78354ad..104de852c17 100644
--- a/naga/tests/out/ir/spv-fetch_depth.ron
+++ b/naga/tests/out/ir/spv-fetch_depth.ron
@@ -266,6 +266,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/spv-shadow.compact.ron b/naga/tests/out/ir/spv-shadow.compact.ron
index b49cd9b55be..bed86a5334d 100644
--- a/naga/tests/out/ir/spv-shadow.compact.ron
+++ b/naga/tests/out/ir/spv-shadow.compact.ron
@@ -974,6 +974,7 @@
                             interpolation: Some(Perspective),
                             sampling: Some(Center),
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                     (
@@ -984,6 +985,7 @@
                             interpolation: Some(Perspective),
                             sampling: Some(Center),
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                 ],
@@ -994,6 +996,7 @@
                         interpolation: None,
                         sampling: None,
                         blend_src: None,
+                        per_primitive: false,
                     )),
                 )),
                 local_variables: [],
@@ -1032,6 +1035,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/spv-shadow.ron b/naga/tests/out/ir/spv-shadow.ron
index e1f0f60b6bb..bdda1d18566 100644
--- a/naga/tests/out/ir/spv-shadow.ron
+++ b/naga/tests/out/ir/spv-shadow.ron
@@ -1252,6 +1252,7 @@
                             interpolation: Some(Perspective),
                             sampling: Some(Center),
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                     (
@@ -1262,6 +1263,7 @@
                             interpolation: Some(Perspective),
                             sampling: Some(Center),
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                 ],
@@ -1272,6 +1274,7 @@
                         interpolation: None,
                         sampling: None,
                         blend_src: None,
+                        per_primitive: false,
                     )),
                 )),
                 local_variables: [],
@@ -1310,6 +1313,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/spv-spec-constants.compact.ron b/naga/tests/out/ir/spv-spec-constants.compact.ron
index 3fa6ffef4ff..67eb29c2475 100644
--- a/naga/tests/out/ir/spv-spec-constants.compact.ron
+++ b/naga/tests/out/ir/spv-spec-constants.compact.ron
@@ -151,6 +151,7 @@
                             interpolation: Some(Perspective),
                             sampling: Some(Center),
                             blend_src: None,
+                            per_primitive: false,
                         )),
                         offset: 0,
                     ),
@@ -510,6 +511,7 @@
                             interpolation: None,
                             sampling: None,
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                     (
@@ -520,6 +522,7 @@
                             interpolation: None,
                             sampling: None,
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                     (
@@ -530,6 +533,7 @@
                             interpolation: None,
                             sampling: None,
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                 ],
@@ -613,6 +617,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/spv-spec-constants.ron b/naga/tests/out/ir/spv-spec-constants.ron
index 94c90aa78f9..51686aa20eb 100644
--- a/naga/tests/out/ir/spv-spec-constants.ron
+++ b/naga/tests/out/ir/spv-spec-constants.ron
@@ -242,6 +242,7 @@
                             interpolation: Some(Perspective),
                             sampling: Some(Center),
                             blend_src: None,
+                            per_primitive: false,
                         )),
                         offset: 0,
                     ),
@@ -616,6 +617,7 @@
                             interpolation: None,
                             sampling: None,
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                     (
@@ -626,6 +628,7 @@
                             interpolation: None,
                             sampling: None,
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                     (
@@ -636,6 +639,7 @@
                             interpolation: None,
                             sampling: None,
                             blend_src: None,
+                            per_primitive: false,
                         )),
                     ),
                 ],
@@ -719,6 +723,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-access.compact.ron b/naga/tests/out/ir/wgsl-access.compact.ron
index 30e88984f3c..c3df0c8c500 100644
--- a/naga/tests/out/ir/wgsl-access.compact.ron
+++ b/naga/tests/out/ir/wgsl-access.compact.ron
@@ -2655,6 +2655,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "foo_frag",
@@ -2672,6 +2674,7 @@
                         interpolation: Some(Perspective),
                         sampling: Some(Center),
                         blend_src: None,
+                        per_primitive: false,
                     )),
                 )),
                 local_variables: [],
@@ -2848,6 +2851,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "foo_compute",
@@ -2907,6 +2912,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-access.ron b/naga/tests/out/ir/wgsl-access.ron
index 30e88984f3c..c3df0c8c500 100644
--- a/naga/tests/out/ir/wgsl-access.ron
+++ b/naga/tests/out/ir/wgsl-access.ron
@@ -2655,6 +2655,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "foo_frag",
@@ -2672,6 +2674,7 @@
                         interpolation: Some(Perspective),
                         sampling: Some(Center),
                         blend_src: None,
+                        per_primitive: false,
                     )),
                 )),
                 local_variables: [],
@@ -2848,6 +2851,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "foo_compute",
@@ -2907,6 +2912,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-collatz.compact.ron b/naga/tests/out/ir/wgsl-collatz.compact.ron
index f72c652d032..fc4daaa1296 100644
--- a/naga/tests/out/ir/wgsl-collatz.compact.ron
+++ b/naga/tests/out/ir/wgsl-collatz.compact.ron
@@ -334,6 +334,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-collatz.ron b/naga/tests/out/ir/wgsl-collatz.ron
index f72c652d032..fc4daaa1296 100644
--- a/naga/tests/out/ir/wgsl-collatz.ron
+++ b/naga/tests/out/ir/wgsl-collatz.ron
@@ -334,6 +334,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-const_assert.compact.ron b/naga/tests/out/ir/wgsl-const_assert.compact.ron
index 2816364f88b..648f4ff9bc9 100644
--- a/naga/tests/out/ir/wgsl-const_assert.compact.ron
+++ b/naga/tests/out/ir/wgsl-const_assert.compact.ron
@@ -34,6 +34,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-const_assert.ron b/naga/tests/out/ir/wgsl-const_assert.ron
index 2816364f88b..648f4ff9bc9 100644
--- a/naga/tests/out/ir/wgsl-const_assert.ron
+++ b/naga/tests/out/ir/wgsl-const_assert.ron
@@ -34,6 +34,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron b/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron
index c5746696d52..9a2bf193f30 100644
--- a/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron
+++ b/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron
@@ -73,6 +73,8 @@
                 ],
                 diagnostic_filter_leaf: Some(0),
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [
diff --git a/naga/tests/out/ir/wgsl-diagnostic-filter.ron b/naga/tests/out/ir/wgsl-diagnostic-filter.ron
index c5746696d52..9a2bf193f30 100644
--- a/naga/tests/out/ir/wgsl-diagnostic-filter.ron
+++ b/naga/tests/out/ir/wgsl-diagnostic-filter.ron
@@ -73,6 +73,8 @@
                 ],
                 diagnostic_filter_leaf: Some(0),
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [
diff --git a/naga/tests/out/ir/wgsl-index-by-value.compact.ron b/naga/tests/out/ir/wgsl-index-by-value.compact.ron
index a4f84a7a6b2..addd0e5871c 100644
--- a/naga/tests/out/ir/wgsl-index-by-value.compact.ron
+++ b/naga/tests/out/ir/wgsl-index-by-value.compact.ron
@@ -465,6 +465,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-index-by-value.ron b/naga/tests/out/ir/wgsl-index-by-value.ron
index a4f84a7a6b2..addd0e5871c 100644
--- a/naga/tests/out/ir/wgsl-index-by-value.ron
+++ b/naga/tests/out/ir/wgsl-index-by-value.ron
@@ -465,6 +465,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-local-const.compact.ron b/naga/tests/out/ir/wgsl-local-const.compact.ron
index 512972657ed..0e4e2e4d40e 100644
--- a/naga/tests/out/ir/wgsl-local-const.compact.ron
+++ b/naga/tests/out/ir/wgsl-local-const.compact.ron
@@ -100,6 +100,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-local-const.ron b/naga/tests/out/ir/wgsl-local-const.ron
index 512972657ed..0e4e2e4d40e 100644
--- a/naga/tests/out/ir/wgsl-local-const.ron
+++ b/naga/tests/out/ir/wgsl-local-const.ron
@@ -100,6 +100,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-must-use.compact.ron b/naga/tests/out/ir/wgsl-must-use.compact.ron
index a701a6805da..16e925f2fb8 100644
--- a/naga/tests/out/ir/wgsl-must-use.compact.ron
+++ b/naga/tests/out/ir/wgsl-must-use.compact.ron
@@ -201,6 +201,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-must-use.ron b/naga/tests/out/ir/wgsl-must-use.ron
index a701a6805da..16e925f2fb8 100644
--- a/naga/tests/out/ir/wgsl-must-use.ron
+++ b/naga/tests/out/ir/wgsl-must-use.ron
@@ -201,6 +201,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron
index 640ee25ca49..28a824bb035 100644
--- a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron
+++ b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron
@@ -128,6 +128,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron
index 640ee25ca49..28a824bb035 100644
--- a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron
+++ b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron
@@ -128,6 +128,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron b/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron
index f65e8f186db..152a45008c5 100644
--- a/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron
+++ b/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron
@@ -263,6 +263,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-overrides-ray-query.ron b/naga/tests/out/ir/wgsl-overrides-ray-query.ron
index f65e8f186db..152a45008c5 100644
--- a/naga/tests/out/ir/wgsl-overrides-ray-query.ron
+++ b/naga/tests/out/ir/wgsl-overrides-ray-query.ron
@@ -263,6 +263,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-overrides.compact.ron b/naga/tests/out/ir/wgsl-overrides.compact.ron
index 81221ff7941..fe136e71e4d 100644
--- a/naga/tests/out/ir/wgsl-overrides.compact.ron
+++ b/naga/tests/out/ir/wgsl-overrides.compact.ron
@@ -221,6 +221,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-overrides.ron b/naga/tests/out/ir/wgsl-overrides.ron
index 81221ff7941..fe136e71e4d 100644
--- a/naga/tests/out/ir/wgsl-overrides.ron
+++ b/naga/tests/out/ir/wgsl-overrides.ron
@@ -221,6 +221,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-storage-textures.compact.ron b/naga/tests/out/ir/wgsl-storage-textures.compact.ron
index ec63fecac27..68c867a19e2 100644
--- a/naga/tests/out/ir/wgsl-storage-textures.compact.ron
+++ b/naga/tests/out/ir/wgsl-storage-textures.compact.ron
@@ -218,6 +218,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "csStore",
@@ -315,6 +317,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-storage-textures.ron b/naga/tests/out/ir/wgsl-storage-textures.ron
index ec63fecac27..68c867a19e2 100644
--- a/naga/tests/out/ir/wgsl-storage-textures.ron
+++ b/naga/tests/out/ir/wgsl-storage-textures.ron
@@ -218,6 +218,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "csStore",
@@ -315,6 +317,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron b/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron
index a8208c09b86..db619dff836 100644
--- a/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron
+++ b/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron
@@ -190,6 +190,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron b/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron
index a8208c09b86..db619dff836 100644
--- a/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron
+++ b/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron
@@ -190,6 +190,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-texture-external.compact.ron b/naga/tests/out/ir/wgsl-texture-external.compact.ron
index 2b9e1c8d5e4..689fe215e36 100644
--- a/naga/tests/out/ir/wgsl-texture-external.compact.ron
+++ b/naga/tests/out/ir/wgsl-texture-external.compact.ron
@@ -394,6 +394,7 @@
                         interpolation: Some(Perspective),
                         sampling: Some(Center),
                         blend_src: None,
+                        per_primitive: false,
                     )),
                 )),
                 local_variables: [],
@@ -416,6 +417,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "vertex_main",
@@ -452,6 +455,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "compute_main",
@@ -483,6 +488,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-texture-external.ron b/naga/tests/out/ir/wgsl-texture-external.ron
index 2b9e1c8d5e4..689fe215e36 100644
--- a/naga/tests/out/ir/wgsl-texture-external.ron
+++ b/naga/tests/out/ir/wgsl-texture-external.ron
@@ -394,6 +394,7 @@
                         interpolation: Some(Perspective),
                         sampling: Some(Center),
                         blend_src: None,
+                        per_primitive: false,
                     )),
                 )),
                 local_variables: [],
@@ -416,6 +417,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "vertex_main",
@@ -452,6 +455,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
         (
             name: "compute_main",
@@ -483,6 +488,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-types_with_comments.compact.ron b/naga/tests/out/ir/wgsl-types_with_comments.compact.ron
index 7186209f00e..7c0d856946f 100644
--- a/naga/tests/out/ir/wgsl-types_with_comments.compact.ron
+++ b/naga/tests/out/ir/wgsl-types_with_comments.compact.ron
@@ -116,6 +116,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/naga/tests/out/ir/wgsl-types_with_comments.ron b/naga/tests/out/ir/wgsl-types_with_comments.ron
index 480b0d2337f..34e44cb9653 100644
--- a/naga/tests/out/ir/wgsl-types_with_comments.ron
+++ b/naga/tests/out/ir/wgsl-types_with_comments.ron
@@ -172,6 +172,8 @@
                 ],
                 diagnostic_filter_leaf: None,
             ),
+            mesh_info: None,
+            task_payload: None,
         ),
     ],
     diagnostic_filters: [],
diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs
index 2c2f4b36c44..ae199f2c703 100644
--- a/wgpu-core/src/validation.rs
+++ b/wgpu-core/src/validation.rs
@@ -1085,6 +1085,8 @@ impl Interface {
             wgt::ShaderStages::VERTEX => naga::ShaderStage::Vertex,
             wgt::ShaderStages::FRAGMENT => naga::ShaderStage::Fragment,
             wgt::ShaderStages::COMPUTE => naga::ShaderStage::Compute,
+            wgt::ShaderStages::MESH => naga::ShaderStage::Mesh,
+            wgt::ShaderStages::TASK => naga::ShaderStage::Task,
             _ => unreachable!(),
         }
     }
@@ -1229,7 +1231,7 @@ impl Interface {
         }
 
         // check workgroup size limits
-        if shader_stage == naga::ShaderStage::Compute {
+        if shader_stage.compute_like() {
             let max_workgroup_size_limits = [
                 self.limits.max_compute_workgroup_size_x,
                 self.limits.max_compute_workgroup_size_y,
diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs
index 71c469806e0..629d9bc937c 100644
--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@@ -2107,6 +2107,9 @@ impl super::Adapter {
             if features.contains(wgt::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN) {
                 capabilities.push(spv::Capability::RayQueryPositionFetchKHR)
             }
+            if features.contains(wgt::Features::EXPERIMENTAL_MESH_SHADER) {
+                capabilities.push(spv::Capability::MeshShadingEXT);
+            }
             if self.private_caps.shader_integer_dot_product {
                 // See <https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VK_KHR_shader_integer_dot_product.html#_new_spir_v_capabilities>.
                 capabilities.extend(&[
diff --git a/wgpu/src/api/render_pass.rs b/wgpu/src/api/render_pass.rs
index 8163b4261f0..a832e380fbf 100644
--- a/wgpu/src/api/render_pass.rs
+++ b/wgpu/src/api/render_pass.rs
@@ -226,7 +226,33 @@ impl RenderPass<'_> {
         self.inner.draw_indexed(indices, base_vertex, instances);
     }
 
-    /// Draws using a mesh shader pipeline
+    /// Draws using a mesh shader pipeline.
+    ///
+    /// The current pipeline must be a mesh shader pipeline.
+    ///
+    /// If the current pipeline has a task shader, run it with an workgroup for
+    /// every `vec3<u32>(i, j, k)` where `i`, `j`, and `k` are between `0` and
+    /// `group_count_x`, `group_count_y`, and `group_count_z`. The invocation with
+    /// index zero in each group is responsible for determining the mesh shader dispatch.
+    /// Its return value indicates the number of workgroups of mesh shaders to invoke. It also
+    /// passes a payload value for them to consume. Because each task workgroup is essentially
+    /// a mesh shader draw call, mesh workgroups dispatched by different task workgroups
+    /// cannot interact in any way, and `workgroup_id` corresponds to its location in the
+    /// calling specific task shader's dispatch group.
+    ///
+    /// If the current pipeline lacks a task shader, run its mesh shader with a
+    /// workgroup for every `vec3<u32>(i, j, k)` where `i`, `j`, and `k` are
+    /// between `0` and `group_count_x`, `group_count_y`, and `group_count_z`.
+    ///
+    /// Each mesh shader workgroup outputs a set of vertices and indices for primitives.
+    /// The indices outputted correspond to the vertices outputted by that same workgroup;
+    /// there is no global vertex buffer. These primitives are passed to the rasterizer and
+    /// essentially treated like a vertex shader output, except that the mesh shader may
+    /// choose to cull specific primitives or pass per-primitive non-interpolated values
+    /// to the mesh shader. As such, each primitive is then rendered with the current
+    /// pipeline's fragment shader, if present. Otherwise, [No Color Output mode] is used.
+    ///
+    /// [No Color Output mode]: https://www.w3.org/TR/webgpu/#no-color-output
     pub fn draw_mesh_tasks(&mut self, group_count_x: u32, group_count_y: u32, group_count_z: u32) {
         self.inner
             .draw_mesh_tasks(group_count_x, group_count_y, group_count_z);
diff --git a/wgpu/src/api/render_pipeline.rs b/wgpu/src/api/render_pipeline.rs
index e887bb4b97e..be16d91f27a 100644
--- a/wgpu/src/api/render_pipeline.rs
+++ b/wgpu/src/api/render_pipeline.rs
@@ -238,7 +238,43 @@ static_assertions::assert_impl_all!(RenderPipelineDescriptor<'_>: Send, Sync);
 
 /// Describes a mesh shader (graphics) pipeline.
 ///
-/// For use with [`Device::create_mesh_pipeline`].
+/// For use with [`Device::create_mesh_pipeline`]. A mesh pipeline is very much
+/// like a render pipeline, except that instead of [`RenderPass::draw`] it is
+/// invoked with [`RenderPass::draw_mesh_tasks`], and instead of a vertex shader
+/// and a fragment shader:
+///
+/// - [`task`] specifies an optional task shader entry point, which determines how
+///   many groups of mesh shaders to dispatch.
+///
+/// - [`mesh`] specifies a mesh shader entry point, which generates groups of
+///   primitives to draw
+///
+/// - [`fragment`] specifies as fragment shader for drawing those primitives,
+///   just like in an ordinary render pipeline.
+///
+/// The key difference is that, whereas a vertex shader is invoked on the
+/// elements of vertex buffers, the task shader gets to decide how many mesh
+/// shader workgroups to make, and then each mesh shader workgroup gets to
+/// decide which primitives it wants to generate, and what their vertex
+/// attributes are. Task and mesh shaders can use whatever they please as
+/// inputs, like a compute shader. However, they cannot use specialized vertex
+/// or index buffers.
+///
+/// A mesh pipeline is invoked by [`RenderPass::draw_mesh_tasks`], which looks
+/// like a compute shader dispatch with [`ComputePass::dispatch_workgroups`]:
+/// you pass `x`, `y`, and `z` values indicating the number of task shaders to
+/// invoke in parallel. The output value of the first thread in a task shader
+/// workgroup determines how many mesh workgroups should be dispatched from there.
+/// Those mesh workgroups also get a special payload passed from the task shader.
+///
+/// If the task shader is omitted, then the (`x`, `y`, `z`) parameters to
+/// `draw_mesh_tasks` are used to decide how many invocations of the mesh shader
+/// to invoke directly, without a task payload.
+///
+/// [vertex formats]: wgpu_types::VertexFormat
+/// [`task`]: Self::task
+/// [`mesh`]: Self::mesh
+/// [`fragment`]: Self::fragment
 #[derive(Clone, Debug)]
 pub struct MeshPipelineDescriptor<'a> {
     /// Debug label of the pipeline. This will show up in graphics debuggers for easy identification.
@@ -263,7 +299,7 @@ pub struct MeshPipelineDescriptor<'a> {
     ///
     /// [default layout]: https://www.w3.org/TR/webgpu/#default-pipeline-layout
     pub layout: Option<&'a PipelineLayout>,
-    /// The compiled task stage, its entry point, and the color targets.
+    /// The compiled task stage and its entry point.
     pub task: Option<TaskState<'a>>,
     /// The compiled mesh stage and its entry point
     pub mesh: MeshState<'a>,