Skip to content

Commit bf9f752

Browse files
inner-daemonsjimblandy
inner-daemons
and
jimblandy
authored
Add mesh shading info to naga IR (#8104)
Co-authored-by: Jim Blandy <[email protected]> Co-authored-by: SupaMaggie70Incorporated <[email protected]>
1 parent e7fcb94 commit bf9f752

File tree

77 files changed

+1419
-112
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

77 files changed

+1419
-112
lines changed

‎docs/api-specs/mesh_shading.md‎

Lines changed: 108 additions & 27 deletions
Large diffs are not rendered by default.

‎naga-cli/src/bin/naga.rs‎

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -64,6 +64,12 @@ struct Args {
6464
#[argh(option)]
6565
shader_model: Option<ShaderModelArg>,
6666

67+
/// the SPIR-V version to use if targeting SPIR-V
68+
///
69+
/// For example, 1.0, 1.4, etc
70+
#[argh(option)]
71+
spirv_version: Option<SpirvVersionArg>,
72+
6773
/// the shader stage, for example 'frag', 'vert', or 'compute'.
6874
/// if the shader stage is unspecified it will be derived from
6975
/// the file extension.
@@ -189,6 +195,22 @@ impl FromStr for ShaderModelArg {
189195
}
190196
}
191197

198+
#[derive(Debug, Clone)]
199+
struct SpirvVersionArg(u8, u8);
200+
201+
impl FromStr for SpirvVersionArg {
202+
type Err = String;
203+
204+
fn from_str(s: &str) -> Result<Self, Self::Err> {
205+
let dot = s
206+
.find(".")
207+
.ok_or_else(|| "Missing dot separator".to_owned())?;
208+
let major = s[..dot].parse::<u8>().map_err(|e| e.to_string())?;
209+
let minor = s[dot + 1..].parse::<u8>().map_err(|e| e.to_string())?;
210+
Ok(Self(major, minor))
211+
}
212+
}
213+
192214
/// Newtype so we can implement [`FromStr`] for `ShaderSource`.
193215
#[derive(Debug, Clone, Copy)]
194216
struct ShaderStage(naga::ShaderStage);
@@ -465,6 +487,9 @@ fn run() -> anyhow::Result<()> {
465487
if let Some(ref version) = args.metal_version {
466488
params.msl.lang_version = version.0;
467489
}
490+
if let Some(ref version) = args.spirv_version {
491+
params.spv_out.lang_version = (version.0, version.1);
492+
}
468493
params.keep_coordinate_space = args.keep_coordinate_space;
469494

470495
params.dot.cfg_only = args.dot_cfg_only;

‎naga/src/back/dot/mod.rs‎

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -307,6 +307,25 @@ impl StatementGraph {
307307
crate::RayQueryFunction::Terminate => "RayQueryTerminate",
308308
}
309309
}
310+
S::MeshFunction(crate::MeshFunction::SetMeshOutputs {
311+
vertex_count,
312+
primitive_count,
313+
}) => {
314+
self.dependencies.push((id, vertex_count, "vertex_count"));
315+
self.dependencies
316+
.push((id, primitive_count, "primitive_count"));
317+
"SetMeshOutputs"
318+
}
319+
S::MeshFunction(crate::MeshFunction::SetVertex { index, value }) => {
320+
self.dependencies.push((id, index, "index"));
321+
self.dependencies.push((id, value, "value"));
322+
"SetVertex"
323+
}
324+
S::MeshFunction(crate::MeshFunction::SetPrimitive { index, value }) => {
325+
self.dependencies.push((id, index, "index"));
326+
self.dependencies.push((id, value, "value"));
327+
"SetPrimitive"
328+
}
310329
S::SubgroupBallot { result, predicate } => {
311330
if let Some(predicate) = predicate {
312331
self.dependencies.push((id, predicate, "predicate"));

‎naga/src/back/glsl/features.rs‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -623,6 +623,7 @@ impl<W> Writer<'_, W> {
623623
interpolation,
624624
sampling,
625625
blend_src,
626+
per_primitive: _,
626627
} => {
627628
if interpolation == Some(Interpolation::Linear) {
628629
self.features.request(Features::NOPERSPECTIVE_QUALIFIER);

‎naga/src/back/glsl/mod.rs‎

Lines changed: 23 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -139,7 +139,8 @@ impl crate::AddressSpace {
139139
| crate::AddressSpace::Uniform
140140
| crate::AddressSpace::Storage { .. }
141141
| crate::AddressSpace::Handle
142-
| crate::AddressSpace::PushConstant => false,
142+
| crate::AddressSpace::PushConstant
143+
| crate::AddressSpace::TaskPayload => false,
143144
}
144145
}
145146
}
@@ -1300,6 +1301,9 @@ impl<'a, W: Write> Writer<'a, W> {
13001301
crate::AddressSpace::Storage { .. } => {
13011302
self.write_interface_block(handle, global)?;
13021303
}
1304+
crate::AddressSpace::TaskPayload => {
1305+
self.write_interface_block(handle, global)?;
1306+
}
13031307
// A global variable in the `Function` address space is a
13041308
// contradiction in terms.
13051309
crate::AddressSpace::Function => unreachable!(),
@@ -1614,6 +1618,7 @@ impl<'a, W: Write> Writer<'a, W> {
16141618
interpolation,
16151619
sampling,
16161620
blend_src,
1621+
per_primitive: _,
16171622
} => (location, interpolation, sampling, blend_src),
16181623
crate::Binding::BuiltIn(built_in) => {
16191624
match built_in {
@@ -1732,6 +1737,7 @@ impl<'a, W: Write> Writer<'a, W> {
17321737
interpolation: None,
17331738
sampling: None,
17341739
blend_src,
1740+
per_primitive: false,
17351741
},
17361742
stage: self.entry_point.stage,
17371743
options: VaryingOptions::from_writer_options(self.options, output),
@@ -1873,7 +1879,7 @@ impl<'a, W: Write> Writer<'a, W> {
18731879
writeln!(self.out, ") {{")?;
18741880

18751881
if self.options.zero_initialize_workgroup_memory
1876-
&& ctx.ty.is_compute_entry_point(self.module)
1882+
&& ctx.ty.is_compute_like_entry_point(self.module)
18771883
{
18781884
self.write_workgroup_variables_initialization(&ctx)?;
18791885
}
@@ -2669,6 +2675,11 @@ impl<'a, W: Write> Writer<'a, W> {
26692675
self.write_image_atomic(ctx, image, coordinate, array_index, fun, value)?
26702676
}
26712677
Statement::RayQuery { .. } => unreachable!(),
2678+
Statement::MeshFunction(
2679+
crate::MeshFunction::SetMeshOutputs { .. }
2680+
| crate::MeshFunction::SetVertex { .. }
2681+
| crate::MeshFunction::SetPrimitive { .. },
2682+
) => unreachable!(),
26722683
Statement::SubgroupBallot { result, predicate } => {
26732684
write!(self.out, "{level}")?;
26742685
let res_name = Baked(result).to_string();
@@ -5248,6 +5259,15 @@ const fn glsl_built_in(built_in: crate::BuiltIn, options: VaryingOptions) -> &'s
52485259
Bi::SubgroupId => "gl_SubgroupID",
52495260
Bi::SubgroupSize => "gl_SubgroupSize",
52505261
Bi::SubgroupInvocationId => "gl_SubgroupInvocationID",
5262+
// mesh
5263+
// TODO: figure out how to map these to glsl things as glsl treats them as arrays
5264+
Bi::CullPrimitive
5265+
| Bi::PointIndex
5266+
| Bi::LineIndices
5267+
| Bi::TriangleIndices
5268+
| Bi::MeshTaskSize => {
5269+
unimplemented!()
5270+
}
52515271
}
52525272
}
52535273

@@ -5263,6 +5283,7 @@ const fn glsl_storage_qualifier(space: crate::AddressSpace) -> Option<&'static s
52635283
As::Handle => Some("uniform"),
52645284
As::WorkGroup => Some("shared"),
52655285
As::PushConstant => Some("uniform"),
5286+
As::TaskPayload => unreachable!(),
52665287
}
52675288
}
52685289

‎naga/src/back/hlsl/conv.rs‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -184,6 +184,9 @@ impl crate::BuiltIn {
184184
Self::PointSize | Self::ViewIndex | Self::PointCoord | Self::DrawID => {
185185
return Err(Error::Custom(format!("Unsupported builtin {self:?}")))
186186
}
187+
Self::CullPrimitive => "SV_CullPrimitive",
188+
Self::PointIndex | Self::LineIndices | Self::TriangleIndices => unimplemented!(),
189+
Self::MeshTaskSize => unreachable!(),
187190
})
188191
}
189192
}

