Quelle mesh_shader.rs
Sprache: Rust
use alloc::{
format,
string::{String, ToString},
vec::Vec,
};
use crate ::{
back::{
self ,
msl::{
writer::{TypeContext, TypedGlobalVariable},
BackendResult, EntryPointArgument, Error, NAMESPACE, WRAPPED_ARRAY_FIELD,
},
},
proc::NameKey,
};
pub (super ) struct MeshOutputInfo {
out_vertex_ty_name: String,
out_primitive_ty_name: String,
out_vertex_member_names: Vec<Option<String>>,
out_primitive_member_names: Vec<Option<String
}
pub (super ) struct NestedFunctionInfo<'a> {
pub (super ) options: &'a super::Options,
pub (super ) ep: &'a crate::EntryPoint,
pub (super ) module: &'a crate::Module,
pub (super ) mod_info: &'a crate::valid::ModuleInfo,
pub (super ) fun_info: &'a crate::valid::FunctionInfo,
pub (super ) args: Vec<EntryPointArgument>,
pub (super ) local_invocation_index: Option<&'a NameKey>,
pub (super ) nested_name: &'a str,
pub (super ) outer_name: &'a str,
pub (super ) out_mesh_info: Option<MeshOutputInfo>,
}
impl <W: core::fmt::Write> super ::Writer<W> {
/// This writes the output vertex and primitive structs given the reflection information about them.
pub (super ) fn write_mesh_output_types(
&mut self ,
mesh_info: &crate ::MeshStageInfo,
fun_name: &str,
module: &crate ::Module,
// See `PipelineOptions::allow_and_force_point_size`
allow_and_force_point_size: bool,
options: &super ::Options,
) -> Result<MeshOutputInfo, Error> {
let mut vertex_member_names = Vec::new();
let mut primitive_member_names = Vec::new();
let vertex_out_name = self .namer.call(&format!("{fun_name}VertexOutput" ));
let primitive_out_name = self .namer.call(&format!("{fun_name}PrimitiveOutput" ));
let mut existing_names = Vec::new();
for (out_name, struct_ty, is_primitive, member_names) in [
(
&vertex_out_name,
mesh_info.vertex_output_type,
false ,
&mut vertex_member_names,
),
(
&primitive_out_name,
mesh_info.primitive_output_type,
true ,
&mut primitive_member_names,
),
] {
writeln!(self .out, "struct {out_name} {{" )?;
// Mesh output types are guaranteed to be user defined structs. This is validated by naga.
let crate ::TypeInner::Struct { ref members, .. } = module.types[struct_ty].inner else {
unreachable!()
};
let mut has_point_size = false ;
for (index, member) in members.iter().enumerate() {
member_names.push(None);
let ty_name = TypeContext {
handle: member.ty,
gctx: module.to_ctx(),
names: &self .names,
access: crate ::StorageAccess::empty(),
first_time: true ,
};
let binding = member
.binding
.clone()
.ok_or_else(|| Error::GenericValidation("Expected binding, got None" .into()))?;
if let crate ::Binding::BuiltIn(crate ::BuiltIn::PointSize) = binding {
has_point_size = true ;
if !allow_and_force_point_size {
continue ;
}
}
if let crate ::Binding::BuiltIn(
crate ::BuiltIn::PointIndex
| crate ::BuiltIn::LineIndices
| crate ::BuiltIn::TriangleIndices,
) = binding
{
continue ;
}
// Names of struct members must be unique across vertex and primitive output.
// Therefore, when writing the primitive output struct, we might need to rename some fields.
let mut name = self .names[&NameKey::StructMember(struct_ty, index as u32)].clone();
if existing_names.contains(&name) {
name = self .namer.call(&name);
} else {
// Let the namer know this is illegal to use again
let _ = self .namer.call(&name);
}
let array_len = match module.types[member.ty].inner {
crate ::TypeInner::Array {
size: crate ::ArraySize::Constant(size),
..
} => Some(size),
_ => None,
};
let resolved =
options.resolve_local_binding(&binding, back::msl::LocationMode::MeshOutput)?;
write!(self .out, "{}{} {}" , back::INDENT, ty_name, name)?;
if let Some(array_len) = array_len {
write!(self .out, " [{array_len}]" )?;
}
resolved.try_fmt(&mut self .out)?;
writeln!(self .out, ";" )?;
*member_names.last_mut().unwrap() = Some(name.clone());
existing_names.push(name);
}
if allow_and_force_point_size && !has_point_size && !is_primitive {
// inject the point size output last
writeln!(
self .out,
"{}float _point_size [[point_size]];" ,
back::INDENT
)?;
}
writeln!(self .out, "}};" )?;
}
Ok(MeshOutputInfo {
out_vertex_ty_name: vertex_out_name,
out_primitive_ty_name: primitive_out_name,
out_vertex_member_names: vertex_member_names,
out_primitive_member_names: primitive_member_names,
})
}
pub (super ) fn write_wrapper_function(&mut self , info: NestedFunctionInfo<'_>) -> BackendResult {
let NestedFunctionInfo {
options,
ep,
module,
mod_info,
fun_info,
args,
local_invocation_index: local_invocation_index_key,
nested_name,
outer_name,
out_mesh_info,
} = info;
let indent = back::INDENT;
let em_str = match ep.stage {
crate ::ShaderStage::Mesh => "[[mesh]]" ,
crate ::ShaderStage::Task => "[[object]]" ,
_ => unreachable!(),
};
writeln!(self .out, "{em_str} void {outer_name}(" )?;
// Arguments
let mut mesh_out_name: Option<String> = None;
let mut mesh_variable_name = None;
let mut task_grid_name = None;
if let Some(ref info) = ep.mesh_info {
let mesh_out = out_mesh_info.as_ref().unwrap();
let mesh_name = self .namer.call("meshOutput" );
let topology_name = match info.topology {
crate ::MeshOutputTopology::Points => "point" ,
crate ::MeshOutputTopology::Lines => "line" ,
crate ::MeshOutputTopology::Triangles => "triangle" ,
};
let num_verts = info.max_vertices;
let num_prims = info.max_primitives;
writeln!(self .out,
" {NAMESPACE}::mesh<{}, {}, {num_verts}, {num_prims}, metal::topology::{topology_name}> {mesh_name}" ,
mesh_out.out_vertex_ty_name,
mesh_out.out_primitive_ty_name,
)?;
mesh_out_name = Some(mesh_name);
mesh_variable_name = Some(
self .names
[&NameKey::GlobalVariable(ep.mesh_info.as_ref().unwrap().output_variable)]
.clone(),
);
} else if ep.stage == crate ::ShaderStage::Task {
let grid_name = self .namer.call("nagaMeshGrid" );
writeln!(self .out, " {NAMESPACE}::mesh_grid_properties {grid_name}" )?;
task_grid_name = Some(grid_name);
}
let local_invocation_index = if let Some(key) = local_invocation_index_key {
self .names[key].clone()
} else {
"__local_invocation_index" .to_string()
};
for arg in &args {
write!(self .out, ", {} {}{}" , arg.ty_name, arg.name, arg.binding)?;
if let Some(init) = arg.init {
write!(self .out, " = " )?;
self .put_const_expression(init, module, mod_info, &module.global_expressions)?;
}
writeln!(self .out)?;
}
writeln!(self .out, ") {{" )?;
// Function body
if ep.stage == crate ::ShaderStage::Mesh {
for (handle, var) in module.global_variables.iter() {
if var.space != crate ::AddressSpace::WorkGroup || fun_info[handle].is_empty() {
continue ;
}
let tyvar = TypedGlobalVariable {
module,
names: &self .names,
handle,
usage: crate ::valid::GlobalUse::WRITE | crate ::valid::GlobalUse::READ,
reference: false ,
};
write!(self .out, "{}" , back::INDENT)?;
tyvar.try_fmt(&mut self .out)?;
writeln!(self .out, ";" )?;
}
}
write!(self .out, "{indent}" )?;
let result_name = if ep.stage == crate ::ShaderStage::Task {
let name = self .namer.call("nagaGridSize" );
write!(self .out, "uint3 {} = " , name)?;
Some(name)
} else {
None
};
write!(self .out, "{nested_name}(" )?;
{
let mut is_first = true ;
for arg in &args {
if !is_first {
write!(self .out, ", " )?;
}
is_first = false ;
write!(self .out, "{}" , arg.name)?;
}
if ep.stage == crate ::ShaderStage::Mesh {
for (handle, var) in module.global_variables.iter() {
if var.space != crate ::AddressSpace::WorkGroup || fun_info[handle].is_empty() {
continue ;
}
if !is_first {
write!(self .out, ", " )?;
}
let name = &self .names[&NameKey::GlobalVariable(handle)];
write!(self .out, "{name}" )?;
}
}
}
writeln!(self .out, ");" )?;
self .write_barrier(crate ::Barrier::WORK_GROUP, back::Level(1 ))?;
if let Some(grid_name) = task_grid_name {
let result_name = result_name.unwrap();
writeln!(self .out, "{indent}if ({local_invocation_index} == 0u) {{" )?;
{
let level2 = back::Level(2 );
if let Some(limits) = options.task_dispatch_limits {
let level3 = back::Level(3 );
let max_per_dim = limits.max_mesh_workgroups_per_dim;
let max_total = limits.max_mesh_workgroups_total;
writeln!(self .out, "{level2}if (" )?;
writeln!(self .out, "{level3}{result_name}.x > {max_per_dim}u ||" )?;
writeln!(self .out, "{level3}{result_name}.y > {max_per_dim}u ||" )?;
writeln!(self .out, "{level3}{result_name}.z > {max_per_dim}u ||" )?;
writeln!(
self .out,
"{level3}{NAMESPACE}::mulhi({result_name}.x, {result_name}.y) != 0u ||"
)?;
writeln!(
self .out,
"{level3}{NAMESPACE}::mulhi({result_name}.x * {result_name}.y, {result_name}.z) != 0u ||"
)?;
writeln!(self .out, "{level3}({result_name}.x * {result_name}.y * {result_name}.z) > {max_total}u" )?;
writeln!(self .out, "{level2}) {{" )?;
writeln!(self .out, "{level3}{result_name} = {NAMESPACE}::uint3(0u);" )?;
writeln!(self .out, "{level2}}}" )?;
}
writeln!(
self .out,
"{level2}{grid_name}.set_threadgroups_per_grid({result_name});"
)?;
}
writeln!(self .out, "{indent}}}" )?;
writeln!(self .out, "{indent}return;" )?;
} else if let Some(ref info) = ep.mesh_info {
let mesh_out = out_mesh_info.as_ref().unwrap();
let out_ty = module.global_variables[info.output_variable].ty;
let mesh_out_name = mesh_out_name.unwrap();
let mesh_variable_name = mesh_variable_name.unwrap();
// The output type is guaranteed to be a struct with exactly 4 members
let crate ::TypeInner::Struct { ref members, .. } = module.types[out_ty].inner else {
unreachable!();
};
let get_out_value = |bi| {
let member_idx = members
.iter()
.position(|a| a.binding == Some(crate ::Binding::BuiltIn(bi)))
.unwrap() as u32;
format!(
"{}.{}" ,
mesh_variable_name,
self .names[&NameKey::StructMember(out_ty, member_idx)]
)
};
let vert_count = format!(
"{NAMESPACE}::min({}, {}u)" ,
get_out_value(crate ::BuiltIn::VertexCount),
info.max_vertices
);
let prim_count = format!(
"{NAMESPACE}::min({}, {}u)" ,
get_out_value(crate ::BuiltIn::PrimitiveCount),
info.max_primitives
);
let workgroup_size: u32 = ep.workgroup_size.iter().product();
{
let vert_index = self .namer.call("vertexIndex" );
let in_array = get_out_value(crate ::BuiltIn::Vertices);
writeln!(
self .out,
"{indent}for(uint {vert_index} = {local_invocation_index}; {vert_index} < {vert_count}; {vert_index} += {workgroup_size}) {{"
)?;
let out_vert = self .namer.call("vertex" );
writeln!(
self .out,
"{indent}{indent}{} {out_vert};" ,
mesh_out.out_vertex_ty_name,
)?;
for (member_idx, new_name) in mesh_out.out_vertex_member_names.iter().enumerate() {
let in_value = format!(
"{in_array}.{WRAPPED_ARRAY_FIELD}[{vert_index}].{}" ,
self .names
[&NameKey::StructMember(info.vertex_output_type, member_idx as u32)]
);
let out_value = format!("{out_vert}.{}" , new_name.as_ref().unwrap());
writeln!(self .out, "{indent}{indent}{out_value} = {in_value};" )?;
}
writeln!(
self .out,
"{indent}{indent}{}.set_vertex({vert_index}, {out_vert});" ,
mesh_out_name
)?;
writeln!(self .out, "{indent}}}" )?;
}
{
let prim_index = self .namer.call("primitiveIndex" );
let in_array = get_out_value(crate ::BuiltIn::Primitives);
writeln!(
self .out,
"{indent}for(uint {prim_index} = {local_invocation_index}; {prim_index} < {prim_count}; {prim_index} += {workgroup_size}) {{"
)?;
let out_prim = self .namer.call("primitive" );
writeln!(
self .out,
"{indent}{indent}{} {out_prim};" ,
mesh_out.out_primitive_ty_name
)?;
for (member_idx, new_name) in mesh_out.out_primitive_member_names.iter().enumerate()
{
let in_value = format!(
"{in_array}.{WRAPPED_ARRAY_FIELD}[{prim_index}].{}" ,
self .names
[&NameKey::StructMember(info.primitive_output_type, member_idx as u32)]
);
if let Some(new_name) = new_name.as_ref() {
let out_value = format!("{out_prim}.{new_name}" );
writeln!(
self .out,
"{indent}{}{out_value} = {in_value};" ,
back::INDENT
)?;
} else {
let num_indices = match info.topology {
crate ::MeshOutputTopology::Points => 1 ,
crate ::MeshOutputTopology::Lines => 2 ,
crate ::MeshOutputTopology::Triangles => 3 ,
};
for i in 0 ..num_indices {
let component = if num_indices == 1 {
"" .to_string()
} else {
format!(".{}" , back::COMPONENTS[i])
};
writeln!(
self .out,
"{indent}{}{}.set_index({prim_index} * {num_indices} + {i}, {in_value}{component});" ,
back::INDENT,
mesh_out_name,
)?;
}
}
}
writeln!(
self .out,
"{indent}{}{}.set_primitive({prim_index}, {out_prim});" ,
back::INDENT,
mesh_out_name
)?;
writeln!(self .out, "{indent}}}" )?;
}
writeln!(self .out, "{indent}if ({local_invocation_index} == 0u) {{" )?;
writeln!(
self .out,
"{indent}{indent}{}.set_primitive_count({prim_count});" ,
mesh_out_name,
)?;
writeln!(self .out, "{indent}}}" )?;
} else {
// Must either have task output grid (task shader) or mesh output info (mesh shader)
unreachable!()
}
writeln!(self .out, "}}" )?;
Ok(())
}
}
Messung V0.5 in Prozent C=94 H=100 G=96
='red' font-size=6>H=100 G=96
¤ Dauer der Verarbeitung: 0.17 Sekunden
(vorverarbeitet am 2026-08-26)
¤
*© Formatika GbR, Deutschland