github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

Update IR to have entry points with inputs/outputs

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-03-07 00:03:05 -05:00
committed by Dzmitry Malyshau
parent 58a4e330dc
commit ebcd815250
50 changed files with 3175 additions and 2367 deletions
Download Patch File Download Diff File Expand all files Collapse all files

266
src/back/glsl/mod.rs
View File

@@ -202,16 +202,14 @@ impl<'a, 'b> FunctionCtx<'a, 'b> {
}
}
/// Helper method that retrieves the name of the argument in the current function
/// Helper method that generates a [`NameKey`](crate::proc::NameKey) for a function argument.
///
/// # Panics
/// - If the function is an entry point
/// - If the function arguments are less or equal to `arg`
/// - If `names` hasn't been filled properly
fn get_arg<'c>(&self, arg: u32, names: &'c FastHashMap<NameKey, String>) -> &'c str {
fn argument_key(&self, arg: u32) -> NameKey {
match self.func {
FunctionType::Function(handle) => &names[&NameKey::FunctionArgument(handle, arg)],
FunctionType::EntryPoint(_) => unreachable!(),
FunctionType::Function(handle) => NameKey::FunctionArgument(handle, arg),
FunctionType::EntryPoint(ep_index) => NameKey::EntryPointArgument(ep_index, arg),
}
}
}
@@ -230,6 +228,33 @@ impl IdGenerator {
}
}
/// Helper wrapper used to get a name for a varying
///
/// Varying have different naming schemes depending on their binding:
/// - Varyings with builtin bindings get the from [`glsl_built_in`](glsl_built_in)
/// - Varyings with location bindings are named `_location_X` where `X` is the location
struct VaryingName<'a> {
binding: &'a Binding,
output: bool,
}
impl fmt::Display for VaryingName<'_> {
fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result {
match *self.binding {
Binding::Location(location, _) => {
write!(
f,
"_{}_location_{}",
if self.output { "out" } else { "in" },
location,
)
}
Binding::BuiltIn(built_in) => {
write!(f, "{}", glsl_built_in(built_in))
}
}
}
}
/// Shorthand result used internally by the backend
type BackendResult = Result<(), Error>;
@@ -450,12 +475,6 @@ impl<'a, W: Write> Writer<'a, W> {
continue;
}
// Skip builtins
// TODO: Write them if they have modifiers
if let Some(crate::Binding::BuiltIn(_)) = global.binding {
continue;
}
match self.module.types[global.ty].inner {
// We treat images separately because they might require
// writing the storage format
@@ -506,6 +525,15 @@ impl<'a, W: Write> Writer<'a, W> {
_ => self.write_global(handle, global)?,
}
}
for arg in self.entry_point.function.arguments.iter() {
self.write_varying(arg.binding.as_ref(), arg.ty, false)?;
}
if let Some(ref result) = self.entry_point.function.result {
self.write_varying(result.binding.as_ref(), result.ty, true)?;
}
writeln!(self.out)?;
// Write all regular functions
for (handle, function) in self.module.functions.iter() {
// Check that the function doesn't use globals that aren't supported
@@ -516,7 +544,6 @@ impl<'a, W: Write> Writer<'a, W> {
// We also `clone` to satisfy the borrow checker
let name = self.names[&NameKey::Function(handle)].clone();
let fun_info = &self.analysis[handle];
// Write the function
@@ -719,22 +746,6 @@ impl<'a, W: Write> Writer<'a, W> {
write!(self.out, "writeonly ")?;
}
// Write the interpolation modifier if needed
//
// We ignore all interpolation modifiers that aren't used in input globals in fragment
// shaders or output globals in vertex shaders
//
// TODO: Should this throw an error?
if let Some(interpolation) = global.interpolation {
match (self.options.shader_stage, global.class) {
(ShaderStage::Fragment, StorageClass::Input)
| (ShaderStage::Vertex, StorageClass::Output) => {
write!(self.out, "{} ", glsl_interpolation(interpolation)?)?;
}
_ => (),
};
}
// Write the storage class
// Trailing space is important
write!(self.out, "{} ", glsl_storage_class(global.class))?;
@@ -755,31 +766,65 @@ impl<'a, W: Write> Writer<'a, W> {
///
/// Globals have different naming schemes depending on their binding:
/// - Globals without bindings use the name from the [`Namer`](crate::proc::Namer)
/// - Globals with builtin bindings get the from [`glsl_built_in`](glsl_built_in)
/// - Globals with location bindings are named `_location_X` where `X` is the location
/// - Globals with resource binding are named `_group_X_binding_Y` where `X`
/// is the group and `Y` is the binding
fn get_global_name(&self, handle: Handle<GlobalVariable>, global: &GlobalVariable) -> String {
match global.binding {
Some(Binding::Location(location)) => {
format!(
"_location_{}{}",
location,
match (self.options.shader_stage, global.class) {
(ShaderStage::Fragment, StorageClass::Input) => "_vs",
(ShaderStage::Vertex, StorageClass::Output) => "_vs",
_ => "",
}
)
Some(ref br) => {
format!("_group_{}_binding_{}", br.group, br.binding)
}
Some(Binding::Resource { group, binding }) => {
format!("_group_{}_binding_{}", group, binding)
}
Some(Binding::BuiltIn(built_in)) => glsl_built_in(built_in).to_string(),
None => self.names[&NameKey::GlobalVariable(handle)].clone(),
}
}
/// Writes the varying declaration.
fn write_varying(
&mut self,
binding: Option<&Binding>,
ty: Handle<Type>,
output: bool,
) -> Result<(), Error> {
match self.module.types[ty].inner {
crate::TypeInner::Struct {
block: _,
ref members,
} => {
for member in members {
self.write_varying(member.binding.as_ref(), member.ty, output)?;
}
}
_ => {
let binding = match binding {
Some(binding @ &Binding::Location(..)) => binding,
_ => return Ok(()),
};
// Write the interpolation modifier if needed
//
// We ignore all interpolation modifiers that aren't used in input globals in fragment
// shaders or output globals in vertex shaders
//
// TODO: Should this throw an error?
if let Binding::Location(_, Some(interp)) = *binding {
if self.options.shader_stage == ShaderStage::Fragment {
write!(self.out, "{} ", glsl_interpolation(interp))?;
}
}
// Write the storage class
write!(self.out, "{} ", if output { "out" } else { "in" })?;
// Write the type
// `write_type` adds no leading or trailing spaces
self.write_type(ty)?;
// Finally write the global name and end the global with a `;` and a newline
// Leading space is important
writeln!(self.out, " {};", VaryingName { binding, output })?;
}
}
Ok(())
}
/// Helper method used to write functions (both entry points and regular functions)
///
/// # Notes
@@ -828,8 +873,10 @@ impl<'a, W: Write> Writer<'a, W> {
// Start by writing the return type if any otherwise write void
// This is the only place where `void` is a valid type
// (though it's more a keyword than a type)
if let Some(ty) = func.return_type {
self.write_type(ty)?;
if let FunctionType::EntryPoint(_) = ctx.func {
write!(self.out, "void")?;
} else if let Some(ref result) = func.result {
self.write_type(result.ty)?;
} else {
write!(self.out, "void")?;
}
@@ -841,14 +888,18 @@ impl<'a, W: Write> Writer<'a, W> {
//
// We need access to `Self` here so we use the reference passed to the closure as an
// argument instead of capturing as that would cause a borrow checker error
self.write_slice(&func.arguments, |this, i, arg| {
let arguments = match ctx.func {
FunctionType::EntryPoint(_) => &[][..],
FunctionType::Function(_) => &func.arguments,
};
self.write_slice(arguments, |this, i, arg| {
// Write the argument type
// `write_type` adds no trailing spaces
this.write_type(arg.ty)?;
// Write the argument name
// The leading space is important
write!(this.out, " {}", ctx.get_arg(i, &this.names))?;
write!(this.out, " {}", &this.names[&ctx.argument_key(i)])?;
Ok(())
})?;
@@ -856,6 +907,44 @@ impl<'a, W: Write> Writer<'a, W> {
// Close the parentheses and open braces to start the function body
writeln!(self.out, ") {{")?;
// Compose the function arguments from globals, in case of an entry point.
if let FunctionType::EntryPoint(ep_index) = ctx.func {
for (index, arg) in func.arguments.iter().enumerate() {
write!(self.out, "{}", INDENT)?;
self.write_type(arg.ty)?;
let name = &self.names[&NameKey::EntryPointArgument(ep_index, index as u32)];
write!(self.out, " {}", name)?;
write!(self.out, " = ")?;
match self.module.types[arg.ty].inner {
crate::TypeInner::Struct {
block: _,
ref members,
} => {
self.write_type(arg.ty)?;
write!(self.out, "(")?;
for (index, member) in members.iter().enumerate() {
let varying_name = VaryingName {
binding: member.binding.as_ref().unwrap(),
output: false,
};
if index != 0 {
write!(self.out, ", ")?;
}
write!(self.out, "{}", varying_name)?;
}
writeln!(self.out, ");")?;
}
_ => {
let varying_name = VaryingName {
binding: arg.binding.as_ref().unwrap(),
output: false,
};
writeln!(self.out, "{};", varying_name)?;
}
}
}
}
// Write all function locals
// Locals are `type name (= init)?;` where the init part (including the =) are optional
//
@@ -1228,13 +1317,54 @@ impl<'a, W: Write> Writer<'a, W> {
// `return expr;`, `expr` is optional
Statement::Return { value } => {
write!(self.out, "{}", INDENT.repeat(indent))?;
write!(self.out, "return")?;
// Write the expression to be returned if needed
if let Some(expr) = value {
write!(self.out, " ")?;
self.write_expr(expr, ctx)?;
match ctx.func {
FunctionType::Function(_) => {
write!(self.out, "return")?;
// Write the expression to be returned if needed
if let Some(expr) = value {
write!(self.out, " ")?;
self.write_expr(expr, ctx)?;
}
writeln!(self.out, ";")?;
}
FunctionType::EntryPoint(ep_index) => {
if let Some(ref result) =
self.module.entry_points[ep_index as usize].function.result
{
let value = value.unwrap();
match self.module.types[result.ty].inner {
crate::TypeInner::Struct {
block: _,
ref members,
} => {
for (index, member) in members.iter().enumerate() {
let varying_name = VaryingName {
binding: member.binding.as_ref().unwrap(),
output: true,
};
write!(self.out, "{} = ", varying_name)?;
self.write_expr(value, ctx)?;
let field_name = &self.names
[&NameKey::StructMember(result.ty, index as u32)];
writeln!(self.out, ".{};", field_name)?;
write!(self.out, "{}", INDENT.repeat(indent))?;
}
}
_ => {
let name = VaryingName {
binding: result.binding.as_ref().unwrap(),
output: true,
};
write!(self.out, "{} = ", name)?;
self.write_expr(value, ctx)?;
writeln!(self.out, ";")?;
write!(self.out, "{}", INDENT.repeat(indent))?;
}
}
}
writeln!(self.out, "return;")?;
}
}
writeln!(self.out, ";")?;
}
// This is one of the places were glsl adds to the syntax of C in this case the discard
// keyword which ceases all further processing in a fragment shader, it's called OpKill
@@ -1282,8 +1412,8 @@ impl<'a, W: Write> Writer<'a, W> {
write!(self.out, "{}", INDENT.repeat(indent))?;
if let Some(expr) = result {
let name = format!("_expr{}", expr.index());
let ty = self.module.functions[function].return_type.unwrap();
self.write_type(ty)?;
let result = self.module.functions[function].result.as_ref().unwrap();
self.write_type(result.ty)?;
write!(self.out, " {} = ", name)?;
self.cached_expressions.insert(expr, name);
}
@@ -1362,7 +1492,7 @@ impl<'a, W: Write> Writer<'a, W> {
}
// Function arguments are written as the argument name
Expression::FunctionArgument(pos) => {
write!(self.out, "{}", ctx.get_arg(pos, &self.names))?
write!(self.out, "{}", &self.names[&ctx.argument_key(pos)])?
}
// Global variables need some special work for their name but
// `get_global_name` does the work for us
@@ -2018,12 +2148,12 @@ fn glsl_built_in(built_in: BuiltIn) -> &'static str {
match built_in {
// vertex
BuiltIn::Position => "gl_Position",
BuiltIn::BaseInstance => "gl_BaseInstance",
BuiltIn::BaseVertex => "gl_BaseVertex",
BuiltIn::BaseInstance => "uint(gl_BaseInstance)",
BuiltIn::BaseVertex => "uint(gl_BaseVertex)",
BuiltIn::ClipDistance => "gl_ClipDistance",
BuiltIn::InstanceIndex => "gl_InstanceID",
BuiltIn::InstanceIndex => "uint(gl_InstanceID)",
BuiltIn::PointSize => "gl_PointSize",
BuiltIn::VertexIndex => "gl_VertexID",
BuiltIn::VertexIndex => "uint(gl_VertexID)",
// fragment
BuiltIn::FragCoord => "gl_FragCoord",
BuiltIn::FragDepth => "gl_FragDepth",
@@ -2044,8 +2174,6 @@ fn glsl_built_in(built_in: BuiltIn) -> &'static str {
fn glsl_storage_class(class: StorageClass) -> &'static str {
match class {
StorageClass::Function => "",
StorageClass::Input => "in",
StorageClass::Output => "out",
StorageClass::Private => "",
StorageClass::Storage => "buffer",
StorageClass::Uniform => "uniform",
@@ -2056,18 +2184,14 @@ fn glsl_storage_class(class: StorageClass) -> &'static str {
}
/// Helper function that returns the string corresponding to the glsl interpolation qualifier
///
/// # Errors
/// If [`Patch`](crate::Interpolation::Patch) is passed, as it isn't supported in glsl
fn glsl_interpolation(interpolation: Interpolation) -> Result<&'static str, Error> {
Ok(match interpolation {
fn glsl_interpolation(interpolation: Interpolation) -> &'static str {
match interpolation {
Interpolation::Perspective => "smooth",
Interpolation::Linear => "noperspective",
Interpolation::Flat => "flat",
Interpolation::Centroid => "centroid",
Interpolation::Sample => "sample",
Interpolation::Patch => return Err(Error::PatchInterpolationNotSupported),
})
}
}
/// Helper function that returns the glsl dimension string of [`ImageDimension`](crate::ImageDimension)

67
src/back/msl/mod.rs
View File

@@ -7,11 +7,20 @@ from SPIR-V's descriptor sets, we require a separate mapping provided in the opt
This mapping may have one or more resource end points for each descriptor set + index
pair.
## Outputs
## Entry points
In Metal, built-in shader outputs can not be nested into structures within
the output struct. If there is a structure in the outputs, and it contains any built-ins,
we move them up to the root output structure that we define ourselves.
Even though MSL and our IR appear to be similar in that the entry points in both can
accept arguments and return values, the restrictions are different.
MSL allows the varyings to be either in separate arguments, or inside a single
`[[stage_in]]` struct. We gather input varyings and form this artificial structure.
We also add all the (non-Private) globals into the arguments.
At the beginning of the entry point, we assign the local constants and re-compose
the arguments as they are declared on IR side, so that the rest of the logic can
pretend that MSL doesn't have all the restrictions it has.
For the result type, if it's a structure, we re-compose it with a temporary value
holding the result.
!*/
use crate::{
@@ -112,15 +121,14 @@ impl Default for Options {
}
impl Options {
fn resolve_binding(
fn resolve_local_binding(
&self,
stage: crate::ShaderStage,
var: &crate::GlobalVariable,
binding: &crate::Binding,
mode: LocationMode,
) -> Result<ResolvedBinding, Error> {
match var.binding {
Some(crate::Binding::BuiltIn(built_in)) => Ok(ResolvedBinding::BuiltIn(built_in)),
Some(crate::Binding::Location(index)) => match mode {
match *binding {
crate::Binding::BuiltIn(built_in) => Ok(ResolvedBinding::BuiltIn(built_in)),
crate::Binding::Location(index, _) => match mode {
LocationMode::VertexInput => Ok(ResolvedBinding::Attribute(index)),
LocationMode::FragmentOutput => Ok(ResolvedBinding::Color(index)),
LocationMode::Intermediate => Ok(ResolvedBinding::User {
@@ -139,25 +147,26 @@ impl Options {
Err(Error::Validation)
}
},
Some(crate::Binding::Resource { group, binding }) => {
let source = BindSource {
stage,
group,
binding,
};
match self.binding_map.get(&source) {
Some(target) => Ok(ResolvedBinding::Resource(target.clone())),
None if self.fake_missing_bindings => Ok(ResolvedBinding::User {
prefix: "fake",
index: 0,
}),
None => Err(Error::MissingBindTarget(source)),
}
}
None => {
log::error!("Missing binding for {:?}", var.name);
Err(Error::Validation)
}
}
}
fn resolve_global_binding(
&self,
stage: crate::ShaderStage,
res_binding: &crate::ResourceBinding,
) -> Result<ResolvedBinding, Error> {
let source = BindSource {
stage,
group: res_binding.group,
binding: res_binding.binding,
};
match self.binding_map.get(&source) {
Some(target) => Ok(ResolvedBinding::Resource(target.clone())),
None if self.fake_missing_bindings => Ok(ResolvedBinding::User {
prefix: "fake",
index: 0,
}),
None => Err(Error::MissingBindTarget(source)),
}
}
}

388
src/back/msl/writer.rs
View File

@@ -92,15 +92,13 @@ fn vector_size_string(size: crate::VectorSize) -> &'static str {
}
}
const OUTPUT_STRUCT_NAME: &str = "output";
const LOCATION_INPUT_STRUCT_NAME: &str = "input";
const COMPONENTS: &[char] = &['x', 'y', 'z', 'w'];
fn separate(is_last: bool) -> &'static str {
if is_last {
""
} else {
fn separate(need_separator: bool) -> &'static str {
if need_separator {
","
} else {
""
}
}
@@ -111,8 +109,7 @@ impl crate::StorageClass {
/// called from the entry point.
fn needs_pass_through(&self) -> bool {
match *self {
crate::StorageClass::Input
| crate::StorageClass::Uniform
crate::StorageClass::Uniform
| crate::StorageClass::Storage
| crate::StorageClass::Handle => true,
_ => false,
@@ -121,7 +118,7 @@ impl crate::StorageClass {
fn get_name(&self, global_use: GlobalUse) -> Option<&'static str> {
match *self {
Self::Input | Self::Output | Self::Handle => None,
Self::Handle => None,
Self::Uniform => Some("constant"),
//TODO: should still be "constant" for read-only buffers
Self::Storage => Some(if global_use.contains(GlobalUse::WRITE) {
@@ -149,7 +146,7 @@ struct ExpressionContext<'a> {
struct StatementContext<'a> {
expression: ExpressionContext<'a>,
fun_info: &'a FunctionInfo,
return_value: Option<&'a str>,
result_struct: Option<&'a str>,
}
impl<W: Write> Writer<W> {
@@ -287,30 +284,16 @@ impl<W: Write> Writer<W> {
}
}
crate::Expression::FunctionArgument(index) => {
let fun_handle = match context.origin {
FunctionOrigin::Handle(handle) => handle,
FunctionOrigin::EntryPoint(_) => unreachable!(),
let name_key = match context.origin {
FunctionOrigin::Handle(handle) => NameKey::FunctionArgument(handle, index),
FunctionOrigin::EntryPoint(ep_index) => {
NameKey::EntryPointArgument(ep_index, index)
}
};
let name = &self.names[&NameKey::FunctionArgument(fun_handle, index)];
let name = &self.names[&name_key];
write!(self.out, "{}", name)?;
}
crate::Expression::GlobalVariable(handle) => {
let var = &context.module.global_variables[handle];
match var.class {
crate::StorageClass::Output => {
if let crate::TypeInner::Struct { .. } = context.module.types[var.ty].inner
{
return Ok(());
}
write!(self.out, "{}.", OUTPUT_STRUCT_NAME)?;
}
crate::StorageClass::Input => {
if let Some(crate::Binding::Location(_)) = var.binding {
write!(self.out, "{}.", LOCATION_INPUT_STRUCT_NAME)?;
}
}
_ => {}
}
let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, "{}", name)?;
}
@@ -771,17 +754,42 @@ impl<W: Write> Writer<W> {
crate::Statement::Return {
value: Some(expr_handle),
} => {
write!(self.out, "{}return ", level)?;
self.put_expression(expr_handle, &context.expression)?;
match context.result_struct {
Some(struct_name) => {
let result_ty = context.expression.function.result.as_ref().unwrap().ty;
match context.expression.module.types[result_ty].inner {
crate::TypeInner::Struct {
block: _,
ref members,
} => {
let tmp = "_tmp";
write!(self.out, "{}const auto {} = ", level, tmp)?;
self.put_expression(expr_handle, &context.expression)?;
writeln!(self.out, ";")?;
write!(self.out, "{}return {} {{", level, struct_name)?;
for index in 0..members.len() as u32 {
let comma = if index == 0 { "" } else { "," };
let name =
&self.names[&NameKey::StructMember(result_ty, index)];
write!(self.out, "{} {}.{}", comma, tmp, name)?;
}
}
_ => {
write!(self.out, "{}return {} {{ ", level, struct_name)?;
self.put_expression(expr_handle, &context.expression)?;
}
}
write!(self.out, " }}")?;
}
None => {
write!(self.out, "{}return ", level)?;
self.put_expression(expr_handle, &context.expression)?;
}
}
writeln!(self.out, ";")?;
}
crate::Statement::Return { value: None } => {
writeln!(
self.out,
"{}return {};",
level,
context.return_value.unwrap_or_default(),
)?;
writeln!(self.out, "{}return;", level)?;
}
crate::Statement::Kill => {
writeln!(self.out, "{}{}::discard_fragment();", level, NAMESPACE)?;
@@ -1107,8 +1115,8 @@ impl<W: Write> Writer<W> {
}
let fun_name = &self.names[&NameKey::Function(fun_handle)];
let result_type_name = match fun.return_type {
Some(ret_ty) => &self.names[&NameKey::Type(ret_ty)],
let result_type_name = match fun.result {
Some(ref result) => &self.names[&NameKey::Type(result.ty)],
None => "void",
};
writeln!(self.out, "{} {}(", result_type_name, fun_name)?;
@@ -1117,7 +1125,7 @@ impl<W: Write> Writer<W> {
let name = &self.names[&NameKey::FunctionArgument(fun_handle, index as u32)];
let param_type_name = &self.names[&NameKey::Type(arg.ty)];
let separator =
separate(pass_through_globals.is_empty() && index + 1 == fun.arguments.len());
separate(!pass_through_globals.is_empty() || index + 1 != fun.arguments.len());
writeln!(
self.out,
"{}{} {}{}",
@@ -1131,7 +1139,7 @@ impl<W: Write> Writer<W> {
handle,
usage: fun_info[handle],
};
let separator = separate(index + 1 == pass_through_globals.len());
let separator = separate(index + 1 != pass_through_globals.len());
write!(self.out, "{}", INDENT)?;
tyvar.try_fmt(&mut self.out)?;
writeln!(self.out, "{}", separator)?;
@@ -1157,7 +1165,7 @@ impl<W: Write> Writer<W> {
analysis,
},
fun_info,
return_value: None,
result_struct: None,
};
self.named_expressions.clear();
self.put_block(Level(1), &fun.body, &context)?;
@@ -1183,27 +1191,11 @@ impl<W: Write> Writer<W> {
},
)?;
// find the entry point(s) and inputs/outputs
let mut last_used_global = None;
for (handle, var) in module.global_variables.iter() {
match var.class {
crate::StorageClass::Input => {
if let Some(crate::Binding::Location(_)) = var.binding {
continue;
}
}
crate::StorageClass::Output => continue,
_ => {}
}
if !fun_info[handle].is_empty() {
last_used_global = Some(handle);
}
}
let fun_name = &self.names[&NameKey::EntryPoint(ep_index as _)];
info.entry_point_names.push(fun_name.clone());
let output_name = format!("{}Output", fun_name);
let location_input_name = format!("{}Input", fun_name);
let stage_out_name = format!("{}Output", fun_name);
let stage_in_name = format!("{}Input", fun_name);
let (em_str, in_mode, out_mode) = match ep.stage {
crate::ShaderStage::Vertex => (
@@ -1221,110 +1213,135 @@ impl<W: Write> Writer<W> {
}
};
let return_value = match ep.stage {
crate::ShaderStage::Vertex | crate::ShaderStage::Fragment => {
// make dedicated input/output structs
writeln!(self.out, "struct {} {{", location_input_name)?;
for (handle, var) in module.global_variables.iter() {
if var.class != crate::StorageClass::Input
|| !fun_info[handle].contains(GlobalUse::READ)
{
continue;
let mut argument_members = Vec::new();
for (arg_index, arg) in fun.arguments.iter().enumerate() {
match module.types[arg.ty].inner {
crate::TypeInner::Struct {
block: _,
ref members,
} => {
for (member_index, member) in members.iter().enumerate() {
argument_members.push((
NameKey::StructMember(arg.ty, member_index as u32),
member.ty,
member.binding.as_ref(),
))
}
if let Some(crate::Binding::BuiltIn(_)) = var.binding {
// MSL disallows built-ins in input structs
continue;
}
_ => argument_members.push((
NameKey::EntryPointArgument(ep_index as _, arg_index as u32),
arg.ty,
arg.binding.as_ref(),
)),
}
}
let varyings_member_name = self.namer.call("varyings");
let mut varying_count = 0;
if !argument_members.is_empty() {
writeln!(self.out, "struct {} {{", stage_in_name)?;
for &(ref name_key, ty, binding) in argument_members.iter() {
let binding = match binding {
Some(ref binding @ &crate::Binding::Location(..)) => binding,
_ => continue,
};
varying_count += 1;
let name = &self.names[&name_key];
let type_name = &self.names[&NameKey::Type(ty)];
let resolved = options.resolve_local_binding(binding, in_mode)?;
write!(self.out, "{}{} {}", INDENT, type_name, name)?;
resolved.try_fmt_decorated(&mut self.out, "")?;
writeln!(self.out, ";")?;
}
writeln!(self.out, "}};")?;
}
let result_member_name = self.namer.call("member");
let result_type_name = match fun.result {
Some(ref result) => {
let mut result_members = Vec::new();
if let crate::TypeInner::Struct {
block: _,
ref members,
} = module.types[result.ty].inner
{
for (member_index, member) in members.iter().enumerate() {
result_members.push((
&self.names[&NameKey::StructMember(result.ty, member_index as u32)],
member.ty,
member.binding.as_ref(),
));
}
let tyvar = TypedGlobalVariable {
module,
names: &self.names,
handle,
usage: GlobalUse::empty(),
};
write!(self.out, "{}", INDENT)?;
tyvar.try_fmt(&mut self.out)?;
let resolved = options.resolve_binding(ep.stage, var, in_mode)?;
resolved.try_fmt_decorated(&mut self.out, ";")?;
writeln!(self.out)?;
} else {
result_members.push((
&result_member_name,
result.ty,
result.binding.as_ref(),
));
}
writeln!(self.out, "struct {} {{", stage_out_name)?;
for (name, ty, binding) in result_members {
let type_name = &self.names[&NameKey::Type(ty)];
let binding = binding.ok_or(Error::Validation)?;
let resolved = options.resolve_local_binding(binding, out_mode)?;
write!(self.out, "{}{} {}", INDENT, type_name, name)?;
resolved.try_fmt_decorated(&mut self.out, "")?;
writeln!(self.out, ";")?;
}
writeln!(self.out, "}};")?;
writeln!(self.out)?;
writeln!(self.out, "struct {} {{", output_name)?;
for (handle, var) in module.global_variables.iter() {
if var.class != crate::StorageClass::Output
|| !fun_info[handle].contains(GlobalUse::WRITE)
{
continue;
}
let tyvar = TypedGlobalVariable {
module,
names: &self.names,
handle,
usage: GlobalUse::empty(),
};
write!(self.out, "{}", INDENT)?;
tyvar.try_fmt(&mut self.out)?;
let resolved = options.resolve_binding(ep.stage, var, out_mode)?;
resolved.try_fmt_decorated(&mut self.out, ";")?;
writeln!(self.out)?;
}
writeln!(self.out, "}};")?;
writeln!(self.out)?;
writeln!(self.out, "{} {} {}(", em_str, output_name, fun_name)?;
let separator = separate(last_used_global.is_none());
writeln!(
self.out,
"{}{} {} [[stage_in]]{}",
INDENT, location_input_name, LOCATION_INPUT_STRUCT_NAME, separator
)?;
Some(OUTPUT_STRUCT_NAME)
}
crate::ShaderStage::Compute => {
writeln!(self.out, "{} void {}(", em_str, fun_name)?;
None
&stage_out_name
}
None => "void",
};
writeln!(self.out, "{} {} {}(", em_str, result_type_name, fun_name)?;
let mut is_first_argument = true;
if varying_count != 0 {
writeln!(
self.out,
" {} {} [[stage_in]]",
stage_in_name, varyings_member_name
)?;
is_first_argument = false;
}
for &(ref name_key, ty, binding) in argument_members.iter() {
let binding = match binding {
Some(ref binding @ &crate::Binding::BuiltIn(..)) => binding,
_ => continue,
};
let name = &self.names[&name_key];
let type_name = &self.names[&NameKey::Type(ty)];
let resolved = options.resolve_local_binding(binding, in_mode)?;
let separator = if is_first_argument {
is_first_argument = false;
' '
} else {
','
};
write!(self.out, "{} {} {}", separator, type_name, name)?;
resolved.try_fmt_decorated(&mut self.out, "\n")?;
}
for (handle, var) in module.global_variables.iter() {
let usage = fun_info[handle];
if usage.is_empty() || var.class == crate::StorageClass::Output {
if usage.is_empty() || var.class == crate::StorageClass::Private {
continue;
}
if var.class == crate::StorageClass::Input {
if let Some(crate::Binding::Location(_)) = var.binding {
// location inputs are put into a separate struct
continue;
}
}
let loc_mode = match (ep.stage, var.class) {
(crate::ShaderStage::Vertex, crate::StorageClass::Input) => {
LocationMode::VertexInput
}
(crate::ShaderStage::Vertex, crate::StorageClass::Output)
| (crate::ShaderStage::Fragment { .. }, crate::StorageClass::Input) => {
LocationMode::Intermediate
}
(crate::ShaderStage::Fragment { .. }, crate::StorageClass::Output) => {
LocationMode::FragmentOutput
}
_ => LocationMode::Uniform,
};
let tyvar = TypedGlobalVariable {
module,
names: &self.names,
handle,
usage,
};
let separator = separate(last_used_global == Some(handle));
write!(self.out, "{}", INDENT)?;
let separator = if is_first_argument {
is_first_argument = false;
' '
} else {
','
};
write!(self.out, "{} ", separator)?;
tyvar.try_fmt(&mut self.out)?;
if var.binding.is_some() {
let resolved = options.resolve_binding(ep.stage, var, loc_mode)?;
resolved.try_fmt_decorated(&mut self.out, separator)?;
if let Some(ref binding) = var.binding {
let resolved = options.resolve_global_binding(ep.stage, binding)?;
resolved.try_fmt_decorated(&mut self.out, "")?;
}
if let Some(value) = var.init {
let value_str = &self.names[&NameKey::Constant(value)];
@@ -1332,18 +1349,72 @@ impl<W: Write> Writer<W> {
}
writeln!(self.out)?;
}
// end of the entry point argument list
writeln!(self.out, ") {{")?;
match ep.stage {
crate::ShaderStage::Vertex | crate::ShaderStage::Fragment => {
writeln!(
self.out,
"{}{} {};",
INDENT, output_name, OUTPUT_STRUCT_NAME
)?;
// Metal doesn't support private mutable variables outside of functions,
// so we put them here, just like the locals.
for (handle, var) in module.global_variables.iter() {
let usage = fun_info[handle];
if usage.is_empty() || var.class != crate::StorageClass::Private {
continue;
}
crate::ShaderStage::Compute => {}
let tyvar = TypedGlobalVariable {
module,
names: &self.names,
handle,
usage,
};
write!(self.out, "{}", INDENT)?;
tyvar.try_fmt(&mut self.out)?;
if let Some(value) = var.init {
let value_str = &self.names[&NameKey::Constant(value)];
write!(self.out, " = {}", value_str)?;
}
writeln!(self.out, ";")?;
}
// Now refactor the inputs in a way that the rest of the code expects
for (arg_index, arg) in fun.arguments.iter().enumerate() {
let arg_name =
&self.names[&NameKey::EntryPointArgument(ep_index as _, arg_index as u32)];
match module.types[arg.ty].inner {
crate::TypeInner::Struct {
block: _,
ref members,
} => {
let struct_name = &self.names[&NameKey::Type(arg.ty)];
write!(
self.out,
"{}const {} {} = {{",
INDENT, struct_name, arg_name
)?;
for member_index in 0..members.len() {
let name =
&self.names[&NameKey::StructMember(arg.ty, member_index as u32)];
let separator = if member_index != 0 { ", " } else { "" };
write!(self.out, "{}{}", INDENT, separator)?;
if let Some(crate::Binding::Location(..)) = arg.binding {
write!(self.out, "{}.", varyings_member_name)?;
}
write!(self.out, "{}", name)?;
}
}
_ => {
if let Some(crate::Binding::Location(..)) = arg.binding {
writeln!(
self.out,
"{}const auto {} = {}.{};",
INDENT, arg_name, varyings_member_name, arg_name
)?;
}
}
}
}
// Finally, declare all the local variables that we need
//TODO: we can postpone this till the relevant expressions are emitted
for (local_handle, local) in fun.local_variables.iter() {
let name = &self.names[&NameKey::EntryPointLocal(ep_index as _, local_handle)];
let ty_name = &self.names[&NameKey::Type(local.ty)];
@@ -1363,13 +1434,12 @@ impl<W: Write> Writer<W> {
analysis,
},
fun_info,
return_value,
result_struct: Some(&stage_out_name),
};
self.named_expressions.clear();
self.put_block(Level(1), &fun.body, &context)?;
writeln!(self.out, "}}")?;
let is_last = ep_index == module.entry_points.len() - 1;
if !is_last {
if ep_index + 1 != module.entry_points.len() {
writeln!(self.out)?;
}
}

16
src/back/spv/layout.rs
View File

@@ -27,6 +27,22 @@ impl PhysicalLayout {
pub(super) fn supports_storage_buffers(&self) -> bool {
self.version >= 0x10300
}
pub(super) fn map_storage_class(&self, class: crate::StorageClass) -> spirv::StorageClass {
match class {
crate::StorageClass::Handle => spirv::StorageClass::UniformConstant,
crate::StorageClass::Function => spirv::StorageClass::Function,
crate::StorageClass::Private => spirv::StorageClass::Private,
crate::StorageClass::Storage if self.supports_storage_buffers() => {
spirv::StorageClass::StorageBuffer
}
crate::StorageClass::Storage | crate::StorageClass::Uniform => {
spirv::StorageClass::Uniform
}
crate::StorageClass::WorkGroup => spirv::StorageClass::Workgroup,
crate::StorageClass::PushConstant => spirv::StorageClass::PushConstant,
}
}
}
impl LogicalLayout {

428
src/back/spv/writer.rs
View File

@@ -46,12 +46,19 @@ struct LocalVariable {
instruction: Instruction,
}
#[derive(Default)]
struct EntryPointContext {
argument_ids: Vec<Word>,
result_ids_typed: Vec<(Word, Word)>,
}
#[derive(Default)]
struct Function {
signature: Option<Instruction>,
parameters: Vec<Instruction>,
variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
blocks: Vec<Block>,
entry_point_context: Option<EntryPointContext>,
}
impl Function {
@@ -86,7 +93,7 @@ enum LocalType {
vector_size: Option<crate::VectorSize>,
kind: crate::ScalarKind,
width: crate::Bytes,
pointer_class: Option<crate::StorageClass>,
pointer_class: Option<spirv::StorageClass>,
},
Matrix {
columns: crate::VectorSize,
@@ -95,15 +102,15 @@ enum LocalType {
},
Pointer {
base: Handle<crate::Type>,
class: crate::StorageClass,
class: spirv::StorageClass,
},
SampledImage {
image_type: Handle<crate::Type>,
},
}
impl LocalType {
fn from_inner(inner: &crate::TypeInner) -> Option<Self> {
impl PhysicalLayout {
fn make_local(&self, inner: &crate::TypeInner) -> Option<LocalType> {
Some(match *inner {
crate::TypeInner::Scalar { kind, width } => LocalType::Value {
vector_size: None,
@@ -126,7 +133,10 @@ impl LocalType {
rows,
width,
},
crate::TypeInner::Pointer { base, class } => LocalType::Pointer { base, class },
crate::TypeInner::Pointer { base, class } => LocalType::Pointer {
base,
class: self.map_storage_class(class),
},
crate::TypeInner::ValuePointer {
size,
kind,
@@ -136,7 +146,7 @@ impl LocalType {
vector_size: size,
kind,
width,
pointer_class: Some(class),
pointer_class: Some(self.map_storage_class(class)),
},
_ => return None,
})
@@ -318,10 +328,12 @@ impl Writer {
Ok(*e.get())
} else {
match lookup_ty {
LookupType::Handle(handle) => match LocalType::from_inner(&arena[handle].inner) {
Some(local) => self.get_type_id(arena, LookupType::Local(local)),
None => self.write_type_declaration_arena(arena, handle),
},
LookupType::Handle(handle) => {
match self.physical_layout.make_local(&arena[handle].inner) {
Some(local) => self.get_type_id(arena, LookupType::Local(local)),
None => self.write_type_declaration_arena(arena, handle),
}
}
LookupType::Local(local_ty) => self.write_type_declaration_local(arena, local_ty),
}
}
@@ -331,36 +343,25 @@ impl Writer {
&mut self,
arena: &Arena<crate::Type>,
handle: Handle<crate::Type>,
class: crate::StorageClass,
class: spirv::StorageClass,
) -> Result<Word, Error> {
let ty_id = self.get_type_id(arena, LookupType::Handle(handle))?;
if let crate::TypeInner::Pointer { .. } = arena[handle].inner {
return Ok(ty_id);
}
Ok(
match self
.lookup_type
.entry(LookupType::Local(LocalType::Pointer {
base: handle,
class,
})) {
Entry::Occupied(e) => *e.get(),
_ => {
let storage_class = self.parse_to_spirv_storage_class(class);
let id = self.generate_id();
let instruction = Instruction::type_pointer(id, storage_class, ty_id);
instruction.to_words(&mut self.logical_layout.declarations);
self.lookup_type.insert(
LookupType::Local(LocalType::Pointer {
base: handle,
class,
}),
id,
);
id
}
},
)
let lookup_type = LookupType::Local(LocalType::Pointer {
base: handle,
class,
});
Ok(if let Some(&id) = self.lookup_type.get(&lookup_type) {
id
} else {
let id = self.generate_id();
let instruction = Instruction::type_pointer(id, class, ty_id);
instruction.to_words(&mut self.logical_layout.declarations);
self.lookup_type.insert(lookup_type, id);
id
})
}
fn create_constant(&mut self, type_id: Word, value: &[Word]) -> Word {
@@ -375,6 +376,7 @@ impl Writer {
ir_function: &crate::Function,
info: &FunctionInfo,
ir_module: &crate::Module,
mut varying_ids: Option<&mut Vec<Word>>,
) -> Result<Word, Error> {
let mut function = Function::default();
@@ -403,7 +405,7 @@ impl Writer {
.init
.map(|constant| self.lookup_constant[&constant]);
let pointer_type_id =
self.get_pointer_id(&ir_module.types, variable.ty, crate::StorageClass::Function)?;
self.get_pointer_id(&ir_module.types, variable.ty, spirv::StorageClass::Function)?;
let instruction = Instruction::variable(
pointer_type_id,
id,
@@ -415,21 +417,101 @@ impl Writer {
.insert(handle, LocalVariable { id, instruction });
}
let return_type_id = match ir_function.return_type {
Some(handle) => self.get_type_id(&ir_module.types, LookupType::Handle(handle))?,
let prelude_id = self.generate_id();
let mut prelude = Block::new(prelude_id);
let mut ep_context = EntryPointContext::default();
let mut parameter_type_ids = Vec::with_capacity(ir_function.arguments.len());
for argument in ir_function.arguments.iter() {
let class = spirv::StorageClass::Input;
let argument_type_id =
self.get_type_id(&ir_module.types, LookupType::Handle(argument.ty))?;
if let Some(ref mut list) = varying_ids {
let id = if let Some(ref binding) = argument.binding {
let name = argument.name.as_ref().map(AsRef::as_ref);
let varying_id =
self.write_varying(ir_module, class, name, argument.ty, binding)?;
list.push(varying_id);
let id = self.generate_id();
prelude
.body
.push(Instruction::load(argument_type_id, id, varying_id, None));
id
} else if let crate::TypeInner::Struct {
block: _,
ref members,
} = ir_module.types[argument.ty].inner
{
let struct_id = self.generate_id();
let mut constituent_ids = Vec::with_capacity(members.len());
for member in members {
let type_id =
self.get_type_id(&ir_module.types, LookupType::Handle(member.ty))?;
let name = member.name.as_ref().map(AsRef::as_ref);
let binding = member.binding.as_ref().unwrap();
let varying_id =
self.write_varying(ir_module, class, name, member.ty, binding)?;
list.push(varying_id);
let id = self.generate_id();
prelude
.body
.push(Instruction::load(type_id, id, varying_id, None));
constituent_ids.push(id);
}
prelude.body.push(Instruction::composite_construct(
argument_type_id,
struct_id,
&constituent_ids,
));
struct_id
} else {
unreachable!("Missing argument binding on an entry point");
};
ep_context.argument_ids.push(id);
} else {
let id = self.generate_id();
let instruction = Instruction::function_parameter(argument_type_id, id);
function.parameters.push(instruction);
parameter_type_ids.push(argument_type_id);
};
}
let return_type_id = match ir_function.result {
Some(ref result) => {
if let Some(ref mut list) = varying_ids {
let class = spirv::StorageClass::Output;
if let Some(ref binding) = result.binding {
let type_id =
self.get_type_id(&ir_module.types, LookupType::Handle(result.ty))?;
let varying_id =
self.write_varying(ir_module, class, None, result.ty, binding)?;
list.push(varying_id);
ep_context.result_ids_typed.push((varying_id, type_id));
} else if let crate::TypeInner::Struct {
block: _,
ref members,
} = ir_module.types[result.ty].inner
{
for member in members {
let type_id =
self.get_type_id(&ir_module.types, LookupType::Handle(member.ty))?;
let name = member.name.as_ref().map(AsRef::as_ref);
let binding = member.binding.as_ref().unwrap();
let varying_id =
self.write_varying(ir_module, class, name, member.ty, binding)?;
list.push(varying_id);
ep_context.result_ids_typed.push((varying_id, type_id));
}
} else {
unreachable!("Missing result binding on an entry point");
}
self.void_type
} else {
self.get_type_id(&ir_module.types, LookupType::Handle(result.ty))?
}
}
None => self.void_type,
};
let mut parameter_type_ids = Vec::with_capacity(ir_function.arguments.len());
for argument in ir_function.arguments.iter() {
let id = self.generate_id();
let parameter_type_id =
self.get_type_id(&ir_module.types, LookupType::Handle(argument.ty))?;
parameter_type_ids.push(parameter_type_id);
function
.parameters
.push(Instruction::function_parameter(parameter_type_id, id));
}
let lookup_function_type = LookupFunctionType {
return_type_id,
@@ -451,8 +533,9 @@ impl Writer {
function_type,
));
let prelude_id = self.generate_id();
let mut prelude = Block::new(prelude_id);
if varying_ids.is_some() {
function.entry_point_context = Some(ep_context);
}
// fill up the `GlobalVariable::handle_id`
for gv in self.global_variables.iter_mut() {
@@ -510,18 +593,13 @@ impl Writer {
info: &FunctionInfo,
ir_module: &crate::Module,
) -> Result<Instruction, Error> {
let function_id = self.write_function(&entry_point.function, info, ir_module)?;
let mut interface_ids = Vec::new();
for (handle, var) in ir_module.global_variables.iter() {
if info[handle].is_empty() {
continue;
}
if let crate::StorageClass::Input | crate::StorageClass::Output = var.class {
let id = self.global_variables[handle.index()].id;
interface_ids.push(id);
}
}
let function_id = self.write_function(
&entry_point.function,
info,
ir_module,
Some(&mut interface_ids),
)?;
let exec_model = match entry_point.stage {
crate::ShaderStage::Vertex => spirv::ExecutionModel::Vertex,
@@ -573,24 +651,6 @@ impl Writer {
}
}
fn parse_to_spirv_storage_class(&self, class: crate::StorageClass) -> spirv::StorageClass {
match class {
crate::StorageClass::Handle => spirv::StorageClass::UniformConstant,
crate::StorageClass::Function => spirv::StorageClass::Function,
crate::StorageClass::Input => spirv::StorageClass::Input,
crate::StorageClass::Output => spirv::StorageClass::Output,
crate::StorageClass::Private => spirv::StorageClass::Private,
crate::StorageClass::Storage if self.physical_layout.supports_storage_buffers() => {
spirv::StorageClass::StorageBuffer
}
crate::StorageClass::Storage | crate::StorageClass::Uniform => {
spirv::StorageClass::Uniform
}
crate::StorageClass::WorkGroup => spirv::StorageClass::Workgroup,
crate::StorageClass::PushConstant => spirv::StorageClass::PushConstant,
}
}
fn write_type_declaration_local(
&mut self,
arena: &Arena<crate::Type>,
@@ -639,7 +699,7 @@ impl Writer {
}
LocalType::Pointer { base, class } => {
let type_id = self.get_type_id(arena, LookupType::Handle(base))?;
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
Instruction::type_pointer(id, class, type_id)
}
LocalType::Value {
vector_size,
@@ -656,7 +716,7 @@ impl Writer {
pointer_class: None,
}),
)?;
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
Instruction::type_pointer(id, class, type_id)
}
LocalType::SampledImage { image_type } => {
let image_type_id = self.get_type_id(arena, LookupType::Handle(image_type))?;
@@ -860,9 +920,15 @@ impl Writer {
}
crate::TypeInner::Pointer { base, class } => {
let type_id = self.get_type_id(arena, LookupType::Handle(base))?;
self.lookup_type
.insert(LookupType::Local(LocalType::Pointer { base, class }), id);
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
let raw_class = self.physical_layout.map_storage_class(class);
self.lookup_type.insert(
LookupType::Local(LocalType::Pointer {
base,
class: raw_class,
}),
id,
);
Instruction::type_pointer(id, raw_class, type_id)
}
crate::TypeInner::ValuePointer {
size,
@@ -870,6 +936,7 @@ impl Writer {
width,
class,
} => {
let raw_class = self.physical_layout.map_storage_class(class);
let type_id = self.get_type_id(
arena,
LookupType::Local(LocalType::Value {
@@ -884,11 +951,11 @@ impl Writer {
vector_size: size,
kind,
width,
pointer_class: Some(class),
pointer_class: Some(raw_class),
}),
id,
);
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
Instruction::type_pointer(id, raw_class, type_id)
}
};
@@ -979,76 +1046,45 @@ impl Writer {
Ok(())
}
fn write_global_variable(
fn write_varying(
&mut self,
ir_module: &crate::Module,
handle: Handle<crate::GlobalVariable>,
) -> Result<(Instruction, Word, spirv::StorageClass), Error> {
let global_variable = &ir_module.global_variables[handle];
class: spirv::StorageClass,
debug_name: Option<&str>,
ty: Handle<crate::Type>,
binding: &crate::Binding,
) -> Result<Word, Error> {
let id = self.generate_id();
let class = self.parse_to_spirv_storage_class(global_variable.class);
self.check(class.required_capabilities())?;
let init_word = global_variable
.init
.map(|constant| self.lookup_constant[&constant]);
let pointer_type_id =
self.get_pointer_id(&ir_module.types, global_variable.ty, global_variable.class)?;
let instruction = Instruction::variable(pointer_type_id, id, class, init_word);
let pointer_type_id = self.get_pointer_id(&ir_module.types, ty, class)?;
Instruction::variable(pointer_type_id, id, class, None)
.to_words(&mut self.logical_layout.declarations);
if self.flags.contains(WriterFlags::DEBUG) {
if let Some(ref name) = global_variable.name {
if let Some(name) = debug_name {
self.debugs.push(Instruction::name(id, name));
}
}
let access_decoration = match global_variable.storage_access {
crate::StorageAccess::LOAD => Some(spirv::Decoration::NonWritable),
crate::StorageAccess::STORE => Some(spirv::Decoration::NonReadable),
_ => None,
};
if let Some(decoration) = access_decoration {
self.annotations
.push(Instruction::decorate(id, decoration, &[]));
}
if let Some(interpolation) = global_variable.interpolation {
let decoration = match interpolation {
crate::Interpolation::Linear => Some(spirv::Decoration::NoPerspective),
crate::Interpolation::Flat => Some(spirv::Decoration::Flat),
crate::Interpolation::Patch => Some(spirv::Decoration::Patch),
crate::Interpolation::Centroid => Some(spirv::Decoration::Centroid),
crate::Interpolation::Sample => Some(spirv::Decoration::Sample),
crate::Interpolation::Perspective => None,
};
if let Some(decoration) = decoration {
self.annotations
.push(Instruction::decorate(id, decoration, &[]));
}
}
match global_variable.binding {
Some(crate::Binding::Location(location)) => {
match *binding {
crate::Binding::Location(location, interpolation) => {
self.annotations.push(Instruction::decorate(
id,
spirv::Decoration::Location,
&[location],
));
let interp_decoration = match interpolation {
Some(crate::Interpolation::Linear) => Some(spirv::Decoration::NoPerspective),
Some(crate::Interpolation::Flat) => Some(spirv::Decoration::Flat),
Some(crate::Interpolation::Centroid) => Some(spirv::Decoration::Centroid),
Some(crate::Interpolation::Sample) => Some(spirv::Decoration::Sample),
Some(crate::Interpolation::Perspective) | None => None,
};
if let Some(decoration) = interp_decoration {
self.annotations
.push(Instruction::decorate(id, decoration, &[]));
}
}
Some(crate::Binding::Resource { group, binding }) => {
self.annotations.push(Instruction::decorate(
id,
spirv::Decoration::DescriptorSet,
&[group],
));
self.annotations.push(Instruction::decorate(
id,
spirv::Decoration::Binding,
&[binding],
));
}
Some(crate::Binding::BuiltIn(built_in)) => {
crate::Binding::BuiltIn(built_in) => {
use crate::BuiltIn as Bi;
let built_in = match built_in {
Bi::BaseInstance => spirv::BuiltIn::BaseInstance,
@@ -1079,7 +1115,57 @@ impl Writer {
&[built_in as u32],
));
}
None => {}
}
Ok(id)
}
fn write_global_variable(
&mut self,
ir_module: &crate::Module,
handle: Handle<crate::GlobalVariable>,
) -> Result<(Instruction, Word, spirv::StorageClass), Error> {
let global_variable = &ir_module.global_variables[handle];
let id = self.generate_id();
let class = self
.physical_layout
.map_storage_class(global_variable.class);
self.check(class.required_capabilities())?;
let init_word = global_variable
.init
.map(|constant| self.lookup_constant[&constant]);
let pointer_type_id = self.get_pointer_id(&ir_module.types, global_variable.ty, class)?;
let instruction = Instruction::variable(pointer_type_id, id, class, init_word);
if self.flags.contains(WriterFlags::DEBUG) {
if let Some(ref name) = global_variable.name {
self.debugs.push(Instruction::name(id, name));
}
}
let access_decoration = match global_variable.storage_access {
crate::StorageAccess::LOAD => Some(spirv::Decoration::NonWritable),
crate::StorageAccess::STORE => Some(spirv::Decoration::NonReadable),
_ => None,
};
if let Some(decoration) = access_decoration {
self.annotations
.push(Instruction::decorate(id, decoration, &[]));
}
if let Some(ref res_binding) = global_variable.binding {
self.annotations.push(Instruction::decorate(
id,
spirv::Decoration::DescriptorSet,
&[res_binding.group],
));
self.annotations.push(Instruction::decorate(
id,
spirv::Decoration::Binding,
&[res_binding.binding],
));
}
// TODO Initializer is optional and not (yet) included in the IR
@@ -1212,7 +1298,7 @@ impl Writer {
) -> Result<(), Error> {
let result_lookup_ty = match self.typifier.get_handle(expr_handle) {
Ok(ty_handle) => LookupType::Handle(ty_handle),
Err(inner) => LookupType::Local(LocalType::from_inner(inner).unwrap()),
Err(inner) => LookupType::Local(self.physical_layout.make_local(inner).unwrap()),
};
let result_type_id = self.get_type_id(&ir_module.types, result_lookup_ty)?;
@@ -1598,9 +1684,10 @@ impl Writer {
.push(Instruction::load(result_type_id, id, pointer_id, None));
id
}
crate::Expression::FunctionArgument(index) => {
function.parameters[index as usize].result_id.unwrap()
}
crate::Expression::FunctionArgument(index) => match function.entry_point_context {
Some(ref context) => context.argument_ids[index as usize],
None => function.parameters[index as usize].result_id.unwrap(),
},
crate::Expression::Call(_function) => self.lookup_function_call[&expr_handle],
crate::Expression::As {
expr,
@@ -1829,7 +1916,7 @@ impl Writer {
) -> Result<(Word, spirv::StorageClass), Error> {
let result_lookup_ty = match self.typifier.get_handle(expr_handle) {
Ok(ty_handle) => LookupType::Handle(ty_handle),
Err(inner) => LookupType::Local(LocalType::from_inner(inner).unwrap()),
Err(inner) => LookupType::Local(self.physical_layout.make_local(inner).unwrap()),
};
let result_type_id = self.get_type_id(&ir_module.types, result_lookup_ty)?;
@@ -2120,8 +2207,40 @@ impl Writer {
Some(Instruction::branch(loop_context.continuing_id.unwrap()));
}
crate::Statement::Return { value: Some(value) } => {
let id = self.cached[value];
block.termination = Some(Instruction::return_value(id));
let value_id = self.cached[value];
let instruction = match function.entry_point_context {
// If this is an entry point, and we need to return anything,
// let's instead store the output variables and return `void`.
Some(ref context) => {
let result = ir_function.result.as_ref().unwrap();
if result.binding.is_none() {
for (index, &(varying_id, type_id)) in
context.result_ids_typed.iter().enumerate()
{
let member_value_id = self.generate_id();
block.body.push(Instruction::composite_extract(
type_id,
member_value_id,
value_id,
&[index as u32],
));
block.body.push(Instruction::store(
varying_id,
member_value_id,
None,
));
}
} else {
let (varying_id, _) = context.result_ids_typed[0];
block
.body
.push(Instruction::store(varying_id, value_id, None));
};
Instruction::return_void()
}
None => Instruction::return_value(value_id),
};
block.termination = Some(instruction);
}
crate::Statement::Return { value: None } => {
block.termination = Some(Instruction::return_void());
@@ -2180,8 +2299,11 @@ impl Writer {
Some(expr) => {
self.cached[expr] = id;
self.lookup_function_call.insert(expr, id);
let ty_handle =
ir_module.functions[local_function].return_type.unwrap();
let ty_handle = ir_module.functions[local_function]
.result
.as_ref()
.unwrap()
.ty;
self.get_type_id(&ir_module.types, LookupType::Handle(ty_handle))?
}
None => self.void_type,
@@ -2253,7 +2375,7 @@ impl Writer {
for (handle, ir_function) in ir_module.functions.iter() {
let info = &analysis[handle];
let id = self.write_function(ir_function, info, ir_module)?;
let id = self.write_function(ir_function, info, ir_module, None)?;
self.lookup_function.insert(handle, id);
}

8
src/front/glsl/ast.rs
View File

@@ -2,8 +2,8 @@ use super::{constants::ConstantSolver, error::ErrorKind};
use crate::{
proc::{ResolveContext, Typifier},
Arena, BinaryOperator, Binding, Constant, Expression, FastHashMap, Function, FunctionArgument,
GlobalVariable, Handle, Interpolation, LocalVariable, Module, RelationalFunction, ShaderStage,
Statement, StorageClass, Type, UnaryOperator,
GlobalVariable, Handle, Interpolation, LocalVariable, Module, RelationalFunction,
ResourceBinding, ShaderStage, Statement, StorageClass, Type, UnaryOperator,
};
#[derive(Debug)]
@@ -224,6 +224,7 @@ impl ExpressionRule {
#[derive(Debug)]
pub enum TypeQualifier {
StorageQualifier(StorageQualifier),
ResourceBinding(ResourceBinding),
Binding(Binding),
Interpolation(Interpolation),
}
@@ -250,11 +251,14 @@ pub struct FunctionCall {
#[derive(Debug, Clone, Copy)]
pub enum StorageQualifier {
StorageClass(StorageClass),
Input,
Output,
Const,
}
#[derive(Debug, Clone)]
pub enum StructLayout {
Binding(Binding),
Resource(ResourceBinding),
PushConstant,
}

63
src/front/glsl/parser.rs
View File

@@ -8,10 +8,11 @@ pomelo! {
BOOL_WIDTH,
Arena, BinaryOperator, Binding, Block, Constant,
ConstantInner, Expression,
Function, GlobalVariable, Handle, Interpolation,
LocalVariable, ScalarValue, ScalarKind,
Function, FunctionArgument, FunctionResult,
GlobalVariable, Handle, Interpolation,
LocalVariable, ResourceBinding, ScalarValue, ScalarKind,
Statement, StorageAccess, StorageClass, StructMember,
SwitchCase, Type, TypeInner, UnaryOperator, FunctionArgument,
SwitchCase, Type, TypeInner, UnaryOperator,
};
use pp_rs::token::PreprocessorError;
}
@@ -602,14 +603,15 @@ pomelo! {
layout_qualifier ::= Layout LeftParen layout_qualifier_id_list(l) RightParen {
if let Some(&(_, loc)) = l.iter().find(|&q| q.0.as_str() == "location") {
StructLayout::Binding(Binding::Location(loc))
let interpolation = None; //TODO
StructLayout::Binding(Binding::Location(loc, interpolation))
} else if let Some(&(_, binding)) = l.iter().find(|&q| q.0.as_str() == "binding") {
let group = if let Some(&(_, set)) = l.iter().find(|&q| q.0.as_str() == "set") {
set
} else {
0
};
StructLayout::Binding(Binding::Resource{ group, binding })
StructLayout::Resource(ResourceBinding{ group, binding })
} else if l.iter().any(|q| q.0.as_str() == "push_constant") {
StructLayout::PushConstant
} else {
@@ -648,6 +650,7 @@ pomelo! {
single_type_qualifier ::= layout_qualifier(l) {
match l {
StructLayout::Binding(b) => TypeQualifier::Binding(b),
StructLayout::Resource(b) => TypeQualifier::ResourceBinding(b),
StructLayout::PushConstant => TypeQualifier::StorageQualifier(StorageQualifier::StorageClass(StorageClass::PushConstant)),
}
}
@@ -663,10 +666,10 @@ pomelo! {
}
// storage_qualifier ::= InOut;
storage_qualifier ::= In {
StorageQualifier::StorageClass(StorageClass::Input)
StorageQualifier::Input
}
storage_qualifier ::= Out {
StorageQualifier::StorageClass(StorageClass::Output)
StorageQualifier::Output
}
// storage_qualifier ::= Centroid;
// storage_qualifier ::= Patch;
@@ -724,6 +727,7 @@ pomelo! {
name: Some(name.clone()),
span: None,
ty,
binding: None, //TODO
}).collect()
} else {
return Err(ErrorKind::SemanticError("Struct member can't be void".into()))
@@ -1005,7 +1009,7 @@ pomelo! {
Function {
name: Some(n.1),
arguments: vec![],
return_type: t.1,
result: t.1.map(|ty| FunctionResult { ty, binding: None }),
local_variables: Arena::<LocalVariable>::new(),
expressions: Arena::<Expression>::new(),
body: vec![],
@@ -1019,12 +1023,12 @@ pomelo! {
(h, args)
}
parameter_declarator ::= parameter_type_specifier(ty) Identifier(n) {
FunctionArgument { name: Some(n.1), ty }
FunctionArgument { name: Some(n.1), ty, binding: None }
}
// parameter_declarator ::= type_specifier(ty) Identifier(ident) array_specifier;
parameter_declaration ::= parameter_declarator;
parameter_declaration ::= parameter_type_specifier(ty) {
FunctionArgument { name: None, ty }
FunctionArgument { name: None, ty, binding: None }
}
parameter_type_specifier ::= type_specifier(t) {
@@ -1077,13 +1081,8 @@ pomelo! {
StorageQualifier::StorageClass(storage_class) => {
// TODO: Check that the storage qualifiers allow for the bindings
let binding = d.type_qualifiers.iter().find_map(|tq| {
if let TypeQualifier::Binding(ref b) = *tq { Some(b.clone()) } else { None }
if let TypeQualifier::ResourceBinding(ref b) = *tq { Some(b.clone()) } else { None }
});
let interpolation = d.type_qualifiers.iter().find_map(|tq| {
if let TypeQualifier::Interpolation(interp) = *tq { Some(interp) } else { None }
});
for (id, initializer) in d.ids_initializers {
let init = initializer.map(|init| extra.solve_constant(init.expression)).transpose()?;
@@ -1104,7 +1103,6 @@ pomelo! {
binding: binding.clone(),
ty: d.ty,
init,
interpolation,
storage_access: StorageAccess::empty(), //TODO
},
);
@@ -1113,6 +1111,37 @@ pomelo! {
}
}
}
StorageQualifier::Input => {
let mut binding = d.type_qualifiers.iter().find_map(|tq| {
if let TypeQualifier::Binding(ref b) = *tq { Some(b.clone()) } else { None }
});
let interpolation = d.type_qualifiers.iter().find_map(|tq| {
if let TypeQualifier::Interpolation(interp) = *tq { Some(interp) } else { None }
});
if let Some(Binding::Location(_, ref mut interp)) = binding {
*interp = interpolation;
}
for (id, _initializer) in d.ids_initializers {
if let Some(id) = id {
//TODO!
let expr = extra.context.expressions.append(Expression::FunctionArgument(0));
extra.context.lookup_global_var_exps.insert(id, expr);
}
}
}
StorageQualifier::Output => {
let _binding = d.type_qualifiers.iter().find_map(|tq| {
if let TypeQualifier::Binding(ref b) = *tq { Some(b.clone()) } else { None }
});
for (id, _initializer) in d.ids_initializers {
if let Some(id) = id {
//TODO!
let expr = extra.context.expressions.append(Expression::FunctionArgument(0));
extra.context.lookup_global_var_exps.insert(id, expr);
}
}
}
StorageQualifier::Const => {
for (id, initializer) in d.ids_initializers {
if let Some(init) = initializer {

25
src/front/glsl/variables.rs
View File

@@ -1,7 +1,4 @@
use crate::{
Binding, BuiltIn, Expression, GlobalVariable, Handle, ScalarKind, StorageAccess, StorageClass,
Type, TypeInner, VectorSize,
};
use crate::{Expression, Handle, Type, TypeInner, VectorSize};
use super::ast::*;
use super::error::ErrorKind;
@@ -20,7 +17,7 @@ impl Program<'_> {
}
match name {
"gl_Position" => {
let h = self.module.global_variables.append(GlobalVariable {
/*let h = self.module.global_variables.append(GlobalVariable {
name: Some(name.into()),
class: StorageClass::Output,
binding: Some(Binding::BuiltIn(BuiltIn::Position)),
@@ -40,12 +37,17 @@ impl Program<'_> {
let exp = self
.context
.expressions
.append(Expression::GlobalVariable(h));
.append(Expression::GlobalVariable(h));*/
let exp = self
.context
.expressions
.append(Expression::FunctionArgument(0)); //TODO
self.context.lookup_global_var_exps.insert(name.into(), exp);
Ok(Some(exp))
}
"gl_VertexIndex" => {
/* TODO
let h = self.module.global_variables.append(GlobalVariable {
name: Some(name.into()),
class: StorageClass::Input,
@@ -71,6 +73,11 @@ impl Program<'_> {
kind: ScalarKind::Sint,
convert: true,
});
*/
let expr = self
.context
.expressions
.append(Expression::FunctionArgument(0)); //TODO
self.context
.lookup_global_var_exps
.insert(name.into(), expr);
@@ -78,6 +85,7 @@ impl Program<'_> {
Ok(Some(expr))
}
"gl_InstanceIndex" => {
/* TODO
let h = self.module.global_variables.append(GlobalVariable {
name: Some(name.into()),
class: StorageClass::Input,
@@ -103,6 +111,11 @@ impl Program<'_> {
kind: ScalarKind::Sint,
convert: true,
});
*/
let expr = self
.context
.expressions
.append(Expression::FunctionArgument(0)); //TODO
self.context
.lookup_global_var_exps
.insert(name.into(), expr);

35
src/front/spv/convert.rs
View File

@@ -2,7 +2,7 @@ use super::error::Error;
use num_traits::cast::FromPrimitive;
use std::convert::TryInto;
pub fn map_binary_operator(word: spirv::Op) -> Result<crate::BinaryOperator, Error> {
pub(super) fn map_binary_operator(word: spirv::Op) -> Result<crate::BinaryOperator, Error> {
use crate::BinaryOperator;
use spirv::Op;
@@ -34,7 +34,7 @@ pub fn map_binary_operator(word: spirv::Op) -> Result<crate::BinaryOperator, Err
}
}
pub fn map_relational_fun(word: spirv::Op) -> Result<crate::RelationalFunction, Error> {
pub(super) fn map_relational_fun(word: spirv::Op) -> Result<crate::RelationalFunction, Error> {
use crate::RelationalFunction as Rf;
use spirv::Op;
@@ -49,7 +49,7 @@ pub fn map_relational_fun(word: spirv::Op) -> Result<crate::RelationalFunction,
}
}
pub fn map_vector_size(word: spirv::Word) -> Result<crate::VectorSize, Error> {
pub(super) fn map_vector_size(word: spirv::Word) -> Result<crate::VectorSize, Error> {
match word {
2 => Ok(crate::VectorSize::Bi),
3 => Ok(crate::VectorSize::Tri),
@@ -58,7 +58,7 @@ pub fn map_vector_size(word: spirv::Word) -> Result<crate::VectorSize, Error> {
}
}
pub fn map_image_dim(word: spirv::Word) -> Result<crate::ImageDimension, Error> {
pub(super) fn map_image_dim(word: spirv::Word) -> Result<crate::ImageDimension, Error> {
use spirv::Dim as D;
match D::from_u32(word) {
Some(D::Dim1D) => Ok(crate::ImageDimension::D1),
@@ -69,7 +69,7 @@ pub fn map_image_dim(word: spirv::Word) -> Result<crate::ImageDimension, Error>
}
}
pub fn map_image_format(word: spirv::Word) -> Result<crate::StorageFormat, Error> {
pub(super) fn map_image_format(word: spirv::Word) -> Result<crate::StorageFormat, Error> {
match spirv::ImageFormat::from_u32(word) {
Some(spirv::ImageFormat::R8) => Ok(crate::StorageFormat::R8Unorm),
Some(spirv::ImageFormat::R8Snorm) => Ok(crate::StorageFormat::R8Snorm),
@@ -107,13 +107,13 @@ pub fn map_image_format(word: spirv::Word) -> Result<crate::StorageFormat, Error
}
}
pub fn map_width(word: spirv::Word) -> Result<crate::Bytes, Error> {
pub(super) fn map_width(word: spirv::Word) -> Result<crate::Bytes, Error> {
(word >> 3) // bits to bytes
.try_into()
.map_err(|_| Error::InvalidTypeWidth(word))
}
pub fn map_builtin(word: spirv::Word, is_output: bool) -> Result<crate::BuiltIn, Error> {
pub(super) fn map_builtin(word: spirv::Word, is_output: bool) -> Result<crate::BuiltIn, Error> {
use spirv::BuiltIn as Bi;
Ok(match spirv::BuiltIn::from_u32(word) {
Some(Bi::BaseInstance) => crate::BuiltIn::BaseInstance,
@@ -145,19 +145,20 @@ pub fn map_builtin(word: spirv::Word, is_output: bool) -> Result<crate::BuiltIn,
})
}
pub fn map_storage_class(word: spirv::Word) -> Result<crate::StorageClass, Error> {
pub(super) fn map_storage_class(word: spirv::Word) -> Result<super::ExtendedClass, Error> {
use super::ExtendedClass as Ec;
use spirv::StorageClass as Sc;
Ok(match Sc::from_u32(word) {
Some(Sc::Function) => crate::StorageClass::Function,
Some(Sc::Input) => crate::StorageClass::Input,
Some(Sc::Output) => crate::StorageClass::Output,
Some(Sc::Private) => crate::StorageClass::Private,
Some(Sc::UniformConstant) => crate::StorageClass::Handle,
Some(Sc::StorageBuffer) => crate::StorageClass::Storage,
Some(Sc::Function) => Ec::Global(crate::StorageClass::Function),
Some(Sc::Input) => Ec::Input,
Some(Sc::Output) => Ec::Output,
Some(Sc::Private) => Ec::Global(crate::StorageClass::Private),
Some(Sc::UniformConstant) => Ec::Global(crate::StorageClass::Handle),
Some(Sc::StorageBuffer) => Ec::Global(crate::StorageClass::Storage),
// we expect the `Storage` case to be filtered out before calling this function.
Some(Sc::Uniform) => crate::StorageClass::Uniform,
Some(Sc::Workgroup) => crate::StorageClass::WorkGroup,
Some(Sc::PushConstant) => crate::StorageClass::PushConstant,
Some(Sc::Uniform) => Ec::Global(crate::StorageClass::Uniform),
Some(Sc::Workgroup) => Ec::Global(crate::StorageClass::WorkGroup),
Some(Sc::PushConstant) => Ec::Global(crate::StorageClass::PushConstant),
_ => return Err(Error::UnsupportedStorageClass(word)),
})
}

131
src/front/spv/function.rs
View File

@@ -63,10 +63,14 @@ impl<I: Iterator<Item = u32>> super::Parser<I> {
crate::Function {
name: self.future_decor.remove(&fun_id).and_then(|dec| dec.name),
arguments: Vec::with_capacity(ft.parameter_type_ids.len()),
return_type: if self.lookup_void_type == Some(result_type_id) {
result: if self.lookup_void_type == Some(result_type_id) {
None
} else {
Some(self.lookup_type.lookup(result_type_id)?.handle)
let lookup_result_ty = self.lookup_type.lookup(result_type_id)?;
Some(crate::FunctionResult {
ty: lookup_result_ty.handle,
binding: None,
})
},
local_variables: Arena::new(),
expressions: self.make_expression_storage(),
@@ -99,8 +103,12 @@ impl<I: Iterator<Item = u32>> super::Parser<I> {
return Err(Error::WrongFunctionArgumentType(type_id));
}
let ty = self.lookup_type.lookup(type_id)?.handle;
fun.arguments
.push(crate::FunctionArgument { name: None, ty });
let decor = self.future_decor.remove(&id).unwrap_or_default();
fun.arguments.push(crate::FunctionArgument {
name: decor.name,
ty,
binding: None,
});
}
Instruction { op, .. } => return Err(Error::InvalidParameter(op)),
}
@@ -157,19 +165,128 @@ impl<I: Iterator<Item = u32>> super::Parser<I> {
fun.body = flow_graph.to_naga()?;
// done
let fun_handle = module.functions.append(fun);
match self.lookup_entry_point.remove(&fun_id) {
Some(ep) => {
// create a wrapping function
let mut function = crate::Function {
name: None,
arguments: Vec::new(),
result: None,
local_variables: Arena::new(),
expressions: Arena::new(),
body: Vec::new(),
};
// 1. copy the inputs from arguments to privates
for &v_id in ep.variable_ids.iter() {
let lvar = self.lookup_variable.lookup(v_id)?;
if let super::Variable::Input(ref arg) = lvar.inner {
function.body.push(crate::Statement::Store {
pointer: function
.expressions
.append(crate::Expression::GlobalVariable(lvar.handle)),
value:
function
.expressions
.append(crate::Expression::FunctionArgument(
function.arguments.len() as u32,
)),
});
let mut arg = arg.clone();
if ep.stage == crate::ShaderStage::Fragment {
if let Some(crate::Binding::Location(_, ref mut interpolation @ None)) =
arg.binding
{
*interpolation = Some(crate::Interpolation::Perspective);
// default
}
}
function.arguments.push(arg);
}
}
// 2. call the wrapped function
function.body.push(crate::Statement::Call {
function: fun_handle,
arguments: Vec::new(),
result: None,
});
// 3. copy the outputs from privates to the result
let mut members = Vec::new();
let mut components = Vec::new();
for &v_id in ep.variable_ids.iter() {
let lvar = self.lookup_variable.lookup(v_id)?;
if let super::Variable::Output(ref result) = lvar.inner {
members.push(crate::StructMember {
name: None,
span: None,
ty: result.ty,
binding: result.binding.clone(),
});
// populate just the globals first, then do `Load` in a
// separate step, so that we can get a range.
components.push(
function
.expressions
.append(crate::Expression::GlobalVariable(lvar.handle)),
);
}
}
let old_len = function.expressions.len();
for component in components.iter_mut() {
*component = function.expressions.append(crate::Expression::Load {
pointer: *component,
});
}
match members.len() {
0 => {}
1 => {
let member = members.remove(0);
function.body.push(crate::Statement::Emit(
function.expressions.range_from(old_len),
));
function.body.push(crate::Statement::Return {
value: components.first().cloned(),
});
function.result = Some(crate::FunctionResult {
ty: member.ty,
binding: member.binding,
});
}
_ => {
let ty = module.types.append(crate::Type {
name: None,
inner: crate::TypeInner::Struct {
block: false,
members,
},
});
let result_expr = function
.expressions
.append(crate::Expression::Compose { ty, components });
function.body.push(crate::Statement::Emit(
function.expressions.range_from(old_len),
));
function.body.push(crate::Statement::Return {
value: Some(result_expr),
});
function.result = Some(crate::FunctionResult { ty, binding: None });
}
}
module.entry_points.push(crate::EntryPoint {
name: ep.name,
stage: ep.stage,
early_depth_test: ep.early_depth_test,
workgroup_size: ep.workgroup_size,
function: fun,
function,
});
}
None => {
let handle = module.functions.append(fun);
self.lookup_function.insert(fun_id, handle);
self.lookup_function.insert(fun_id, fun_handle);
}
};

236
src/front/spv/mod.rs
View File

@@ -8,6 +8,14 @@ There map `spv::Word` into a specific IR handle, plus potentially a bit of
extra info, such as the related SPIR-V type ID.
TODO: would be nice to find ways that avoid looking up as much
## Inputs/Outputs
We create a private variable for each input/output. The relevant inputs are
populated at the start of an entry point. The outputs are saved at the end.
The function associated with an entry point is wrapped in another function,
such that we can handle any `Return` statements without problems.
!*/
#![allow(dead_code)]
@@ -203,37 +211,31 @@ impl Decoration {
}
}
fn get_binding(&self, is_output: bool) -> Option<crate::Binding> {
//TODO: validate this better
fn resource_binding(&self) -> Option<crate::ResourceBinding> {
match *self {
Decoration {
desc_set: Some(group),
desc_index: Some(binding),
..
} => Some(crate::ResourceBinding { group, binding }),
_ => None,
}
}
fn io_binding(&self, is_output: bool) -> Result<crate::Binding, Error> {
match *self {
Decoration {
built_in: Some(built_in),
location: None,
desc_set: None,
desc_index: None,
..
} => match map_builtin(built_in, is_output) {
Ok(built_in) => Some(crate::Binding::BuiltIn(built_in)),
Err(e) => {
log::warn!("{:?}", e);
None
}
},
} => map_builtin(built_in, is_output).map(crate::Binding::BuiltIn),
Decoration {
built_in: None,
location: Some(loc),
desc_set: None,
desc_index: None,
interpolation,
..
} => Some(crate::Binding::Location(loc)),
Decoration {
built_in: None,
location: None,
desc_set: Some(group),
desc_index: Some(binding),
..
} => Some(crate::Binding::Resource { group, binding }),
_ => None,
} => Ok(crate::Binding::Location(loc, interpolation)),
_ => Err(Error::MissingDecoration(spirv::Decoration::Location)),
}
}
}
@@ -266,8 +268,16 @@ struct LookupConstant {
type_id: spirv::Word,
}
#[derive(Debug)]
enum Variable {
Global,
Input(crate::FunctionArgument),
Output(crate::FunctionResult),
}
#[derive(Debug)]
struct LookupVariable {
inner: Variable,
handle: Handle<crate::GlobalVariable>,
type_id: spirv::Word,
}
@@ -279,11 +289,17 @@ struct LookupExpression {
}
#[derive(Clone, Debug)]
pub struct Assignment {
struct Assignment {
to: Handle<crate::Expression>,
value: Handle<crate::Expression>,
}
enum ExtendedClass {
Global(crate::StorageClass),
Input,
Output,
}
#[derive(Clone, Debug, Default)]
pub struct Options {
pub flow_graph_dump_prefix: Option<PathBuf>,
@@ -429,7 +445,7 @@ impl<I: Iterator<Item = u32>> Parser<I> {
dec.interpolation = Some(crate::Interpolation::Flat);
}
spirv::Decoration::Patch => {
dec.interpolation = Some(crate::Interpolation::Patch);
// skip
}
spirv::Decoration::Centroid => {
dec.interpolation = Some(crate::Interpolation::Centroid);
@@ -1706,6 +1722,9 @@ impl<I: Iterator<Item = u32>> Parser<I> {
self.index_constants.push(handle);
}
self.dummy_functions = Arena::new();
self.lookup_function.clear();
loop {
use spirv::Op;
@@ -1764,10 +1783,8 @@ impl<I: Iterator<Item = u32>> Parser<I> {
for (_, fun) in module.functions.iter_mut() {
self.patch_function_calls(fun)?;
}
for ep in module.entry_points.iter_mut() {
self.patch_function_calls(&mut ep.function)?;
}
self.lookup_function.clear();
// Note: we aren't patching the entry point functions, because they are simply
// wrappers behind real functions, and are already resolved.
// Check all the images and samplers to have consistent comparison property.
for (handle, flags) in self.handle_sampling.drain() {
@@ -2215,7 +2232,10 @@ impl<I: Iterator<Item = u32>> Parser<I> {
{
crate::StorageClass::Storage
}
_ => map_storage_class(storage_class)?,
_ => match map_storage_class(storage_class)? {
ExtendedClass::Global(class) => class,
ExtendedClass::Input | ExtendedClass::Output => crate::StorageClass::Private,
},
};
// Don't bother with pointer stuff for `Handle` types.
@@ -2325,6 +2345,7 @@ impl<I: Iterator<Item = u32>> Parser<I> {
name: decor.name,
span: None, //TODO
ty,
binding: None,
});
}
@@ -2695,78 +2716,119 @@ impl<I: Iterator<Item = u32>> Parser<I> {
} => true,
_ => false,
};
let class = if self.lookup_storage_buffer_types.contains(&effective_ty) {
crate::StorageClass::Storage
let ext_class = if self.lookup_storage_buffer_types.contains(&effective_ty) {
ExtendedClass::Global(crate::StorageClass::Storage)
} else {
map_storage_class(storage_class)?
};
let storage_access = if is_storage {
let mut access = crate::StorageAccess::all();
if dec.flags.contains(DecorationFlags::NON_READABLE) {
access ^= crate::StorageAccess::LOAD;
let (inner, var) = match ext_class {
ExtendedClass::Global(class) => {
let storage_access = if is_storage {
let mut access = crate::StorageAccess::all();
if dec.flags.contains(DecorationFlags::NON_READABLE) {
access ^= crate::StorageAccess::LOAD;
}
if dec.flags.contains(DecorationFlags::NON_WRITABLE) {
access ^= crate::StorageAccess::STORE;
}
access
} else {
crate::StorageAccess::empty()
};
let var = crate::GlobalVariable {
binding: dec.resource_binding(),
name: dec.name,
class,
ty: effective_ty,
init,
storage_access,
};
(Variable::Global, var)
}
if dec.flags.contains(DecorationFlags::NON_WRITABLE) {
access ^= crate::StorageAccess::STORE;
ExtendedClass::Input => {
let binding = dec.io_binding(false)?;
if let crate::Binding::BuiltIn(built_in) = binding {
let needs_inner_uint = match built_in {
crate::BuiltIn::BaseInstance
| crate::BuiltIn::BaseVertex
| crate::BuiltIn::InstanceIndex
| crate::BuiltIn::SampleIndex
| crate::BuiltIn::VertexIndex
| crate::BuiltIn::LocalInvocationIndex => Some(crate::TypeInner::Scalar {
kind: crate::ScalarKind::Uint,
width: 4,
}),
crate::BuiltIn::GlobalInvocationId
| crate::BuiltIn::LocalInvocationId
| crate::BuiltIn::WorkGroupId
| crate::BuiltIn::WorkGroupSize => Some(crate::TypeInner::Vector {
size: crate::VectorSize::Tri,
kind: crate::ScalarKind::Uint,
width: 4,
}),
_ => None,
};
if let (Some(inner), Some(crate::ScalarKind::Sint)) = (
needs_inner_uint,
module.types[effective_ty].inner.scalar_kind(),
) {
log::warn!("Treating {:?} as unsigned", built_in);
effective_ty = module
.types
.fetch_or_append(crate::Type { name: None, inner });
}
}
let var = crate::GlobalVariable {
name: dec.name.clone(),
class: crate::StorageClass::Private,
binding: None,
ty: effective_ty,
init: None,
storage_access: crate::StorageAccess::empty(),
};
let inner = Variable::Input(crate::FunctionArgument {
name: dec.name,
ty: effective_ty,
binding: Some(binding),
});
(inner, var)
}
ExtendedClass::Output => {
let binding = dec.io_binding(true)?;
let var = crate::GlobalVariable {
name: dec.name,
class: crate::StorageClass::Private,
binding: None,
ty: effective_ty,
init: None,
storage_access: crate::StorageAccess::empty(),
};
let inner = Variable::Output(crate::FunctionResult {
ty: effective_ty,
binding: Some(binding),
});
(inner, var)
}
access
} else {
crate::StorageAccess::empty()
};
let binding = dec.get_binding(class == crate::StorageClass::Output);
if let Some(crate::Binding::BuiltIn(built_in)) = binding {
// SPIR-V only cares about some of the built-in types being integer.
// Naga requires them to be strictly unsigned, so we have to patch it.
let needs_inner_uint = match built_in {
crate::BuiltIn::BaseInstance
| crate::BuiltIn::BaseVertex
| crate::BuiltIn::InstanceIndex
| crate::BuiltIn::SampleIndex
| crate::BuiltIn::VertexIndex
| crate::BuiltIn::LocalInvocationIndex => Some(crate::TypeInner::Scalar {
kind: crate::ScalarKind::Uint,
width: 4,
}),
crate::BuiltIn::GlobalInvocationId
| crate::BuiltIn::LocalInvocationId
| crate::BuiltIn::WorkGroupId
| crate::BuiltIn::WorkGroupSize => Some(crate::TypeInner::Vector {
size: crate::VectorSize::Tri,
kind: crate::ScalarKind::Uint,
width: 4,
}),
_ => None,
};
if let (Some(inner), Some(crate::ScalarKind::Sint)) = (
needs_inner_uint,
module.types[effective_ty].inner.scalar_kind(),
) {
log::warn!("Treating {:?} as unsigned", built_in);
effective_ty = module
.types
.fetch_or_append(crate::Type { name: None, inner });
}
}
let var = crate::GlobalVariable {
name: dec.name,
class,
binding,
ty: effective_ty,
init,
interpolation: dec.interpolation,
storage_access,
};
let handle = module.global_variables.append(var);
self.lookup_variable
.insert(id, LookupVariable { handle, type_id });
if module.types[effective_ty].inner.can_comparison_sample() {
log::debug!("\t\ttracking {:?} for sampling properties", handle);
self.handle_sampling
.insert(handle, image::SamplingFlags::empty());
}
self.lookup_variable.insert(
id,
LookupVariable {
inner,
handle,
type_id,
},
);
Ok(())
}
}

2
src/front/wgsl/conv.rs
View File

@@ -2,8 +2,6 @@ use super::Error;
pub fn map_storage_class(word: &str) -> Result<crate::StorageClass, Error<'_>> {
match word {
"in" => Ok(crate::StorageClass::Input),
"out" => Ok(crate::StorageClass::Output),
"private" => Ok(crate::StorageClass::Private),
"uniform" => Ok(crate::StorageClass::Uniform),
"storage" => Ok(crate::StorageClass::Storage),

186
src/front/wgsl/mod.rs
View File

@@ -109,6 +109,12 @@ pub enum Error<'a> {
ZeroStride,
#[error("not a composite type: {0:?}")]
NotCompositeType(Handle<crate::Type>),
#[error("Input/output binding is not consistent: location {0:?}, built-in {1:?} and interpolation {2:?}")]
InconsistentBinding(
Option<u32>,
Option<crate::BuiltIn>,
Option<crate::Interpolation>,
),
#[error("call to local `{0}(..)` can't be resolved")]
UnknownLocalFunction(&'a str),
#[error("builtin {0:?} is not implemented")]
@@ -416,6 +422,54 @@ pub enum Scope {
type LocalFunctionCall = (Handle<crate::Function>, Vec<Handle<crate::Expression>>);
#[derive(Default)]
struct BindingParser {
location: Option<u32>,
built_in: Option<crate::BuiltIn>,
interpolation: Option<crate::Interpolation>,
}
impl BindingParser {
fn parse<'a>(&mut self, lexer: &mut Lexer<'a>, name: &'a str) -> Result<(), Error<'a>> {
match name {
"location" => {
lexer.expect(Token::Paren('('))?;
self.location = Some(lexer.next_uint_literal()?);
lexer.expect(Token::Paren(')'))?;
}
"builtin" => {
lexer.expect(Token::Paren('('))?;
let raw = lexer.next_ident()?;
self.built_in = Some(conv::map_built_in(raw)?);
lexer.expect(Token::Paren(')'))?;
}
"interpolate" => {
lexer.expect(Token::Paren('('))?;
let raw = lexer.next_ident()?;
self.interpolation = Some(conv::map_interpolation(raw)?);
lexer.expect(Token::Paren(')'))?;
}
_ => return Err(Error::UnknownDecoration(name)),
}
Ok(())
}
fn finish<'a>(self) -> Result<Option<crate::Binding>, Error<'a>> {
match (self.location, self.built_in, self.interpolation) {
(None, None, None) => Ok(None),
(Some(loc), None, interpolation) => {
Ok(Some(crate::Binding::Location(loc, interpolation)))
}
(None, Some(bi), None) => Ok(Some(crate::Binding::BuiltIn(bi))),
(location, built_in, interpolation) => Err(Error::InconsistentBinding(
location,
built_in,
interpolation,
)),
}
}
}
struct ParsedVariable<'a> {
name: &'a str,
class: Option<crate::StorageClass>,
@@ -1447,6 +1501,7 @@ impl Parser {
lexer.expect(Token::Paren('{'))?;
loop {
let mut span = 0;
let mut bind_parser = BindingParser::default();
if lexer.skip(Token::DoubleParen('[')) {
self.scopes.push(Scope::Decoration);
let mut ready = true;
@@ -1458,17 +1513,22 @@ impl Parser {
(Token::Separator(','), _) if !ready => {
ready = true;
}
(Token::Word("span"), _) if ready => {
lexer.expect(Token::Paren('('))?;
//Note: 0 is not handled
span = lexer.next_uint_literal()?;
lexer.expect(Token::Paren(')'))?;
ready = false;
}
(Token::Word("offset"), _) if ready => {
lexer.expect(Token::Paren('('))?;
let _offset = lexer.next_uint_literal()?;
lexer.expect(Token::Paren(')'))?;
(Token::Word(word), _) if ready => {
match word {
"span" => {
lexer.expect(Token::Paren('('))?;
//Note: 0 is not handled
span = lexer.next_uint_literal()?;
lexer.expect(Token::Paren(')'))?;
}
"offset" => {
// skip - only here for parsing compatibility
lexer.expect(Token::Paren('('))?;
let _offset = lexer.next_uint_literal()?;
lexer.expect(Token::Paren(')'))?;
}
_ => bind_parser.parse(lexer, word)?,
}
ready = false;
}
other => return Err(Error::Unexpected(other, "decoration separator")),
@@ -1476,6 +1536,7 @@ impl Parser {
}
self.scopes.pop();
}
let name = match lexer.next() {
(Token::Word(word), _) => word,
(Token::Paren('}'), _) => return Ok(members),
@@ -1484,10 +1545,12 @@ impl Parser {
lexer.expect(Token::Separator(':'))?;
let (ty, _access) = self.parse_type_decl(lexer, None, type_arena, const_arena)?;
lexer.expect(Token::Separator(';'))?;
members.push(crate::StructMember {
name: Some(name.to_owned()),
span: NonZeroU32::new(span),
ty,
binding: bind_parser.finish()?,
});
}
}
@@ -2259,6 +2322,31 @@ impl Parser {
Ok(block)
}
fn parse_varying_binding<'a>(
&mut self,
lexer: &mut Lexer<'a>,
) -> Result<Option<crate::Binding>, Error<'a>> {
if !lexer.skip(Token::DoubleParen('[')) {
return Ok(None);
}
let mut bind_parser = BindingParser::default();
self.scopes.push(Scope::Decoration);
loop {
let word = lexer.next_ident()?;
bind_parser.parse(lexer, word)?;
match lexer.next() {
(Token::DoubleParen(']'), _) => {
break;
}
(Token::Separator(','), _) => {}
other => return Err(Error::Unexpected(other, "decoration separator")),
}
}
self.scopes.pop();
bind_parser.finish()
}
fn parse_function_decl<'a>(
&mut self,
lexer: &mut Lexer<'a>,
@@ -2278,10 +2366,12 @@ impl Parser {
// read parameter list
let mut arguments = Vec::new();
lexer.expect(Token::Paren('('))?;
let mut ready = true;
while !lexer.skip(Token::Paren(')')) {
if !arguments.is_empty() {
lexer.expect(Token::Separator(','))?;
if !ready {
return Err(Error::Unexpected(lexer.next(), "comma"));
}
let binding = self.parse_varying_binding(lexer)?;
let (param_name, param_type, _access) =
self.parse_variable_ident_decl(lexer, &mut module.types, &mut module.constants)?;
let param_index = arguments.len() as u32;
@@ -2291,13 +2381,16 @@ impl Parser {
arguments.push(crate::FunctionArgument {
name: Some(param_name.to_string()),
ty: param_type,
binding,
});
ready = lexer.skip(Token::Separator(','));
}
// read return type
let return_type = if lexer.skip(Token::Arrow) && !lexer.skip(Token::Word("void")) {
let (handle, _access) =
let result = if lexer.skip(Token::Arrow) && !lexer.skip(Token::Word("void")) {
let binding = self.parse_varying_binding(lexer)?;
let (ty, _access) =
self.parse_type_decl(lexer, None, &mut module.types, &mut module.constants)?;
Some(handle)
Some(crate::FunctionResult { ty, binding })
} else {
None
};
@@ -2305,7 +2398,7 @@ impl Parser {
let mut fun = crate::Function {
name: Some(fun_name.to_string()),
arguments,
return_type,
result,
local_variables: Arena::new(),
expressions,
body: Vec::new(),
@@ -2345,7 +2438,6 @@ impl Parser {
// read decorations
let mut binding = None;
// Perspective is the default qualifier.
let mut interpolation = None;
let mut stage = None;
let mut is_block = false;
let mut workgroup_size = [0u32; 3];
@@ -2356,18 +2448,6 @@ impl Parser {
self.scopes.push(Scope::Decoration);
loop {
match lexer.next_ident()? {
"location" => {
lexer.expect(Token::Paren('('))?;
let loc = lexer.next_uint_literal()?;
lexer.expect(Token::Paren(')'))?;
binding = Some(crate::Binding::Location(loc));
}
"builtin" => {
lexer.expect(Token::Paren('('))?;
let builtin = conv::map_built_in(lexer.next_ident()?)?;
lexer.expect(Token::Paren(')'))?;
binding = Some(crate::Binding::BuiltIn(builtin));
}
"binding" => {
lexer.expect(Token::Paren('('))?;
bind_index = Some(lexer.next_uint_literal()?);
@@ -2381,11 +2461,6 @@ impl Parser {
bind_group = Some(lexer.next_uint_literal()?);
lexer.expect(Token::Paren(')'))?;
}
"interpolate" => {
lexer.expect(Token::Paren('('))?;
interpolation = Some(conv::map_interpolation(lexer.next_ident()?)?);
lexer.expect(Token::Paren(')'))?;
}
"stage" => {
lexer.expect(Token::Paren('('))?;
stage = Some(conv::map_shader_stage(lexer.next_ident()?)?);
@@ -2433,7 +2508,7 @@ impl Parser {
}
}
if let (Some(group), Some(index)) = (bind_group, bind_index) {
binding = Some(crate::Binding::Resource {
binding = Some(crate::ResourceBinding {
group,
binding: index,
});
@@ -2494,33 +2569,21 @@ impl Parser {
self.parse_variable_decl(lexer, &mut module.types, &mut module.constants)?;
let class = match pvar.class {
Some(c) => c,
None => match binding {
Some(crate::Binding::BuiltIn(builtin)) => match builtin {
crate::BuiltIn::GlobalInvocationId => crate::StorageClass::Input,
crate::BuiltIn::Position => crate::StorageClass::Output,
_ => return Err(Error::UnimplementedBuiltin(builtin)),
},
Some(crate::Binding::Resource { .. }) => {
match module.types[pvar.ty].inner {
crate::TypeInner::Struct { .. } if pvar.access.is_empty() => {
crate::StorageClass::Uniform
}
crate::TypeInner::Struct { .. }
| crate::TypeInner::Array { .. } => crate::StorageClass::Storage,
crate::TypeInner::Image { .. }
| crate::TypeInner::Sampler { .. } => crate::StorageClass::Handle,
ref other => {
log::error!("Resource type {:?}", other);
return Err(Error::InvalidResourceType(pvar.ty));
}
None => match module.types[pvar.ty].inner {
crate::TypeInner::Struct { .. } if binding.is_some() => {
if pvar.access.is_empty() {
crate::StorageClass::Uniform
} else {
crate::StorageClass::Storage
}
}
_ => match module.types[pvar.ty].inner {
crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } => {
crate::StorageClass::Handle
}
_ => crate::StorageClass::Private,
},
crate::TypeInner::Array { .. } if binding.is_some() => {
crate::StorageClass::Storage
}
crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } => {
crate::StorageClass::Handle
}
_ => crate::StorageClass::Private,
},
};
let var_handle = module.global_variables.append(crate::GlobalVariable {
@@ -2529,7 +2592,6 @@ impl Parser {
binding: binding.take(),
ty: pvar.ty,
init: pvar.init,
interpolation,
storage_access: pvar.access,
});
lookup_global_expression

54
src/lib.rs
View File

@@ -123,10 +123,6 @@ pub enum ShaderStage {
pub enum StorageClass {
/// Function locals.
Function,
/// Pipeline input, per invocation.
Input,
/// Pipeline output, per invocation, mutable.
Output,
/// Private data, per invocation, mutable.
Private,
/// Workgroup shared data, mutable.
@@ -228,8 +224,6 @@ pub enum Interpolation {
Linear,
/// Indicates that no interpolation will be performed.
Flat,
/// Indicates a tessellation patch.
Patch,
/// When used with multi-sampling rasterization, allow
/// a single interpolation location for an entire pixel.
Centroid,
@@ -247,6 +241,8 @@ pub struct StructMember {
pub name: Option<String>,
pub span: Option<NonZeroU32>,
pub ty: Handle<Type>,
/// For I/O structs, defines the binding.
pub binding: Option<Binding>,
}
/// The number of dimensions an image has.
@@ -452,9 +448,18 @@ pub enum Binding {
/// Built-in shader variable.
BuiltIn(BuiltIn),
/// Indexed location.
Location(u32),
/// Binding within a resource group.
Resource { group: u32, binding: u32 },
Location(u32, Option<Interpolation>),
}
/// Pipeline binding information for global resources.
#[derive(Clone, Debug, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
pub struct ResourceBinding {
/// The bind group index.
pub group: u32,
/// Binding number within the group.
pub binding: u32,
}
/// Variable defined at module level.
@@ -466,19 +471,12 @@ pub struct GlobalVariable {
pub name: Option<String>,
/// How this variable is to be stored.
pub class: StorageClass,
/// How this variable is to be bound.
pub binding: Option<Binding>,
/// For resources, defines the binding point.
pub binding: Option<ResourceBinding>,
/// The type of this variable.
pub ty: Handle<Type>,
/// Initial value for this variable.
pub init: Option<Handle<Constant>>,
//TODO: require fragment input interpolation once the entry point I/O
// is refactored.
/// The interpolation qualifier, if any.
/// If the this `GlobalVariable` is a vertex output
/// or fragment input, `None` corresponds to the
/// `smooth`/`perspective` interpolation qualifier.
pub interpolation: Option<Interpolation>,
/// Access bit for storage types of images and buffers.
pub storage_access: StorageAccess,
}
@@ -841,7 +839,7 @@ pub enum Statement {
}
/// A function argument.
#[derive(Debug)]
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
pub struct FunctionArgument {
@@ -849,6 +847,20 @@ pub struct FunctionArgument {
pub name: Option<String>,
/// Type of the argument.
pub ty: Handle<Type>,
/// For entry points, an argument has to have a binding
/// unless it's a structure.
pub binding: Option<Binding>,
}
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
pub struct FunctionResult {
/// Type of the result.
pub ty: Handle<Type>,
/// For entry points, the result has to have a binding
/// unless it's a structure.
pub binding: Option<Binding>,
}
/// A function defined in the module.
@@ -860,8 +872,8 @@ pub struct Function {
pub name: Option<String>,
/// Information about function argument.
pub arguments: Vec<FunctionArgument>,
/// The return type of this function, if any.
pub return_type: Option<Handle<Type>>,
/// The result of this function, if any.
pub result: Option<FunctionResult>,
/// Local variables defined and used in the function.
pub local_variables: Arena<LocalVariable>,
/// Expressions used inside this function.

71
src/proc/analyzer.rs
View File

@@ -288,6 +288,7 @@ impl FunctionInfo {
&mut self,
handle: Handle<crate::Expression>,
expression_arena: &Arena<crate::Expression>,
arguments: &[crate::FunctionArgument],
global_var_arena: &Arena<crate::GlobalVariable>,
other_functions: &[FunctionInfo],
) -> Result<(), AnalysisError> {
@@ -316,33 +317,39 @@ impl FunctionInfo {
requirement: None,
}
}
// same as `LocalVariable` - generally non-uniform
E::FunctionArgument(_) => Uniformity::non_uniform_result(handle),
// depends on the builtin and storage class
E::GlobalVariable(gh) => {
assignable_global = Some(gh);
let var = &global_var_arena[gh];
let uniform = if let Some(crate::Binding::BuiltIn(built_in)) = var.binding {
match built_in {
// depends on the builtin or interpolation
E::FunctionArgument(index) => {
let arg = &arguments[index as usize];
let uniform = match arg.binding {
Some(crate::Binding::BuiltIn(built_in)) => match built_in {
// per-polygon built-ins are uniform
crate::BuiltIn::FrontFacing
// per-work-group built-ins are uniform
| crate::BuiltIn::WorkGroupId
| crate::BuiltIn::WorkGroupSize => true,
_ => false,
}
} else {
use crate::StorageClass as Sc;
match var.class {
// only flat inputs are uniform
Sc::Input => var.interpolation == Some(crate::Interpolation::Flat),
Sc::Output | Sc::Function | Sc::Private | Sc::WorkGroup => false,
// uniform data
Sc::Uniform | Sc::PushConstant => true,
// storage data is only uniform when read-only
Sc::Handle | Sc::Storage => {
!var.storage_access.contains(crate::StorageAccess::STORE)
}
},
// only flat inputs are uniform
Some(crate::Binding::Location(_, Some(crate::Interpolation::Flat))) => true,
_ => false,
};
Uniformity {
non_uniform_result: if uniform { None } else { Some(handle) },
requirement: None,
}
}
// depends on the storage class
E::GlobalVariable(gh) => {
use crate::StorageClass as Sc;
assignable_global = Some(gh);
let var = &global_var_arena[gh];
let uniform = match var.class {
Sc::Function | Sc::Private | Sc::WorkGroup => false,
// uniform data
Sc::Uniform | Sc::PushConstant => true,
// storage data is only uniform when read-only
Sc::Handle | Sc::Storage => {
!var.storage_access.contains(crate::StorageAccess::STORE)
}
};
Uniformity {
@@ -660,7 +667,13 @@ impl Analysis {
};
for (handle, _) in fun.expressions.iter() {
info.process_expression(handle, &fun.expressions, global_var_arena, &self.functions)?;
info.process_expression(
handle,
&fun.expressions,
&fun.arguments,
global_var_arena,
&self.functions,
)?;
}
let uniformity = info.process_block(&fun.body, &self.functions, None)?;
@@ -727,18 +740,16 @@ fn uniform_control_flow() {
name: None,
init: None,
ty,
binding: Some(crate::Binding::BuiltIn(crate::BuiltIn::VertexIndex)),
class: crate::StorageClass::Input,
interpolation: None,
storage_access: crate::StorageAccess::empty(),
class: crate::StorageClass::Handle,
binding: None,
storage_access: crate::StorageAccess::STORE,
});
let uniform_global = global_var_arena.append(crate::GlobalVariable {
name: None,
init: None,
ty,
binding: Some(crate::Binding::Location(0)),
class: crate::StorageClass::Input,
interpolation: Some(crate::Interpolation::Flat),
binding: None,
class: crate::StorageClass::Uniform,
storage_access: crate::StorageAccess::empty(),
});
@@ -772,7 +783,7 @@ fn uniform_control_flow() {
expressions: vec![ExpressionInfo::default(); expressions.len()].into_boxed_slice(),
};
for (handle, _) in expressions.iter() {
info.process_expression(handle, &expressions, &global_var_arena, &[])
info.process_expression(handle, &expressions, &[], &global_var_arena, &[])
.unwrap();
}
assert_eq!(info[non_uniform_global_expr].ref_count, 1);

8
src/proc/namer.rs
View File

@@ -14,6 +14,7 @@ pub enum NameKey {
FunctionLocal(Handle<crate::Function>, Handle<crate::LocalVariable>),
EntryPoint(EntryPointIndex),
EntryPointLocal(EntryPointIndex, Handle<crate::LocalVariable>),
EntryPointArgument(EntryPointIndex, u32),
}
/// This processor assigns names to all the things in a module
@@ -160,6 +161,13 @@ impl Namer {
for (ep_index, ep) in module.entry_points.iter().enumerate() {
let ep_name = self.call(&ep.name);
output.insert(NameKey::EntryPoint(ep_index as _), ep_name);
for (index, arg) in ep.function.arguments.iter().enumerate() {
let name = self.call_or(&arg.name, "param");
output.insert(
NameKey::EntryPointArgument(ep_index as _, index as u32),
name,
);
}
for (handle, var) in ep.function.local_variables.iter() {
let name = self.call_or(&var.name, "local");
output.insert(NameKey::EntryPointLocal(ep_index as _, handle), name);

7
src/proc/typifier.rs
View File

@@ -646,10 +646,11 @@ impl Typifier {
}
},
crate::Expression::Call(function) => {
let ty = ctx.functions[function]
.return_type
let result = ctx.functions[function]
.result
.as_ref()
.ok_or(ResolveError::FunctionReturnsVoid)?;
Resolution::Handle(ty)
Resolution::Handle(result.ty)
}
crate::Expression::ArrayLength(_) => Resolution::Value(Ti::Scalar {
kind: crate::ScalarKind::Uint,

470
src/proc/validator.rs
View File

@@ -70,8 +70,7 @@ pub struct Validator {
// already have to use the typifier, so the work here is redundant in a way.
typifier: Typifier,
type_flags: Vec<TypeFlags>,
location_in_mask: BitSet,
location_out_mask: BitSet,
location_mask: BitSet,
bind_group_masks: Vec<BitSet>,
select_cases: FastHashSet<i32>,
valid_expression_list: Vec<Handle<crate::Expression>>,
@@ -112,8 +111,6 @@ pub enum GlobalVariableError {
InvalidUsage,
#[error("Type isn't compatible with the storage class")]
InvalidType,
#[error("Interpolation is not valid")]
InvalidInterpolation,
#[error("Storage access {seen:?} exceeds the allowed {allowed:?}")]
InvalidStorageAccess {
allowed: crate::StorageAccess,
@@ -126,8 +123,6 @@ pub enum GlobalVariableError {
},
#[error("Binding decoration is missing or not applicable")]
InvalidBinding,
#[error("BuiltIn type for {0:?} is invalid")]
InvalidBuiltInType(crate::BuiltIn),
}
#[derive(Clone, Debug, Error)]
@@ -136,6 +131,22 @@ pub enum LocalVariableError {
InitializerType,
}
#[derive(Clone, Debug, Error)]
pub enum VaryingError {
#[error("The type does not match the varying")]
InvalidType(Handle<crate::Type>),
#[error("Interpolation is not valid")]
InvalidInterpolation,
#[error("BuiltIn {0:?} is not available at this stage")]
InvalidBuiltInStage(crate::BuiltIn),
#[error("BuiltIn type for {0:?} is invalid")]
InvalidBuiltInType(crate::BuiltIn),
#[error("Struct member {0} is missing a binding")]
MemberMissingBinding(u32),
#[error("Multiple bindings at location {location} are present")]
BindingCollision { location: u32 },
}
#[derive(Clone, Debug, Error)]
pub enum ExpressionError {
#[error("Is invalid")]
@@ -241,16 +252,14 @@ pub enum EntryPointError {
UnexpectedWorkgroupSize,
#[error("Workgroup size is out of range")]
OutOfRangeWorkgroupSize,
#[error("Can't have arguments")]
UnexpectedArguments,
#[error("Can't have a return value")]
UnexpectedReturnValue,
#[error("Global variable {0:?} is used incorrectly as {1:?}")]
InvalidGlobalUsage(Handle<crate::GlobalVariable>, GlobalUse),
#[error("Bindings for {0:?} conflict with other global variables")]
#[error("Bindings for {0:?} conflict with other resource")]
BindingCollision(Handle<crate::GlobalVariable>),
#[error("Built-in {0:?} is not applicable to this entry point")]
InvalidBuiltIn(crate::BuiltIn),
#[error("Argument {0} varying error")]
Argument(u32, #[source] VaryingError),
#[error("Result varying error")]
Result(#[source] VaryingError),
#[error("Location {location} onterpolation of an integer has to be flat")]
InvalidIntegerInterpolation { location: u32 },
#[error(transparent)]
@@ -300,111 +309,6 @@ pub enum ValidationError {
Corrupted,
}
impl crate::GlobalVariable {
fn forbid_interpolation(&self) -> Result<(), GlobalVariableError> {
match self.interpolation {
Some(_) => Err(GlobalVariableError::InvalidInterpolation),
None => Ok(()),
}
}
fn check_resource(&self) -> Result<(), GlobalVariableError> {
match self.binding {
Some(crate::Binding::Resource { .. }) => {}
Some(crate::Binding::BuiltIn(_)) | Some(crate::Binding::Location(_)) | None => {
return Err(GlobalVariableError::InvalidBinding)
}
}
self.forbid_interpolation()
}
fn check_varying(&self, types: &Arena<crate::Type>) -> Result<(), GlobalVariableError> {
match self.binding {
Some(crate::Binding::BuiltIn(built_in)) => {
use crate::{BuiltIn as Bi, ScalarKind as Sk, TypeInner as Ti, VectorSize as Vs};
// Only validate the type here. Whether or not it's legal to access
// this builtin is up to the entry point.
let width = 4;
let expected_ty_inner = match built_in {
Bi::BaseInstance
| Bi::BaseVertex
| Bi::InstanceIndex
| Bi::VertexIndex
| Bi::SampleIndex
| Bi::SampleMaskIn
| Bi::SampleMaskOut
| Bi::LocalInvocationIndex => Some(Ti::Scalar {
kind: Sk::Uint,
width,
}),
Bi::PointSize | Bi::FragDepth => Some(Ti::Scalar {
kind: Sk::Float,
width,
}),
Bi::Position | Bi::FragCoord => Some(Ti::Vector {
size: Vs::Quad,
kind: Sk::Float,
width,
}),
Bi::FrontFacing => Some(Ti::Scalar {
kind: Sk::Bool,
width: crate::BOOL_WIDTH,
}),
Bi::GlobalInvocationId
| Bi::LocalInvocationId
| Bi::WorkGroupId
| Bi::WorkGroupSize => Some(Ti::Vector {
size: Vs::Tri,
kind: Sk::Uint,
width,
}),
Bi::ClipDistance => None,
};
let ty_inner = &types[self.ty].inner;
if Some(ty_inner) != expected_ty_inner.as_ref() {
match (built_in, &types[self.ty].inner) {
(Bi::ClipDistance, &Ti::Array { base, .. }) => match types[base].inner {
Ti::Scalar {
kind: Sk::Float, ..
} => {}
ref other => {
log::warn!("Wrong array base type: {:?}", other);
return Err(GlobalVariableError::InvalidBuiltInType(built_in));
}
},
(_, other) => {
log::warn!("Wrong builtin type: {:?}", other);
return Err(GlobalVariableError::InvalidBuiltInType(built_in));
}
}
}
self.forbid_interpolation()?
}
Some(crate::Binding::Location(_)) => match types[self.ty].inner {
crate::TypeInner::Scalar { .. }
| crate::TypeInner::Vector { .. }
| crate::TypeInner::Matrix { .. } => {}
_ => return Err(GlobalVariableError::InvalidType),
},
Some(crate::Binding::Resource { .. }) => {
return Err(GlobalVariableError::InvalidBinding)
}
None => {
match types[self.ty].inner {
//TODO: check the member types
crate::TypeInner::Struct {
block: _,
members: _,
} => self.forbid_interpolation()?,
_ => return Err(GlobalVariableError::InvalidType),
}
}
}
Ok(())
}
}
fn storage_usage(access: crate::StorageAccess) -> GlobalUse {
let mut storage_usage = GlobalUse::QUERY;
if access.contains(crate::StorageAccess::LOAD) {
@@ -416,27 +320,184 @@ fn storage_usage(access: crate::StorageAccess) -> GlobalUse {
storage_usage
}
fn built_in_usage(built_in: crate::BuiltIn) -> (crate::ShaderStage, GlobalUse) {
use crate::{BuiltIn as Bi, ShaderStage as Ss};
match built_in {
Bi::BaseInstance => (Ss::Vertex, GlobalUse::READ),
Bi::BaseVertex => (Ss::Vertex, GlobalUse::READ),
Bi::ClipDistance => (Ss::Vertex, GlobalUse::WRITE),
Bi::InstanceIndex => (Ss::Vertex, GlobalUse::READ),
Bi::PointSize => (Ss::Vertex, GlobalUse::WRITE),
Bi::Position => (Ss::Vertex, GlobalUse::WRITE),
Bi::VertexIndex => (Ss::Vertex, GlobalUse::READ),
Bi::FragCoord => (Ss::Fragment, GlobalUse::READ),
Bi::FragDepth => (Ss::Fragment, GlobalUse::WRITE),
Bi::FrontFacing => (Ss::Fragment, GlobalUse::READ),
Bi::SampleIndex => (Ss::Fragment, GlobalUse::READ),
Bi::SampleMaskIn => (Ss::Fragment, GlobalUse::READ),
Bi::SampleMaskOut => (Ss::Fragment, GlobalUse::WRITE),
Bi::GlobalInvocationId => (Ss::Compute, GlobalUse::READ),
Bi::LocalInvocationId => (Ss::Compute, GlobalUse::READ),
Bi::LocalInvocationIndex => (Ss::Compute, GlobalUse::READ),
Bi::WorkGroupId => (Ss::Compute, GlobalUse::READ),
Bi::WorkGroupSize => (Ss::Compute, GlobalUse::READ),
struct VaryingContext<'a> {
ty: Handle<crate::Type>,
stage: crate::ShaderStage,
output: bool,
types: &'a Arena<crate::Type>,
location_mask: &'a mut BitSet,
}
impl VaryingContext<'_> {
fn validate_impl(&mut self, binding: &crate::Binding) -> Result<(), VaryingError> {
use crate::{
BuiltIn as Bi, ScalarKind as Sk, ShaderStage as St, TypeInner as Ti, VectorSize as Vs,
};
let ty_inner = &self.types[self.ty].inner;
match *binding {
crate::Binding::BuiltIn(built_in) => {
let width = 4;
let (visible, type_good) = match built_in {
Bi::BaseInstance | Bi::BaseVertex | Bi::InstanceIndex | Bi::VertexIndex => (
self.stage == St::Vertex && !self.output,
*ty_inner
== Ti::Scalar {
kind: Sk::Uint,
width,
},
),
Bi::ClipDistance => (
self.stage == St::Vertex && self.output,
match *ty_inner {
Ti::Array { base, .. } => {
self.types[base].inner
== Ti::Scalar {
kind: Sk::Float,
width,
}
}
_ => false,
},
),
Bi::PointSize => (
self.stage == St::Vertex && self.output,
*ty_inner
== Ti::Scalar {
kind: Sk::Float,
width,
},
),
Bi::Position => (
self.stage == St::Vertex && self.output,
*ty_inner
== Ti::Vector {
size: Vs::Quad,
kind: Sk::Float,
width,
},
),
Bi::FragCoord => (
self.stage == St::Fragment && !self.output,
*ty_inner
== Ti::Vector {
size: Vs::Quad,
kind: Sk::Float,
width,
},
),
Bi::FragDepth => (
self.stage == St::Fragment && self.output,
*ty_inner
== Ti::Scalar {
kind: Sk::Float,
width,
},
),
Bi::FrontFacing => (
self.stage == St::Fragment && !self.output,
*ty_inner
== Ti::Scalar {
kind: Sk::Bool,
width: crate::BOOL_WIDTH,
},
),
Bi::SampleIndex | Bi::SampleMaskIn => (
self.stage == St::Fragment && !self.output,
*ty_inner
== Ti::Scalar {
kind: Sk::Uint,
width,
},
),
Bi::SampleMaskOut => (
self.stage == St::Fragment && self.output,
*ty_inner
== Ti::Scalar {
kind: Sk::Uint,
width,
},
),
Bi::LocalInvocationIndex => (
self.stage == St::Compute && !self.output,
*ty_inner
== Ti::Scalar {
kind: Sk::Uint,
width,
},
),
Bi::GlobalInvocationId
| Bi::LocalInvocationId
| Bi::WorkGroupId
| Bi::WorkGroupSize => (
self.stage == St::Compute && !self.output,
*ty_inner
== Ti::Vector {
size: Vs::Tri,
kind: Sk::Uint,
width,
},
),
};
if !visible {
return Err(VaryingError::InvalidBuiltInStage(built_in));
}
if !type_good {
log::warn!("Wrong builtin type: {:?}", ty_inner);
return Err(VaryingError::InvalidBuiltInType(built_in));
}
}
crate::Binding::Location(location, interpolation) => {
if !self.location_mask.insert(location as usize) {
return Err(VaryingError::BindingCollision { location });
}
let needs_interpolation =
self.stage == crate::ShaderStage::Fragment && !self.output;
if !needs_interpolation && interpolation.is_some() {
return Err(VaryingError::InvalidInterpolation);
}
match ty_inner.scalar_kind() {
Some(crate::ScalarKind::Float) => {}
Some(_)
if needs_interpolation
&& interpolation != Some(crate::Interpolation::Flat) =>
{
return Err(VaryingError::InvalidInterpolation);
}
_ => return Err(VaryingError::InvalidType(self.ty)),
}
}
}
Ok(())
}
fn validate(mut self, binding: Option<&crate::Binding>) -> Result<(), VaryingError> {
match binding {
Some(binding) => self.validate_impl(binding),
None => {
match self.types[self.ty].inner {
//TODO: check the member types
crate::TypeInner::Struct {
block: false,
ref members,
} => {
for (index, member) in members.iter().enumerate() {
self.ty = member.ty;
match member.binding {
None => {
return Err(VaryingError::MemberMissingBinding(index as u32))
}
Some(ref binding) => self.validate_impl(binding)?,
}
}
}
_ => return Err(VaryingError::InvalidType(self.ty)),
}
Ok(())
}
}
}
}
@@ -446,8 +507,7 @@ impl Validator {
Validator {
typifier: Typifier::new(),
type_flags: Vec::new(),
location_in_mask: BitSet::new(),
location_out_mask: BitSet::new(),
location_mask: BitSet::new(),
bind_group_masks: Vec::new(),
select_cases: FastHashSet::default(),
valid_expression_list: Vec::new(),
@@ -619,17 +679,9 @@ impl Validator {
types: &Arena<crate::Type>,
) -> Result<(), GlobalVariableError> {
log::debug!("var {:?}", var);
let (allowed_storage_access, required_type_flags) = match var.class {
let (allowed_storage_access, required_type_flags, is_resource) = match var.class {
crate::StorageClass::Function => return Err(GlobalVariableError::InvalidUsage),
crate::StorageClass::Input | crate::StorageClass::Output => {
var.check_varying(types)?;
(
crate::StorageAccess::empty(),
TypeFlags::DATA | TypeFlags::INTERFACE,
)
}
crate::StorageClass::Storage => {
var.check_resource()?;
match types[var.ty].inner {
crate::TypeInner::Struct { .. } => (),
_ => return Err(GlobalVariableError::InvalidType),
@@ -637,10 +689,10 @@ impl Validator {
(
crate::StorageAccess::all(),
TypeFlags::DATA | TypeFlags::HOST_SHARED,
true,
)
}
crate::StorageClass::Uniform => {
var.check_resource()?;
match types[var.ty].inner {
crate::TypeInner::Struct { .. } => (),
_ => return Err(GlobalVariableError::InvalidType),
@@ -648,10 +700,10 @@ impl Validator {
(
crate::StorageAccess::empty(),
TypeFlags::DATA | TypeFlags::SIZED | TypeFlags::HOST_SHARED,
true,
)
}
crate::StorageClass::Handle => {
var.check_resource()?;
let access = match types[var.ty].inner {
crate::TypeInner::Image {
class: crate::ImageClass::Storage(_),
@@ -662,18 +714,15 @@ impl Validator {
}
_ => return Err(GlobalVariableError::InvalidType),
};
(access, TypeFlags::empty())
(access, TypeFlags::empty(), true)
}
crate::StorageClass::Private | crate::StorageClass::WorkGroup => {
if var.binding.is_some() {
return Err(GlobalVariableError::InvalidBinding);
}
var.forbid_interpolation()?;
(crate::StorageAccess::empty(), TypeFlags::DATA)
(crate::StorageAccess::empty(), TypeFlags::DATA, false)
}
crate::StorageClass::PushConstant => (
crate::StorageAccess::LOAD,
TypeFlags::DATA | TypeFlags::HOST_SHARED,
false,
),
};
@@ -692,6 +741,10 @@ impl Validator {
});
}
if is_resource != var.binding.is_some() {
return Err(GlobalVariableError::InvalidBinding);
}
Ok(())
}
@@ -765,7 +818,7 @@ impl Validator {
.map(|expr| self.resolve_type_impl(expr, context.types))
.transpose()
.map_err(CallError::ResultValue)?;
let expected_ty = fun.return_type.map(|ty| &context.types[ty].inner);
let expected_ty = fun.result.as_ref().map(|fr| &context.types[fr.ty].inner);
if result_ty != expected_ty {
log::error!(
"Called function returns {:?} where {:?} is expected",
@@ -773,7 +826,7 @@ impl Validator {
expected_ty
);
return Err(CallError::ResultType {
required: fun.return_type,
required: fun.result.as_ref().map(|fr| fr.ty),
seen_expression: result,
});
}
@@ -1054,7 +1107,7 @@ impl Validator {
expressions: &fun.expressions,
types: &module.types,
functions: &module.functions,
return_type: fun.return_type,
return_type: fun.result.as_ref().map(|fr| fr.ty),
},
)
}
@@ -1080,73 +1133,43 @@ impl Validator {
return Err(EntryPointError::UnexpectedWorkgroupSize);
}
self.location_in_mask.clear();
self.location_out_mask.clear();
self.location_mask.clear();
for (index, fa) in ep.function.arguments.iter().enumerate() {
let ctx = VaryingContext {
ty: fa.ty,
stage: ep.stage,
output: false,
types: &module.types,
location_mask: &mut self.location_mask,
};
ctx.validate(fa.binding.as_ref())
.map_err(|e| EntryPointError::Argument(index as u32, e))?;
}
self.location_mask.clear();
if let Some(ref fr) = ep.function.result {
let ctx = VaryingContext {
ty: fr.ty,
stage: ep.stage,
output: true,
types: &module.types,
location_mask: &mut self.location_mask,
};
ctx.validate(fr.binding.as_ref())
.map_err(EntryPointError::Result)?;
}
for bg in self.bind_group_masks.iter_mut() {
bg.clear();
}
for (var_handle, var) in module.global_variables.iter() {
let usage = info[var_handle];
if usage.is_empty() {
continue;
}
if let Some(crate::Binding::Location(location)) = var.binding {
if ep.stage == crate::ShaderStage::Fragment
&& var.class == crate::StorageClass::Input
{
match module.types[var.ty].inner.scalar_kind() {
Some(crate::ScalarKind::Float) => {}
Some(_) if var.interpolation != Some(crate::Interpolation::Flat) => {
return Err(EntryPointError::InvalidIntegerInterpolation { location });
}
_ => {}
}
}
}
let allowed_usage = match var.class {
crate::StorageClass::Function => unreachable!(),
crate::StorageClass::Input => {
match var.binding {
Some(crate::Binding::BuiltIn(built_in)) => {
let (allowed_stage, allowed_usage) = built_in_usage(built_in);
if allowed_stage != ep.stage || !allowed_usage.contains(GlobalUse::READ)
{
return Err(EntryPointError::InvalidBuiltIn(built_in));
}
}
Some(crate::Binding::Location(loc)) => {
if !self.location_in_mask.insert(loc as usize) {
return Err(EntryPointError::BindingCollision(var_handle));
}
}
Some(crate::Binding::Resource { .. }) => unreachable!(),
None => (),
}
GlobalUse::READ
}
crate::StorageClass::Output => {
match var.binding {
Some(crate::Binding::BuiltIn(built_in)) => {
let (allowed_stage, allowed_usage) = built_in_usage(built_in);
if allowed_stage != ep.stage
|| !allowed_usage.contains(GlobalUse::WRITE)
{
return Err(EntryPointError::InvalidBuiltIn(built_in));
}
}
Some(crate::Binding::Location(loc)) => {
if !self.location_out_mask.insert(loc as usize) {
return Err(EntryPointError::BindingCollision(var_handle));
}
}
Some(crate::Binding::Resource { .. }) => unreachable!(),
None => (),
}
GlobalUse::READ | GlobalUse::WRITE
}
crate::StorageClass::Uniform => GlobalUse::READ | GlobalUse::QUERY,
crate::StorageClass::Storage => storage_usage(var.storage_access),
crate::StorageClass::Handle => match module.types[var.ty].inner {
@@ -1169,23 +1192,16 @@ impl Validator {
return Err(EntryPointError::InvalidGlobalUsage(var_handle, usage));
}
if let Some(crate::Binding::Resource { group, binding }) = var.binding {
while self.bind_group_masks.len() <= group as usize {
if let Some(ref bind) = var.binding {
while self.bind_group_masks.len() <= bind.group as usize {
self.bind_group_masks.push(BitSet::new());
}
if !self.bind_group_masks[group as usize].insert(binding as usize) {
if !self.bind_group_masks[bind.group as usize].insert(bind.binding as usize) {
return Err(EntryPointError::BindingCollision(var_handle));
}
}
}
if !ep.function.arguments.is_empty() {
return Err(EntryPointError::UnexpectedArguments);
}
if ep.function.return_type.is_some() {
return Err(EntryPointError::UnexpectedReturnValue);
}
self.validate_function(&ep.function, info, module)?;
Ok(())
}