‎naga/src/back/hlsl/writer.rs‎

Lines changed: 18 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -507,7 +507,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
507507

508508
self.write_wrapped_functions(module, &ctx)?;
509509

510-
if ep.stage == ShaderStage::Compute {
510+
if ep.stage.compute_like() {
511511
// HLSL is calling workgroup size "num threads"
512512
let num_threads = ep.workgroup_size;
513513
writeln!(
@@ -967,6 +967,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
967967
self.write_type(module, global.ty)?;
968968
""
969969
}
970+
crate::AddressSpace::TaskPayload => unimplemented!(),
970971
crate::AddressSpace::Uniform => {
971972
// constant buffer declarations are expected to be inlined, e.g.
972973
// `cbuffer foo: register(b0) { field1: type1; }`
@@ -1764,7 +1765,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
17641765
module: &Module,
17651766
) -> bool {
17661767
self.options.zero_initialize_workgroup_memory
1767-
&& func_ctx.ty.is_compute_entry_point(module)
1768+
&& func_ctx.ty.is_compute_like_entry_point(module)
17681769
&& module.global_variables.iter().any(|(handle, var)| {
17691770
!func_ctx.info[handle].is_empty() && var.space == crate::AddressSpace::WorkGroup
17701771
})
@@ -2599,6 +2600,19 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
25992600
writeln!(self.out, ".Abort();")?;
26002601
}
26012602
},
2603+
Statement::MeshFunction(crate::MeshFunction::SetMeshOutputs {
2604+
vertex_count,
2605+
primitive_count,
2606+
}) => {
2607+
write!(self.out, "{level}SetMeshOutputCounts(")?;
2608+
self.write_expr(module, vertex_count, func_ctx)?;
2609+
write!(self.out, ", ")?;
2610+
self.write_expr(module, primitive_count, func_ctx)?;
2611+
write!(self.out, ");")?;
2612+
}
2613+
Statement::MeshFunction(
2614+
crate::MeshFunction::SetVertex { .. } | crate::MeshFunction::SetPrimitive { .. },
2615+
) => unimplemented!(),
26022616
Statement::SubgroupBallot { result, predicate } => {
26032617
write!(self.out, "{level}")?;
26042618
let name = Baked(result).to_string();
@@ -3076,7 +3090,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
30763090
crate::AddressSpace::Function
30773091
| crate::AddressSpace::Private
30783092
| crate::AddressSpace::WorkGroup
3079-
| crate::AddressSpace::PushConstant,
3093+
| crate::AddressSpace::PushConstant
3094+
| crate::AddressSpace::TaskPayload,
30803095
)
30813096
| None => true,
30823097
Some(crate::AddressSpace::Uniform) => {

‎naga/src/back/mod.rs‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -139,11 +139,11 @@ pub enum FunctionType {
139139
}
140140

141141
impl FunctionType {
142-
/// Returns true if the function is an entry point for a compute shader.
143-
pub fn is_compute_entry_point(&self, module: &crate::Module) -> bool {
142+
/// Returns true if the function is an entry point for a compute-like shader.
143+
pub fn is_compute_like_entry_point(&self, module: &crate::Module) -> bool {
144144
match *self {
145145
FunctionType::EntryPoint(index) => {
146-
module.entry_points[index as usize].stage == crate::ShaderStage::Compute
146+
module.entry_points[index as usize].stage.compute_like()
147147
}
148148
FunctionType::Function(_) => false,
149149
}

‎naga/src/back/msl/mod.rs‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -545,6 +545,7 @@ impl Options {
545545
interpolation,
546546
sampling,
547547
blend_src,
548+
per_primitive: _,
548549
} => match mode {
549550
LocationMode::VertexInput => Ok(ResolvedBinding::Attribute(location)),
550551
LocationMode::FragmentOutput => {
@@ -703,6 +704,10 @@ impl ResolvedBinding {
703704
Bi::CullDistance | Bi::ViewIndex | Bi::DrawID => {
704705
return Err(Error::UnsupportedBuiltIn(built_in))
705706
}
707+
Bi::CullPrimitive => "primitive_culled",
708+
// TODO: figure out how to make this written as a function call
709+
Bi::PointIndex | Bi::LineIndices | Bi::TriangleIndices => unimplemented!(),
710+
Bi::MeshTaskSize => unreachable!(),
706711
};
707712
write!(out, "{name}")?;
708713
}

‎naga/src/back/msl/writer.rs‎

Lines changed: 17 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -594,7 +594,8 @@ impl crate::AddressSpace {
594594
| Self::Private
595595
| Self::WorkGroup
596596
| Self::PushConstant
597-
| Self::Handle => true,
597+
| Self::Handle
598+
| Self::TaskPayload => true,
598599
Self::Function => false,
599600
}
600601
}
@@ -607,6 +608,7 @@ impl crate::AddressSpace {
607608
// may end up with "const" even if the binding is read-write,
608609
// and that should be OK.
609610
Self::Storage { .. } => true,
611+
Self::TaskPayload => unimplemented!(),
610612
// These should always be read-write.
611613
Self::Private | Self::WorkGroup => false,
612614
// These translate to `constant` address space, no need for qualifiers.
@@ -623,6 +625,7 @@ impl crate::AddressSpace {
623625
Self::Storage { .. } => Some("device"),
624626
Self::Private | Self::Function => Some("thread"),
625627
Self::WorkGroup => Some("threadgroup"),
628+
Self::TaskPayload => Some("object_data"),
626629
}
627630
}
628631
}
@@ -4060,6 +4063,14 @@ impl<W: Write> Writer<W> {
40604063
}
40614064
}
40624065
}
4066+
// TODO: write emitters for these
4067+
crate::Statement::MeshFunction(crate::MeshFunction::SetMeshOutputs { .. }) => {
4068+
unimplemented!()
4069+
}
4070+
crate::Statement::MeshFunction(
4071+
crate::MeshFunction::SetVertex { .. }
4072+
| crate::MeshFunction::SetPrimitive { .. },
4073+
) => unimplemented!(),
40634074
crate::Statement::SubgroupBallot { result, predicate } => {
40644075
write!(self.out, "{level}")?;
40654076
let name = self.namer.call("");
@@ -6619,7 +6630,7 @@ template <typename A>
66196630
LocationMode::Uniform,
66206631
false,
66216632
),
6622-
crate::ShaderStage::Task | crate::ShaderStage::Mesh => unreachable!(),
6633+
crate::ShaderStage::Task | crate::ShaderStage::Mesh => unimplemented!(),
66236634
};
66246635

66256636
// Should this entry point be modified to do vertex pulling?
@@ -6686,6 +6697,9 @@ template <typename A>
66866697
break;
66876698
}
66886699
}
6700+
crate::AddressSpace::TaskPayload => {
6701+
unimplemented!()
6702+
}
66896703
crate::AddressSpace::Function
66906704
| crate::AddressSpace::Private
66916705
| crate::AddressSpace::WorkGroup => {}
@@ -7683,7 +7697,7 @@ mod workgroup_mem_init {
76837697
fun_info: &valid::FunctionInfo,
76847698
) -> bool {
76857699
options.zero_initialize_workgroup_memory
7686-
&& ep.stage == crate::ShaderStage::Compute
7700+
&& ep.stage.compute_like()
76877701
&& module.global_variables.iter().any(|(handle, var)| {
76887702
!fun_info[handle].is_empty() && var.space == crate::AddressSpace::WorkGroup
76897703
})

0 commit comments

Comments
 (0)