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

implement const-expressions (phase 1)

Browse Source
This commit is contained in:
teoxoy
2023-04-05 17:38:03 +02:00
committed by Jim Blandy
parent dfd88ee485
commit d348d843e1
145 changed files with 7055 additions and 7685 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
src/back/dot/mod.rs
View File

@@ -405,6 +405,10 @@ fn write_function_expressions(
E::Literal(_) => ("Literal".into(), 2),
E::Constant(_) => ("Constant".into(), 2),
E::ZeroValue(_) => ("ZeroValue".into(), 2),
E::Compose { ref components, .. } => {
payload = Some(Payload::Arguments(components));
("Compose".into(), 3)
}
E::Access { base, index } => {
edges.insert("base", base);
edges.insert("index", index);
@@ -426,10 +430,6 @@ fn write_function_expressions(
edges.insert("vector", vector);
(format!("Swizzle{:?}", &pattern[..size as usize]).into(), 3)
}
E::Compose { ref components, .. } => {
payload = Some(Payload::Arguments(components));
("Compose".into(), 3)
}
E::FunctionArgument(index) => (format!("Argument[{index}]").into(), 1),
E::GlobalVariable(h) => {
payload = Some(Payload::Global(h));

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

@@ -604,8 +604,6 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);
// Write struct types.
//
// This are always ordered because the IR is structured in a way that
@@ -627,6 +625,23 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
// Write all named constants
let mut constants = self
.module
.constants
.iter()
.filter(|&(_, c)| c.name.is_some())
.peekable();
while let Some((handle, _)) = constants.next() {
self.write_global_constant(handle)?;
// Add extra newline for readability on last iteration
if constants.peek().is_none() {
writeln!(self.out)?;
}
}
let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);
// Write the globals
//
// Unless explicitly disabled with WriterFlags::INCLUDE_UNUSED_ITEMS,
@@ -722,31 +737,6 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
if include_unused {
// write named constants
for (handle, constant) in self.module.constants.iter() {
if let Some(name) = constant.name.as_ref() {
write!(self.out, "const ")?;
match constant.inner {
crate::ConstantInner::Scalar { width, value } => {
// create a TypeInner to write
let inner = TypeInner::Scalar {
width,
kind: value.scalar_kind(),
};
self.write_value_type(&inner)?;
}
crate::ConstantInner::Composite { ty, .. } => {
self.write_type(ty)?;
}
};
write!(self.out, " {name} = ")?;
self.write_constant(handle)?;
writeln!(self.out, ";")?;
}
}
}
for arg in self.entry_point.function.arguments.iter() {
self.write_varying(arg.binding.as_ref(), arg.ty, false)?;
}
@@ -818,12 +808,6 @@ impl<'a, W: Write> Writer<'a, W> {
///
/// # Notes
/// Adds no trailing or leading whitespace
///
/// # Panics
/// - If type is either a image, a sampler, a pointer, or a struct
/// - If it's an Array with a [`ArraySize::Constant`](crate::ArraySize::Constant) with a
/// constant that isn't a [`Scalar`](crate::ConstantInner::Scalar) or if the
/// scalar value isn't an [`Sint`](crate::ScalarValue::Sint) or [`Uint`](crate::ScalarValue::Uint)
fn write_value_type(&mut self, inner: &TypeInner) -> BackendResult {
match *inner {
// Scalars are simple we just get the full name from `glsl_scalar`
@@ -865,11 +849,9 @@ impl<'a, W: Write> Writer<'a, W> {
rows as u8
)?,
// GLSL arrays are written as `type name[size]`
// Current code is written arrays only as `[size]`
// Here we only write the size of the array i.e. `[size]`
// Base `type` and `name` should be written outside
TypeInner::Array { base, size, .. } => self.write_array_size(base, size)?,
// Panic if either Image, Sampler, Pointer, or a Struct is being written
//
// Write all variants instead of `_` so that if new variants are added a
// no exhaustiveness error is thrown
TypeInner::Pointer { .. }
@@ -890,12 +872,6 @@ impl<'a, W: Write> Writer<'a, W> {
///
/// # Notes
/// Adds no trailing or leading whitespace
///
/// # Panics
/// - If type is either a image or sampler
/// - If it's an Array with a [`ArraySize::Constant`](crate::ArraySize::Constant) with a
/// constant that isn't a [`Scalar`](crate::ConstantInner::Scalar) or if the
/// scalar value isn't an [`Sint`](crate::ScalarValue::Sint) or [`Uint`](crate::ScalarValue::Uint)
fn write_type(&mut self, ty: Handle<crate::Type>) -> BackendResult {
match self.module.types[ty].inner {
// glsl has no pointer types so just write types as normal and loads are skipped
@@ -1050,7 +1026,7 @@ impl<'a, W: Write> Writer<'a, W> {
if global.space.initializable() && is_value_init_supported(self.module, global.ty) {
write!(self.out, " = ")?;
if let Some(init) = global.init {
self.write_constant(init)?;
self.write_const_expr(init)?;
} else {
self.write_zero_init_value(global.ty)?;
}
@@ -1550,7 +1526,7 @@ impl<'a, W: Write> Writer<'a, W> {
// Write the constant
// `write_constant` adds no trailing or leading space/newline
self.write_constant(init)?;
self.write_const_expr(init)?;
} else if is_value_init_supported(self.module, local.ty) {
write!(self.out, " = ")?;
self.write_zero_init_value(local.ty)?;
@@ -1621,60 +1597,29 @@ impl<'a, W: Write> Writer<'a, W> {
mut f: F,
) -> BackendResult {
// Loop trough `data` invoking `f` for each element
for (i, item) in data.iter().enumerate() {
f(self, i as u32, item)?;
// Only write a comma if isn't the last element
if i != data.len().saturating_sub(1) {
// The leading space is for readability only
for (index, item) in data.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
f(self, index as u32, item)?;
}
Ok(())
}
/// Helper method used to write constants
///
/// # Notes
/// Adds no newlines or leading/trailing whitespace
fn write_constant(&mut self, handle: Handle<crate::Constant>) -> BackendResult {
use crate::ScalarValue as Sv;
match self.module.constants[handle].inner {
crate::ConstantInner::Scalar {
width: _,
ref value,
} => match *value {
// Signed integers don't need anything special
Sv::Sint(int) => write!(self.out, "{int}")?,
// Unsigned integers need a `u` at the end
//
// While `core` doesn't necessarily need it, it's allowed and since `es` needs it we
// always write it as the extra branch wouldn't have any benefit in readability
Sv::Uint(int) => write!(self.out, "{int}u")?,
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero which is needed for a valid glsl float constant
Sv::Float(float) => write!(self.out, "{float:?}")?,
// Booleans are either `true` or `false` so nothing special needs to be done
Sv::Bool(boolean) => write!(self.out, "{boolean}")?,
},
// Composite constant are created using the same syntax as compose
// `type(components)` where `components` is a comma separated list of constants
crate::ConstantInner::Composite { ty, ref components } => {
self.write_type(ty)?;
if let TypeInner::Array { base, size, .. } = self.module.types[ty].inner {
self.write_array_size(base, size)?;
}
write!(self.out, "(")?;
// Write the comma separated constants
self.write_slice(components, |this, _, arg| this.write_constant(*arg))?;
write!(self.out, ")")?
}
/// Helper method used to write global constants
fn write_global_constant(&mut self, handle: Handle<crate::Constant>) -> BackendResult {
write!(self.out, "const ")?;
let constant = &self.module.constants[handle];
self.write_type(constant.ty)?;
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, " {name}")?;
if let TypeInner::Array { base, size, .. } = self.module.types[constant.ty].inner {
self.write_array_size(base, size)?;
}
write!(self.out, " = ")?;
self.write_const_expr(constant.init)?;
writeln!(self.out, ";")?;
Ok(())
}
@@ -2219,6 +2164,80 @@ impl<'a, W: Write> Writer<'a, W> {
Ok(())
}
/// Helper method used to write constant expressions
///
/// # Notes
/// Adds no newlines or leading/trailing whitespace
fn write_const_expr(&mut self, expr: Handle<crate::Expression>) -> BackendResult {
self.write_possibly_const_expr(expr, &self.module.const_expressions, |writer, expr| {
writer.write_const_expr(expr)
})
}
/// Helper method used to write possibly constant expressions
///
/// # Notes
/// Adds no newlines or leading/trailing whitespace
fn write_possibly_const_expr<E>(
&mut self,
expr: Handle<crate::Expression>,
expressions: &crate::Arena<crate::Expression>,
write_expression: E,
) -> BackendResult
where
E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
{
use crate::Expression;
match expressions[expr] {
Expression::Literal(literal) => {
match literal {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero which is needed for a valid glsl float constant
crate::Literal::F64(value) => write!(self.out, "{:?}LF", value)?,
crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
// Unsigned integers need a `u` at the end
//
// While `core` doesn't necessarily need it, it's allowed and since `es` needs it we
// always write it as the extra branch wouldn't have any benefit in readability
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
}
}
Expression::Constant(handle) => {
let constant = &self.module.constants[handle];
if constant.name.is_some() {
write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
} else {
self.write_const_expr(constant.init)?;
}
}
Expression::ZeroValue(ty) => {
self.write_zero_init_value(ty)?;
}
Expression::Compose { ty, ref components } => {
self.write_type(ty)?;
if let TypeInner::Array { base, size, .. } = self.module.types[ty].inner {
self.write_array_size(base, size)?;
}
write!(self.out, "(")?;
for (index, component) in components.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
write_expression(self, *component)?;
}
write!(self.out, ")")?
}
_ => unreachable!(),
}
Ok(())
}
/// Helper method to write expressions
///
/// # Notes
@@ -2236,6 +2255,14 @@ impl<'a, W: Write> Writer<'a, W> {
}
match ctx.expressions[expr] {
Expression::Literal(_)
| Expression::Constant(_)
| Expression::ZeroValue(_)
| Expression::Compose { .. } => {
self.write_possibly_const_expr(expr, ctx.expressions, |writer, expr| {
writer.write_expr(expr, ctx)
})?;
}
// `Access` is applied to arrays, vectors and matrices and is written as indexing
Expression::Access { base, index } => {
self.write_expr(base, ctx)?;
@@ -2281,26 +2308,6 @@ impl<'a, W: Write> Writer<'a, W> {
ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))),
}
}
// Constants are delegated to `write_constant`
Expression::Constant(constant) => self.write_constant(constant)?,
Expression::ZeroValue(ty) => {
self.write_zero_init_value(ty)?;
}
Expression::Literal(literal) => {
match literal {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero which is needed for a valid glsl float constant
crate::Literal::F64(value) => write!(self.out, "{:?}LF", value)?,
crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
// Unsigned integers need a `u` at the end
//
// While `core` doesn't necessarily need it, it's allowed and since `es` needs it we
// always write it as the extra branch wouldn't have any benefit in readability
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
}
}
// `Splat` needs to actually write down a vector, it's not always inferred in GLSL.
Expression::Splat { size: _, value } => {
let resolved = ctx.info[expr].ty.inner_with(&self.module.types);
@@ -2321,20 +2328,6 @@ impl<'a, W: Write> Writer<'a, W> {
self.out.write_char(back::COMPONENTS[sc as usize])?;
}
}
// `Compose` is pretty simple we just write `type(components)` where `components` is a
// comma separated list of expressions
Expression::Compose { ty, ref components } => {
self.write_type(ty)?;
let resolved = ctx.info[expr].ty.inner_with(&self.module.types);
if let TypeInner::Array { base, size, .. } = *resolved {
self.write_array_size(base, size)?;
}
write!(self.out, "(")?;
self.write_slice(components, |this, _, arg| this.write_expr(*arg, ctx))?;
write!(self.out, ")")?
}
// Function arguments are written as the argument name
Expression::FunctionArgument(pos) => {
write!(self.out, "{}", &self.names[&ctx.argument_key(pos)])?
@@ -2524,7 +2517,7 @@ impl<'a, W: Write> Writer<'a, W> {
if tex_1d_hack {
write!(self.out, "ivec2(")?;
}
self.write_constant(constant)?;
self.write_const_expr(constant)?;
if tex_1d_hack {
write!(self.out, ", 0)")?;
}
@@ -3857,11 +3850,11 @@ impl<'a, W: Write> Writer<'a, W> {
TypeInner::Struct { ref members, .. } => {
let name = &self.names[&NameKey::Type(ty)];
write!(self.out, "{name}(")?;
for (i, member) in members.iter().enumerate() {
self.write_zero_init_value(member.ty)?;
if i != members.len().saturating_sub(1) {
for (index, member) in members.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
self.write_zero_init_value(member.ty)?;
}
write!(self.out, ")")?;
}

98
src/back/hlsl/help.rs
View File

@@ -159,10 +159,7 @@ impl<'a, W: Write> super::Writer<'a, W> {
/// <https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/sm5-object-rwbyteaddressbuffer-getdimensions>
pub(super) fn write_wrapped_array_length_function(
&mut self,
module: &crate::Module,
wal: WrappedArrayLength,
expr_handle: Handle<crate::Expression>,
func_ctx: &FunctionCtx,
) -> BackendResult {
use crate::back::INDENT;
@@ -170,9 +167,7 @@ impl<'a, W: Write> super::Writer<'a, W> {
const RETURN_VARIABLE_NAME: &str = "ret";
// Write function return type and name
let ret_ty = func_ctx.info[expr_handle].ty.inner_with(&module.types);
self.write_value_type(module, ret_ty)?;
write!(self.out, " ")?;
write!(self.out, "uint ")?;
self.write_wrapped_array_length_function_name(wal)?;
// Write function parameters
@@ -786,14 +781,36 @@ impl<'a, W: Write> super::Writer<'a, W> {
Ok(())
}
/// Helper function that write wrapped function for `Expression::ImageQuery` and `Expression::ArrayLength`
///
/// <https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions>
/// Helper function that writes compose wrapped functions
pub(super) fn write_wrapped_compose_functions(
&mut self,
module: &crate::Module,
expressions: &crate::Arena<crate::Expression>,
) -> BackendResult {
for (handle, _) in expressions.iter() {
if let crate::Expression::Compose { ty, .. } = expressions[handle] {
match module.types[ty].inner {
crate::TypeInner::Struct { .. } | crate::TypeInner::Array { .. } => {
let constructor = WrappedConstructor { ty };
if self.wrapped.constructors.insert(constructor) {
self.write_wrapped_constructor_function(module, constructor)?;
}
}
_ => {}
};
}
}
Ok(())
}
/// Helper function that writes various wrapped functions
pub(super) fn write_wrapped_functions(
&mut self,
module: &crate::Module,
func_ctx: &FunctionCtx,
) -> BackendResult {
self.write_wrapped_compose_functions(module, func_ctx.expressions)?;
for (handle, _) in func_ctx.expressions.iter() {
match func_ctx.expressions[handle] {
crate::Expression::ArrayLength(expr) => {
@@ -816,9 +833,8 @@ impl<'a, W: Write> super::Writer<'a, W> {
writable: storage_access.contains(crate::StorageAccess::STORE),
};
if !self.wrapped.array_lengths.contains(&wal) {
self.write_wrapped_array_length_function(module, wal, handle, func_ctx)?;
self.wrapped.array_lengths.insert(wal);
if self.wrapped.array_lengths.insert(wal) {
self.write_wrapped_array_length_function(wal)?;
}
}
crate::Expression::ImageQuery { image, query } => {
@@ -836,9 +852,8 @@ impl<'a, W: Write> super::Writer<'a, W> {
_ => unreachable!("we only query images"),
};
if !self.wrapped.image_queries.contains(&wiq) {
if self.wrapped.image_queries.insert(wiq) {
self.write_wrapped_image_query_function(module, wiq, handle, func_ctx)?;
self.wrapped.image_queries.insert(wiq);
}
}
// Write `WrappedConstructor` for structs that are loaded from `AddressSpace::Storage`
@@ -867,19 +882,18 @@ impl<'a, W: Write> super::Writer<'a, W> {
}
let constructor = WrappedConstructor { ty };
if !writer.wrapped.constructors.contains(&constructor) {
if writer.wrapped.constructors.insert(constructor) {
writer
.write_wrapped_constructor_function(module, constructor)?;
writer.wrapped.constructors.insert(constructor);
}
}
crate::TypeInner::Array { base, .. } => {
write_wrapped_constructor(writer, base, module)?;
let constructor = WrappedConstructor { ty };
if !writer.wrapped.constructors.contains(&constructor) {
if writer.wrapped.constructors.insert(constructor) {
writer
.write_wrapped_constructor_function(module, constructor)?;
writer.wrapped.constructors.insert(constructor);
}
}
_ => {}
@@ -888,18 +902,6 @@ impl<'a, W: Write> super::Writer<'a, W> {
Ok(())
}
}
crate::Expression::Compose { ty, components: _ } => {
let constructor = match module.types[ty].inner {
crate::TypeInner::Struct { .. } | crate::TypeInner::Array { .. } => {
WrappedConstructor { ty }
}
_ => continue,
};
if !self.wrapped.constructors.contains(&constructor) {
self.write_wrapped_constructor_function(module, constructor)?;
self.wrapped.constructors.insert(constructor);
}
}
// We treat matrices of the form `matCx2` as a sequence of C `vec2`s
// (see top level module docs for details).
//
@@ -925,7 +927,7 @@ impl<'a, W: Write> super::Writer<'a, W> {
let ty = base_ty_handle.unwrap();
let access = WrappedStructMatrixAccess { ty, index };
if !self.wrapped.struct_matrix_access.contains(&access) {
if self.wrapped.struct_matrix_access.insert(access) {
self.write_wrapped_struct_matrix_get_function(module, access)?;
self.write_wrapped_struct_matrix_set_function(module, access)?;
self.write_wrapped_struct_matrix_set_vec_function(
@@ -934,7 +936,6 @@ impl<'a, W: Write> super::Writer<'a, W> {
self.write_wrapped_struct_matrix_set_scalar_function(
module, access,
)?;
self.wrapped.struct_matrix_access.insert(access);
}
}
_ => {}
@@ -948,33 +949,6 @@ impl<'a, W: Write> super::Writer<'a, W> {
Ok(())
}
pub(super) fn write_wrapped_constructor_function_for_constant(
&mut self,
module: &crate::Module,
constant: &crate::Constant,
) -> BackendResult {
if let crate::ConstantInner::Composite { ty, ref components } = constant.inner {
match module.types[ty].inner {
crate::TypeInner::Struct { .. } | crate::TypeInner::Array { .. } => {
let constructor = WrappedConstructor { ty };
if !self.wrapped.constructors.contains(&constructor) {
self.write_wrapped_constructor_function(module, constructor)?;
self.wrapped.constructors.insert(constructor);
}
}
_ => {}
}
for constant in components {
self.write_wrapped_constructor_function_for_constant(
module,
&module.constants[*constant],
)?;
}
}
Ok(())
}
pub(super) fn write_texture_coordinates(
&mut self,
kind: &str,
@@ -1085,9 +1059,8 @@ impl<'a, W: Write> super::Writer<'a, W> {
}) = super::writer::get_inner_matrix_data(module, global.ty)
{
let entry = WrappedMatCx2 { columns };
if !self.wrapped.mat_cx2s.contains(&entry) {
if self.wrapped.mat_cx2s.insert(entry) {
self.write_mat_cx2_typedef_and_functions(entry)?;
self.wrapped.mat_cx2s.insert(entry);
}
}
}
@@ -1104,9 +1077,8 @@ impl<'a, W: Write> super::Writer<'a, W> {
}) = super::writer::get_inner_matrix_data(module, member.ty)
{
let entry = WrappedMatCx2 { columns };
if !self.wrapped.mat_cx2s.contains(&entry) {
if self.wrapped.mat_cx2s.insert(entry) {
self.write_mat_cx2_typedef_and_functions(entry)?;
self.wrapped.mat_cx2s.insert(entry);
}
}
}

297
src/back/hlsl/writer.rs
View File

@@ -199,31 +199,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
write!(self.out, ", space{}", bt.space)?;
}
writeln!(self.out, ");")?;
}
// Write all constants
// For example, input wgsl shader:
// ```wgsl
// let c_scale: f32 = 1.2;
// return VertexOutput(uv, vec4<f32>(c_scale * pos, 0.0, 1.0));
// ```
//
// Output shader:
// ```hlsl
// static const float c_scale = 1.2;
// const VertexOutput vertexoutput1 = { vertexinput.uv3, float4((c_scale * vertexinput.pos1), 0.0, 1.0) };
// ```
//
// If we remove `write_global_constant` `c_scale` will be inlined.
for (handle, constant) in module.constants.iter() {
if constant.name.is_some() {
self.write_global_constant(module, &constant.inner, handle)?;
}
// Extra newline for readability
writeln!(self.out)?;
}
// Extra newline for readability
writeln!(self.out)?;
// Save all entry point output types
let ep_results = module
.entry_points
@@ -264,9 +244,20 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
}
}
// Write wrapped constructor functions used in constants
for (_, constant) in module.constants.iter() {
self.write_wrapped_constructor_function_for_constant(module, constant)?;
self.write_wrapped_compose_functions(module, &module.const_expressions)?;
// Write all named constants
let mut constants = module
.constants
.iter()
.filter(|&(_, c)| c.name.is_some())
.peekable();
while let Some((handle, _)) = constants.next() {
self.write_global_constant(module, handle)?;
// Add extra newline for readability on last iteration
if constants.peek().is_none() {
writeln!(self.out)?;
}
}
// Write all globals
@@ -321,8 +312,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
};
let name = self.names[&NameKey::Function(handle)].clone();
// Write wrapped function for `Expression::ImageQuery` and `Expressions::ArrayLength`
// before writing all statements and expressions.
self.write_wrapped_functions(module, &ctx)?;
self.write_function(module, name.as_str(), function, &ctx, info)?;
@@ -362,8 +351,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
named_expressions: &ep.function.named_expressions,
};
// Write wrapped function for `Expression::ImageQuery` and `Expressions::ArrayLength`
// before writing all statements and expressions.
self.write_wrapped_functions(module, &ctx)?;
if ep.stage == ShaderStage::Compute {
@@ -769,7 +756,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
if global.space == crate::AddressSpace::Private {
write!(self.out, " = ")?;
if let Some(init) = global.init {
self.write_constant(module, init)?;
self.write_const_expression(module, init)?;
} else {
self.write_default_init(module, global.ty)?;
}
@@ -807,44 +794,19 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
fn write_global_constant(
&mut self,
module: &Module,
inner: &crate::ConstantInner,
handle: Handle<crate::Constant>,
) -> BackendResult {
write!(self.out, "static const ")?;
match *inner {
crate::ConstantInner::Scalar {
width: _,
ref value,
} => {
// Write type
let ty_str = match *value {
crate::ScalarValue::Sint(_) => "int",
crate::ScalarValue::Uint(_) => "uint",
crate::ScalarValue::Float(_) => "float",
crate::ScalarValue::Bool(_) => "bool",
};
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, "{ty_str} {name} = ")?;
// Second match required to avoid heap allocation by `format!()`
match *value {
crate::ScalarValue::Sint(value) => write!(self.out, "{value}")?,
crate::ScalarValue::Uint(value) => write!(self.out, "{value}")?,
crate::ScalarValue::Float(value) => {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
write!(self.out, "{value:?}")?
}
crate::ScalarValue::Bool(value) => write!(self.out, "{value}")?,
};
}
crate::ConstantInner::Composite { ty, ref components } => {
self.write_type(module, ty)?;
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, " {name} = ")?;
self.write_composite_constant(module, ty, components)?;
}
let constant = &module.constants[handle];
self.write_type(module, constant.ty)?;
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, " {}", name)?;
// Write size for array type
if let TypeInner::Array { base, size, .. } = module.types[constant.ty].inner {
self.write_array_size(module, base, size)?;
}
write!(self.out, " = ")?;
self.write_const_expression(module, constant.init)?;
writeln!(self.out, ";")?;
Ok(())
}
@@ -1261,12 +1223,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
write!(self.out, " = ")?;
// Write the local initializer if needed
if let Some(init) = local.init {
// Put the equal signal only if there's a initializer
// The leading and trailing spaces aren't needed but help with readability
// Write the constant
// `write_constant` adds no trailing or leading space/newline
self.write_constant(module, init)?;
self.write_const_expression(module, init)?;
} else {
// Zero initialize local variables
self.write_default_init(module, local.ty)?;
@@ -1859,12 +1816,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
let func_name = &self.names[&NameKey::Function(function)];
write!(self.out, "{func_name}(")?;
for (index, argument) in arguments.iter().enumerate() {
self.write_expr(module, *argument, func_ctx)?;
// Only write a comma if isn't the last element
if index != arguments.len().saturating_sub(1) {
// The leading space is for readability only
if index != 0 {
write!(self.out, ", ")?;
}
self.write_expr(module, *argument, func_ctx)?;
}
writeln!(self.out, ");")?
}
@@ -2043,6 +1998,77 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
Ok(())
}
fn write_const_expression(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
) -> BackendResult {
self.write_possibly_const_expression(
module,
expr,
&module.const_expressions,
|writer, expr| writer.write_const_expression(module, expr),
)
}
fn write_possibly_const_expression<E>(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
expressions: &crate::Arena<crate::Expression>,
write_expression: E,
) -> BackendResult
where
E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
{
use crate::Expression;
match expressions[expr] {
Expression::Literal(literal) => match literal {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
crate::Literal::F64(value) => write!(self.out, "{value:?}L")?,
crate::Literal::F32(value) => write!(self.out, "{value:?}")?,
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
},
Expression::Constant(handle) => {
let constant = &module.constants[handle];
if constant.name.is_some() {
write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
} else {
self.write_const_expression(module, constant.init)?;
}
}
Expression::ZeroValue(ty) => self.write_default_init(module, ty)?,
Expression::Compose { ty, ref components } => {
match module.types[ty].inner {
TypeInner::Struct { .. } | TypeInner::Array { .. } => {
self.write_wrapped_constructor_function_name(
module,
WrappedConstructor { ty },
)?;
}
_ => {
self.write_type(module, ty)?;
}
};
write!(self.out, "(")?;
for (index, component) in components.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
write_expression(self, *component)?;
}
write!(self.out, ")")?;
}
_ => unreachable!(),
}
Ok(())
}
/// Helper method to write expressions
///
/// # Notes
@@ -2091,41 +2117,16 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
let expression = &func_ctx.expressions[expr];
match *expression {
Expression::Constant(constant) => self.write_constant(module, constant)?,
Expression::ZeroValue(ty) => self.write_default_init(module, ty)?,
Expression::Literal(literal) => match literal {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
crate::Literal::F64(value) => write!(self.out, "{value:?}L")?,
crate::Literal::F32(value) => write!(self.out, "{value:?}")?,
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
},
Expression::Compose { ty, ref components } => {
match module.types[ty].inner {
TypeInner::Struct { .. } | TypeInner::Array { .. } => {
self.write_wrapped_constructor_function_name(
module,
WrappedConstructor { ty },
)?;
}
_ => {
self.write_type(module, ty)?;
}
};
write!(self.out, "(")?;
for (index, &component) in components.iter().enumerate() {
if index != 0 {
// The leading space is for readability only
write!(self.out, ", ")?;
}
self.write_expr(module, component, func_ctx)?;
}
write!(self.out, ")")?;
Expression::Literal(_)
| Expression::Constant(_)
| Expression::ZeroValue(_)
| Expression::Compose { .. } => {
self.write_possibly_const_expression(
module,
expr,
func_ctx.expressions,
|writer, expr| writer.write_expr(module, expr, func_ctx),
)?;
}
// All of the multiplication can be expressed as `mul`,
// except vector * vector, which needs to use the "*" operator.
@@ -2412,7 +2413,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
if let Some(offset) = offset {
write!(self.out, ", ")?;
self.write_constant(module, offset)?;
self.write_const_expression(module, offset)?;
}
write!(self.out, ")")?;
@@ -3184,82 +3185,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
Ok(())
}
/// Helper method used to write constants
///
/// # Notes
/// Doesn't add any newlines or leading/trailing spaces
fn write_constant(
&mut self,
module: &Module,
handle: Handle<crate::Constant>,
) -> BackendResult {
let constant = &module.constants[handle];
match constant.inner {
crate::ConstantInner::Scalar {
width: _,
ref value,
} => {
if constant.name.is_some() {
write!(self.out, "{}", &self.names[&NameKey::Constant(handle)])?;
} else {
self.write_scalar_value(*value)?;
}
}
crate::ConstantInner::Composite { ty, ref components } => {
self.write_composite_constant(module, ty, components)?;
}
}
Ok(())
}
fn write_composite_constant(
&mut self,
module: &Module,
ty: Handle<crate::Type>,
components: &[Handle<crate::Constant>],
) -> BackendResult {
match module.types[ty].inner {
TypeInner::Struct { .. } | TypeInner::Array { .. } => {
self.write_wrapped_constructor_function_name(module, WrappedConstructor { ty })?;
}
_ => {
self.write_type(module, ty)?;
}
};
write!(self.out, "(")?;
for (index, constant) in components.iter().enumerate() {
self.write_constant(module, *constant)?;
// Only write a comma if isn't the last element
if index != components.len().saturating_sub(1) {
// The leading space is for readability only
write!(self.out, ", ")?;
}
}
write!(self.out, ")")?;
Ok(())
}
/// Helper method used to write [`ScalarValue`](crate::ScalarValue)
///
/// # Notes
/// Adds no trailing or leading whitespace
fn write_scalar_value(&mut self, value: crate::ScalarValue) -> BackendResult {
use crate::ScalarValue as Sv;
match value {
Sv::Sint(value) => write!(self.out, "{value}")?,
Sv::Uint(value) => write!(self.out, "{value}u")?,
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
Sv::Float(value) => write!(self.out, "{value:?}")?,
Sv::Bool(value) => write!(self.out, "{value}")?,
}
Ok(())
}
fn write_named_expr(
&mut self,
module: &Module,

8
src/back/mod.rs
View File

@@ -71,6 +71,14 @@ struct FunctionCtx<'a> {
}
impl FunctionCtx<'_> {
fn resolve_type<'a>(
&'a self,
handle: crate::Handle<crate::Expression>,
types: &'a crate::UniqueArena<crate::Type>,
) -> &'a crate::TypeInner {
self.info[handle].ty.inner_with(types)
}
/// Helper method that generates a [`NameKey`](crate::proc::NameKey) for a local in the current function
const fn name_key(&self, local: crate::Handle<crate::LocalVariable>) -> crate::proc::NameKey {
match self.ty {

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

@@ -300,50 +300,6 @@ impl<'a> TypedGlobalVariable<'a> {
}
}
struct ConstantContext<'a> {
handle: Handle<crate::Constant>,
arena: &'a crate::Arena<crate::Constant>,
names: &'a FastHashMap<NameKey, String>,
first_time: bool,
}
impl<'a> Display for ConstantContext<'a> {
fn fmt(&self, out: &mut Formatter<'_>) -> Result<(), FmtError> {
let con = &self.arena[self.handle];
if con.needs_alias() && !self.first_time {
let name = &self.names[&NameKey::Constant(self.handle)];
return write!(out, "{name}");
}
match con.inner {
crate::ConstantInner::Scalar { value, width: _ } => match value {
crate::ScalarValue::Sint(value) => {
write!(out, "{value}")
}
crate::ScalarValue::Uint(value) => {
write!(out, "{value}u")
}
crate::ScalarValue::Float(value) => {
if value.is_infinite() {
let sign = if value.is_sign_negative() { "-" } else { "" };
write!(out, "{sign}INFINITY")
} else if value.is_nan() {
write!(out, "NAN")
} else {
let suffix = if value.fract() == 0.0 { ".0" } else { "" };
write!(out, "{value}{suffix}")
}
}
crate::ScalarValue::Bool(value) => {
write!(out, "{value}")
}
},
crate::ConstantInner::Composite { .. } => unreachable!("should be aliased"),
}
}
}
pub struct Writer<W> {
out: W,
names: FastHashMap<NameKey, String>,
@@ -501,16 +457,6 @@ impl crate::Type {
}
}
impl crate::Constant {
// Returns `true` if we need to emit an alias for this constant.
const fn needs_alias(&self) -> bool {
match self.inner {
crate::ConstantInner::Scalar { .. } => self.name.is_some(),
crate::ConstantInner::Composite { .. } => true,
}
}
}
enum FunctionOrigin {
Handle(Handle<crate::Function>),
EntryPoint(proc::EntryPointIndex),
@@ -655,12 +601,25 @@ impl<W: Write> Writer<W> {
parameters: impl Iterator<Item = Handle<crate::Expression>>,
context: &ExpressionContext,
) -> BackendResult {
self.put_call_parameters_impl(parameters, |writer, expr| {
writer.put_expression(expr, context, true)
})
}
fn put_call_parameters_impl<E>(
&mut self,
parameters: impl Iterator<Item = Handle<crate::Expression>>,
put_expression: E,
) -> BackendResult
where
E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
{
write!(self.out, "(")?;
for (i, handle) in parameters.enumerate() {
if i != 0 {
write!(self.out, ", ")?;
}
self.put_expression(handle, context, true)?;
put_expression(self, handle)?;
}
write!(self.out, ")")?;
Ok(())
@@ -1077,44 +1036,6 @@ impl<W: Write> Writer<W> {
Ok(())
}
fn put_compose(
&mut self,
ty: Handle<crate::Type>,
components: &[Handle<crate::Expression>],
context: &ExpressionContext,
) -> BackendResult {
match context.module.types[ty].inner {
crate::TypeInner::Scalar { width: 4, kind } if components.len() == 1 => {
write!(self.out, "{}", kind.to_msl_name())?;
self.put_call_parameters(components.iter().cloned(), context)?;
}
crate::TypeInner::Vector { size, kind, .. } => {
put_numeric_type(&mut self.out, kind, &[size])?;
self.put_call_parameters(components.iter().cloned(), context)?;
}
crate::TypeInner::Matrix { columns, rows, .. } => {
put_numeric_type(&mut self.out, crate::ScalarKind::Float, &[rows, columns])?;
self.put_call_parameters(components.iter().cloned(), context)?;
}
crate::TypeInner::Array { .. } | crate::TypeInner::Struct { .. } => {
write!(self.out, "{} {{", &self.names[&NameKey::Type(ty)])?;
for (index, &component) in components.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
// insert padding initialization, if needed
if self.struct_member_pads.contains(&(ty, index as u32)) {
write!(self.out, "{{}}, ")?;
}
self.put_expression(component, context, true)?;
}
write!(self.out, "}}")?;
}
_ => return Err(Error::UnsupportedCompose(ty)),
}
Ok(())
}
/// Write the maximum valid index of the dynamically sized array at the end of `handle`.
///
/// The 'maximum valid index' is simply one less than the array's length.
@@ -1260,6 +1181,113 @@ impl<W: Write> Writer<W> {
Ok(())
}
fn put_const_expression(
&mut self,
expr_handle: Handle<crate::Expression>,
module: &crate::Module,
) -> BackendResult {
self.put_possibly_const_expression(
expr_handle,
&module.const_expressions,
module,
|writer, expr| writer.put_const_expression(expr, module),
)
}
fn put_possibly_const_expression<E>(
&mut self,
expr_handle: Handle<crate::Expression>,
expressions: &crate::Arena<crate::Expression>,
module: &crate::Module,
put_expression: E,
) -> BackendResult
where
E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
{
match expressions[expr_handle] {
crate::Expression::Literal(literal) => match literal {
crate::Literal::F64(_) => {
return Err(Error::CapabilityNotSupported(valid::Capabilities::FLOAT64))
}
crate::Literal::F32(value) => {
if value.is_infinite() {
let sign = if value.is_sign_negative() { "-" } else { "" };
write!(self.out, "{sign}INFINITY")?;
} else if value.is_nan() {
write!(self.out, "NAN")?;
} else {
let suffix = if value.fract() == 0.0 { ".0" } else { "" };
write!(self.out, "{value}{suffix}")?;
}
}
crate::Literal::U32(value) => {
write!(self.out, "{value}u")?;
}
crate::Literal::I32(value) => {
write!(self.out, "{value}")?;
}
crate::Literal::Bool(value) => {
write!(self.out, "{value}")?;
}
},
crate::Expression::Constant(handle) => {
let constant = &module.constants[handle];
if constant.name.is_some() {
write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
} else {
self.put_const_expression(constant.init, module)?;
}
}
crate::Expression::ZeroValue(ty) => {
let ty_name = TypeContext {
handle: ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
binding: None,
first_time: false,
};
write!(self.out, "{ty_name} {{}}")?;
}
crate::Expression::Compose { ty, ref components } => {
let ty_name = TypeContext {
handle: ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
binding: None,
first_time: false,
};
write!(self.out, "{ty_name}")?;
match module.types[ty].inner {
crate::TypeInner::Scalar { .. }
| crate::TypeInner::Vector { .. }
| crate::TypeInner::Matrix { .. } => {
self.put_call_parameters_impl(components.iter().copied(), put_expression)?;
}
crate::TypeInner::Array { .. } | crate::TypeInner::Struct { .. } => {
write!(self.out, " {{")?;
for (index, &component) in components.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
// insert padding initialization, if needed
if self.struct_member_pads.contains(&(ty, index as u32)) {
write!(self.out, "{{}}, ")?;
}
put_expression(self, component)?;
}
write!(self.out, "}}")?;
}
_ => return Err(Error::UnsupportedCompose(ty)),
}
}
_ => unreachable!(),
}
Ok(())
}
/// Emit code for the expression `expr_handle`.
///
/// The `is_scoped` argument is true if the surrounding operators have the
@@ -1291,6 +1319,17 @@ impl<W: Write> Writer<W> {
let expression = &context.function.expressions[expr_handle];
log::trace!("expression {:?} = {:?}", expr_handle, expression);
match *expression {
crate::Expression::Literal(_)
| crate::Expression::Constant(_)
| crate::Expression::ZeroValue(_)
| crate::Expression::Compose { .. } => {
self.put_possibly_const_expression(
expr_handle,
&context.function.expressions,
context.module,
|writer, expr| writer.put_expression(expr, context, true),
)?;
}
crate::Expression::Access { base, .. }
| crate::Expression::AccessIndex { base, .. } => {
// This is an acceptable place to generate a `ReadZeroSkipWrite` check.
@@ -1318,51 +1357,6 @@ impl<W: Write> Writer<W> {
self.put_access_chain(expr_handle, policy, context)?;
}
}
crate::Expression::Constant(handle) => {
let coco = ConstantContext {
handle,
arena: &context.module.constants,
names: &self.names,
first_time: false,
};
write!(self.out, "{coco}")?;
}
crate::Expression::ZeroValue(ty) => {
let ty_name = TypeContext {
handle: ty,
gctx: context.module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
binding: None,
first_time: false,
};
write!(self.out, "{ty_name} {{}}")?;
}
crate::Expression::Literal(literal) => match literal {
crate::Literal::F64(_) => {
return Err(Error::CapabilityNotSupported(valid::Capabilities::FLOAT64))
}
crate::Literal::F32(value) => {
if value.is_infinite() {
let sign = if value.is_sign_negative() { "-" } else { "" };
write!(self.out, "{sign}INFINITY")?;
} else if value.is_nan() {
write!(self.out, "NAN")?;
} else {
let suffix = if value.fract() == 0.0 { ".0" } else { "" };
write!(self.out, "{value}{suffix}")?;
}
}
crate::Literal::U32(value) => {
write!(self.out, "{value}u")?;
}
crate::Literal::I32(value) => {
write!(self.out, "{value}")?;
}
crate::Literal::Bool(value) => {
write!(self.out, "{value}")?;
}
},
crate::Expression::Splat { size, value } => {
let scalar_kind = match *context.resolve_type(value) {
crate::TypeInner::Scalar { kind, .. } => kind,
@@ -1384,9 +1378,6 @@ impl<W: Write> Writer<W> {
write!(self.out, "{}", back::COMPONENTS[sc as usize])?;
}
}
crate::Expression::Compose { ty, ref components } => {
self.put_compose(ty, components, context)?;
}
crate::Expression::FunctionArgument(index) => {
let name_key = match context.origin {
FunctionOrigin::Handle(handle) => NameKey::FunctionArgument(handle, index),
@@ -1448,15 +1439,11 @@ impl<W: Write> Writer<W> {
self.put_image_sample_level(image, level, context)?;
if let Some(constant) = offset {
let coco = ConstantContext {
handle: constant,
arena: &context.module.constants,
names: &self.names,
first_time: false,
};
write!(self.out, ", {coco}")?;
if let Some(offset) = offset {
write!(self.out, ", ")?;
self.put_const_expression(offset, context.module)?;
}
match gather {
None | Some(crate::SwizzleComponent::X) => {}
Some(component) => {
@@ -2270,15 +2257,11 @@ impl<W: Write> Writer<W> {
policy: index::BoundsCheckPolicy,
context: &ExpressionContext,
) -> BackendResult {
let is_atomic = match *context.resolve_type(pointer) {
crate::TypeInner::Pointer { base, .. } => match context.module.types[base].inner {
crate::TypeInner::Atomic { .. } => true,
_ => false,
},
_ => false,
};
let is_atomic_pointer = context
.resolve_type(pointer)
.is_atomic_pointer(&context.module.types);
if is_atomic {
if is_atomic_pointer {
write!(
self.out,
"{NAMESPACE}::atomic_load_explicit({ATOMIC_REFERENCE}"
@@ -2420,9 +2403,8 @@ impl<W: Write> Writer<W> {
// check what kind of product this is depending
// on the resolve type of the Dot function itself
if let crate::TypeInner::Scalar { kind, .. } =
*context.resolve_type(expr_handle)
{
let inner = context.resolve_type(expr_handle);
if let crate::TypeInner::Scalar { kind, .. } = *inner {
match kind {
crate::ScalarKind::Sint | crate::ScalarKind::Uint => {
self.need_bake_expressions.insert(arg);
@@ -2991,18 +2973,12 @@ impl<W: Write> Writer<W> {
level: back::Level,
context: &StatementContext,
) -> BackendResult {
let pointer_inner = context.expression.resolve_type(pointer);
let is_atomic = match *pointer_inner {
crate::TypeInner::Pointer { base, .. } => {
match context.expression.module.types[base].inner {
crate::TypeInner::Atomic { .. } => true,
_ => false,
}
}
_ => false,
};
let is_atomic_pointer = context
.expression
.resolve_type(pointer)
.is_atomic_pointer(&context.expression.module.types);
if is_atomic {
if is_atomic_pointer {
write!(
self.out,
"{level}{NAMESPACE}::atomic_store_explicit({ATOMIC_REFERENCE}"
@@ -3107,9 +3083,8 @@ impl<W: Write> Writer<W> {
}
};
self.write_scalar_constants(module)?;
self.write_type_defs(module)?;
self.write_composite_constants(module)?;
self.write_global_constants(module)?;
self.write_functions(module, info, options, pipeline_options)
}
@@ -3264,78 +3239,25 @@ impl<W: Write> Writer<W> {
Ok(())
}
fn write_scalar_constants(&mut self, module: &crate::Module) -> BackendResult {
for (handle, constant) in module.constants.iter() {
match constant.inner {
crate::ConstantInner::Scalar {
width: _,
ref value,
} if constant.name.is_some() => {
debug_assert!(constant.needs_alias());
write!(self.out, "constexpr constant ")?;
match *value {
crate::ScalarValue::Sint(_) => {
write!(self.out, "int")?;
}
crate::ScalarValue::Uint(_) => {
write!(self.out, "unsigned")?;
}
crate::ScalarValue::Float(_) => {
write!(self.out, "float")?;
}
crate::ScalarValue::Bool(_) => {
write!(self.out, "bool")?;
}
}
let name = &self.names[&NameKey::Constant(handle)];
let coco = ConstantContext {
handle,
arena: &module.constants,
names: &self.names,
first_time: true,
};
writeln!(self.out, " {name} = {coco};")?;
}
_ => {}
}
}
Ok(())
}
/// Writes all named constants
fn write_global_constants(&mut self, module: &crate::Module) -> BackendResult {
let constants = module.constants.iter().filter(|&(_, c)| c.name.is_some());
fn write_composite_constants(&mut self, module: &crate::Module) -> BackendResult {
for (handle, constant) in module.constants.iter() {
match constant.inner {
crate::ConstantInner::Scalar { .. } => {}
crate::ConstantInner::Composite { ty, ref components } => {
debug_assert!(constant.needs_alias());
let name = &self.names[&NameKey::Constant(handle)];
let ty_name = TypeContext {
handle: ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
binding: None,
first_time: false,
};
write!(self.out, "constant {ty_name} {name} = {{",)?;
for (i, &sub_handle) in components.iter().enumerate() {
// insert padding initialization, if needed
if self.struct_member_pads.contains(&(ty, i as u32)) {
write!(self.out, ", {{}}")?;
}
let separator = if i != 0 { ", " } else { "" };
let coco = ConstantContext {
handle: sub_handle,
arena: &module.constants,
names: &self.names,
first_time: false,
};
write!(self.out, "{separator}{coco}")?;
}
writeln!(self.out, "}};")?;
}
}
for (handle, constant) in constants {
let ty_name = TypeContext {
handle: constant.ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
binding: None,
first_time: false,
};
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, "constant {ty_name} {name} = ")?;
self.put_const_expression(constant.init, module)?;
writeln!(self.out, ";")?;
}
Ok(())
}
@@ -3528,13 +3450,8 @@ impl<W: Write> Writer<W> {
write!(self.out, "{}{} {}", back::INDENT, ty_name, local_name)?;
match local.init {
Some(value) => {
let coco = ConstantContext {
handle: value,
arena: &module.constants,
names: &self.names,
first_time: false,
};
write!(self.out, " = {coco}")?;
write!(self.out, " = ")?;
self.put_const_expression(value, module)?;
}
None => {
write!(self.out, " = {{}}")?;
@@ -3935,13 +3852,8 @@ impl<W: Write> Writer<W> {
resolved.try_fmt(&mut self.out)?;
}
if let Some(value) = var.init {
let coco = ConstantContext {
handle: value,
arena: &module.constants,
names: &self.names,
first_time: false,
};
write!(self.out, " = {coco}")?;
write!(self.out, " = ")?;
self.put_const_expression(value, module)?;
}
writeln!(self.out)?;
}
@@ -3996,13 +3908,9 @@ impl<W: Write> Writer<W> {
tyvar.try_fmt(&mut self.out)?;
match var.init {
Some(value) => {
let coco = ConstantContext {
handle: value,
arena: &module.constants,
names: &self.names,
first_time: false,
};
writeln!(self.out, " = {coco};")?;
write!(self.out, " = ")?;
self.put_const_expression(value, module)?;
writeln!(self.out, ";")?;
}
None => {
writeln!(self.out, " = {{}};")?;
@@ -4100,13 +4008,8 @@ impl<W: Write> Writer<W> {
write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?;
match local.init {
Some(value) => {
let coco = ConstantContext {
handle: value,
arena: &module.constants,
names: &self.names,
first_time: false,
};
write!(self.out, " = {coco}")?;
write!(self.out, " = ")?;
self.put_const_expression(value, module)?;
}
None => {
write!(self.out, " = {{}}")?;
@@ -4373,21 +4276,11 @@ fn test_stack_size() {
use crate::valid::{Capabilities, ValidationFlags};
// create a module with at least one expression nested
let mut module = crate::Module::default();
let constant = module.constants.append(
crate::Constant {
name: None,
specialization: None,
inner: crate::ConstantInner::Scalar {
value: crate::ScalarValue::Float(1.0),
width: 4,
},
},
let mut fun = crate::Function::default();
let const_expr = fun.expressions.append(
crate::Expression::Literal(crate::Literal::F32(1.0)),
Default::default(),
);
let mut fun = crate::Function::default();
let const_expr = fun
.expressions
.append(crate::Expression::Constant(constant), Default::default());
let nested_expr = fun.expressions.append(
crate::Expression::Unary {
op: crate::UnaryOperator::Negate,

43
src/back/spv/block.rs
View File

@@ -237,6 +237,29 @@ impl<'w> BlockContext<'w> {
let result_type_id = self.get_expression_type_id(&self.fun_info[expr_handle].ty);
let id = match self.ir_function.expressions[expr_handle] {
crate::Expression::Literal(literal) => self.writer.get_constant_scalar(literal),
crate::Expression::Constant(handle) => {
let init = self.ir_module.constants[handle].init;
self.writer.constant_ids[init.index()]
}
crate::Expression::ZeroValue(_) => self.writer.write_constant_null(result_type_id),
crate::Expression::Compose {
ty: _,
ref components,
} => {
self.temp_list.clear();
for &component in components {
self.temp_list.push(self.cached[component]);
}
let id = self.gen_id();
block.body.push(Instruction::composite_construct(
result_type_id,
id,
&self.temp_list,
));
id
}
crate::Expression::Access { base, index: _ } if self.is_intermediate(base) => {
// See `is_intermediate`; we'll handle this later in
// `write_expression_pointer`.
@@ -372,9 +395,6 @@ impl<'w> BlockContext<'w> {
crate::Expression::GlobalVariable(handle) => {
self.writer.global_variables[handle.index()].access_id
}
crate::Expression::Constant(handle) => self.writer.constant_ids[handle.index()],
crate::Expression::ZeroValue(_) => self.writer.write_constant_null(result_type_id),
crate::Expression::Literal(literal) => self.writer.get_constant_scalar(literal),
crate::Expression::Splat { size, value } => {
let value_id = self.cached[value];
let components = [value_id; 4];
@@ -406,23 +426,6 @@ impl<'w> BlockContext<'w> {
));
id
}
crate::Expression::Compose {
ty: _,
ref components,
} => {
self.temp_list.clear();
for &component in components {
self.temp_list.push(self.cached[component]);
}
let id = self.gen_id();
block.body.push(Instruction::composite_construct(
result_type_id,
id,
&self.temp_list,
));
id
}
crate::Expression::Unary { op, expr } => {
let id = self.gen_id();
let expr_id = self.cached[expr];

2
src/back/spv/image.rs
View File

@@ -823,7 +823,7 @@ impl<'w> BlockContext<'w> {
gather: Option<crate::SwizzleComponent>,
coordinate: Handle<crate::Expression>,
array_index: Option<Handle<crate::Expression>>,
offset: Option<Handle<crate::Constant>>,
offset: Option<Handle<crate::Expression>>,
level: crate::SampleLevel,
depth_ref: Option<Handle<crate::Expression>>,
block: &mut Block,

76
src/back/spv/index.rs
View File

@@ -176,17 +176,19 @@ impl<'w> BlockContext<'w> {
// done the bounds check.
let max_index_id = match self.write_sequence_max_index(sequence, block)? {
MaybeKnown::Known(known_max_index) => {
if let crate::Expression::Constant(index_k) = self.ir_function.expressions[index] {
if let Some(known_index) = self.ir_module.constants[index_k].to_array_length() {
// Both the index and length are known at compile time.
//
// In strict WGSL compliance mode, out-of-bounds indices cannot be
// reported at shader translation time, and must be replaced with
// in-bounds indices at run time. So we cannot assume that
// validation ensured the index was in bounds. Restrict now.
let restricted = std::cmp::min(known_index, known_max_index);
return Ok(BoundsCheckResult::KnownInBounds(restricted));
}
if let Ok(known_index) = self
.ir_module
.to_ctx()
.eval_expr_to_u32_from(index, &self.ir_function.expressions)
{
// Both the index and length are known at compile time.
//
// In strict WGSL compliance mode, out-of-bounds indices cannot be
// reported at shader translation time, and must be replaced with
// in-bounds indices at run time. So we cannot assume that
// validation ensured the index was in bounds. Restrict now.
let restricted = std::cmp::min(known_index, known_max_index);
return Ok(BoundsCheckResult::KnownInBounds(restricted));
}
self.get_index_constant(known_max_index)
@@ -236,31 +238,33 @@ impl<'w> BlockContext<'w> {
// bounds check.
let length_id = match self.write_sequence_length(sequence, block)? {
MaybeKnown::Known(known_length) => {
if let crate::Expression::Constant(index_k) = self.ir_function.expressions[index] {
if let Some(known_index) = self.ir_module.constants[index_k].to_array_length() {
// Both the index and length are known at compile time.
//
// It would be nice to assume that, since we are using the
// `ReadZeroSkipWrite` policy, we are not in strict WGSL
// compliance mode, and thus we can count on the validator to have
// rejected any programs with known out-of-bounds indices, and
// thus just return `KnownInBounds` here without actually
// checking.
//
// But it's also reasonable to expect that bounds check policies
// and error reporting policies should be able to vary
// independently without introducing security holes. So, we should
// support the case where bad indices do not cause validation
// errors, and are handled via `ReadZeroSkipWrite`.
//
// In theory, when `known_index` is bad, we could return a new
// `KnownOutOfBounds` variant here. But it's simpler just to fall
// through and let the bounds check take place. The shader is
// broken anyway, so it doesn't make sense to invest in emitting
// the ideal code for it.
if known_index < known_length {
return Ok(BoundsCheckResult::KnownInBounds(known_index));
}
if let Ok(known_index) = self
.ir_module
.to_ctx()
.eval_expr_to_u32_from(index, &self.ir_function.expressions)
{
// Both the index and length are known at compile time.
//
// It would be nice to assume that, since we are using the
// `ReadZeroSkipWrite` policy, we are not in strict WGSL
// compliance mode, and thus we can count on the validator to have
// rejected any programs with known out-of-bounds indices, and
// thus just return `KnownInBounds` here without actually
// checking.
//
// But it's also reasonable to expect that bounds check policies
// and error reporting policies should be able to vary
// independently without introducing security holes. So, we should
// support the case where bad indices do not cause validation
// errors, and are handled via `ReadZeroSkipWrite`.
//
// In theory, when `known_index` is bad, we could return a new
// `KnownOutOfBounds` variant here. But it's simpler just to fall
// through and let the bounds check take place. The shader is
// broken anyway, so it doesn't make sense to invest in emitting
// the ideal code for it.
if known_index < known_length {
return Ok(BoundsCheckResult::KnownInBounds(known_index));
}
}

1
src/back/spv/mod.rs
View File

@@ -616,6 +616,7 @@ pub struct Writer {
lookup_type: crate::FastHashMap<LookupType, Word>,
lookup_function: crate::FastHashMap<Handle<crate::Function>, Word>,
lookup_function_type: crate::FastHashMap<LookupFunctionType, Word>,
/// Indexed by const-expression handle indexes
constant_ids: Vec<Word>,
cached_constants: crate::FastHashMap<CachedConstant, Word>,
global_variables: Vec<GlobalVariable>,

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

@@ -1213,6 +1213,36 @@ impl Writer {
null_id
}
fn write_constant_expr(
&mut self,
handle: Handle<crate::Expression>,
ir_module: &crate::Module,
) -> Result<Word, Error> {
let id = match ir_module.const_expressions[handle] {
crate::Expression::Literal(literal) => self.get_constant_scalar(literal),
crate::Expression::Constant(constant) => {
let constant = &ir_module.constants[constant];
self.constant_ids[constant.init.index()]
}
crate::Expression::ZeroValue(ty) => {
let type_id = self.get_type_id(LookupType::Handle(ty));
self.write_constant_null(type_id)
}
crate::Expression::Compose { ty, ref components } => {
let component_ids: Vec<_> = components
.iter()
.map(|component| self.constant_ids[component.index()])
.collect();
self.get_constant_composite(LookupType::Handle(ty), component_ids.as_slice())
}
_ => unreachable!(),
};
self.constant_ids[handle.index()] = id;
Ok(id)
}
pub(super) fn write_barrier(&mut self, flags: crate::Barrier, block: &mut Block) {
let memory_scope = if flags.contains(crate::Barrier::STORAGE) {
spirv::Scope::Device
@@ -1796,59 +1826,30 @@ impl Writer {
}
}
self.constant_ids.resize(ir_module.constants.len(), 0);
// first, output all the scalar constants
for (handle, constant) in ir_module.constants.iter() {
match constant.inner {
crate::ConstantInner::Composite { .. } => continue,
crate::ConstantInner::Scalar { width, ref value } => {
let literal = crate::Literal::from_scalar(*value, width).ok_or(
Error::Validation("Unexpected kind and/or width for Literal"),
)?;
self.constant_ids[handle.index()] = match constant.name {
Some(ref name) => {
let id = self.id_gen.next();
self.write_constant_scalar(id, &literal, Some(name));
id
}
None => self.get_constant_scalar(literal),
};
}
}
}
// then all types, some of them may rely on constants and struct type set
// write all types
for (handle, _) in ir_module.types.iter() {
self.write_type_declaration_arena(&ir_module.types, handle)?;
}
// then all the composite constants, they rely on types
for (handle, constant) in ir_module.constants.iter() {
match constant.inner {
crate::ConstantInner::Scalar { .. } => continue,
crate::ConstantInner::Composite { ty, ref components } => {
let ty = LookupType::Handle(ty);
// write all const-expressions as constants
self.constant_ids
.resize(ir_module.const_expressions.len(), 0);
for (handle, _) in ir_module.const_expressions.iter() {
self.write_constant_expr(handle, ir_module)?;
}
debug_assert!(self.constant_ids.iter().all(|&id| id != 0));
let mut constituent_ids = Vec::with_capacity(components.len());
for constituent in components.iter() {
let constituent_id = self.constant_ids[constituent.index()];
constituent_ids.push(constituent_id);
}
self.constant_ids[handle.index()] = match constant.name {
Some(ref name) => {
let id = self.id_gen.next();
self.write_constant_composite(id, ty, &constituent_ids, Some(name));
id
}
None => self.get_constant_composite(ty, &constituent_ids),
};
// write the name of constants on their respective const-expression initializer
if self.flags.contains(WriterFlags::DEBUG) {
for (_, constant) in ir_module.constants.iter() {
if let Some(ref name) = constant.name {
let id = self.constant_ids[constant.init.index()];
self.debugs.push(Instruction::name(id, name));
}
}
}
debug_assert_eq!(self.constant_ids.iter().position(|&id| id == 0), None);
// now write all globals
// write all global variables
for (handle, var) in ir_module.global_variables.iter() {
// If a single entry point was specified, only write `OpVariable` instructions
// for the globals it actually uses. Emit dummies for the others,
@@ -1865,7 +1866,7 @@ impl Writer {
self.global_variables.push(gvar);
}
// all functions
// write all functions
for (handle, ir_function) in ir_module.functions.iter() {
let info = &mod_info[handle];
if let Some(index) = ep_index {
@@ -1882,7 +1883,7 @@ impl Writer {
self.lookup_function.insert(handle, id);
}
// and entry points
// write all or one entry points
for (index, ir_ep) in ir_module.entry_points.iter().enumerate() {
if ep_index.is_some() && ep_index != Some(index) {
continue;

298
src/back/wgsl/writer.rs
View File

@@ -118,10 +118,17 @@ impl<W: Write> Writer<W> {
}
}
// Write all constants
for (handle, constant) in module.constants.iter() {
if constant.name.is_some() {
self.write_global_constant(module, &constant.inner, handle)?;
// Write all named constants
let mut constants = module
.constants
.iter()
.filter(|&(_, c)| c.name.is_some())
.peekable();
while let Some((handle, _)) = constants.next() {
self.write_global_constant(module, handle)?;
// Add extra newline for readability on last iteration
if constants.peek().is_none() {
writeln!(self.out)?;
}
}
@@ -182,25 +189,6 @@ impl<W: Write> Writer<W> {
Ok(())
}
/// Helper method used to write [`ScalarValue`](crate::ScalarValue)
///
/// # Notes
/// Adds no trailing or leading whitespace
fn write_scalar_value(&mut self, value: crate::ScalarValue) -> BackendResult {
use crate::ScalarValue as Sv;
match value {
Sv::Sint(value) => write!(self.out, "{value}")?,
Sv::Uint(value) => write!(self.out, "{value}u")?,
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
Sv::Float(value) => write!(self.out, "{value:?}")?,
Sv::Bool(value) => write!(self.out, "{value}")?,
}
Ok(())
}
/// Helper method used to write struct name
///
/// # Notes
@@ -302,7 +290,7 @@ impl<W: Write> Writer<W> {
// Write the constant
// `write_constant` adds no trailing or leading space/newline
self.write_constant(module, init)?;
self.write_const_expression(module, init)?;
}
// Finish the local with `;` and add a newline (only for readability)
@@ -708,14 +696,11 @@ impl<W: Write> Writer<W> {
Statement::Store { pointer, value } => {
write!(self.out, "{level}")?;
let is_atomic = match *func_ctx.info[pointer].ty.inner_with(&module.types) {
crate::TypeInner::Pointer { base, .. } => match module.types[base].inner {
crate::TypeInner::Atomic { .. } => true,
_ => false,
},
_ => false,
};
if is_atomic {
let is_atomic_pointer = func_ctx
.resolve_type(pointer, &module.types)
.is_atomic_pointer(&module.types);
if is_atomic_pointer {
write!(self.out, "atomicStore(")?;
self.write_expr(module, pointer, func_ctx)?;
write!(self.out, ", ")?;
@@ -747,12 +732,10 @@ impl<W: Write> Writer<W> {
let func_name = &self.names[&NameKey::Function(function)];
write!(self.out, "{func_name}(")?;
for (index, &argument) in arguments.iter().enumerate() {
self.write_expr(module, argument, func_ctx)?;
// Only write a comma if isn't the last element
if index != arguments.len().saturating_sub(1) {
// The leading space is for readability only
if index != 0 {
write!(self.out, ", ")?;
}
self.write_expr(module, argument, func_ctx)?;
}
writeln!(self.out, ");")?
}
@@ -987,7 +970,7 @@ impl<W: Write> Writer<W> {
}
}
Ex::Access { base, .. } | Ex::AccessIndex { base, .. } => {
let base_ty = func_ctx.info[base].ty.inner_with(&module.types);
let base_ty = func_ctx.resolve_type(base, &module.types);
match *base_ty {
crate::TypeInner::Pointer { .. } | crate::TypeInner::ValuePointer { .. } => {
Indirection::Reference
@@ -1077,6 +1060,74 @@ impl<W: Write> Writer<W> {
Ok(())
}
fn write_const_expression(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
) -> BackendResult {
self.write_possibly_const_expression(
module,
expr,
&module.const_expressions,
|writer, expr| writer.write_const_expression(module, expr),
)
}
fn write_possibly_const_expression<E>(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
expressions: &crate::Arena<crate::Expression>,
write_expression: E,
) -> BackendResult
where
E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
{
use crate::Expression;
match expressions[expr] {
Expression::Literal(literal) => {
match literal {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
crate::Literal::F64(_) => {
return Err(Error::Custom("unsupported f64 literal".to_string()));
}
crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
}
}
Expression::Constant(handle) => {
let constant = &module.constants[handle];
if constant.name.is_some() {
write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
} else {
self.write_const_expression(module, constant.init)?;
}
}
Expression::ZeroValue(ty) => {
self.write_type(module, ty)?;
write!(self.out, "()")?;
}
Expression::Compose { ty, ref components } => {
self.write_type(module, ty)?;
write!(self.out, "(")?;
for (index, component) in components.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
write_expression(self, *component)?;
}
write!(self.out, ")")?
}
_ => unreachable!(),
}
Ok(())
}
/// Write the 'plain form' of `expr`.
///
/// An expression's 'plain form' is the most general rendition of that
@@ -1109,36 +1160,16 @@ impl<W: Write> Writer<W> {
// `postfix_expression` forms for member/component access and
// subscripting.
match *expression {
Expression::Literal(literal) => {
match literal {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
crate::Literal::F64(_) => {
return Err(Error::Custom("unsupported f64 literal".to_string()));
}
crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
}
}
Expression::Constant(constant) => self.write_constant(module, constant)?,
Expression::ZeroValue(ty) => {
self.write_type(module, ty)?;
write!(self.out, "()")?;
}
Expression::Compose { ty, ref components } => {
self.write_type(module, ty)?;
write!(self.out, "(")?;
for (index, component) in components.iter().enumerate() {
self.write_expr(module, *component, func_ctx)?;
// Only write a comma if isn't the last element
if index != components.len().saturating_sub(1) {
// The leading space is for readability only
write!(self.out, ", ")?;
}
}
write!(self.out, ")")?
Expression::Literal(_)
| Expression::Constant(_)
| Expression::ZeroValue(_)
| Expression::Compose { .. } => {
self.write_possibly_const_expression(
module,
expr,
func_ctx.expressions,
|writer, expr| writer.write_expr(module, expr, func_ctx),
)?;
}
Expression::FunctionArgument(pos) => {
let name_key = func_ctx.argument_key(pos);
@@ -1261,11 +1292,12 @@ impl<W: Write> Writer<W> {
if let Some(offset) = offset {
write!(self.out, ", ")?;
self.write_constant(module, offset)?;
self.write_const_expression(module, offset)?;
}
write!(self.out, ")")?;
}
Expression::ImageSample {
image,
sampler,
@@ -1282,7 +1314,7 @@ impl<W: Write> Writer<W> {
};
write!(self.out, "textureGather{suffix_cmp}(")?;
match *func_ctx.info[image].ty.inner_with(&module.types) {
match *func_ctx.resolve_type(image, &module.types) {
TypeInner::Image {
class: crate::ImageClass::Depth { multi: _ },
..
@@ -1309,7 +1341,7 @@ impl<W: Write> Writer<W> {
if let Some(offset) = offset {
write!(self.out, ", ")?;
self.write_constant(module, offset)?;
self.write_const_expression(module, offset)?;
}
write!(self.out, ")")?;
@@ -1332,6 +1364,7 @@ impl<W: Write> Writer<W> {
};
write!(self.out, ")")?;
}
Expression::ImageLoad {
image,
coordinate,
@@ -1357,12 +1390,13 @@ impl<W: Write> Writer<W> {
let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, "{name}")?;
}
Expression::As {
expr,
kind,
convert,
} => {
let inner = func_ctx.info[expr].ty.inner_with(&module.types);
let inner = func_ctx.resolve_type(expr, &module.types);
match *inner {
TypeInner::Matrix {
columns,
@@ -1407,7 +1441,7 @@ impl<W: Write> Writer<W> {
write!(self.out, ")")?;
}
Expression::Splat { size, value } => {
let inner = func_ctx.info[value].ty.inner_with(&module.types);
let inner = func_ctx.resolve_type(value, &module.types);
let (scalar_kind, scalar_width) = match *inner {
crate::TypeInner::Scalar { kind, width } => (kind, width),
_ => {
@@ -1424,15 +1458,11 @@ impl<W: Write> Writer<W> {
write!(self.out, ")")?;
}
Expression::Load { pointer } => {
let is_atomic = match *func_ctx.info[pointer].ty.inner_with(&module.types) {
crate::TypeInner::Pointer { base, .. } => match module.types[base].inner {
crate::TypeInner::Atomic { .. } => true,
_ => false,
},
_ => false,
};
let is_atomic_pointer = func_ctx
.resolve_type(pointer, &module.types)
.is_atomic_pointer(&module.types);
if is_atomic {
if is_atomic_pointer {
write!(self.out, "atomicLoad(")?;
self.write_expr(module, pointer, func_ctx)?;
write!(self.out, ")")?;
@@ -1453,6 +1483,7 @@ impl<W: Write> Writer<W> {
self.write_expr(module, expr, func_ctx)?;
write!(self.out, ")")?;
}
Expression::Math {
fun,
arg,
@@ -1561,6 +1592,7 @@ impl<W: Write> Writer<W> {
}
}
}
Expression::Swizzle {
size,
vector,
@@ -1576,7 +1608,7 @@ impl<W: Write> Writer<W> {
let unary = match op {
crate::UnaryOperator::Negate => "-",
crate::UnaryOperator::Not => {
match *func_ctx.info[expr].ty.inner_with(&module.types) {
match *func_ctx.resolve_type(expr, &module.types) {
TypeInner::Scalar {
kind: crate::ScalarKind::Bool,
..
@@ -1592,6 +1624,7 @@ impl<W: Write> Writer<W> {
write!(self.out, ")")?
}
Expression::Select {
condition,
accept,
@@ -1688,7 +1721,7 @@ impl<W: Write> Writer<W> {
// Write initializer
if let Some(init) = global.init {
write!(self.out, " = ")?;
self.write_constant(module, init)?;
self.write_const_expression(module, init)?;
}
// End with semicolon
@@ -1697,47 +1730,6 @@ impl<W: Write> Writer<W> {
Ok(())
}
/// Helper method used to write constants
///
/// # Notes
/// Doesn't add any newlines or leading/trailing spaces
fn write_constant(
&mut self,
module: &Module,
handle: Handle<crate::Constant>,
) -> BackendResult {
let constant = &module.constants[handle];
match constant.inner {
crate::ConstantInner::Scalar {
width: _,
ref value,
} => {
if constant.name.is_some() {
write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
} else {
self.write_scalar_value(*value)?;
}
}
crate::ConstantInner::Composite { ty, ref components } => {
self.write_type(module, ty)?;
write!(self.out, "(")?;
// Write the comma separated constants
for (index, constant) in components.iter().enumerate() {
self.write_constant(module, *constant)?;
// Only write a comma if isn't the last element
if index != components.len().saturating_sub(1) {
// The leading space is for readability only
write!(self.out, ", ")?;
}
}
write!(self.out, ")")?
}
}
Ok(())
}
/// Helper method used to write global constants
///
/// # Notes
@@ -1745,61 +1737,17 @@ impl<W: Write> Writer<W> {
fn write_global_constant(
&mut self,
module: &Module,
inner: &crate::ConstantInner,
handle: Handle<crate::Constant>,
) -> BackendResult {
match *inner {
crate::ConstantInner::Scalar {
width: _,
ref value,
} => {
let name = &self.names[&NameKey::Constant(handle)];
// First write only constant name
write!(self.out, "const {name}: ")?;
// Next write constant type and value
match *value {
crate::ScalarValue::Sint(value) => {
write!(self.out, "i32 = {value}")?;
}
crate::ScalarValue::Uint(value) => {
write!(self.out, "u32 = {value}u")?;
}
crate::ScalarValue::Float(value) => {
// Floats are written using `Debug` instead of `Display` because it always appends the
// decimal part even it's zero
write!(self.out, "f32 = {value:?}")?;
}
crate::ScalarValue::Bool(value) => {
write!(self.out, "bool = {value}")?;
}
};
// End with semicolon
writeln!(self.out, ";")?;
}
crate::ConstantInner::Composite { ty, ref components } => {
let name = &self.names[&NameKey::Constant(handle)];
// First write only constant name
write!(self.out, "const {name}: ")?;
// Next write constant type
self.write_type(module, ty)?;
let name = &self.names[&NameKey::Constant(handle)];
// First write only constant name
write!(self.out, "const {name}: ")?;
self.write_type(module, module.constants[handle].ty)?;
write!(self.out, " = ")?;
let init = module.constants[handle].init;
self.write_const_expression(module, init)?;
writeln!(self.out, ";")?;
write!(self.out, " = ")?;
self.write_type(module, ty)?;
write!(self.out, "(")?;
for (index, constant) in components.iter().enumerate() {
self.write_constant(module, *constant)?;
// Only write a comma if isn't the last element
if index != components.len().saturating_sub(1) {
// The leading space is for readability only
write!(self.out, ", ")?;
}
}
write!(self.out, ");")?;
}
}
// End with extra newline for readability
writeln!(self.out)?;
Ok(())
}

9
src/front/glsl/builtins.rs
View File

@@ -7,10 +7,9 @@ use super::{
Error, ErrorKind, Frontend, Result,
};
use crate::{
BinaryOperator, Block, Constant, DerivativeAxis as Axis, DerivativeControl as Ctrl, Expression,
Handle, ImageClass, ImageDimension as Dim, ImageQuery, MathFunction, Module,
RelationalFunction, SampleLevel, ScalarKind as Sk, Span, Type, TypeInner, UnaryOperator,
VectorSize,
BinaryOperator, Block, DerivativeAxis as Axis, DerivativeControl as Ctrl, Expression, Handle,
ImageClass, ImageDimension as Dim, ImageQuery, MathFunction, Module, RelationalFunction,
SampleLevel, ScalarKind as Sk, Span, Type, TypeInner, UnaryOperator, VectorSize,
};
impl crate::ScalarKind {
@@ -2183,7 +2182,7 @@ fn texture_call(
image: Handle<Expression>,
level: SampleLevel,
comps: CoordComponents,
offset: Option<Handle<Constant>>,
offset: Option<Handle<Expression>>,
body: &mut Block,
meta: Span,
) -> Result<Handle<Expression>> {

1129
src/front/glsl/constants.rs
View File

File diff suppressed because it is too large Load Diff

41
src/front/glsl/context.rs
View File

@@ -10,10 +10,10 @@ use super::{
use crate::{
front::{Emitter, Typifier},
AddressSpace, Arena, BinaryOperator, Block, Expression, FastHashMap, FunctionArgument, Handle,
Literal, LocalVariable, RelationalFunction, ScalarKind, ScalarValue, Span, Statement, Type,
TypeInner, VectorSize,
Literal, LocalVariable, RelationalFunction, ScalarKind, Span, Statement, Type, TypeInner,
VectorSize,
};
use std::{convert::TryFrom, ops::Index};
use std::ops::Index;
/// The position at which an expression is, used while lowering
#[derive(Clone, Copy, PartialEq, Eq, Debug)]
@@ -519,21 +519,12 @@ impl Context {
.0;
let pointer = maybe_constant_index
.and_then(|constant| {
.and_then(|const_expr| {
Some(self.add_expression(
Expression::AccessIndex {
base,
index: match frontend.module.constants[constant].inner {
crate::ConstantInner::Scalar {
value: ScalarValue::Uint(i),
..
} => u32::try_from(i).ok()?,
crate::ConstantInner::Scalar {
value: ScalarValue::Sint(i),
..
} => u32::try_from(i).ok()?,
_ => return None,
},
index:
frontend.module.to_ctx().eval_expr_to_u32(const_expr).ok()?,
},
meta,
body,
@@ -1039,14 +1030,18 @@ impl Context {
// pointer type which is required for dynamic indexing
if !constant_index {
if let Some((constant, ty)) = var.constant {
let local = self.locals.append(
LocalVariable {
name: None,
ty,
init: Some(constant),
},
Span::default(),
);
let local =
self.locals.append(
LocalVariable {
name: None,
ty,
init: Some(frontend.module.const_expressions.append(
Expression::Constant(constant),
Span::default(),
)),
},
Span::default(),
);
self.add_expression(
Expression::LocalVariable(local),

109
src/front/glsl/functions.rs
View File

@@ -8,8 +8,8 @@ use super::{
};
use crate::{
front::glsl::types::type_power, proc::ensure_block_returns, AddressSpace, Arena, Block,
Constant, ConstantInner, EntryPoint, Expression, Function, FunctionArgument, FunctionResult,
Handle, LocalVariable, ScalarKind, ScalarValue, Span, Statement, StructMember, Type, TypeInner,
EntryPoint, Expression, Function, FunctionArgument, FunctionResult, Handle, Literal,
LocalVariable, ScalarKind, Span, Statement, StructMember, Type, TypeInner,
};
use std::iter;
@@ -24,29 +24,6 @@ struct ProxyWrite {
}
impl Frontend {
fn add_constant_value(
&mut self,
scalar_kind: ScalarKind,
value: u64,
meta: Span,
) -> Handle<Constant> {
let value = match scalar_kind {
ScalarKind::Uint => ScalarValue::Uint(value),
ScalarKind::Sint => ScalarValue::Sint(value as i64),
ScalarKind::Float => ScalarValue::Float(value as f64),
ScalarKind::Bool => unreachable!(),
};
self.module.constants.fetch_or_append(
Constant {
name: None,
specialization: None,
inner: ConstantInner::Scalar { width: 4, value },
},
meta,
)
}
pub(crate) fn function_or_constructor_call(
&mut self,
ctx: &mut Context,
@@ -97,10 +74,10 @@ impl Frontend {
if expr_type.scalar_kind() == Some(ScalarKind::Bool)
&& result_scalar_kind != ScalarKind::Bool =>
{
let c0 = self.add_constant_value(result_scalar_kind, 0u64, meta);
let c1 = self.add_constant_value(result_scalar_kind, 1u64, meta);
let mut reject = ctx.add_expression(Expression::Constant(c0), expr_meta, body);
let mut accept = ctx.add_expression(Expression::Constant(c1), expr_meta, body);
let l0 = Literal::zero(result_scalar_kind, 4).unwrap();
let l1 = Literal::one(result_scalar_kind, 4).unwrap();
let mut reject = ctx.add_expression(Expression::Literal(l0), expr_meta, body);
let mut accept = ctx.add_expression(Expression::Literal(l1), expr_meta, body);
ctx.implicit_splat(self, &mut reject, meta, vector_size)?;
ctx.implicit_splat(self, &mut accept, meta, vector_size)?;
@@ -282,18 +259,9 @@ impl Frontend {
},
meta,
);
let zero_constant = self.module.constants.fetch_or_append(
Constant {
name: None,
specialization: None,
inner: ConstantInner::Scalar {
width,
value: ScalarValue::Float(0.0),
},
},
meta,
);
let zero = ctx.add_expression(Expression::Constant(zero_constant), meta, body);
let zero_literal = Literal::zero(ScalarKind::Float, width).unwrap();
let zero = ctx.add_expression(Expression::Literal(zero_literal), meta, body);
for i in 0..columns as u32 {
components.push(
@@ -323,30 +291,12 @@ impl Frontend {
// (column i, row j) in the argument will be initialized from there. All
// other components will be initialized to the identity matrix.
let zero_constant = self.module.constants.fetch_or_append(
Constant {
name: None,
specialization: None,
inner: ConstantInner::Scalar {
width,
value: ScalarValue::Float(0.0),
},
},
meta,
);
let zero = ctx.add_expression(Expression::Constant(zero_constant), meta, body);
let one_constant = self.module.constants.fetch_or_append(
Constant {
name: None,
specialization: None,
inner: ConstantInner::Scalar {
width,
value: ScalarValue::Float(1.0),
},
},
meta,
);
let one = ctx.add_expression(Expression::Constant(one_constant), meta, body);
let zero_literal = Literal::zero(ScalarKind::Float, width).unwrap();
let one_literal = Literal::one(ScalarKind::Float, width).unwrap();
let zero = ctx.add_expression(Expression::Literal(zero_literal), meta, body);
let one = ctx.add_expression(Expression::Literal(one_literal), meta, body);
let vector_ty = self.module.types.insert(
Type {
name: None,
@@ -406,24 +356,17 @@ impl Frontend {
Ordering::Greater => ctx.vector_resize(rows, vector, meta, body),
})
} else {
let vec_constant = self.module.constants.fetch_or_append(
Constant {
name: None,
specialization: None,
inner: ConstantInner::Composite {
ty: vector_ty,
components: (0..rows as u32)
.map(|r| match r == i {
true => one_constant,
false => zero_constant,
})
.collect(),
},
},
meta,
);
let vec =
ctx.add_expression(Expression::Constant(vec_constant), meta, body);
let compose_expr = Expression::Compose {
ty: vector_ty,
components: (0..rows as u32)
.map(|r| match r == i {
true => one,
false => zero,
})
.collect(),
};
let vec = ctx.add_expression(compose_expr, meta, body);
components.push(vec)
}

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

@@ -9,8 +9,7 @@ use super::{
variables::{GlobalOrConstant, VarDeclaration},
Frontend, Result,
};
use crate::{arena::Handle, Block, Constant, ConstantInner, Expression, ScalarValue, Span, Type};
use core::convert::TryFrom;
use crate::{arena::Handle, proc::U32EvalError, Block, Expression, Span, Type};
use pp_rs::token::{PreprocessorError, Token as PPToken, TokenValue as PPTokenValue};
use std::iter::Peekable;
@@ -192,30 +191,21 @@ impl<'source> ParsingContext<'source> {
}
fn parse_uint_constant(&mut self, frontend: &mut Frontend) -> Result<(u32, Span)> {
let (value, meta) = self.parse_constant_expression(frontend)?;
let (const_expr, meta) = self.parse_constant_expression(frontend)?;
let int = match frontend.module.constants[value].inner {
ConstantInner::Scalar {
value: ScalarValue::Uint(int),
..
} => u32::try_from(int).map_err(|_| Error {
let res = frontend.module.to_ctx().eval_expr_to_u32(const_expr);
let int = match res {
Ok(value) => Ok(value),
Err(U32EvalError::Negative) => Err(Error {
kind: ErrorKind::SemanticError("int constant overflows".into()),
meta,
})?,
ConstantInner::Scalar {
value: ScalarValue::Sint(int),
..
} => u32::try_from(int).map_err(|_| Error {
kind: ErrorKind::SemanticError("int constant overflows".into()),
}),
Err(U32EvalError::NonConst) => Err(Error {
kind: ErrorKind::SemanticError("Expected a uint constant".into()),
meta,
})?,
_ => {
return Err(Error {
kind: ErrorKind::SemanticError("Expected a uint constant".into()),
meta,
})
}
};
}),
}?;
Ok((int, meta))
}
@@ -223,7 +213,7 @@ impl<'source> ParsingContext<'source> {
fn parse_constant_expression(
&mut self,
frontend: &mut Frontend,
) -> Result<(Handle<Constant>, Span)> {
) -> Result<(Handle<Expression>, Span)> {
let mut block = Block::new();
let mut ctx = Context::new(frontend, &mut block);
@@ -412,7 +402,7 @@ impl<'ctx, 'qualifiers> DeclarationContext<'ctx, 'qualifiers> {
frontend: &mut Frontend,
ty: Handle<Type>,
name: String,
init: Option<Handle<Constant>>,
init: Option<Handle<Expression>>,
meta: Span,
) -> Result<Handle<Expression>> {
let decl = VarDeclaration {

6
src/front/glsl/parser/declarations.rs
View File

@@ -244,7 +244,7 @@ impl<'source> ParsingContext<'source> {
.transpose()?;
let is_const = ctx.qualifiers.storage.0 == StorageQualifier::Const;
let maybe_constant = if ctx.external {
let maybe_const_expr = if ctx.external {
if let Some((root, meta)) = init {
match frontend.solve_constant(ctx.ctx, root, meta) {
Ok(res) => Some(res),
@@ -260,9 +260,9 @@ impl<'source> ParsingContext<'source> {
None
};
let pointer = ctx.add_var(frontend, ty, name, maybe_constant, meta)?;
let pointer = ctx.add_var(frontend, ty, name, maybe_const_expr, meta)?;
if let Some((value, _)) = init.filter(|_| maybe_constant.is_none()) {
if let Some((value, _)) = init.filter(|_| maybe_const_expr.is_none()) {
ctx.flush_expressions();
ctx.body.push(Statement::Store { pointer, value }, meta);
}

23
src/front/glsl/parser/functions.rs
View File

@@ -1,5 +1,6 @@
use crate::front::glsl::context::ExprPos;
use crate::front::glsl::Span;
use crate::Literal;
use crate::{
front::glsl::{
ast::ParameterQualifier,
@@ -9,7 +10,7 @@ use crate::{
variables::VarDeclaration,
Error, ErrorKind, Frontend, Result,
},
Block, ConstantInner, Expression, ScalarValue, Statement, SwitchCase, UnaryOperator,
Block, Expression, Statement, SwitchCase, UnaryOperator,
};
impl<'source> ParsingContext<'source> {
@@ -194,20 +195,16 @@ impl<'source> ParsingContext<'source> {
let expr = self.parse_expression(frontend, ctx, &mut stmt, body)?;
let (root, meta) =
ctx.lower_expect(stmt, frontend, expr, ExprPos::Rhs, body)?;
let constant = frontend.solve_constant(ctx, root, meta)?;
let const_expr = frontend.solve_constant(ctx, root, meta)?;
match frontend.module.constants[constant].inner {
ConstantInner::Scalar {
value: ScalarValue::Sint(int),
..
} => match uint {
true => crate::SwitchValue::U32(int as u32),
false => crate::SwitchValue::I32(int as i32),
match frontend.module.const_expressions[const_expr] {
Expression::Literal(Literal::I32(value)) => match uint {
true => crate::SwitchValue::U32(value as u32),
false => crate::SwitchValue::I32(value),
},
ConstantInner::Scalar {
value: ScalarValue::Uint(int),
..
} => crate::SwitchValue::U32(int as u32),
Expression::Literal(Literal::U32(value)) => {
crate::SwitchValue::U32(value)
}
_ => {
frontend.errors.push(Error {
kind: ErrorKind::SemanticError(

56
src/front/glsl/parser_tests.rs
View File

@@ -509,7 +509,8 @@ fn functions() {
#[test]
fn constants() {
use crate::{Constant, ConstantInner, ScalarValue};
use crate::{Constant, Expression, ScalarKind, Type, TypeInner};
let mut frontend = Frontend::default();
let module = frontend
@@ -526,28 +527,49 @@ fn constants() {
)
.unwrap();
let mut types = module.types.iter();
let mut constants = module.constants.iter();
let mut const_expressions = module.const_expressions.iter();
let (ty_handle, ty) = types.next().unwrap();
assert_eq!(
ty,
&Type {
name: None,
inner: TypeInner::Scalar {
kind: ScalarKind::Float,
width: 4
}
}
);
let (init_a_handle, init_a) = const_expressions.next().unwrap();
assert_eq!(init_a, &Expression::Literal(crate::Literal::F32(1.0)));
let (constant_a_handle, constant_a) = constants.next().unwrap();
assert_eq!(
constant_a,
&Constant {
name: Some("a".to_owned()),
r#override: crate::Override::None,
ty: ty_handle,
init: init_a_handle
}
);
// skip const expr that was inserted for `global` var
const_expressions.next().unwrap();
let (init_b_handle, init_b) = const_expressions.next().unwrap();
assert_eq!(init_b, &Expression::Constant(constant_a_handle));
assert_eq!(
constants.next().unwrap().1,
&Constant {
name: Some("a".to_owned()),
specialization: None,
inner: ConstantInner::Scalar {
width: 4,
value: ScalarValue::Float(1.0)
}
}
);
assert_eq!(
constants.next().unwrap().1,
&Constant {
name: Some("b".to_owned()),
specialization: None,
inner: ConstantInner::Scalar {
width: 4,
value: ScalarValue::Float(1.0)
}
r#override: crate::Override::None,
ty: ty_handle,
init: init_b_handle
}
);

7
src/front/glsl/types.rs
View File

@@ -2,8 +2,8 @@ use super::{
constants::ConstantSolver, context::Context, Error, ErrorKind, Frontend, Result, Span,
};
use crate::{
proc::ResolveContext, Bytes, Constant, Expression, Handle, ImageClass, ImageDimension,
ScalarKind, Type, TypeInner, VectorSize,
proc::ResolveContext, Bytes, Expression, Handle, ImageClass, ImageDimension, ScalarKind, Type,
TypeInner, VectorSize,
};
pub fn parse_type(type_name: &str) -> Option<Type> {
@@ -320,11 +320,12 @@ impl Frontend {
ctx: &Context,
root: Handle<Expression>,
meta: Span,
) -> Result<Handle<Constant>> {
) -> Result<Handle<Expression>> {
let mut solver = ConstantSolver {
types: &mut self.module.types,
expressions: &ctx.expressions,
constants: &mut self.module.constants,
const_expressions: &mut self.module.const_expressions,
};
solver.solve(root).map_err(|e| Error {

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

@@ -14,7 +14,7 @@ pub struct VarDeclaration<'a, 'key> {
pub qualifiers: &'a mut TypeQualifiers<'key>,
pub ty: Handle<Type>,
pub name: Option<String>,
pub init: Option<Handle<Constant>>,
pub init: Option<Handle<Expression>>,
pub meta: Span,
}
@@ -494,29 +494,21 @@ impl Frontend {
meta,
})?;
let constant = Constant {
name: name.clone(),
r#override: crate::Override::None,
ty,
init,
};
let handle = self.module.constants.fetch_or_append(constant, meta);
let lookup = GlobalLookup {
kind: GlobalLookupKind::Constant(init, ty),
kind: GlobalLookupKind::Constant(handle, ty),
entry_arg: None,
mutable: false,
};
if let Some(name) = name.as_ref() {
let constant = self.module.constants.get_mut(init);
if constant.name.is_none() {
// set the name of the constant
constant.name = Some(name.clone())
} else {
// add a copy of the constant with the new name
let new_const = Constant {
name: Some(name.clone()),
specialization: constant.specialization,
inner: constant.inner.clone(),
};
self.module.constants.fetch_or_append(new_const, meta);
}
}
(GlobalOrConstant::Constant(init), lookup)
(GlobalOrConstant::Constant(handle), lookup)
}
StorageQualifier::AddressSpace(mut space) => {
match space {

10
src/front/mod.rs
View File

@@ -55,16 +55,6 @@ impl Emitter {
}
}
#[allow(dead_code)]
impl super::ConstantInner {
const fn boolean(value: bool) -> Self {
Self::Scalar {
width: super::BOOL_WIDTH,
value: super::ScalarValue::Bool(value),
}
}
}
/// A table of types for an `Arena<Expression>`.
///
/// A front end can use a `Typifier` to get types for an arena's expressions

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

@@ -128,6 +128,7 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
expressions: &mut fun.expressions,
local_arena: &mut fun.local_variables,
const_arena: &mut module.constants,
const_expressions: &mut module.const_expressions,
type_arena: &module.types,
global_arena: &module.global_variables,
arguments: &fun.arguments,
@@ -572,6 +573,7 @@ impl<'function> BlockContext<'function> {
crate::proc::GlobalCtx {
types: self.type_arena,
constants: self.const_arena,
const_expressions: self.const_expressions,
}
}

4
src/front/spv/image.rs
View File

@@ -503,6 +503,10 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
spirv::ImageOperands::CONST_OFFSET => {
let offset_constant = self.next()?;
let offset_handle = self.lookup_constant.lookup(offset_constant)?.handle;
let offset_handle = ctx.const_expressions.append(
crate::Expression::Constant(offset_handle),
Default::default(),
);
offset = Some(offset_handle);
words_left -= 1;
}

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

@@ -221,6 +221,11 @@ impl Decoration {
}
}
fn specialization(&self) -> crate::Override {
self.specialization
.map_or(crate::Override::None, crate::Override::ByNameOrId)
}
const fn resource_binding(&self) -> Option<crate::ResourceBinding> {
match *self {
Decoration {
@@ -536,6 +541,7 @@ struct BlockContext<'function> {
local_arena: &'function mut Arena<crate::LocalVariable>,
/// Constants arena of the module being processed
const_arena: &'function mut Arena<crate::Constant>,
const_expressions: &'function mut Arena<crate::Expression>,
/// Type arena of the module being processed
type_arena: &'function UniqueArena<crate::Type>,
/// Global arena of the module being processed
@@ -584,8 +590,6 @@ pub struct Frontend<I> {
// so that in the IR any called function is already known.
function_call_graph: GraphMap<spirv::Word, (), petgraph::Directed>,
options: Options,
index_constants: Vec<Handle<crate::Constant>>,
index_constant_expressions: Vec<Handle<crate::Expression>>,
/// Maps for a switch from a case target to the respective body and associated literals that
/// use that target block id.
@@ -636,8 +640,6 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
dummy_functions: Arena::new(),
function_call_graph: GraphMap::new(),
options: options.clone(),
index_constants: Vec::new(),
index_constant_expressions: Vec::new(),
switch_cases: indexmap::IndexMap::default(),
gl_per_vertex_builtin_access: FastHashSet::default(),
}
@@ -1395,7 +1397,10 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
inst.expect(5)?;
let init_id = self.next()?;
let lconst = self.lookup_constant.lookup(init_id)?;
Some(lconst.handle)
Some(
ctx.const_expressions
.append(crate::Expression::Constant(lconst.handle), span),
)
} else {
None
};
@@ -1531,11 +1536,15 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
let index_expr_handle = get_expr_handle!(access_id, &index_expr);
let index_expr_data = &ctx.expressions[index_expr.handle];
let index_maybe = match *index_expr_data {
crate::Expression::Constant(const_handle) => {
Some(ctx.const_arena[const_handle].to_array_length().ok_or(
Error::InvalidAccess(crate::Expression::Constant(const_handle)),
)?)
}
crate::Expression::Constant(const_handle) => Some(
ctx.gctx()
.eval_expr_to_u32(ctx.const_arena[const_handle].init)
.map_err(|_| {
Error::InvalidAccess(crate::Expression::Constant(
const_handle,
))
})?,
),
_ => None,
};
@@ -3731,13 +3740,6 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
},
);
}
// register special constants
self.index_constant_expressions.clear();
for &con_handle in self.index_constants.iter() {
let span = constants.get_span(con_handle);
let handle = expressions.append(crate::Expression::Constant(con_handle), span);
self.index_constant_expressions.push(handle);
}
// register constants
for (&id, con) in self.lookup_constant.iter() {
let span = constants.get_span(con.handle);
@@ -3902,23 +3904,6 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
crate::Module::default()
};
// register indexing constants
self.index_constants.clear();
for i in 0..4 {
let handle = module.constants.append(
crate::Constant {
name: None,
specialization: None,
inner: crate::ConstantInner::Scalar {
width: 4,
value: crate::ScalarValue::Sint(i),
},
},
Default::default(),
);
self.index_constants.push(handle);
}
self.layouter.clear();
self.dummy_functions = Arena::new();
self.lookup_function.clear();
@@ -3965,9 +3950,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
Op::TypeSampler => self.parse_type_sampler(inst, &mut module),
Op::Constant | Op::SpecConstant => self.parse_constant(inst, &mut module),
Op::ConstantComposite => self.parse_composite_constant(inst, &mut module),
Op::ConstantNull | Op::Undef => self
.parse_null_constant(inst, &module.types, &mut module.constants)
.map(|_| ()),
Op::ConstantNull | Op::Undef => self.parse_null_constant(inst, &mut module),
Op::ConstantTrue => self.parse_bool_constant(inst, true, &mut module),
Op::ConstantFalse => self.parse_bool_constant(inst, false, &mut module),
Op::Variable => self.parse_global_variable(inst, &mut module),
@@ -4892,21 +4875,15 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
let type_lookup = self.lookup_type.lookup(type_id)?;
let ty = type_lookup.handle;
let inner = match module.types[ty].inner {
let literal = match module.types[ty].inner {
crate::TypeInner::Scalar {
kind: crate::ScalarKind::Uint,
width,
} => {
let low = self.next()?;
let high = if width > 4 {
inst.expect(5)?;
self.next()?
} else {
0
};
crate::ConstantInner::Scalar {
width,
value: crate::ScalarValue::Uint((u64::from(high) << 32) | u64::from(low)),
match width {
4 => crate::Literal::U32(low),
_ => return Err(Error::InvalidTypeWidth(width as u32)),
}
}
crate::TypeInner::Scalar {
@@ -4914,17 +4891,9 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
width,
} => {
let low = self.next()?;
let high = if width > 4 {
inst.expect(5)?;
self.next()?
} else {
0
};
crate::ConstantInner::Scalar {
width,
value: crate::ScalarValue::Sint(
(i64::from(high as i32) << 32) | ((i64::from(low as i32) << 32) >> 32),
),
match width {
4 => crate::Literal::I32(low as i32),
_ => return Err(Error::InvalidTypeWidth(width as u32)),
}
}
crate::TypeInner::Scalar {
@@ -4932,18 +4901,16 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
width,
} => {
let low = self.next()?;
let extended = match width {
4 => f64::from(f32::from_bits(low)),
match width {
4 => crate::Literal::F32(f32::from_bits(low)),
8 => {
inst.expect(5)?;
let high = self.next()?;
f64::from_bits((u64::from(high) << 32) | u64::from(low))
crate::Literal::F64(f64::from_bits(
(u64::from(high) << 32) | u64::from(low),
))
}
_ => return Err(Error::InvalidTypeWidth(width as u32)),
};
crate::ConstantInner::Scalar {
width,
value: crate::ScalarValue::Float(extended),
}
}
_ => return Err(Error::UnsupportedType(type_lookup.handle)),
@@ -4951,16 +4918,22 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
let decor = self.future_decor.remove(&id).unwrap_or_default();
let span = self.span_from_with_op(start);
let init = module
.const_expressions
.append(crate::Expression::Literal(literal), span);
self.lookup_constant.insert(
id,
LookupConstant {
handle: module.constants.append(
crate::Constant {
specialization: decor.specialization,
r#override: decor.specialization(),
name: decor.name,
inner,
ty,
init,
},
self.span_from_with_op(start),
span,
),
type_id,
},
@@ -4977,27 +4950,41 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
self.switch(ModuleState::Type, inst.op)?;
inst.expect_at_least(3)?;
let type_id = self.next()?;
let id = self.next()?;
let type_lookup = self.lookup_type.lookup(type_id)?;
let ty = type_lookup.handle;
let id = self.next()?;
let mut components = Vec::with_capacity(inst.wc as usize - 3);
for _ in 0..components.capacity() {
let start = self.data_offset;
let component_id = self.next()?;
let span = self.span_from_with_op(start);
let constant = self.lookup_constant.lookup(component_id)?;
components.push(constant.handle);
let expr = module
.const_expressions
.append(crate::Expression::Constant(constant.handle), span);
components.push(expr);
}
let decor = self.future_decor.remove(&id).unwrap_or_default();
let span = self.span_from_with_op(start);
let init = module
.const_expressions
.append(crate::Expression::Compose { ty, components }, span);
self.lookup_constant.insert(
id,
LookupConstant {
handle: module.constants.append(
crate::Constant {
name: self.future_decor.remove(&id).and_then(|dec| dec.name),
specialization: None,
inner: crate::ConstantInner::Composite { ty, components },
r#override: decor.specialization(),
name: decor.name,
ty,
init,
},
self.span_from_with_op(start),
span,
),
type_id,
},
@@ -5008,30 +4995,35 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
fn parse_null_constant(
&mut self,
inst: Instruction,
types: &UniqueArena<crate::Type>,
constants: &mut Arena<crate::Constant>,
) -> Result<(u32, u32, Handle<crate::Constant>), Error> {
module: &mut crate::Module,
) -> Result<(), Error> {
let start = self.data_offset;
self.switch(ModuleState::Type, inst.op)?;
inst.expect(3)?;
let type_id = self.next()?;
let id = self.next()?;
let span = self.span_from_with_op(start);
let type_lookup = self.lookup_type.lookup(type_id)?;
let ty = type_lookup.handle;
let inner = null::generate_null_constant(ty, types, constants, span)?;
let handle = constants.append(
let decor = self.future_decor.remove(&id).unwrap_or_default();
let init = module
.const_expressions
.append(crate::Expression::ZeroValue(ty), span);
let handle = module.constants.append(
crate::Constant {
name: self.future_decor.remove(&id).and_then(|dec| dec.name),
specialization: None, //TODO
inner,
r#override: decor.specialization(),
name: decor.name,
ty,
init,
},
span,
);
self.lookup_constant
.insert(id, LookupConstant { handle, type_id });
Ok((type_id, id, handle))
Ok(())
}
fn parse_bool_constant(
@@ -5045,17 +5037,28 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
inst.expect(3)?;
let type_id = self.next()?;
let id = self.next()?;
let span = self.span_from_with_op(start);
let type_lookup = self.lookup_type.lookup(type_id)?;
let ty = type_lookup.handle;
let decor = self.future_decor.remove(&id).unwrap_or_default();
let init = module.const_expressions.append(
crate::Expression::Literal(crate::Literal::Bool(value)),
span,
);
self.lookup_constant.insert(
id,
LookupConstant {
handle: module.constants.append(
crate::Constant {
name: self.future_decor.remove(&id).and_then(|dec| dec.name),
specialization: None, //TODO
inner: crate::ConstantInner::boolean(value),
r#override: decor.specialization(),
name: decor.name,
ty,
init,
},
self.span_from_with_op(start),
span,
),
type_id,
},
@@ -5076,9 +5079,14 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
let storage_class = self.next()?;
let init = if inst.wc > 4 {
inst.expect(5)?;
let start = self.data_offset;
let init_id = self.next()?;
let span = self.span_from_with_op(start);
let lconst = self.lookup_constant.lookup(init_id)?;
Some(lconst.handle)
let expr = module
.const_expressions
.append(crate::Expression::Constant(lconst.handle), span);
Some(expr)
} else {
None
};
@@ -5211,8 +5219,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
match null::generate_default_built_in(
Some(built_in),
ty,
&module.types,
&mut module.constants,
&mut module.const_expressions,
span,
) {
Ok(handle) => Some(handle),
@@ -5225,37 +5232,25 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
Some(crate::Binding::Location { .. }) => None,
None => match module.types[ty].inner {
crate::TypeInner::Struct { ref members, .. } => {
// A temporary to avoid borrowing `module.types`
let pairs = members
.iter()
.map(|member| {
let built_in = match member.binding {
Some(crate::Binding::BuiltIn(built_in)) => Some(built_in),
_ => None,
};
(built_in, member.ty)
})
.collect::<Vec<_>>();
let mut components = Vec::with_capacity(members.len());
for (built_in, member_ty) in pairs {
for member in members.iter() {
let built_in = match member.binding {
Some(crate::Binding::BuiltIn(built_in)) => Some(built_in),
_ => None,
};
let handle = null::generate_default_built_in(
built_in,
member_ty,
&module.types,
&mut module.constants,
member.ty,
&mut module.const_expressions,
span,
)?;
components.push(handle);
}
Some(module.constants.append(
crate::Constant {
name: None,
specialization: None,
inner: crate::ConstantInner::Composite { ty, components },
},
span,
))
Some(
module
.const_expressions
.append(crate::Expression::Compose { ty, components }, span),
)
}
_ => None,
},
@@ -5322,7 +5317,11 @@ fn resolve_constant(
gctx: crate::proc::GlobalCtx,
constant: Handle<crate::Constant>,
) -> Option<u32> {
gctx.constants[constant].to_array_length()
match gctx.const_expressions[gctx.constants[constant].init] {
crate::Expression::Literal(crate::Literal::U32(id)) => Some(id),
crate::Expression::Literal(crate::Literal::I32(id)) => Some(id as u32),
_ => None,
}
}
pub fn parse_u8_slice(data: &[u8], options: &Options) -> Result<crate::Module, Error> {

175
src/front/spv/null.rs
View File

@@ -1,172 +1,31 @@
use super::Error;
use crate::arena::{Arena, Handle, UniqueArena};
const fn make_scalar_inner(kind: crate::ScalarKind, width: crate::Bytes) -> crate::ConstantInner {
crate::ConstantInner::Scalar {
width,
value: match kind {
crate::ScalarKind::Uint => crate::ScalarValue::Uint(0),
crate::ScalarKind::Sint => crate::ScalarValue::Sint(0),
crate::ScalarKind::Float => crate::ScalarValue::Float(0.0),
crate::ScalarKind::Bool => crate::ScalarValue::Bool(false),
},
}
}
pub fn generate_null_constant(
ty: Handle<crate::Type>,
type_arena: &UniqueArena<crate::Type>,
constant_arena: &mut Arena<crate::Constant>,
span: crate::Span,
) -> Result<crate::ConstantInner, Error> {
let inner = match type_arena[ty].inner {
crate::TypeInner::Scalar { kind, width } => make_scalar_inner(kind, width),
crate::TypeInner::Vector { size, kind, width } => {
let mut components = Vec::with_capacity(size as usize);
for _ in 0..size as usize {
components.push(constant_arena.fetch_or_append(
crate::Constant {
name: None,
specialization: None,
inner: make_scalar_inner(kind, width),
},
span,
));
}
crate::ConstantInner::Composite { ty, components }
}
crate::TypeInner::Matrix {
columns,
rows,
width,
} => {
// If we successfully declared a matrix type, we have declared a vector type for it too.
let vector_ty = type_arena
.get(&crate::Type {
name: None,
inner: crate::TypeInner::Vector {
kind: crate::ScalarKind::Float,
size: rows,
width,
},
})
.unwrap();
let vector_inner = generate_null_constant(vector_ty, type_arena, constant_arena, span)?;
let vector_handle = constant_arena.fetch_or_append(
crate::Constant {
name: None,
specialization: None,
inner: vector_inner,
},
span,
);
crate::ConstantInner::Composite {
ty,
components: vec![vector_handle; columns as usize],
}
}
crate::TypeInner::Struct { ref members, .. } => {
let mut components = Vec::with_capacity(members.len());
// copy out the types to avoid borrowing `members`
let member_tys = members.iter().map(|member| member.ty).collect::<Vec<_>>();
for member_ty in member_tys {
let inner = generate_null_constant(member_ty, type_arena, constant_arena, span)?;
components.push(constant_arena.fetch_or_append(
crate::Constant {
name: None,
specialization: None,
inner,
},
span,
));
}
crate::ConstantInner::Composite { ty, components }
}
crate::TypeInner::Array {
base,
size: crate::ArraySize::Constant(size),
..
} => {
let inner = generate_null_constant(base, type_arena, constant_arena, span)?;
let value = constant_arena.fetch_or_append(
crate::Constant {
name: None,
specialization: None,
inner,
},
span,
);
crate::ConstantInner::Composite {
ty,
components: vec![value; size.get() as usize],
}
}
ref other => {
log::warn!("null constant type {:?}", other);
return Err(Error::UnsupportedType(ty));
}
};
Ok(inner)
}
use crate::arena::{Arena, Handle};
/// Create a default value for an output built-in.
pub fn generate_default_built_in(
built_in: Option<crate::BuiltIn>,
ty: Handle<crate::Type>,
type_arena: &UniqueArena<crate::Type>,
constant_arena: &mut Arena<crate::Constant>,
const_expressions: &mut Arena<crate::Expression>,
span: crate::Span,
) -> Result<Handle<crate::Constant>, Error> {
let inner = match built_in {
) -> Result<Handle<crate::Expression>, Error> {
let expr = match built_in {
Some(crate::BuiltIn::Position { .. }) => {
let zero = constant_arena.fetch_or_append(
crate::Constant {
name: None,
specialization: None,
inner: crate::ConstantInner::Scalar {
value: crate::ScalarValue::Float(0.0),
width: 4,
},
},
span,
);
let one = constant_arena.fetch_or_append(
crate::Constant {
name: None,
specialization: None,
inner: crate::ConstantInner::Scalar {
value: crate::ScalarValue::Float(1.0),
width: 4,
},
},
span,
);
crate::ConstantInner::Composite {
let zero = const_expressions
.append(crate::Expression::Literal(crate::Literal::F32(0.0)), span);
let one = const_expressions
.append(crate::Expression::Literal(crate::Literal::F32(1.0)), span);
crate::Expression::Compose {
ty,
components: vec![zero, zero, zero, one],
}
}
Some(crate::BuiltIn::PointSize) => crate::ConstantInner::Scalar {
value: crate::ScalarValue::Float(1.0),
width: 4,
},
Some(crate::BuiltIn::FragDepth) => crate::ConstantInner::Scalar {
value: crate::ScalarValue::Float(0.0),
width: 4,
},
Some(crate::BuiltIn::SampleMask) => crate::ConstantInner::Scalar {
value: crate::ScalarValue::Uint(u64::MAX),
width: 4,
},
//Note: `crate::BuiltIn::ClipDistance` is intentionally left for the default path
_ => generate_null_constant(ty, type_arena, constant_arena, span)?,
Some(crate::BuiltIn::PointSize) => crate::Expression::Literal(crate::Literal::F32(1.0)),
Some(crate::BuiltIn::FragDepth) => crate::Expression::Literal(crate::Literal::F32(0.0)),
Some(crate::BuiltIn::SampleMask) => {
crate::Expression::Literal(crate::Literal::U32(u32::MAX))
}
// Note: `crate::BuiltIn::ClipDistance` is intentionally left for the default path
_ => crate::Expression::ZeroValue(ty),
};
Ok(constant_arena.fetch_or_append(
crate::Constant {
name: None,
specialization: None,
inner,
},
span,
))
Ok(const_expressions.append(expr, span))
}

20
src/front/wgsl/error.rs
View File

@@ -100,8 +100,6 @@ pub enum ExpectedToken<'a> {
Identifier,
Number,
Integer,
/// A compile-time constant expression.
Constant,
/// Expected: constant, parenthesized expression, identifier
PrimaryExpression,
/// Expected: assignment, increment/decrement expression
@@ -141,6 +139,7 @@ pub enum InvalidAssignmentType {
pub enum Error<'a> {
Unexpected(Span, ExpectedToken<'a>),
UnexpectedComponents(Span),
UnexpectedOperationInConstContext(Span),
BadNumber(Span, NumberError),
/// A negative signed integer literal where both signed and unsigned,
/// but only non-negative literals are allowed.
@@ -227,7 +226,6 @@ pub enum Error<'a> {
/// the same identifier as `ident`, above.
path: Vec<(Span, Span)>,
},
ConstExprUnsupported(Span),
InvalidSwitchValue {
uint: bool,
span: Span,
@@ -273,7 +271,6 @@ impl<'a> Error<'a> {
ExpectedToken::Identifier => "identifier".to_string(),
ExpectedToken::Number => "32-bit signed integer literal".to_string(),
ExpectedToken::Integer => "unsigned/signed integer literal".to_string(),
ExpectedToken::Constant => "compile-time constant".to_string(),
ExpectedToken::PrimaryExpression => "expression".to_string(),
ExpectedToken::Assignment => "assignment or increment/decrement".to_string(),
ExpectedToken::SwitchItem => "switch item ('case' or 'default') or a closing curly bracket to signify the end of the switch statement ('}')".to_string(),
@@ -297,6 +294,11 @@ impl<'a> Error<'a> {
labels: vec![(bad_span, "unexpected components".into())],
notes: vec![],
},
Error::UnexpectedOperationInConstContext(span) => ParseError {
message: "this operation is not supported in a const context".to_string(),
labels: vec![(span, "operation not supported here".into())],
notes: vec![],
},
Error::BadNumber(bad_span, ref err) => ParseError {
message: format!("{}: `{}`", err, &source[bad_span],),
labels: vec![(bad_span, err.to_string().into())],
@@ -625,14 +627,6 @@ impl<'a> Error<'a> {
.collect(),
notes: vec![],
},
Error::ConstExprUnsupported(span) => ParseError {
message: "this constant expression is not supported".to_string(),
labels: vec![(span, "expression is not supported".into())],
notes: vec![
"this should be fixed in a future version of Naga".into(),
"https://github.com/gfx-rs/naga/issues/1829".into(),
],
},
Error::InvalidSwitchValue { uint, span } => ParseError {
message: "invalid switch value".to_string(),
labels: vec![(
@@ -701,7 +695,7 @@ impl<'a> Error<'a> {
},
Error::NonPositiveArrayLength(span) => ParseError {
message: "array element count must be greater than zero".to_string(),
labels: vec![(span, "must be positive".into())],
labels: vec![(span, "must be greater than zero".into())],
notes: vec![],
},
}

87
src/front/wgsl/lower/construction.rs
View File

@@ -4,7 +4,7 @@ use crate::front::wgsl::parse::ast;
use crate::{Handle, Span};
use crate::front::wgsl::error::Error;
use crate::front::wgsl::lower::{ExpressionContext, Lowerer, OutputContext};
use crate::front::wgsl::lower::{ExpressionContext, Lowerer};
use crate::proc::TypeResolution;
enum ConcreteConstructorHandle {
@@ -145,7 +145,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
components: &[Handle<ast::Expression<'source>>],
mut ctx: ExpressionContext<'source, '_, '_>,
) -> Result<Handle<crate::Expression>, Error<'source>> {
let constructor_h = self.constructor(constructor, ctx.as_output())?;
let constructor_h = self.constructor(constructor, ctx.reborrow())?;
let components_h = match *components {
[] => ComponentsHandle::None,
@@ -153,7 +153,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let span = ctx.ast_expressions.get_span(component);
let component = self.expression(component, ctx.reborrow())?;
ctx.grow_types(component)?;
let ty = &ctx.typifier[component];
let ty = &ctx.typifier()[component];
ComponentsHandle::One {
component,
@@ -179,7 +179,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
.collect();
ctx.grow_types(component)?;
let ty = &ctx.typifier[component];
let ty = &ctx.typifier()[component];
ComponentsHandle::Many {
components,
@@ -195,14 +195,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
);
let expr = match (components, constructor) {
// Empty constructor
(Components::None, dst_ty) => {
let ty = match dst_ty {
ConcreteConstructor::Type(ty, _) => ty,
_ => return Err(Error::TypeNotInferrable(ty_span)),
};
return Ok(ctx.interrupt_emitter(crate::Expression::ZeroValue(ty), span));
}
(Components::None, dst_ty) => match dst_ty {
ConcreteConstructor::Type(ty, _) => {
return Ok(ctx.interrupt_emitter(crate::Expression::ZeroValue(ty), span))
}
_ => return Err(Error::TypeNotInferrable(ty_span)),
},
// Scalar constructor & conversion (scalar -> scalar)
(
@@ -402,7 +400,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let components = components
.chunks(rows as usize)
.map(|vec_components| {
ctx.naga_expressions.append(
ctx.append_expression(
crate::Expression::Compose {
ty: vec_ty,
components: Vec::from(vec_components),
@@ -525,55 +523,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
_ => return Err(Error::TypeNotConstructible(ty_span)),
};
let expr = ctx.naga_expressions.append(expr, span);
let expr = ctx.append_expression(expr, span);
Ok(expr)
}
/// Build a Naga IR [`ConstantInner`] given a WGSL construction expression.
///
/// Given `constructor`, representing the head of a WGSL [`type constructor
/// expression`], and a slice of [`ast::Expression`] handles representing
/// the constructor's arguments, build a Naga [`ConstantInner`] value
/// representing the given value.
///
/// If `constructor` is for a composite type, this may entail adding new
/// [`Type`]s and [`Constant`]s to [`ctx.module`], if it doesn't already
/// have what we need.
///
/// If the arguments cannot be evaluated at compile time, return an error.
///
/// [`ConstantInner`]: crate::ConstantInner
/// [`type constructor expression`]: https://gpuweb.github.io/gpuweb/wgsl/#type-constructor-expr
/// [`Function::expressions`]: ast::Function::expressions
/// [`TranslationUnit::global_expressions`]: ast::TranslationUnit::global_expressions
/// [`Type`]: crate::Type
/// [`Constant`]: crate::Constant
/// [`ctx.module`]: OutputContext::module
pub fn const_construct(
&mut self,
span: Span,
constructor: &ast::ConstructorType<'source>,
components: &[Handle<ast::Expression<'source>>],
mut ctx: OutputContext<'source, '_, '_>,
) -> Result<crate::ConstantInner, Error<'source>> {
// TODO: Support zero values, splatting and inference.
let constructor = self.constructor(constructor, ctx.reborrow())?;
let c = match constructor {
ConcreteConstructorHandle::Type(ty) => {
let components = components
.iter()
.map(|&expr| self.constant(expr, ctx.reborrow()))
.collect::<Result<_, _>>()?;
crate::ConstantInner::Composite { ty, components }
}
_ => return Err(Error::ConstExprUnsupported(span)),
};
Ok(c)
}
/// Build a Naga IR [`Type`] for `constructor` if there is enough
/// information to do so.
///
@@ -591,13 +544,13 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
/// array's length.
///
/// [`Type`]: crate::Type
/// [`ctx.module`]: OutputContext::module
/// [`ctx.module`]: ExpressionContext::module
/// [`Array`]: crate::TypeInner::Array
/// [`Constant`]: crate::Constant
fn constructor<'out>(
&mut self,
constructor: &ast::ConstructorType<'source>,
mut ctx: OutputContext<'source, '_, 'out>,
mut ctx: ExpressionContext<'source, '_, 'out>,
) -> Result<ConcreteConstructorHandle, Error<'source>> {
let c = match *constructor {
ast::ConstructorType::Scalar { width, kind } => {
@@ -628,17 +581,11 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
}
ast::ConstructorType::PartialArray => ConcreteConstructorHandle::PartialArray,
ast::ConstructorType::Array { base, size } => {
let base = self.resolve_ast_type(base, ctx.reborrow())?;
let base = self.resolve_ast_type(base, ctx.as_global())?;
let size = match size {
ast::ArraySize::Constant(expr) => {
let span = ctx.ast_expressions.get_span(expr);
let constant = self.constant(expr, ctx.reborrow())?;
let size = ctx.module.constants[constant]
.to_array_length()
.ok_or(Error::ExpectedArraySize(span))?;
let size =
NonZeroU32::new(size).ok_or(Error::NonPositiveArrayLength(span))?;
crate::ArraySize::Constant(size)
let const_expr = self.expression(expr, ctx.as_const())?;
crate::ArraySize::Constant(ctx.array_length(const_expr)?)
}
ast::ArraySize::Dynamic => crate::ArraySize::Dynamic,
};

710
src/front/wgsl/lower/mod.rs
View File

File diff suppressed because it is too large Load Diff

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

@@ -198,6 +198,7 @@ impl crate::TypeInner {
}
mod type_inner_tests {
#[test]
fn to_wgsl() {
use std::num::NonZeroU32;
@@ -228,6 +229,7 @@ mod type_inner_tests {
let gctx = crate::proc::GlobalCtx {
types: &types,
constants: &crate::Arena::new(),
const_expressions: &crate::Arena::new(),
};
let array = crate::TypeInner::Array {
base: mytype1,

6
src/front/wgsl/parse/ast.rs
View File

@@ -107,7 +107,7 @@ pub struct EntryPoint {
}
#[cfg(doc)]
use crate::front::wgsl::lower::{ExpressionContext, StatementContext};
use crate::front::wgsl::lower::{RuntimeExpressionContext, StatementContext};
#[derive(Debug)]
pub struct Function<'a> {
@@ -132,14 +132,14 @@ pub struct Function<'a> {
/// During lowering, [`LocalDecl`] statements add entries to a per-function
/// table that maps `Handle<Local>` values to their Naga representations,
/// accessed via [`StatementContext::local_table`] and
/// [`ExpressionContext::local_table`]. This table is then consulted when
/// [`RuntimeExpressionContext::local_table`]. This table is then consulted when
/// lowering subsequent [`Ident`] expressions.
///
/// [`LocalDecl`]: StatementKind::LocalDecl
/// [`arguments`]: Function::arguments
/// [`Ident`]: Expression::Ident
/// [`StatementContext::local_table`]: StatementContext::local_table
/// [`ExpressionContext::local_table`]: ExpressionContext::local_table
/// [`RuntimeExpressionContext::local_table`]: RuntimeExpressionContext::local_table
pub locals: Arena<Local>,
pub body: Block<'a>,

75
src/lib.rs
View File

@@ -796,18 +796,6 @@ pub enum TypeInner {
BindingArray { base: Handle<Type>, size: ArraySize },
}
/// Constant value.
#[derive(Debug, PartialEq)]
#[cfg_attr(feature = "clone", derive(Clone))]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct Constant {
pub name: Option<String>,
pub specialization: Option<u32>,
pub inner: ConstantInner,
}
#[derive(Debug, Clone, Copy, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
@@ -820,32 +808,29 @@ pub enum Literal {
Bool(bool),
}
/// A literal scalar value, used in constants.
#[derive(Debug, Clone, Copy, PartialOrd)]
#[derive(Debug, PartialEq)]
#[cfg_attr(feature = "clone", derive(Clone))]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ScalarValue {
Sint(i64),
Uint(u64),
Float(f64),
Bool(bool),
pub enum Override {
None,
ByName,
ByNameOrId(u32),
}
/// Additional information, dependent on the kind of constant.
#[derive(Clone, Debug, PartialEq)]
/// Constant value.
#[derive(Debug, PartialEq)]
#[cfg_attr(feature = "clone", derive(Clone))]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ConstantInner {
Scalar {
width: Bytes,
value: ScalarValue,
},
Composite {
ty: Handle<Type>,
components: Vec<Handle<Constant>>,
},
pub struct Constant {
pub name: Option<String>,
pub r#override: Override,
pub ty: Handle<Type>,
/// Expression handle lives in const_expressions
pub init: Handle<Expression>,
}
/// Describes how an input/output variable is to be bound.
@@ -907,7 +892,9 @@ pub struct GlobalVariable {
/// The type of this variable.
pub ty: Handle<Type>,
/// Initial value for this variable.
pub init: Option<Handle<Constant>>,
///
/// Expression handle lives in const_expressions
pub init: Option<Handle<Expression>>,
}
/// Variable defined at function level.
@@ -921,7 +908,9 @@ pub struct LocalVariable {
/// The type of this variable.
pub ty: Handle<Type>,
/// Initial value for this variable.
pub init: Option<Handle<Constant>>,
///
/// Expression handle lives in const_expressions
pub init: Option<Handle<Expression>>,
}
/// Operation that can be applied on a single value.
@@ -1224,6 +1213,11 @@ pub enum Expression {
Constant(Handle<Constant>),
/// Zero value of a type.
ZeroValue(Handle<Type>),
/// Composite expression.
Compose {
ty: Handle<Type>,
components: Vec<Handle<Expression>>,
},
/// Array access with a computed index.
///
@@ -1294,11 +1288,6 @@ pub enum Expression {
vector: Handle<Expression>,
pattern: [SwizzleComponent; 4],
},
/// Composite expression.
Compose {
ty: Handle<Type>,
components: Vec<Handle<Expression>>,
},
/// Reference a function parameter, by its index.
///
@@ -1347,7 +1336,8 @@ pub enum Expression {
gather: Option<SwizzleComponent>,
coordinate: Handle<Expression>,
array_index: Option<Handle<Expression>>,
offset: Option<Handle<Constant>>,
/// Expression handle lives in const_expressions
offset: Option<Handle<Expression>>,
level: SampleLevel,
depth_ref: Option<Handle<Expression>>,
},
@@ -1448,7 +1438,6 @@ pub enum Expression {
Derivative {
axis: DerivativeAxis,
ctrl: DerivativeControl,
//modifier,
expr: Handle<Expression>,
},
/// Call a relational function.
@@ -1510,7 +1499,6 @@ pub enum Expression {
pub use block::Block;
/// The value of the switch case.
// Clone is used only for error reporting and is not intended for end users
#[derive(Clone, Copy, Debug, Eq, Hash, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
@@ -1619,7 +1607,7 @@ pub enum Statement {
/// [`body`]: SwitchCase::body
/// [`Default`]: SwitchValue::Default
Switch {
selector: Handle<Expression>, //int
selector: Handle<Expression>,
cases: Vec<SwitchCase>,
},
@@ -1939,6 +1927,11 @@ pub struct Module {
pub constants: Arena<Constant>,
/// Arena for the global variables defined in this module.
pub global_variables: Arena<GlobalVariable>,
/// const/override-expressions used by this module.
///
/// An `Expression` must occur before all other `Expression`s that use its
/// value.
pub const_expressions: Arena<Expression>,
/// Arena for the functions defined in this module.
///
/// Each function must appear in this arena strictly before all its callers.

20
src/proc/index.rs
View File

@@ -340,22 +340,16 @@ pub fn access_needs_check(
}
impl GuardedIndex {
/// Make A `GuardedIndex::Known` from a `GuardedIndex::Expression` if possible.
///
/// If the expression is a [`Constant`] whose value is a non-specialized, scalar
/// integer constant that can be converted to a `u32`, do so and return a
/// `GuardedIndex::Known`. Otherwise, return the `GuardedIndex::Expression`
/// unchanged.
/// Make a `GuardedIndex::Known` from a `GuardedIndex::Expression` if possible.
///
/// Return values that are already `Known` unchanged.
///
/// [`Constant`]: crate::Expression::Constant
fn try_resolve_to_constant(&mut self, function: &crate::Function, module: &crate::Module) {
if let GuardedIndex::Expression(expr) = *self {
if let crate::Expression::Constant(handle) = function.expressions[expr] {
if let Some(value) = module.constants[handle].to_array_length() {
*self = GuardedIndex::Known(value);
}
if let Ok(value) = module
.to_ctx()
.eval_expr_to_u32_from(expr, &function.expressions)
{
*self = GuardedIndex::Known(value);
}
}
}
@@ -366,7 +360,7 @@ pub enum IndexableLengthError {
#[error("Type is not indexable, and has no length (validation error)")]
TypeNotIndexable,
#[error("Array length constant {0:?} is invalid")]
InvalidArrayLength(Handle<crate::Constant>),
InvalidArrayLength(Handle<crate::Expression>),
}
impl crate::TypeInner {

180
src/proc/mod.rs
View File

@@ -8,8 +8,6 @@ mod namer;
mod terminator;
mod typifier;
use std::cmp::PartialEq;
pub use index::{BoundsCheckPolicies, BoundsCheckPolicy, IndexableLength, IndexableLengthError};
pub use layouter::{Alignment, LayoutError, LayoutErrorInner, Layouter, TypeLayout};
pub use namer::{EntryPointIndex, NameKey, Namer};
@@ -62,17 +60,6 @@ impl From<super::StorageFormat> for super::ScalarKind {
}
}
impl super::ScalarValue {
pub const fn scalar_kind(&self) -> super::ScalarKind {
match *self {
Self::Uint(_) => super::ScalarKind::Uint,
Self::Sint(_) => super::ScalarKind::Sint,
Self::Float(_) => super::ScalarKind::Float,
Self::Bool(_) => super::ScalarKind::Bool,
}
}
}
impl super::ScalarKind {
pub const fn is_numeric(self) -> bool {
match self {
@@ -135,17 +122,6 @@ impl crate::Literal {
}
}
pub const fn from_scalar(scalar: crate::ScalarValue, width: crate::Bytes) -> Option<Self> {
match (scalar, width) {
(crate::ScalarValue::Sint(n), 4) => Some(Self::I32(n as _)),
(crate::ScalarValue::Uint(n), 4) => Some(Self::U32(n as _)),
(crate::ScalarValue::Float(n), 4) => Some(Self::F32(n as _)),
(crate::ScalarValue::Float(n), 8) => Some(Self::F64(n)),
(crate::ScalarValue::Bool(b), _) => Some(Self::Bool(b)),
_ => None,
}
}
pub const fn zero(kind: crate::ScalarKind, width: crate::Bytes) -> Option<Self> {
Self::new(0, kind, width)
}
@@ -209,6 +185,16 @@ impl super::TypeInner {
}
}
pub fn is_atomic_pointer(&self, types: &crate::UniqueArena<crate::Type>) -> bool {
match *self {
crate::TypeInner::Pointer { base, .. } => match types[base].inner {
crate::TypeInner::Atomic { .. } => true,
_ => false,
},
_ => false,
}
}
/// Get the size of this type.
pub fn size(&self, _gctx: GlobalCtx) -> u32 {
match *self {
@@ -308,6 +294,39 @@ impl super::TypeInner {
_ => false,
}
}
pub fn components(&self) -> Option<u32> {
Some(match *self {
Self::Vector { size, .. } => size as u32,
Self::Matrix { columns, .. } => columns as u32,
Self::Array {
size: crate::ArraySize::Constant(len),
..
} => len.get(),
Self::Struct { ref members, .. } => members.len() as u32,
_ => return None,
})
}
pub fn component_type(&self, index: usize) -> Option<TypeResolution> {
Some(match *self {
Self::Vector { kind, width, .. } => {
TypeResolution::Value(crate::TypeInner::Scalar { kind, width })
}
Self::Matrix { rows, width, .. } => TypeResolution::Value(crate::TypeInner::Vector {
size: rows,
kind: crate::ScalarKind::Float,
width,
}),
Self::Array {
base,
size: crate::ArraySize::Constant(_),
..
} => TypeResolution::Handle(base),
Self::Struct { ref members, .. } => TypeResolution::Handle(members[index].ty),
_ => return None,
})
}
}
impl super::AddressSpace {
@@ -441,7 +460,7 @@ impl crate::Expression {
pub fn is_dynamic_index(&self, module: &crate::Module) -> bool {
if let Self::Constant(handle) = *self {
let constant = &module.constants[handle];
constant.specialization.is_some()
!matches!(constant.r#override, crate::Override::None)
} else {
true
}
@@ -484,36 +503,6 @@ impl crate::SampleLevel {
}
}
impl crate::Constant {
/// Interpret this constant as an array length, and return it as a `u32`.
///
/// Ignore any specialization available for this constant; return its
/// unspecialized value.
///
/// If the constant has an inappropriate kind (non-scalar or non-integer) or
/// value (negative, out of range for u32), return `None`. This usually
/// indicates an error, but only the caller has enough information to report
/// the error helpfully: in back ends, it's a validation error, but in front
/// ends, it may indicate ill-formed input (for example, a SPIR-V
/// `OpArrayType` referring to an inappropriate `OpConstant`). So we return
/// `Option` and let the caller sort things out.
pub(crate) fn to_array_length(&self) -> Option<u32> {
match self.inner {
crate::ConstantInner::Scalar { value, width: _ } => match value {
crate::ScalarValue::Uint(value) => value.try_into().ok(),
// Accept a signed integer size to avoid
// requiring an explicit uint
// literal. Type inference should make
// this unnecessary.
crate::ScalarValue::Sint(value) => value.try_into().ok(),
_ => None,
},
// caught by type validation
crate::ConstantInner::Composite { .. } => None,
}
}
}
impl crate::Binding {
pub const fn to_built_in(&self) -> Option<crate::BuiltIn> {
match *self {
@@ -523,30 +512,6 @@ impl crate::Binding {
}
}
//TODO: should we use an existing crate for hashable floats?
impl PartialEq for crate::ScalarValue {
fn eq(&self, other: &Self) -> bool {
match (*self, *other) {
(Self::Uint(a), Self::Uint(b)) => a == b,
(Self::Sint(a), Self::Sint(b)) => a == b,
(Self::Float(a), Self::Float(b)) => a.to_bits() == b.to_bits(),
(Self::Bool(a), Self::Bool(b)) => a == b,
_ => false,
}
}
}
impl Eq for crate::ScalarValue {}
impl std::hash::Hash for crate::ScalarValue {
fn hash<H: std::hash::Hasher>(&self, hasher: &mut H) {
match *self {
Self::Sint(v) => v.hash(hasher),
Self::Uint(v) => v.hash(hasher),
Self::Float(v) => v.to_bits().hash(hasher),
Self::Bool(v) => v.hash(hasher),
}
}
}
impl super::SwizzleComponent {
pub const XYZW: [Self; 4] = [Self::X, Self::Y, Self::Z, Self::W];
@@ -589,14 +554,71 @@ impl crate::Module {
GlobalCtx {
types: &self.types,
constants: &self.constants,
const_expressions: &self.const_expressions,
}
}
}
#[derive(Debug)]
pub(super) enum U32EvalError {
NonConst,
Negative,
}
#[derive(Clone, Copy)]
pub struct GlobalCtx<'a> {
pub types: &'a crate::UniqueArena<crate::Type>,
pub constants: &'a crate::Arena<crate::Constant>,
pub const_expressions: &'a crate::Arena<crate::Expression>,
}
impl GlobalCtx<'_> {
/// Try to evaluate the expression in `self.const_expressions` using its `handle` and return it as a `u32`.
#[allow(dead_code)]
pub(super) fn eval_expr_to_u32(
&self,
handle: crate::Handle<crate::Expression>,
) -> Result<u32, U32EvalError> {
self.eval_expr_to_u32_from(handle, self.const_expressions)
}
/// Try to evaluate the expression in the `arena` using its `handle` and return it as a `u32`.
pub(super) fn eval_expr_to_u32_from(
&self,
handle: crate::Handle<crate::Expression>,
arena: &crate::Arena<crate::Expression>,
) -> Result<u32, U32EvalError> {
fn get(
gctx: GlobalCtx,
handle: crate::Handle<crate::Expression>,
arena: &crate::Arena<crate::Expression>,
) -> Result<u32, U32EvalError> {
match arena[handle] {
crate::Expression::Literal(crate::Literal::U32(value)) => Ok(value),
crate::Expression::Literal(crate::Literal::I32(value)) => {
value.try_into().map_err(|_| U32EvalError::Negative)
}
crate::Expression::ZeroValue(ty)
if matches!(
gctx.types[ty].inner,
crate::TypeInner::Scalar {
kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
width: _
}
) =>
{
Ok(0)
}
_ => Err(U32EvalError::NonConst),
}
}
match arena[handle] {
crate::Expression::Constant(c) => {
get(*self, self.constants[c].init, self.const_expressions)
}
_ => get(*self, handle, arena),
}
}
}
#[test]

38
src/proc/namer.rs
View File

@@ -225,43 +225,7 @@ impl Namer {
use std::fmt::Write;
// Try to be more descriptive about the constant values
temp.clear();
match constant.inner {
crate::ConstantInner::Scalar {
width: _,
value: crate::ScalarValue::Sint(v),
} => write!(temp, "const_{v}i"),
crate::ConstantInner::Scalar {
width: _,
value: crate::ScalarValue::Uint(v),
} => write!(temp, "const_{v}u"),
crate::ConstantInner::Scalar {
width: _,
value: crate::ScalarValue::Float(v),
} => {
let abs = v.abs();
write!(
temp,
"const_{}{}",
if v < 0.0 { "n" } else { "" },
abs.trunc(),
)
.unwrap();
let fract = abs.fract();
if fract == 0.0 {
write!(temp, "f")
} else {
write!(temp, "_{:02}f", (fract * 100.0) as i8)
}
}
crate::ConstantInner::Scalar {
width: _,
value: crate::ScalarValue::Bool(v),
} => write!(temp, "const_{v}"),
crate::ConstantInner::Composite { ty, components: _ } => {
write!(temp, "const_{}", output[&NameKey::Type(ty)])
}
}
.unwrap();
write!(temp, "const_{}", output[&NameKey::Type(constant.ty)]).unwrap();
&temp
}
};

26
src/proc/typifier.rs
View File

@@ -148,18 +148,6 @@ impl Clone for TypeResolution {
}
}
impl crate::ConstantInner {
pub const fn resolve_type(&self) -> TypeResolution {
match *self {
Self::Scalar { width, ref value } => TypeResolution::Value(crate::TypeInner::Scalar {
kind: value.scalar_kind(),
width,
}),
Self::Composite { ty, components: _ } => TypeResolution::Handle(ty),
}
}
}
#[derive(Clone, Debug, Error, PartialEq)]
pub enum ResolveError {
#[error("Index {index} is out of bounds for expression {expr:?}")]
@@ -421,17 +409,6 @@ impl<'a> ResolveContext<'a> {
}
}
}
crate::Expression::Literal(lit) => TypeResolution::Value(lit.ty_inner()),
crate::Expression::Constant(h) => match self.constants[h].inner {
crate::ConstantInner::Scalar { width, ref value } => {
TypeResolution::Value(Ti::Scalar {
kind: value.scalar_kind(),
width,
})
}
crate::ConstantInner::Composite { ty, components: _ } => TypeResolution::Handle(ty),
},
crate::Expression::ZeroValue(ty) => TypeResolution::Handle(ty),
crate::Expression::Splat { size, value } => match *past(value)?.inner_with(types) {
Ti::Scalar { kind, width } => {
TypeResolution::Value(Ti::Vector { size, kind, width })
@@ -456,6 +433,9 @@ impl<'a> ResolveContext<'a> {
return Err(ResolveError::InvalidVector(vector));
}
},
crate::Expression::Literal(lit) => TypeResolution::Value(lit.ty_inner()),
crate::Expression::Constant(h) => TypeResolution::Handle(self.constants[h].ty),
crate::Expression::ZeroValue(ty) => TypeResolution::Handle(ty),
crate::Expression::Compose { ty, .. } => TypeResolution::Handle(ty),
crate::Expression::FunctionArgument(index) => {
let arg = self

30
src/valid/analyzer.rs
View File

@@ -495,7 +495,6 @@ impl FunctionInfo {
requirements: UniformityRequirements::empty(),
},
// always uniform
E::Literal(_) | E::Constant(_) | E::ZeroValue(_) => Uniformity::new(),
E::Splat { size: _, value } => Uniformity {
non_uniform_result: self.add_ref(value),
requirements: UniformityRequirements::empty(),
@@ -504,6 +503,7 @@ impl FunctionInfo {
non_uniform_result: self.add_ref(vector),
requirements: UniformityRequirements::empty(),
},
E::Literal(_) | E::Constant(_) | E::ZeroValue(_) => Uniformity::new(),
E::Compose { ref components, .. } => {
let non_uniform_result = components
.iter()
@@ -969,6 +969,18 @@ impl FunctionInfo {
}
impl ModuleInfo {
/// Populates `self.const_expression_types`
pub(super) fn process_const_expression(
&mut self,
handle: Handle<crate::Expression>,
resolve_context: &ResolveContext,
gctx: crate::proc::GlobalCtx,
) -> Result<(), super::ConstExpressionError> {
self.const_expression_types[handle.index()] =
resolve_context.resolve(&gctx.const_expressions[handle], |h| Ok(&self[h]))?;
Ok(())
}
/// Builds the `FunctionInfo` based on the function, and validates the
/// uniform control flow if required by the expressions of this function.
pub(super) fn process_function(
@@ -1022,18 +1034,6 @@ impl ModuleInfo {
fn uniform_control_flow() {
use crate::{Expression as E, Statement as S};
let mut constant_arena = Arena::new();
let constant = constant_arena.append(
crate::Constant {
name: None,
specialization: None,
inner: crate::ConstantInner::Scalar {
width: 4,
value: crate::ScalarValue::Uint(0),
},
},
Default::default(),
);
let mut type_arena = crate::UniqueArena::new();
let ty = type_arena.insert(
crate::Type {
@@ -1070,7 +1070,7 @@ fn uniform_control_flow() {
let mut expressions = Arena::new();
// checks the uniform control flow
let constant_expr = expressions.append(E::Constant(constant), Default::default());
let constant_expr = expressions.append(E::Literal(crate::Literal::U32(0)), Default::default());
// checks the non-uniform control flow
let derivative_expr = expressions.append(
E::Derivative {
@@ -1110,7 +1110,7 @@ fn uniform_control_flow() {
sampling: crate::FastHashSet::default(),
};
let resolve_context = ResolveContext {
constants: &constant_arena,
constants: &Arena::new(),
types: &type_arena,
special_types: &crate::SpecialTypes::default(),
global_vars: &global_var_arena,

107
src/valid/expression.rs
View File

@@ -1,7 +1,7 @@
#[cfg(feature = "validate")]
use super::{
compose::validate_compose, validate_atomic_compare_exchange_struct, FunctionInfo, ShaderStages,
TypeFlags,
compose::validate_compose, validate_atomic_compare_exchange_struct, FunctionInfo, ModuleInfo,
ShaderStages, TypeFlags,
};
#[cfg(feature = "validate")]
use crate::arena::UniqueArena;
@@ -22,8 +22,10 @@ pub enum ExpressionError {
InvalidBaseType(Handle<crate::Expression>),
#[error("Accessing with index {0:?} can't be done")]
InvalidIndexType(Handle<crate::Expression>),
#[error("Accessing index {1:?} is out of {0:?} bounds")]
IndexOutOfBounds(Handle<crate::Expression>, crate::ScalarValue),
#[error("Accessing {0:?} via a negative index is invalid")]
NegativeIndex(Handle<crate::Expression>),
#[error("Accessing index {1} is out of {0:?} bounds")]
IndexOutOfBounds(Handle<crate::Expression>, u32),
#[error("The expression {0:?} may only be indexed by a constant")]
IndexMustBeConstant(Handle<crate::Expression>),
#[error("Function argument {0:?} doesn't exist")]
@@ -91,7 +93,7 @@ pub enum ExpressionError {
has_ref: bool,
},
#[error("Sample offset constant {1:?} doesn't match the image dimension {0:?}")]
InvalidSampleOffset(crate::ImageDimension, Handle<crate::Constant>),
InvalidSampleOffset(crate::ImageDimension, Handle<crate::Expression>),
#[error("Depth reference {0:?} is not a scalar float")]
InvalidDepthReference(Handle<crate::Expression>),
#[error("Depth sample level can only be Auto or Zero")]
@@ -151,6 +153,29 @@ impl<'a> std::ops::Index<Handle<crate::Expression>> for ExpressionTypeResolver<'
#[cfg(feature = "validate")]
impl super::Validator {
pub(super) fn validate_const_expression(
&self,
handle: Handle<crate::Expression>,
gctx: crate::proc::GlobalCtx,
mod_info: &mut ModuleInfo,
) -> Result<(), super::ConstExpressionError> {
use crate::Expression as E;
match gctx.const_expressions[handle] {
E::Literal(_) | E::Constant(_) | E::ZeroValue(_) => {}
E::Compose { ref components, ty } => {
validate_compose(
ty,
gctx,
components.iter().map(|&handle| mod_info[handle].clone()),
)?;
}
_ => return Err(super::ConstExpressionError::NonConst),
}
Ok(())
}
pub(super) fn validate_expression(
&self,
root: Handle<crate::Expression>,
@@ -158,7 +183,7 @@ impl super::Validator {
function: &crate::Function,
module: &crate::Module,
info: &FunctionInfo,
other_infos: &[FunctionInfo],
mod_info: &ModuleInfo,
) -> Result<ShaderStages, ExpressionError> {
use crate::{Expression as E, ScalarKind as Sk, TypeInner as Ti};
@@ -205,27 +230,19 @@ impl super::Validator {
if let crate::proc::IndexableLength::Known(known_length) =
base_type.indexable_length(module)?
{
if let E::Constant(k) = function.expressions[index] {
if let crate::Constant {
// We must treat specializable constants as unknown.
specialization: None,
// Non-scalar indices should have been caught above.
inner: crate::ConstantInner::Scalar { value, .. },
..
} = module.constants[k]
{
match value {
crate::ScalarValue::Uint(u) if u >= known_length as u64 => {
return Err(ExpressionError::IndexOutOfBounds(base, value));
}
crate::ScalarValue::Sint(s)
if s < 0 || s >= known_length as i64 =>
{
return Err(ExpressionError::IndexOutOfBounds(base, value));
}
_ => (),
match module
.to_ctx()
.eval_expr_to_u32_from(index, &function.expressions)
{
Ok(value) => {
if value >= known_length {
return Err(ExpressionError::IndexOutOfBounds(base, value));
}
}
Err(crate::proc::U32EvalError::Negative) => {
return Err(ExpressionError::NegativeIndex(base))
}
Err(crate::proc::U32EvalError::NonConst) => {}
}
}
@@ -263,16 +280,10 @@ impl super::Validator {
let limit = resolve_index_limit(module, base, &resolver[base], true)?;
if index >= limit {
return Err(ExpressionError::IndexOutOfBounds(
base,
crate::ScalarValue::Uint(limit as _),
));
return Err(ExpressionError::IndexOutOfBounds(base, limit));
}
ShaderStages::all()
}
E::Literal(_value) => ShaderStages::all(),
E::Constant(_handle) => ShaderStages::all(),
E::ZeroValue(_type) => ShaderStages::all(),
E::Splat { size: _, value } => match resolver[value] {
Ti::Scalar { .. } => ShaderStages::all(),
ref other => {
@@ -299,6 +310,7 @@ impl super::Validator {
}
ShaderStages::all()
}
E::Literal(_) | E::Constant(_) | E::ZeroValue(_) => ShaderStages::all(),
E::Compose { ref components, ty } => {
validate_compose(
ty,
@@ -412,26 +424,17 @@ impl super::Validator {
}
// check constant offset
if let Some(const_handle) = offset {
let good = match module.constants[const_handle].inner {
crate::ConstantInner::Scalar {
width: _,
value: crate::ScalarValue::Sint(_),
} => num_components == 1,
crate::ConstantInner::Scalar { .. } => false,
crate::ConstantInner::Composite { ty, .. } => {
match module.types[ty].inner {
Ti::Vector {
size,
kind: Sk::Sint,
..
} => size as u32 == num_components,
_ => false,
}
if let Some(const_expr) = offset {
match *mod_info[const_expr].inner_with(&module.types) {
Ti::Scalar { kind: Sk::Sint, .. } if num_components == 1 => {}
Ti::Vector {
size,
kind: Sk::Sint,
..
} if size as u32 == num_components => {}
_ => {
return Err(ExpressionError::InvalidSampleOffset(dim, const_expr));
}
};
if !good {
return Err(ExpressionError::InvalidSampleOffset(dim, const_handle));
}
}
@@ -1388,7 +1391,7 @@ impl super::Validator {
}
ShaderStages::all()
}
E::CallResult(function) => other_infos[function.index()].available_stages,
E::CallResult(function) => mod_info.functions[function.index()].available_stages,
E::AtomicResult { ty, comparison } => {
let scalar_predicate = |ty: &crate::TypeInner| match ty {
&crate::TypeInner::Scalar {

34
src/valid/function.rs
View File

@@ -941,6 +941,7 @@ impl super::Validator {
&self,
var: &crate::LocalVariable,
gctx: crate::proc::GlobalCtx,
mod_info: &ModuleInfo,
) -> Result<(), LocalVariableError> {
log::debug!("var {:?}", var);
let type_info = self
@@ -954,24 +955,14 @@ impl super::Validator {
return Err(LocalVariableError::InvalidType(var.ty));
}
if let Some(const_handle) = var.init {
match gctx.constants[const_handle].inner {
crate::ConstantInner::Scalar { width, ref value } => {
let ty_inner = crate::TypeInner::Scalar {
width,
kind: value.scalar_kind(),
};
if gctx.types[var.ty].inner != ty_inner {
return Err(LocalVariableError::InitializerType);
}
}
crate::ConstantInner::Composite { ty, components: _ } => {
if ty != var.ty {
return Err(LocalVariableError::InitializerType);
}
}
if let Some(init) = var.init {
let decl_ty = &gctx.types[var.ty].inner;
let init_ty = mod_info[init].inner_with(gctx.types);
if !decl_ty.equivalent(init_ty, gctx.types) {
return Err(LocalVariableError::InitializerType);
}
}
Ok(())
}
@@ -987,7 +978,7 @@ impl super::Validator {
#[cfg(feature = "validate")]
for (var_handle, var) in fun.local_variables.iter() {
self.validate_local_var(var, module.to_ctx())
self.validate_local_var(var, module.to_ctx(), mod_info)
.map_err(|source| {
FunctionError::LocalVariable {
handle: var_handle,
@@ -1061,14 +1052,7 @@ impl super::Validator {
}
#[cfg(feature = "validate")]
if self.flags.contains(super::ValidationFlags::EXPRESSIONS) {
match self.validate_expression(
handle,
expr,
fun,
module,
&info,
&mod_info.functions,
) {
match self.validate_expression(handle, expr, fun, module, &info, mod_info) {
Ok(stages) => info.available_stages &= stages,
Err(source) => {
return Err(FunctionError::Expression { handle, source }

83
src/valid/handles.rs
View File

@@ -40,6 +40,7 @@ impl super::Validator {
ref global_variables,
ref types,
ref special_types,
ref const_expressions,
} = module;
// NOTE: Types being first is important. All other forms of validation depend on this.
@@ -70,24 +71,24 @@ impl super::Validator {
}
}
let validate_type = |handle| Self::validate_type_handle(handle, types);
for (this_handle, constant) in constants.iter() {
let &crate::Constant {
name: _,
specialization: _,
ref inner,
} = constant;
match *inner {
crate::ConstantInner::Scalar { .. } => (),
crate::ConstantInner::Composite { ty, ref components } => {
validate_type(ty)?;
this_handle.check_dep_iter(components.iter().copied())?;
}
}
for handle_and_expr in const_expressions.iter() {
Self::validate_const_expression_handles(handle_and_expr, constants, types)?;
}
let validate_constant = |handle| Self::validate_constant_handle(handle, constants);
let validate_type = |handle| Self::validate_type_handle(handle, types);
let validate_const_expr =
|handle| Self::validate_expression_handle(handle, const_expressions);
for (_handle, constant) in constants.iter() {
let &crate::Constant {
name: _,
r#override: _,
ty,
init,
} = constant;
validate_type(ty)?;
validate_const_expr(init)?;
}
for (_handle, global_variable) in global_variables.iter() {
let &crate::GlobalVariable {
@@ -99,7 +100,7 @@ impl super::Validator {
} = global_variable;
validate_type(ty)?;
if let Some(init_expr) = init {
validate_constant(init_expr)?;
validate_const_expr(init_expr)?;
}
}
@@ -131,7 +132,7 @@ impl super::Validator {
let &crate::LocalVariable { name: _, ty, init } = local_variable;
validate_type(ty)?;
if let Some(init_constant) = init {
validate_constant(init_constant)?;
validate_const_expr(init_constant)?;
}
}
@@ -143,6 +144,7 @@ impl super::Validator {
Self::validate_expression_handles(
handle_and_expr,
constants,
const_expressions,
types,
local_variables,
global_variables,
@@ -202,9 +204,36 @@ impl super::Validator {
handle.check_valid_for(functions).map(|_| ())
}
fn validate_const_expression_handles(
(handle, expression): (Handle<crate::Expression>, &crate::Expression),
constants: &Arena<crate::Constant>,
types: &UniqueArena<crate::Type>,
) -> Result<(), InvalidHandleError> {
let validate_constant = |handle| Self::validate_constant_handle(handle, constants);
let validate_type = |handle| Self::validate_type_handle(handle, types);
match *expression {
crate::Expression::Literal(_) => {}
crate::Expression::Constant(constant) => {
validate_constant(constant)?;
}
crate::Expression::ZeroValue(ty) => {
validate_type(ty)?;
}
crate::Expression::Compose { ty, ref components } => {
validate_type(ty)?;
handle.check_dep_iter(components.iter().copied())?;
}
_ => {}
}
Ok(())
}
#[allow(clippy::too_many_arguments)]
fn validate_expression_handles(
(handle, expression): (Handle<crate::Expression>, &crate::Expression),
constants: &Arena<crate::Constant>,
const_expressions: &Arena<crate::Expression>,
types: &UniqueArena<crate::Type>,
local_variables: &Arena<crate::LocalVariable>,
global_variables: &Arena<crate::GlobalVariable>,
@@ -213,6 +242,8 @@ impl super::Validator {
current_function: Option<Handle<crate::Function>>,
) -> Result<(), InvalidHandleError> {
let validate_constant = |handle| Self::validate_constant_handle(handle, constants);
let validate_const_expr =
|handle| Self::validate_expression_handle(handle, const_expressions);
let validate_type = |handle| Self::validate_type_handle(handle, types);
match *expression {
@@ -222,19 +253,19 @@ impl super::Validator {
crate::Expression::AccessIndex { base, .. } => {
handle.check_dep(base)?;
}
crate::Expression::Literal(_value) => {}
crate::Expression::Constant(constant) => {
validate_constant(constant)?;
}
crate::Expression::ZeroValue(ty) => {
validate_type(ty)?;
}
crate::Expression::Splat { value, .. } => {
handle.check_dep(value)?;
}
crate::Expression::Swizzle { vector, .. } => {
handle.check_dep(vector)?;
}
crate::Expression::Literal(_) => {}
crate::Expression::Constant(constant) => {
validate_constant(constant)?;
}
crate::Expression::ZeroValue(ty) => {
validate_type(ty)?;
}
crate::Expression::Compose { ty, ref components } => {
validate_type(ty)?;
handle.check_dep_iter(components.iter().copied())?;
@@ -260,7 +291,7 @@ impl super::Validator {
depth_ref,
} => {
if let Some(offset) = offset {
validate_constant(offset)?;
validate_const_expr(offset)?;
}
handle

11
src/valid/interface.rs
View File

@@ -31,6 +31,8 @@ pub enum GlobalVariableError {
Handle<crate::Type>,
#[source] Disalignment,
),
#[error("Initializer doesn't match the variable type")]
InitializerType,
}
#[derive(Clone, Debug, thiserror::Error)]
@@ -396,6 +398,7 @@ impl super::Validator {
&self,
var: &crate::GlobalVariable,
gctx: crate::proc::GlobalCtx,
mod_info: &ModuleInfo,
) -> Result<(), GlobalVariableError> {
use super::TypeFlags;
@@ -509,6 +512,14 @@ impl super::Validator {
}
}
if let Some(init) = var.init {
let decl_ty = &gctx.types[var.ty].inner;
let init_ty = mod_info[init].inner_with(gctx.types);
if !decl_ty.equivalent(init_ty, gctx.types) {
return Err(GlobalVariableError::InitializerType);
}
}
Ok(())
}

119
src/valid/mod.rs
View File

@@ -10,12 +10,9 @@ mod handles;
mod interface;
mod r#type;
#[cfg(feature = "validate")]
use crate::arena::UniqueArena;
use crate::{
arena::Handle,
proc::{LayoutError, Layouter},
proc::{LayoutError, Layouter, TypeResolution},
FastHashSet,
};
use bit_set::BitSet;
@@ -143,6 +140,7 @@ pub struct ModuleInfo {
type_flags: Vec<TypeFlags>,
functions: Vec<FunctionInfo>,
entry_points: Vec<FunctionInfo>,
const_expression_types: Box<[TypeResolution]>,
}
impl ops::Index<Handle<crate::Type>> for ModuleInfo {
@@ -159,6 +157,13 @@ impl ops::Index<Handle<crate::Function>> for ModuleInfo {
}
}
impl ops::Index<Handle<crate::Expression>> for ModuleInfo {
type Output = TypeResolution;
fn index(&self, handle: Handle<crate::Expression>) -> &Self::Output {
&self.const_expression_types[handle.index()]
}
}
#[derive(Debug)]
pub struct Validator {
flags: ValidationFlags,
@@ -173,12 +178,22 @@ pub struct Validator {
valid_expression_set: BitSet,
}
#[derive(Clone, Debug, thiserror::Error)]
pub enum ConstExpressionError {
#[error("The expression is not a constant expression")]
NonConst,
#[error(transparent)]
Compose(#[from] ComposeError),
#[error("Type resolution failed")]
Type(#[from] crate::proc::ResolveError),
}
#[derive(Clone, Debug, thiserror::Error)]
pub enum ConstantError {
#[error("The type doesn't match the constant")]
InvalidType,
#[error(transparent)]
Compose(#[from] ComposeError),
#[error("The type is not constructible")]
NonConstructibleType,
}
#[derive(Clone, Debug, thiserror::Error)]
@@ -193,6 +208,11 @@ pub enum ValidationError {
name: String,
source: TypeError,
},
#[error("Constant expression {handle:?} is invalid")]
ConstExpression {
handle: Handle<crate::Expression>,
source: ConstExpressionError,
},
#[error("Constant {handle:?} '{name}' is invalid")]
Constant {
handle: Handle<crate::Constant>,
@@ -300,24 +320,21 @@ impl Validator {
&self,
handle: Handle<crate::Constant>,
gctx: crate::proc::GlobalCtx,
mod_info: &ModuleInfo,
) -> Result<(), ConstantError> {
let con = &gctx.constants[handle];
match con.inner {
crate::ConstantInner::Scalar { width, ref value } => {
if self.check_width(value.scalar_kind(), width).is_err() {
return Err(ConstantError::InvalidType);
}
}
crate::ConstantInner::Composite { ty, ref components } => {
compose::validate_compose(
ty,
gctx,
components
.iter()
.map(|&component| gctx.constants[component].inner.resolve_type()),
)?;
}
let type_info = &self.types[con.ty.index()];
if !type_info.flags.contains(TypeFlags::CONSTRUCTIBLE) {
return Err(ConstantError::NonConstructibleType);
}
let decl_ty = &gctx.types[con.ty].inner;
let init_ty = mod_info[con.init].inner_with(gctx.types);
if !decl_ty.equivalent(init_ty, gctx.types) {
return Err(ConstantError::InvalidType);
}
Ok(())
}
@@ -337,25 +354,17 @@ impl Validator {
ValidationError::from(e).with_span_handle(handle, &module.types)
})?;
#[cfg(feature = "validate")]
if self.flags.contains(ValidationFlags::CONSTANTS) {
for (handle, constant) in module.constants.iter() {
self.validate_constant(handle, module.to_ctx())
.map_err(|source| {
ValidationError::Constant {
handle,
name: constant.name.clone().unwrap_or_default(),
source,
}
.with_span_handle(handle, &module.constants)
})?
}
}
let placeholder = TypeResolution::Value(crate::TypeInner::Scalar {
kind: crate::ScalarKind::Bool,
width: 0,
});
let mut mod_info = ModuleInfo {
type_flags: Vec::with_capacity(module.types.len()),
functions: Vec::with_capacity(module.functions.len()),
entry_points: Vec::with_capacity(module.entry_points.len()),
const_expression_types: vec![placeholder; module.const_expressions.len()]
.into_boxed_slice(),
};
for (handle, ty) in module.types.iter() {
@@ -373,9 +382,45 @@ impl Validator {
self.types[handle.index()] = ty_info;
}
{
let t = crate::Arena::new();
let resolve_context = crate::proc::ResolveContext::with_locals(module, &t, &[]);
for (handle, _) in module.const_expressions.iter() {
mod_info
.process_const_expression(handle, &resolve_context, module.to_ctx())
.map_err(|source| {
ValidationError::ConstExpression { handle, source }
.with_span_handle(handle, &module.const_expressions)
})?
}
}
#[cfg(feature = "validate")]
if self.flags.contains(ValidationFlags::CONSTANTS) {
for (handle, _) in module.const_expressions.iter() {
self.validate_const_expression(handle, module.to_ctx(), &mut mod_info)
.map_err(|source| {
ValidationError::ConstExpression { handle, source }
.with_span_handle(handle, &module.const_expressions)
})?
}
for (handle, constant) in module.constants.iter() {
self.validate_constant(handle, module.to_ctx(), &mod_info)
.map_err(|source| {
ValidationError::Constant {
handle,
name: constant.name.clone().unwrap_or_default(),
source,
}
.with_span_handle(handle, &module.constants)
})?
}
}
#[cfg(feature = "validate")]
for (var_handle, var) in module.global_variables.iter() {
self.validate_global_var(var, module.to_ctx())
self.validate_global_var(var, module.to_ctx(), &mod_info)
.map_err(|source| {
ValidationError::GlobalVariable {
handle: var_handle,
@@ -434,7 +479,7 @@ impl Validator {
#[cfg(feature = "validate")]
fn validate_atomic_compare_exchange_struct(
types: &UniqueArena<crate::Type>,
types: &crate::UniqueArena<crate::Type>,
members: &[crate::StructMember],
scalar_predicate: impl FnOnce(&crate::TypeInner) -> bool,
) -> bool {

156
tests/out/analysis/access.info.ron
View File

@@ -77,15 +77,6 @@
space: Function,
)),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: None,
@@ -107,6 +98,15 @@
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: None,
@@ -708,15 +708,6 @@
space: Function,
)),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: None,
@@ -738,6 +729,15 @@
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: Some(55),
@@ -1294,15 +1294,6 @@
space: Function,
)),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: None,
@@ -1324,6 +1315,15 @@
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: None,
@@ -2051,15 +2051,6 @@
space: Function,
)),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: None,
@@ -2081,6 +2072,15 @@
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: Some(2),
requirements: (""),
),
ref_count: 1,
assignable_global: None,
ty: Handle(3),
),
(
uniformity: (
non_uniform_result: Some(66),
@@ -4290,4 +4290,88 @@
sampling: [],
),
],
const_expression_types: [
Value(Scalar(
kind: Uint,
width: 4,
)),
Value(Scalar(
kind: Uint,
width: 4,
)),
Value(Scalar(
kind: Uint,
width: 4,
)),
Value(Scalar(
kind: Uint,
width: 4,
)),
Handle(2),
Value(Scalar(
kind: Sint,
width: 4,
)),
Handle(4),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
Value(Scalar(
kind: Sint,
width: 4,
)),
],
)

1
tests/out/analysis/collatz.info.ron
View File

@@ -428,4 +428,5 @@
sampling: [],
),
],
const_expression_types: [],
)

787
tests/out/analysis/shadow.info.ron
View File

File diff suppressed because it is too large Load Diff

16
tests/out/glsl/access.foo_vert.Vertex.glsl
View File

@@ -37,8 +37,8 @@ void test_matrix_within_struct_accesses() {
int idx = 0;
Baz t = Baz(mat3x2(0.0));
idx = 1;
int _e2 = idx;
idx = (_e2 - 1);
int _e3 = idx;
idx = (_e3 - 1);
mat3x2 l0_ = _group_0_binding_1_vs.m;
vec2 l1_ = _group_0_binding_1_vs.m[0];
int _e15 = idx;
@@ -52,8 +52,8 @@ void test_matrix_within_struct_accesses() {
int _e43 = idx;
float l6_ = _group_0_binding_1_vs.m[_e41][_e43];
t = Baz(mat3x2(vec2(1.0), vec2(2.0), vec2(3.0)));
int _e55 = idx;
idx = (_e55 + 1);
int _e56 = idx;
idx = (_e56 + 1);
t.m = mat3x2(vec2(6.0), vec2(5.0), vec2(4.0));
t.m[0] = vec2(9.0);
int _e72 = idx;
@@ -73,8 +73,8 @@ void test_matrix_within_array_within_struct_accesses() {
int idx_1 = 0;
MatCx2InArray t_1 = MatCx2InArray(mat4x2[2](mat4x2(0.0), mat4x2(0.0)));
idx_1 = 1;
int _e2 = idx_1;
idx_1 = (_e2 - 1);
int _e3 = idx_1;
idx_1 = (_e3 - 1);
mat4x2 l0_1[2] = _group_0_binding_3_vs.am;
mat4x2 l1_1 = _group_0_binding_3_vs.am[0];
vec2 l2_1 = _group_0_binding_3_vs.am[0][0];
@@ -89,8 +89,8 @@ void test_matrix_within_array_within_struct_accesses() {
int _e60 = idx_1;
float l7_ = _group_0_binding_3_vs.am[0][_e58][_e60];
t_1 = MatCx2InArray(mat4x2[2](mat4x2(0.0), mat4x2(0.0)));
int _e66 = idx_1;
idx_1 = (_e66 + 1);
int _e67 = idx_1;
idx_1 = (_e67 + 1);
t_1.am = mat4x2[2](mat4x2(0.0), mat4x2(0.0));
t_1.am[0] = mat4x2(vec2(8.0), vec2(7.0), vec2(6.0), vec2(5.0));
t_1.am[0][0] = vec2(9.0);

6
tests/out/glsl/boids.main.Compute.glsl
View File

@@ -18,6 +18,8 @@ struct SimParams {
float rule2Scale;
float rule3Scale;
};
const uint NUM_PARTICLES = 1500u;
uniform SimParams_block_0Compute { SimParams _group_0_binding_0_cs; };
layout(std430) readonly buffer Particles_block_1Compute {
@@ -42,7 +44,7 @@ void main() {
vec2 vel = vec2(0.0);
uint i = 0u;
uint index = global_invocation_id.x;
if ((index >= 1500u)) {
if ((index >= NUM_PARTICLES)) {
return;
}
vec2 _e8 = _group_0_binding_1_cs.particles[index].pos;
@@ -63,7 +65,7 @@ void main() {
}
loop_init = false;
uint _e36 = i;
if ((_e36 >= 1500u)) {
if ((_e36 >= NUM_PARTICLES)) {
break;
}
uint _e39 = i;

4
tests/out/glsl/do-while.main.Fragment.glsl
View File

@@ -8,8 +8,8 @@ void fb1_(inout bool cond) {
bool loop_init = true;
while(true) {
if (!loop_init) {
bool _e6 = cond;
if (!(_e6)) {
bool _e2 = cond;
if (!(_e2)) {
break;
}
}

2
tests/out/glsl/globals.main.Compute.glsl
View File

@@ -9,6 +9,8 @@ struct FooStruct {
vec3 v3_;
float v1_;
};
const bool Foo_1 = true;
shared float wg[10];
shared uint at_1;

37
tests/out/glsl/operators.main.Compute.glsl
View File

@@ -9,16 +9,21 @@ struct Foo {
vec4 a;
int b;
};
const vec4 v_f32_one = vec4(1.0, 1.0, 1.0, 1.0);
const vec4 v_f32_zero = vec4(0.0, 0.0, 0.0, 0.0);
const vec4 v_f32_half = vec4(0.5, 0.5, 0.5, 0.5);
const ivec4 v_i32_one = ivec4(1, 1, 1, 1);
vec4 builtins() {
int s1_ = (true ? 1 : 0);
vec4 s2_ = (true ? vec4(1.0, 1.0, 1.0, 1.0) : vec4(0.0, 0.0, 0.0, 0.0));
vec4 s3_ = mix(vec4(1.0, 1.0, 1.0, 1.0), vec4(0.0, 0.0, 0.0, 0.0), bvec4(false, false, false, false));
vec4 m1_ = mix(vec4(0.0, 0.0, 0.0, 0.0), vec4(1.0, 1.0, 1.0, 1.0), vec4(0.5, 0.5, 0.5, 0.5));
vec4 m2_ = mix(vec4(0.0, 0.0, 0.0, 0.0), vec4(1.0, 1.0, 1.0, 1.0), 0.1);
float b1_ = intBitsToFloat(ivec4(1, 1, 1, 1).x);
vec4 b2_ = intBitsToFloat(ivec4(1, 1, 1, 1));
ivec4 v_i32_zero = ivec4(vec4(0.0, 0.0, 0.0, 0.0));
vec4 s2_ = (true ? v_f32_one : v_f32_zero);
vec4 s3_ = mix(v_f32_one, v_f32_zero, bvec4(false, false, false, false));
vec4 m1_ = mix(v_f32_zero, v_f32_one, v_f32_half);
vec4 m2_ = mix(v_f32_zero, v_f32_one, 0.1);
float b1_ = intBitsToFloat(v_i32_one.x);
vec4 b2_ = intBitsToFloat(v_i32_one);
ivec4 v_i32_zero = ivec4(v_f32_zero);
return (((((vec4((ivec4(s1_) + v_i32_zero)) + s2_) + m1_) + m2_) + vec4(b1_)) + b2_);
}
@@ -243,15 +248,15 @@ void assignment() {
a_1 = (_e27 << 2u);
int _e30 = a_1;
a_1 = (_e30 >> 1u);
int _e32 = a_1;
a_1 = (_e32 + 1);
int _e35 = a_1;
a_1 = (_e35 - 1);
int _e33 = a_1;
a_1 = (_e33 + 1);
int _e36 = a_1;
a_1 = (_e36 - 1);
vec0_ = ivec3(0);
int _e42 = vec0_.y;
vec0_.y = (_e42 + 1);
int _e47 = vec0_.y;
vec0_.y = (_e47 - 1);
int _e43 = vec0_.y;
vec0_.y = (_e43 + 1);
int _e48 = vec0_.y;
vec0_.y = (_e48 - 1);
return;
}
@@ -268,7 +273,7 @@ void negation_avoids_prefix_decrement() {
void main() {
vec4 _e0 = builtins();
vec4 _e1 = splat();
vec3 _e4 = bool_cast(vec4(1.0, 1.0, 1.0, 1.0).xyz);
vec3 _e4 = bool_cast(v_f32_one.xyz);
float _e5 = constructors();
logical();
arithmetic();

8
tests/out/glsl/quad-vert.main.Vertex.glsl
View File

@@ -26,10 +26,10 @@ layout(location = 0) in vec2 _p2vs_location0;
layout(location = 0) smooth out vec2 _vs2fs_location0;
void main_1() {
vec2 _e12 = a_uv_1;
v_uv = _e12;
vec2 _e13 = a_pos_1;
perVertexStruct.gen_gl_Position = vec4(_e13.x, _e13.y, 0.0, 1.0);
vec2 _e8 = a_uv_1;
v_uv = _e8;
vec2 _e9 = a_pos_1;
perVertexStruct.gen_gl_Position = vec4(_e9.x, _e9.y, 0.0, 1.0);
return;
}

2
tests/out/glsl/quad.frag_main.Fragment.glsl
View File

@@ -7,6 +7,8 @@ struct VertexOutput {
vec2 uv;
vec4 position;
};
const float c_scale = 1.2;
uniform highp sampler2D _group_0_binding_0_fs;
smooth in vec2 _vs2fs_location0;

2
tests/out/glsl/quad.fs_extra.Fragment.glsl
View File

@@ -7,6 +7,8 @@ struct VertexOutput {
vec2 uv;
vec4 position;
};
const float c_scale = 1.2;
layout(location = 0) out vec4 _fs2p_location0;
void main() {

4
tests/out/glsl/quad.vert_main.Vertex.glsl
View File

@@ -7,6 +7,8 @@ struct VertexOutput {
vec2 uv;
vec4 position;
};
const float c_scale = 1.2;
layout(location = 0) in vec2 _p2vs_location0;
layout(location = 1) in vec2 _p2vs_location1;
smooth out vec2 _vs2fs_location0;
@@ -14,7 +16,7 @@ smooth out vec2 _vs2fs_location0;
void main() {
vec2 pos = _p2vs_location0;
vec2 uv = _p2vs_location1;
VertexOutput _tmp_return = VertexOutput(uv, vec4((1.2000000476837158 * pos), 0.0, 1.0));
VertexOutput _tmp_return = VertexOutput(uv, vec4((c_scale * pos), 0.0, 1.0));
_vs2fs_location0 = _tmp_return.uv;
gl_Position = _tmp_return.position;
return;

11
tests/out/glsl/shadow.fs_main.Fragment.glsl
View File

@@ -21,6 +21,9 @@ struct Light {
vec4 pos;
vec4 color;
};
const vec3 c_ambient = vec3(0.05, 0.05, 0.05);
const uint c_max_lights = 10u;
uniform Globals_block_0Fragment { Globals _group_0_binding_0_fs; };
uniform Entity_block_1Fragment { Entity _group_1_binding_0_fs; };
@@ -49,18 +52,18 @@ void main() {
vec3 color = vec3(0.0);
uint i = 0u;
vec3 normal_1 = normalize(in_.world_normal);
color = vec3(0.05000000074505806, 0.05000000074505806, 0.05000000074505806);
color = c_ambient;
i = 0u;
bool loop_init = true;
while(true) {
if (!loop_init) {
uint _e39 = i;
i = (_e39 + 1u);
uint _e40 = i;
i = (_e40 + 1u);
}
loop_init = false;
uint _e7 = i;
uint _e11 = _group_0_binding_0_fs.num_lights.x;
if ((_e7 < min(_e11, 10u))) {
if ((_e7 < min(_e11, c_max_lights))) {
} else {
break;
}

11
tests/out/glsl/shadow.fs_main_without_storage.Fragment.glsl
View File

@@ -21,6 +21,9 @@ struct Light {
vec4 pos;
vec4 color;
};
const vec3 c_ambient = vec3(0.05, 0.05, 0.05);
const uint c_max_lights = 10u;
uniform Globals_block_0Fragment { Globals _group_0_binding_0_fs; };
uniform Entity_block_1Fragment { Entity _group_1_binding_0_fs; };
@@ -49,18 +52,18 @@ void main() {
vec3 color_1 = vec3(0.0);
uint i_1 = 0u;
vec3 normal_1 = normalize(in_1.world_normal);
color_1 = vec3(0.05000000074505806, 0.05000000074505806, 0.05000000074505806);
color_1 = c_ambient;
i_1 = 0u;
bool loop_init = true;
while(true) {
if (!loop_init) {
uint _e39 = i_1;
i_1 = (_e39 + 1u);
uint _e40 = i_1;
i_1 = (_e40 + 1u);
}
loop_init = false;
uint _e7 = i_1;
uint _e11 = _group_0_binding_0_fs.num_lights.x;
if ((_e7 < min(_e11, 10u))) {
if ((_e7 < min(_e11, c_max_lights))) {
} else {
break;
}

3
tests/out/glsl/shadow.vs_main.Vertex.glsl
View File

@@ -21,6 +21,9 @@ struct Light {
vec4 pos;
vec4 color;
};
const vec3 c_ambient = vec3(0.05, 0.05, 0.05);
const uint c_max_lights = 10u;
uniform Globals_block_0Vertex { Globals _group_0_binding_0_vs; };
uniform Entity_block_1Vertex { Entity _group_1_binding_0_vs; };

2
tests/out/glsl/workgroup-uniform-load.test_workgroupUniformLoad.Compute.glsl
View File

@@ -5,6 +5,8 @@ precision highp int;
layout(local_size_x = 4, local_size_y = 1, local_size_z = 1) in;
const uint SIZE = 128u;
shared int arr_i32_[128];

45
tests/out/hlsl/access.hlsl
View File

@@ -1,4 +1,3 @@
typedef struct { float2 _0; float2 _1; } __mat2x2;
float2 __get_col_of_mat2x2(__mat2x2 mat, uint idx) {
switch(idx) {
@@ -84,6 +83,14 @@ RWByteAddressBuffer qux : register(u2);
cbuffer nested_mat_cx2_ : register(b3) { MatCx2InArray nested_mat_cx2_; }
groupshared uint val;
Baz ConstructBaz(float3x2 arg0) {
Baz ret = (Baz)0;
ret.m_0 = arg0[0];
ret.m_1 = arg0[1];
ret.m_2 = arg0[2];
return ret;
}
float3x2 GetMatmOnBaz(Baz obj) {
return float3x2(obj.m_0, obj.m_1, obj.m_2);
}
@@ -110,22 +117,14 @@ void SetMatScalarmOnBaz(Baz obj, float scalar, uint mat_idx, uint vec_idx) {
}
}
Baz ConstructBaz(float3x2 arg0) {
Baz ret = (Baz)0;
ret.m_0 = arg0[0];
ret.m_1 = arg0[1];
ret.m_2 = arg0[2];
return ret;
}
void test_matrix_within_struct_accesses()
{
int idx = (int)0;
Baz t = (Baz)0;
idx = 1;
int _expr2 = idx;
idx = (_expr2 - 1);
int _expr3 = idx;
idx = (_expr3 - 1);
float3x2 l0_ = GetMatmOnBaz(baz);
float2 l1_ = GetMatmOnBaz(baz)[0];
int _expr15 = idx;
@@ -139,8 +138,8 @@ void test_matrix_within_struct_accesses()
int _expr43 = idx;
float l6_ = GetMatmOnBaz(baz)[_expr41][_expr43];
t = ConstructBaz(float3x2((1.0).xx, (2.0).xx, (3.0).xx));
int _expr55 = idx;
idx = (_expr55 + 1);
int _expr56 = idx;
idx = (_expr56 + 1);
SetMatmOnBaz(t, float3x2((6.0).xx, (5.0).xx, (4.0).xx));
t.m_0 = (9.0).xx;
int _expr72 = idx;
@@ -168,8 +167,8 @@ void test_matrix_within_array_within_struct_accesses()
MatCx2InArray t_1 = (MatCx2InArray)0;
idx_1 = 1;
int _expr2 = idx_1;
idx_1 = (_expr2 - 1);
int _expr3 = idx_1;
idx_1 = (_expr3 - 1);
float4x2 l0_1[2] = ((float4x2[2])nested_mat_cx2_.am);
float4x2 l1_1 = ((float4x2)nested_mat_cx2_.am[0]);
float2 l2_1 = nested_mat_cx2_.am[0]._0;
@@ -184,8 +183,8 @@ void test_matrix_within_array_within_struct_accesses()
int _expr60 = idx_1;
float l7_ = __get_col_of_mat4x2(nested_mat_cx2_.am[0], _expr58)[_expr60];
t_1 = ConstructMatCx2InArray((float4x2[2])0);
int _expr66 = idx_1;
idx_1 = (_expr66 + 1);
int _expr67 = idx_1;
idx_1 = (_expr67 + 1);
t_1.am = (__mat4x2[2])(float4x2[2])0;
t_1.am[0] = (__mat4x2)float4x2((8.0).xx, (7.0).xx, (6.0).xx, (5.0).xx);
t_1.am[0]._0 = (9.0).xx;
@@ -231,6 +230,12 @@ void assign_array_through_ptr_fn(inout float4 foo_2[2])
return;
}
typedef int ret_Constructarray5_int_[5];
ret_Constructarray5_int_ Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4) {
int ret[5] = { arg0, arg1, arg2, arg3, arg4 };
return ret;
}
typedef uint2 ret_Constructarray2_uint2_[2];
ret_Constructarray2_uint2_ Constructarray2_uint2_(uint2 arg0, uint2 arg1) {
uint2 ret[2] = { arg0, arg1 };
@@ -244,12 +249,6 @@ uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
return ret;
}
typedef int ret_Constructarray5_int_[5];
ret_Constructarray5_int_ Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4) {
int ret[5] = { arg0, arg1, arg2, arg3, arg4 };
return ret;
}
float4 foo_vert(uint vi : SV_VertexID) : SV_Position
{
float foo = (float)0;

1
tests/out/hlsl/array-in-ctor.hlsl
View File

@@ -1,4 +1,3 @@
struct Ah {
float inner[2];
};

1
tests/out/hlsl/atomicOps.hlsl
View File

@@ -1,4 +1,3 @@
struct Struct {
uint atomic_scalar;
int atomic_arr[2];

219
tests/out/hlsl/binding-arrays.hlsl
View File

@@ -1,4 +1,3 @@
struct UniformIndex {
uint index;
};
@@ -71,115 +70,115 @@ float4 main(FragmentInput_main fragmentinput_main) : SV_Target0
u2_ = (_expr28 + NagaDimensions2D(texture_array_unbounded[uniform_index]));
uint2 _expr33 = u2_;
u2_ = (_expr33 + NagaDimensions2D(texture_array_unbounded[NonUniformResourceIndex(non_uniform_index)]));
float4 _expr41 = texture_array_bounded[0].Gather(samp[0], uv);
float4 _expr42 = v4_;
v4_ = (_expr42 + _expr41);
float4 _expr48 = texture_array_bounded[uniform_index].Gather(samp[uniform_index], uv);
float4 _expr49 = v4_;
v4_ = (_expr49 + _expr48);
float4 _expr55 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].Gather(samp[NonUniformResourceIndex(non_uniform_index)], uv);
float4 _expr56 = v4_;
v4_ = (_expr56 + _expr55);
float4 _expr65 = texture_array_depth[0].GatherCmp(samp_comp[0], uv, 0.0);
float4 _expr66 = v4_;
v4_ = (_expr66 + _expr65);
float4 _expr73 = texture_array_depth[uniform_index].GatherCmp(samp_comp[uniform_index], uv, 0.0);
float4 _expr74 = v4_;
v4_ = (_expr74 + _expr73);
float4 _expr81 = texture_array_depth[NonUniformResourceIndex(non_uniform_index)].GatherCmp(samp_comp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float4 _expr82 = v4_;
v4_ = (_expr82 + _expr81);
float4 _expr88 = texture_array_unbounded[0].Load(int3(pix, 0));
float4 _expr89 = v4_;
v4_ = (_expr89 + _expr88);
float4 _expr94 = texture_array_unbounded[uniform_index].Load(int3(pix, 0));
float4 _expr95 = v4_;
v4_ = (_expr95 + _expr94);
float4 _expr100 = texture_array_unbounded[NonUniformResourceIndex(non_uniform_index)].Load(int3(pix, 0));
float4 _expr101 = v4_;
v4_ = (_expr101 + _expr100);
uint _expr107 = u1_;
u1_ = (_expr107 + NagaNumLayers2DArray(texture_array_2darray[0]));
uint _expr112 = u1_;
u1_ = (_expr112 + NagaNumLayers2DArray(texture_array_2darray[uniform_index]));
uint _expr117 = u1_;
u1_ = (_expr117 + NagaNumLayers2DArray(texture_array_2darray[NonUniformResourceIndex(non_uniform_index)]));
uint _expr123 = u1_;
u1_ = (_expr123 + NagaNumLevels2D(texture_array_bounded[0]));
uint _expr128 = u1_;
u1_ = (_expr128 + NagaNumLevels2D(texture_array_bounded[uniform_index]));
uint _expr133 = u1_;
u1_ = (_expr133 + NagaNumLevels2D(texture_array_bounded[NonUniformResourceIndex(non_uniform_index)]));
uint _expr139 = u1_;
u1_ = (_expr139 + NagaMSNumSamples2D(texture_array_multisampled[0]));
uint _expr144 = u1_;
u1_ = (_expr144 + NagaMSNumSamples2D(texture_array_multisampled[uniform_index]));
uint _expr149 = u1_;
u1_ = (_expr149 + NagaMSNumSamples2D(texture_array_multisampled[NonUniformResourceIndex(non_uniform_index)]));
float4 _expr157 = texture_array_bounded[0].Sample(samp[0], uv);
float4 _expr158 = v4_;
v4_ = (_expr158 + _expr157);
float4 _expr164 = texture_array_bounded[uniform_index].Sample(samp[uniform_index], uv);
float4 _expr165 = v4_;
v4_ = (_expr165 + _expr164);
float4 _expr171 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].Sample(samp[NonUniformResourceIndex(non_uniform_index)], uv);
float4 _expr172 = v4_;
v4_ = (_expr172 + _expr171);
float4 _expr181 = texture_array_bounded[0].SampleBias(samp[0], uv, 0.0);
float4 _expr182 = v4_;
v4_ = (_expr182 + _expr181);
float4 _expr189 = texture_array_bounded[uniform_index].SampleBias(samp[uniform_index], uv, 0.0);
float4 _expr190 = v4_;
v4_ = (_expr190 + _expr189);
float4 _expr197 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].SampleBias(samp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float4 _expr198 = v4_;
v4_ = (_expr198 + _expr197);
float _expr207 = texture_array_depth[0].SampleCmp(samp_comp[0], uv, 0.0);
float _expr208 = v1_;
v1_ = (_expr208 + _expr207);
float _expr215 = texture_array_depth[uniform_index].SampleCmp(samp_comp[uniform_index], uv, 0.0);
float _expr216 = v1_;
v1_ = (_expr216 + _expr215);
float _expr223 = texture_array_depth[NonUniformResourceIndex(non_uniform_index)].SampleCmp(samp_comp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float _expr224 = v1_;
v1_ = (_expr224 + _expr223);
float _expr233 = texture_array_depth[0].SampleCmpLevelZero(samp_comp[0], uv, 0.0);
float _expr234 = v1_;
v1_ = (_expr234 + _expr233);
float _expr241 = texture_array_depth[uniform_index].SampleCmpLevelZero(samp_comp[uniform_index], uv, 0.0);
float _expr242 = v1_;
v1_ = (_expr242 + _expr241);
float _expr249 = texture_array_depth[NonUniformResourceIndex(non_uniform_index)].SampleCmpLevelZero(samp_comp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float _expr250 = v1_;
v1_ = (_expr250 + _expr249);
float4 _expr258 = texture_array_bounded[0].SampleGrad(samp[0], uv, uv, uv);
float4 _expr259 = v4_;
v4_ = (_expr259 + _expr258);
float4 _expr265 = texture_array_bounded[uniform_index].SampleGrad(samp[uniform_index], uv, uv, uv);
float4 _expr266 = v4_;
v4_ = (_expr266 + _expr265);
float4 _expr272 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].SampleGrad(samp[NonUniformResourceIndex(non_uniform_index)], uv, uv, uv);
float4 _expr273 = v4_;
v4_ = (_expr273 + _expr272);
float4 _expr282 = texture_array_bounded[0].SampleLevel(samp[0], uv, 0.0);
float4 _expr283 = v4_;
v4_ = (_expr283 + _expr282);
float4 _expr290 = texture_array_bounded[uniform_index].SampleLevel(samp[uniform_index], uv, 0.0);
float4 _expr291 = v4_;
v4_ = (_expr291 + _expr290);
float4 _expr298 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].SampleLevel(samp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float4 _expr299 = v4_;
v4_ = (_expr299 + _expr298);
float4 _expr304 = v4_;
texture_array_storage[0][pix] = _expr304;
float4 _expr42 = texture_array_bounded[0].Gather(samp[0], uv);
float4 _expr43 = v4_;
v4_ = (_expr43 + _expr42);
float4 _expr50 = texture_array_bounded[uniform_index].Gather(samp[uniform_index], uv);
float4 _expr51 = v4_;
v4_ = (_expr51 + _expr50);
float4 _expr58 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].Gather(samp[NonUniformResourceIndex(non_uniform_index)], uv);
float4 _expr59 = v4_;
v4_ = (_expr59 + _expr58);
float4 _expr68 = texture_array_depth[0].GatherCmp(samp_comp[0], uv, 0.0);
float4 _expr69 = v4_;
v4_ = (_expr69 + _expr68);
float4 _expr76 = texture_array_depth[uniform_index].GatherCmp(samp_comp[uniform_index], uv, 0.0);
float4 _expr77 = v4_;
v4_ = (_expr77 + _expr76);
float4 _expr84 = texture_array_depth[NonUniformResourceIndex(non_uniform_index)].GatherCmp(samp_comp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float4 _expr85 = v4_;
v4_ = (_expr85 + _expr84);
float4 _expr91 = texture_array_unbounded[0].Load(int3(pix, 0));
float4 _expr92 = v4_;
v4_ = (_expr92 + _expr91);
float4 _expr97 = texture_array_unbounded[uniform_index].Load(int3(pix, 0));
float4 _expr98 = v4_;
v4_ = (_expr98 + _expr97);
float4 _expr103 = texture_array_unbounded[NonUniformResourceIndex(non_uniform_index)].Load(int3(pix, 0));
float4 _expr104 = v4_;
v4_ = (_expr104 + _expr103);
uint _expr110 = u1_;
u1_ = (_expr110 + NagaNumLayers2DArray(texture_array_2darray[0]));
uint _expr115 = u1_;
u1_ = (_expr115 + NagaNumLayers2DArray(texture_array_2darray[uniform_index]));
uint _expr120 = u1_;
u1_ = (_expr120 + NagaNumLayers2DArray(texture_array_2darray[NonUniformResourceIndex(non_uniform_index)]));
uint _expr126 = u1_;
u1_ = (_expr126 + NagaNumLevels2D(texture_array_bounded[0]));
uint _expr131 = u1_;
u1_ = (_expr131 + NagaNumLevels2D(texture_array_bounded[uniform_index]));
uint _expr136 = u1_;
u1_ = (_expr136 + NagaNumLevels2D(texture_array_bounded[NonUniformResourceIndex(non_uniform_index)]));
uint _expr142 = u1_;
u1_ = (_expr142 + NagaMSNumSamples2D(texture_array_multisampled[0]));
uint _expr147 = u1_;
u1_ = (_expr147 + NagaMSNumSamples2D(texture_array_multisampled[uniform_index]));
uint _expr152 = u1_;
u1_ = (_expr152 + NagaMSNumSamples2D(texture_array_multisampled[NonUniformResourceIndex(non_uniform_index)]));
float4 _expr160 = texture_array_bounded[0].Sample(samp[0], uv);
float4 _expr161 = v4_;
v4_ = (_expr161 + _expr160);
float4 _expr167 = texture_array_bounded[uniform_index].Sample(samp[uniform_index], uv);
float4 _expr168 = v4_;
v4_ = (_expr168 + _expr167);
float4 _expr174 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].Sample(samp[NonUniformResourceIndex(non_uniform_index)], uv);
float4 _expr175 = v4_;
v4_ = (_expr175 + _expr174);
float4 _expr184 = texture_array_bounded[0].SampleBias(samp[0], uv, 0.0);
float4 _expr185 = v4_;
v4_ = (_expr185 + _expr184);
float4 _expr192 = texture_array_bounded[uniform_index].SampleBias(samp[uniform_index], uv, 0.0);
float4 _expr193 = v4_;
v4_ = (_expr193 + _expr192);
float4 _expr200 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].SampleBias(samp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float4 _expr201 = v4_;
v4_ = (_expr201 + _expr200);
float _expr210 = texture_array_depth[0].SampleCmp(samp_comp[0], uv, 0.0);
float _expr211 = v1_;
v1_ = (_expr211 + _expr210);
float _expr218 = texture_array_depth[uniform_index].SampleCmp(samp_comp[uniform_index], uv, 0.0);
float _expr219 = v1_;
v1_ = (_expr219 + _expr218);
float _expr226 = texture_array_depth[NonUniformResourceIndex(non_uniform_index)].SampleCmp(samp_comp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float _expr227 = v1_;
v1_ = (_expr227 + _expr226);
float _expr236 = texture_array_depth[0].SampleCmpLevelZero(samp_comp[0], uv, 0.0);
float _expr237 = v1_;
v1_ = (_expr237 + _expr236);
float _expr244 = texture_array_depth[uniform_index].SampleCmpLevelZero(samp_comp[uniform_index], uv, 0.0);
float _expr245 = v1_;
v1_ = (_expr245 + _expr244);
float _expr252 = texture_array_depth[NonUniformResourceIndex(non_uniform_index)].SampleCmpLevelZero(samp_comp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float _expr253 = v1_;
v1_ = (_expr253 + _expr252);
float4 _expr261 = texture_array_bounded[0].SampleGrad(samp[0], uv, uv, uv);
float4 _expr262 = v4_;
v4_ = (_expr262 + _expr261);
float4 _expr268 = texture_array_bounded[uniform_index].SampleGrad(samp[uniform_index], uv, uv, uv);
float4 _expr269 = v4_;
v4_ = (_expr269 + _expr268);
float4 _expr275 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].SampleGrad(samp[NonUniformResourceIndex(non_uniform_index)], uv, uv, uv);
float4 _expr276 = v4_;
v4_ = (_expr276 + _expr275);
float4 _expr285 = texture_array_bounded[0].SampleLevel(samp[0], uv, 0.0);
float4 _expr286 = v4_;
v4_ = (_expr286 + _expr285);
float4 _expr293 = texture_array_bounded[uniform_index].SampleLevel(samp[uniform_index], uv, 0.0);
float4 _expr294 = v4_;
v4_ = (_expr294 + _expr293);
float4 _expr301 = texture_array_bounded[NonUniformResourceIndex(non_uniform_index)].SampleLevel(samp[NonUniformResourceIndex(non_uniform_index)], uv, 0.0);
float4 _expr302 = v4_;
v4_ = (_expr302 + _expr301);
float4 _expr307 = v4_;
texture_array_storage[uniform_index][pix] = _expr307;
texture_array_storage[0][pix] = _expr307;
float4 _expr310 = v4_;
texture_array_storage[NonUniformResourceIndex(non_uniform_index)][pix] = _expr310;
uint2 _expr311 = u2_;
uint _expr312 = u1_;
float2 v2_ = float2((_expr311 + (_expr312).xx));
float4 _expr316 = v4_;
float _expr323 = v1_;
return ((_expr316 + float4(v2_.x, v2_.y, v2_.x, v2_.y)) + (_expr323).xxxx);
texture_array_storage[uniform_index][pix] = _expr310;
float4 _expr313 = v4_;
texture_array_storage[NonUniformResourceIndex(non_uniform_index)][pix] = _expr313;
uint2 _expr314 = u2_;
uint _expr315 = u1_;
float2 v2_ = float2((_expr314 + (_expr315).xx));
float4 _expr319 = v4_;
float _expr326 = v1_;
return ((_expr319 + float4(v2_.x, v2_.y, v2_.x, v2_.y)) + (_expr326).xxxx);
}

1
tests/out/hlsl/bitcast.hlsl
View File

@@ -1,4 +1,3 @@
[numthreads(1, 1, 1)]
void main()
{

1
tests/out/hlsl/bits.hlsl
View File

@@ -1,4 +1,3 @@
[numthreads(1, 1, 1)]
void main()
{

4
tests/out/hlsl/boids.hlsl
View File

@@ -1,5 +1,3 @@
static const uint NUM_PARTICLES = 1500;
struct Particle {
float2 pos;
float2 vel;
@@ -15,6 +13,8 @@ struct SimParams {
float rule3Scale;
};
static const uint NUM_PARTICLES = 1500u;
cbuffer params : register(b0) { SimParams params; }
ByteAddressBuffer particlesSrc : register(t1);
RWByteAddressBuffer particlesDst : register(u2);

1
tests/out/hlsl/break-if.hlsl
View File

@@ -1,4 +1,3 @@
void breakIfEmpty()
{
bool loop_init = true;

1
tests/out/hlsl/collatz.hlsl
View File

@@ -1,4 +1,3 @@
RWByteAddressBuffer v_indices : register(u0);
uint collatz_iterations(uint n_base)

1
tests/out/hlsl/control-flow.hlsl
View File

@@ -1,4 +1,3 @@
void switch_default_break(int i)
{
switch(i) {

5
tests/out/hlsl/do-while.hlsl
View File

@@ -1,11 +1,10 @@
void fb1_(inout bool cond)
{
bool loop_init = true;
while(true) {
if (!loop_init) {
bool _expr6 = cond;
if (!(_expr6)) {
bool _expr2 = cond;
if (!(_expr2)) {
break;
}
}

5
tests/out/hlsl/empty-global-name.hlsl
View File

@@ -1,4 +1,3 @@
struct type_1 {
int member;
};
@@ -7,8 +6,8 @@ RWByteAddressBuffer unnamed : register(u0);
void function()
{
int _expr8 = asint(unnamed.Load(0));
unnamed.Store(0, asuint((_expr8 + 1)));
int _expr4 = asint(unnamed.Load(0));
unnamed.Store(0, asuint((_expr4 + 1)));
return;
}

1
tests/out/hlsl/empty.hlsl
View File

@@ -1,4 +1,3 @@
[numthreads(1, 1, 1)]
void main()
{

1
tests/out/hlsl/fragment-output.hlsl
View File

@@ -1,4 +1,3 @@
struct FragmentOutputVec4Vec3_ {
float4 vec4f : SV_Target0;
nointerpolation int4 vec4i : SV_Target1;

1
tests/out/hlsl/functions.hlsl
View File

@@ -1,4 +1,3 @@
float2 test_fma()
{
float2 a = float2(2.0, 2.0);

4
tests/out/hlsl/globals.hlsl
View File

@@ -1,5 +1,3 @@
static const bool Foo_1 = true;
typedef struct { float2 _0; float2 _1; float2 _2; } __mat3x2;
float2 __get_col_of_mat3x2(__mat3x2 mat, uint idx) {
switch(idx) {
@@ -56,6 +54,8 @@ struct FooStruct {
float v1_;
};
static const bool Foo_1 = true;
groupshared float wg[10];
groupshared uint at_1;
RWByteAddressBuffer alignment : register(u1);

1
tests/out/hlsl/hlsl-keyword.hlsl
View File

@@ -1,4 +1,3 @@
float4 fs_main() : SV_Target0
{
float4 Pass_ = (float4)0;

1
tests/out/hlsl/image.hlsl
View File

@@ -1,4 +1,3 @@
Texture2D<uint4> image_mipmapped_src : register(t0);
Texture2DMS<uint4> image_multisampled_src : register(t3);
Texture2DMS<float> image_depth_multisampled_src : register(t4);

1
tests/out/hlsl/interpolate.hlsl
View File

@@ -1,4 +1,3 @@
struct FragmentInput {
float4 position : SV_Position;
nointerpolation uint _flat : LOC0;

11
tests/out/hlsl/inv-hyperbolic-trig-functions.hlsl
View File

@@ -1,4 +1,3 @@
static float a = (float)0;
void main_1()
@@ -7,12 +6,12 @@ void main_1()
float c = (float)0;
float d = (float)0;
float _expr4 = a;
b = log(_expr4 + sqrt(_expr4 * _expr4 + 1.0));
float _expr6 = a;
c = log(_expr6 + sqrt(_expr6 * _expr6 - 1.0));
float _expr8 = a;
b = log(_expr8 + sqrt(_expr8 * _expr8 + 1.0));
float _expr10 = a;
c = log(_expr10 + sqrt(_expr10 * _expr10 - 1.0));
float _expr12 = a;
d = 0.5 * log((1.0 + _expr12) / (1.0 - _expr12));
d = 0.5 * log((1.0 + _expr8) / (1.0 - _expr8));
return;
}

1
tests/out/hlsl/math-functions.hlsl
View File

@@ -1,4 +1,3 @@
void main()
{
float4 v = (0.0).xxxx;

42
tests/out/hlsl/operators.hlsl
View File

@@ -1,8 +1,3 @@
static const float4 v_f32_one = float4(1.0, 1.0, 1.0, 1.0);
static const float4 v_f32_zero = float4(0.0, 0.0, 0.0, 0.0);
static const float4 v_f32_half = float4(0.5, 0.5, 0.5, 0.5);
static const int4 v_i32_one = int4(1, 1, 1, 1);
struct Foo {
float4 a;
int b;
@@ -11,16 +6,21 @@ struct Foo {
int _end_pad_2;
};
static const float4 v_f32_one = float4(1.0, 1.0, 1.0, 1.0);
static const float4 v_f32_zero = float4(0.0, 0.0, 0.0, 0.0);
static const float4 v_f32_half = float4(0.5, 0.5, 0.5, 0.5);
static const int4 v_i32_one = int4(1, 1, 1, 1);
float4 builtins()
{
int s1_ = (true ? 1 : 0);
float4 s2_ = (true ? float4(1.0, 1.0, 1.0, 1.0) : float4(0.0, 0.0, 0.0, 0.0));
float4 s3_ = (bool4(false, false, false, false) ? float4(0.0, 0.0, 0.0, 0.0) : float4(1.0, 1.0, 1.0, 1.0));
float4 m1_ = lerp(float4(0.0, 0.0, 0.0, 0.0), float4(1.0, 1.0, 1.0, 1.0), float4(0.5, 0.5, 0.5, 0.5));
float4 m2_ = lerp(float4(0.0, 0.0, 0.0, 0.0), float4(1.0, 1.0, 1.0, 1.0), 0.1);
float b1_ = asfloat(int4(1, 1, 1, 1).x);
float4 b2_ = asfloat(int4(1, 1, 1, 1));
int4 v_i32_zero = int4(float4(0.0, 0.0, 0.0, 0.0));
float4 s2_ = (true ? v_f32_one : v_f32_zero);
float4 s3_ = (bool4(false, false, false, false) ? v_f32_zero : v_f32_one);
float4 m1_ = lerp(v_f32_zero, v_f32_one, v_f32_half);
float4 m2_ = lerp(v_f32_zero, v_f32_one, 0.1);
float b1_ = asfloat(v_i32_one.x);
float4 b2_ = asfloat(v_i32_one);
int4 v_i32_zero = int4(v_f32_zero);
return (((((float4(((s1_).xxxx + v_i32_zero)) + s2_) + m1_) + m2_) + (b1_).xxxx) + b2_);
}
@@ -270,15 +270,15 @@ void assignment()
a_1 = (_expr27 << 2u);
int _expr30 = a_1;
a_1 = (_expr30 >> 1u);
int _expr32 = a_1;
a_1 = (_expr32 + 1);
int _expr35 = a_1;
a_1 = (_expr35 - 1);
int _expr33 = a_1;
a_1 = (_expr33 + 1);
int _expr36 = a_1;
a_1 = (_expr36 - 1);
vec0_ = (int3)0;
int _expr42 = vec0_.y;
vec0_.y = (_expr42 + 1);
int _expr47 = vec0_.y;
vec0_.y = (_expr47 - 1);
int _expr43 = vec0_.y;
vec0_.y = (_expr43 + 1);
int _expr48 = vec0_.y;
vec0_.y = (_expr48 - 1);
return;
}
@@ -298,7 +298,7 @@ void main()
{
const float4 _e0 = builtins();
const float4 _e1 = splat();
const float3 _e4 = bool_cast(float4(1.0, 1.0, 1.0, 1.0).xyz);
const float3 _e4 = bool_cast(v_f32_one.xyz);
const float _e5 = constructors();
logical();
arithmetic();

1
tests/out/hlsl/padding.hlsl
View File

@@ -1,4 +1,3 @@
struct S {
float3 a;
int _end_pad_0;

17
tests/out/hlsl/quad-vert.hlsl
View File

@@ -1,4 +1,3 @@
struct gl_PerVertex {
float4 gl_Position : SV_Position;
float gl_PointSize;
@@ -12,12 +11,6 @@ struct type_9 {
float4 gl_Position : SV_Position;
};
typedef float ret_Constructarray1_float_[1];
ret_Constructarray1_float_ Constructarray1_float_(float arg0) {
float ret[1] = { arg0 };
return ret;
}
gl_PerVertex Constructgl_PerVertex(float4 arg0, float arg1, float arg2[1], float arg3[1]) {
gl_PerVertex ret = (gl_PerVertex)0;
ret.gl_Position = arg0;
@@ -29,7 +22,7 @@ gl_PerVertex Constructgl_PerVertex(float4 arg0, float arg1, float arg2[1], float
static float2 v_uv = (float2)0;
static float2 a_uv_1 = (float2)0;
static gl_PerVertex perVertexStruct = Constructgl_PerVertex(float4(0.0, 0.0, 0.0, 1.0), 1.0, Constructarray1_float_(0.0), Constructarray1_float_(0.0));
static gl_PerVertex perVertexStruct = Constructgl_PerVertex(float4(0.0, 0.0, 0.0, 1.0), 1.0, (float[1])0, (float[1])0);
static float2 a_pos_1 = (float2)0;
struct VertexOutput_main {
@@ -39,10 +32,10 @@ struct VertexOutput_main {
void main_1()
{
float2 _expr12 = a_uv_1;
v_uv = _expr12;
float2 _expr13 = a_pos_1;
perVertexStruct.gl_Position = float4(_expr13.x, _expr13.y, 0.0, 1.0);
float2 _expr8 = a_uv_1;
v_uv = _expr8;
float2 _expr9 = a_pos_1;
perVertexStruct.gl_Position = float4(_expr9.x, _expr9.y, 0.0, 1.0);
return;
}

4
tests/out/hlsl/quad.hlsl
View File

@@ -1,10 +1,10 @@
static const float c_scale = 1.2000000476837158;
struct VertexOutput {
float2 uv : LOC0;
float4 position : SV_Position;
};
static const float c_scale = 1.2;
Texture2D<float4> u_texture : register(t0);
SamplerState u_sampler : register(s1);

18
tests/out/hlsl/shadow.hlsl
View File

@@ -1,6 +1,3 @@
static const float3 c_ambient = float3(0.05000000074505806, 0.05000000074505806, 0.05000000074505806);
static const uint c_max_lights = 10;
struct Globals {
row_major float4x4 view_proj;
uint4 num_lights;
@@ -23,6 +20,9 @@ struct Light {
float4 color;
};
static const float3 c_ambient = float3(0.05, 0.05, 0.05);
static const uint c_max_lights = 10u;
cbuffer u_globals : register(b0) { Globals u_globals; }
cbuffer u_entity : register(b0, space1) { Entity u_entity; }
ByteAddressBuffer s_lights : register(t1);
@@ -92,13 +92,13 @@ float4 fs_main(FragmentInput_fs_main fragmentinput_fs_main) : SV_Target0
uint i = (uint)0;
float3 normal_1 = normalize(in_.world_normal);
color = float3(0.05000000074505806, 0.05000000074505806, 0.05000000074505806);
color = c_ambient;
i = 0u;
bool loop_init = true;
while(true) {
if (!loop_init) {
uint _expr39 = i;
i = (_expr39 + 1u);
uint _expr40 = i;
i = (_expr40 + 1u);
}
loop_init = false;
uint _expr7 = i;
@@ -130,13 +130,13 @@ float4 fs_main_without_storage(FragmentInput_fs_main_without_storage fragmentinp
uint i_1 = (uint)0;
float3 normal_2 = normalize(in_1.world_normal);
color_1 = float3(0.05000000074505806, 0.05000000074505806, 0.05000000074505806);
color_1 = c_ambient;
i_1 = 0u;
bool loop_init_1 = true;
while(true) {
if (!loop_init_1) {
uint _expr39 = i_1;
i_1 = (_expr39 + 1u);
uint _expr40 = i_1;
i_1 = (_expr40 + 1u);
}
loop_init_1 = false;
uint _expr7 = i_1;

1
tests/out/hlsl/standard.hlsl
View File

@@ -1,4 +1,3 @@
struct FragmentInput_derivatives {
float4 foo_1 : SV_Position;
};

1
tests/out/hlsl/texture-arg.hlsl
View File

@@ -1,4 +1,3 @@
Texture2D<float4> Texture : register(t0);
SamplerState Sampler : register(s1);

2
tests/out/hlsl/workgroup-uniform-load.hlsl
View File

@@ -1,4 +1,4 @@
static const uint SIZE = 128;
static const uint SIZE = 128u;
groupshared int arr_i32_[128];

1
tests/out/hlsl/workgroup-var-init.hlsl
View File

@@ -1,4 +1,3 @@
struct WStruct {
uint arr[512];
int atom;

163
tests/out/ir/access.ron
View File

@@ -335,87 +335,14 @@
ray_desc: None,
ray_intersection: None,
),
constants: [
(
name: None,
specialization: None,
inner: Scalar(
width: 4,
value: Uint(0),
),
),
(
name: None,
specialization: None,
inner: Composite(
ty: 2,
components: [
1,
1,
1,
],
),
),
(
name: None,
specialization: None,
inner: Scalar(
width: 4,
value: Sint(0),
),
),
(
name: None,
specialization: None,
inner: Composite(
ty: 4,
components: [
1,
2,
3,
],
),
),
(
name: None,
specialization: None,
inner: Scalar(
width: 4,
value: Sint(2),
),
),
(
name: None,
specialization: None,
inner: Scalar(
width: 4,
value: Sint(10),
),
),
(
name: None,
specialization: None,
inner: Scalar(
width: 4,
value: Sint(1),
),
),
(
name: None,
specialization: None,
inner: Scalar(
width: 4,
value: Sint(5),
),
),
],
constants: [],
global_variables: [
(
name: Some("global_const"),
space: Private,
binding: None,
ty: 4,
init: Some(4),
init: Some(7),
),
(
name: Some("bar"),
@@ -469,6 +396,44 @@
init: None,
),
],
const_expressions: [
Literal(U32(0)),
Literal(U32(0)),
Literal(U32(0)),
Literal(U32(0)),
Compose(
ty: 2,
components: [
2,
3,
4,
],
),
Literal(I32(0)),
Compose(
ty: 4,
components: [
1,
5,
6,
],
),
Literal(I32(2)),
Literal(I32(10)),
Literal(I32(2)),
Literal(I32(2)),
Literal(I32(2)),
Literal(I32(2)),
Literal(I32(10)),
Literal(I32(5)),
Literal(I32(5)),
Literal(I32(10)),
Literal(I32(5)),
Literal(I32(2)),
Literal(I32(2)),
Literal(I32(2)),
Literal(I32(2)),
],
functions: [
(
name: Some("test_matrix_within_struct_accesses"),
@@ -489,14 +454,14 @@
expressions: [
Literal(I32(1)),
LocalVariable(1),
Literal(I32(1)),
Load(
pointer: 2,
),
Constant(7),
Binary(
op: Subtract,
left: 3,
right: 4,
left: 4,
right: 3,
),
GlobalVariable(3),
AccessIndex(
@@ -644,14 +609,14 @@
],
),
LocalVariable(2),
Literal(I32(1)),
Load(
pointer: 2,
),
Constant(7),
Binary(
op: Add,
left: 56,
right: 57,
left: 57,
right: 56,
),
AccessIndex(
base: 55,
@@ -794,11 +759,7 @@
value: 1,
),
Emit((
start: 2,
end: 3,
)),
Emit((
start: 4,
start: 3,
end: 5,
)),
Store(
@@ -870,11 +831,7 @@
value: 54,
),
Emit((
start: 55,
end: 56,
)),
Emit((
start: 57,
start: 56,
end: 58,
)),
Store(
@@ -1001,14 +958,14 @@
expressions: [
Literal(I32(1)),
LocalVariable(1),
Literal(I32(1)),
Load(
pointer: 2,
),
Constant(7),
Binary(
op: Subtract,
left: 3,
right: 4,
left: 4,
right: 3,
),
GlobalVariable(5),
AccessIndex(
@@ -1177,14 +1134,14 @@
],
),
LocalVariable(2),
Literal(I32(1)),
Load(
pointer: 2,
),
Constant(7),
Binary(
op: Add,
left: 67,
right: 68,
left: 68,
right: 67,
),
AccessIndex(
base: 66,
@@ -1374,11 +1331,7 @@
value: 1,
),
Emit((
start: 2,
end: 3,
)),
Emit((
start: 4,
start: 3,
end: 5,
)),
Store(
@@ -1474,11 +1427,7 @@
value: 65,
),
Emit((
start: 66,
end: 67,
)),
Emit((
start: 68,
start: 67,
end: 69,
)),
Store(

1
tests/out/ir/collatz.ron
View File

@@ -57,6 +57,7 @@
init: None,
),
],
const_expressions: [],
functions: [
(
name: Some("collatz_iterations"),

875
tests/out/ir/shadow.ron
View File

File diff suppressed because it is too large Load Diff

18
tests/out/msl/access.msl
View File

@@ -57,8 +57,6 @@ struct type_22 {
struct type_26 {
metal::float4 inner[2];
};
constant metal::uint3 const_type_1_ = {0u, 0u, 0u};
constant GlobalConst const_GlobalConst = {0u, {}, const_type_1_, 0};
void test_matrix_within_struct_accesses(
constant Baz& baz
@@ -66,8 +64,8 @@ void test_matrix_within_struct_accesses(
int idx = {};
Baz t = {};
idx = 1;
int _e2 = idx;
idx = _e2 - 1;
int _e3 = idx;
idx = _e3 - 1;
metal::float3x2 l0_ = baz.m;
metal::float2 l1_ = baz.m[0];
int _e15 = idx;
@@ -81,8 +79,8 @@ void test_matrix_within_struct_accesses(
int _e43 = idx;
float l6_ = baz.m[_e41][_e43];
t = Baz {metal::float3x2(metal::float2(1.0), metal::float2(2.0), metal::float2(3.0))};
int _e55 = idx;
idx = _e55 + 1;
int _e56 = idx;
idx = _e56 + 1;
t.m = metal::float3x2(metal::float2(6.0), metal::float2(5.0), metal::float2(4.0));
t.m[0] = metal::float2(9.0);
int _e72 = idx;
@@ -104,8 +102,8 @@ void test_matrix_within_array_within_struct_accesses(
int idx_1 = {};
MatCx2InArray t_1 = {};
idx_1 = 1;
int _e2 = idx_1;
idx_1 = _e2 - 1;
int _e3 = idx_1;
idx_1 = _e3 - 1;
type_15 l0_1 = nested_mat_cx2_.am;
metal::float4x2 l1_1 = nested_mat_cx2_.am.inner[0];
metal::float2 l2_1 = nested_mat_cx2_.am.inner[0][0];
@@ -120,8 +118,8 @@ void test_matrix_within_array_within_struct_accesses(
int _e60 = idx_1;
float l7_ = nested_mat_cx2_.am.inner[0][_e58][_e60];
t_1 = MatCx2InArray {type_15 {}};
int _e66 = idx_1;
idx_1 = _e66 + 1;
int _e67 = idx_1;
idx_1 = _e67 + 1;
t_1.am = type_15 {};
t_1.am.inner[0] = metal::float4x2(metal::float2(8.0), metal::float2(7.0), metal::float2(6.0), metal::float2(5.0));
t_1.am.inner[0][0] = metal::float2(9.0);

226
tests/out/msl/binding-arrays.msl
View File

@@ -54,121 +54,121 @@ fragment main_Output main_(
u2_ = _e28 + metal::uint2(texture_array_unbounded[uniform_index].get_width(), texture_array_unbounded[uniform_index].get_height());
metal::uint2 _e33 = u2_;
u2_ = _e33 + metal::uint2(texture_array_unbounded[non_uniform_index].get_width(), texture_array_unbounded[non_uniform_index].get_height());
metal::float4 _e41 = texture_array_bounded[0].gather(samp[0], uv);
metal::float4 _e42 = v4_;
v4_ = _e42 + _e41;
metal::float4 _e48 = texture_array_bounded[uniform_index].gather(samp[uniform_index], uv);
metal::float4 _e49 = v4_;
v4_ = _e49 + _e48;
metal::float4 _e55 = texture_array_bounded[non_uniform_index].gather(samp[non_uniform_index], uv);
metal::float4 _e56 = v4_;
v4_ = _e56 + _e55;
metal::float4 _e65 = texture_array_depth[0].gather_compare(samp_comp[0], uv, 0.0);
metal::float4 _e66 = v4_;
v4_ = _e66 + _e65;
metal::float4 _e73 = texture_array_depth[uniform_index].gather_compare(samp_comp[uniform_index], uv, 0.0);
metal::float4 _e74 = v4_;
v4_ = _e74 + _e73;
metal::float4 _e81 = texture_array_depth[non_uniform_index].gather_compare(samp_comp[non_uniform_index], uv, 0.0);
metal::float4 _e82 = v4_;
v4_ = _e82 + _e81;
metal::float4 _e88 = (uint(0) < texture_array_unbounded[0].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[0].get_width(0), texture_array_unbounded[0].get_height(0))) ? texture_array_unbounded[0].read(metal::uint2(pix), 0): DefaultConstructible());
metal::float4 _e89 = v4_;
v4_ = _e89 + _e88;
metal::float4 _e94 = (uint(0) < texture_array_unbounded[uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[uniform_index].get_width(0), texture_array_unbounded[uniform_index].get_height(0))) ? texture_array_unbounded[uniform_index].read(metal::uint2(pix), 0): DefaultConstructible());
metal::float4 _e95 = v4_;
v4_ = _e95 + _e94;
metal::float4 _e100 = (uint(0) < texture_array_unbounded[non_uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[non_uniform_index].get_width(0), texture_array_unbounded[non_uniform_index].get_height(0))) ? texture_array_unbounded[non_uniform_index].read(metal::uint2(pix), 0): DefaultConstructible());
metal::float4 _e101 = v4_;
v4_ = _e101 + _e100;
uint _e107 = u1_;
u1_ = _e107 + texture_array_2darray[0].get_array_size();
uint _e112 = u1_;
u1_ = _e112 + texture_array_2darray[uniform_index].get_array_size();
uint _e117 = u1_;
u1_ = _e117 + texture_array_2darray[non_uniform_index].get_array_size();
uint _e123 = u1_;
u1_ = _e123 + texture_array_bounded[0].get_num_mip_levels();
uint _e128 = u1_;
u1_ = _e128 + texture_array_bounded[uniform_index].get_num_mip_levels();
uint _e133 = u1_;
u1_ = _e133 + texture_array_bounded[non_uniform_index].get_num_mip_levels();
uint _e139 = u1_;
u1_ = _e139 + texture_array_multisampled[0].get_num_samples();
uint _e144 = u1_;
u1_ = _e144 + texture_array_multisampled[uniform_index].get_num_samples();
uint _e149 = u1_;
u1_ = _e149 + texture_array_multisampled[non_uniform_index].get_num_samples();
metal::float4 _e157 = texture_array_bounded[0].sample(samp[0], uv);
metal::float4 _e158 = v4_;
v4_ = _e158 + _e157;
metal::float4 _e164 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv);
metal::float4 _e165 = v4_;
v4_ = _e165 + _e164;
metal::float4 _e171 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv);
metal::float4 _e172 = v4_;
v4_ = _e172 + _e171;
metal::float4 _e181 = texture_array_bounded[0].sample(samp[0], uv, metal::bias(0.0));
metal::float4 _e182 = v4_;
v4_ = _e182 + _e181;
metal::float4 _e189 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::bias(0.0));
metal::float4 _e190 = v4_;
v4_ = _e190 + _e189;
metal::float4 _e197 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::bias(0.0));
metal::float4 _e198 = v4_;
v4_ = _e198 + _e197;
float _e207 = texture_array_depth[0].sample_compare(samp_comp[0], uv, 0.0);
float _e208 = v1_;
v1_ = _e208 + _e207;
float _e215 = texture_array_depth[uniform_index].sample_compare(samp_comp[uniform_index], uv, 0.0);
float _e216 = v1_;
v1_ = _e216 + _e215;
float _e223 = texture_array_depth[non_uniform_index].sample_compare(samp_comp[non_uniform_index], uv, 0.0);
float _e224 = v1_;
v1_ = _e224 + _e223;
float _e233 = texture_array_depth[0].sample_compare(samp_comp[0], uv, 0.0);
float _e234 = v1_;
v1_ = _e234 + _e233;
float _e241 = texture_array_depth[uniform_index].sample_compare(samp_comp[uniform_index], uv, 0.0);
float _e242 = v1_;
v1_ = _e242 + _e241;
float _e249 = texture_array_depth[non_uniform_index].sample_compare(samp_comp[non_uniform_index], uv, 0.0);
float _e250 = v1_;
v1_ = _e250 + _e249;
metal::float4 _e258 = texture_array_bounded[0].sample(samp[0], uv, metal::gradient2d(uv, uv));
metal::float4 _e259 = v4_;
v4_ = _e259 + _e258;
metal::float4 _e265 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::gradient2d(uv, uv));
metal::float4 _e266 = v4_;
v4_ = _e266 + _e265;
metal::float4 _e272 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::gradient2d(uv, uv));
metal::float4 _e273 = v4_;
v4_ = _e273 + _e272;
metal::float4 _e282 = texture_array_bounded[0].sample(samp[0], uv, metal::level(0.0));
metal::float4 _e283 = v4_;
v4_ = _e283 + _e282;
metal::float4 _e290 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::level(0.0));
metal::float4 _e291 = v4_;
v4_ = _e291 + _e290;
metal::float4 _e298 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::level(0.0));
metal::float4 _e299 = v4_;
v4_ = _e299 + _e298;
metal::float4 _e304 = v4_;
if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[0].get_width(), texture_array_storage[0].get_height()))) {
texture_array_storage[0].write(_e304, metal::uint2(pix));
}
metal::float4 _e42 = texture_array_bounded[0].gather(samp[0], uv);
metal::float4 _e43 = v4_;
v4_ = _e43 + _e42;
metal::float4 _e50 = texture_array_bounded[uniform_index].gather(samp[uniform_index], uv);
metal::float4 _e51 = v4_;
v4_ = _e51 + _e50;
metal::float4 _e58 = texture_array_bounded[non_uniform_index].gather(samp[non_uniform_index], uv);
metal::float4 _e59 = v4_;
v4_ = _e59 + _e58;
metal::float4 _e68 = texture_array_depth[0].gather_compare(samp_comp[0], uv, 0.0);
metal::float4 _e69 = v4_;
v4_ = _e69 + _e68;
metal::float4 _e76 = texture_array_depth[uniform_index].gather_compare(samp_comp[uniform_index], uv, 0.0);
metal::float4 _e77 = v4_;
v4_ = _e77 + _e76;
metal::float4 _e84 = texture_array_depth[non_uniform_index].gather_compare(samp_comp[non_uniform_index], uv, 0.0);
metal::float4 _e85 = v4_;
v4_ = _e85 + _e84;
metal::float4 _e91 = (uint(0) < texture_array_unbounded[0].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[0].get_width(0), texture_array_unbounded[0].get_height(0))) ? texture_array_unbounded[0].read(metal::uint2(pix), 0): DefaultConstructible());
metal::float4 _e92 = v4_;
v4_ = _e92 + _e91;
metal::float4 _e97 = (uint(0) < texture_array_unbounded[uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[uniform_index].get_width(0), texture_array_unbounded[uniform_index].get_height(0))) ? texture_array_unbounded[uniform_index].read(metal::uint2(pix), 0): DefaultConstructible());
metal::float4 _e98 = v4_;
v4_ = _e98 + _e97;
metal::float4 _e103 = (uint(0) < texture_array_unbounded[non_uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[non_uniform_index].get_width(0), texture_array_unbounded[non_uniform_index].get_height(0))) ? texture_array_unbounded[non_uniform_index].read(metal::uint2(pix), 0): DefaultConstructible());
metal::float4 _e104 = v4_;
v4_ = _e104 + _e103;
uint _e110 = u1_;
u1_ = _e110 + texture_array_2darray[0].get_array_size();
uint _e115 = u1_;
u1_ = _e115 + texture_array_2darray[uniform_index].get_array_size();
uint _e120 = u1_;
u1_ = _e120 + texture_array_2darray[non_uniform_index].get_array_size();
uint _e126 = u1_;
u1_ = _e126 + texture_array_bounded[0].get_num_mip_levels();
uint _e131 = u1_;
u1_ = _e131 + texture_array_bounded[uniform_index].get_num_mip_levels();
uint _e136 = u1_;
u1_ = _e136 + texture_array_bounded[non_uniform_index].get_num_mip_levels();
uint _e142 = u1_;
u1_ = _e142 + texture_array_multisampled[0].get_num_samples();
uint _e147 = u1_;
u1_ = _e147 + texture_array_multisampled[uniform_index].get_num_samples();
uint _e152 = u1_;
u1_ = _e152 + texture_array_multisampled[non_uniform_index].get_num_samples();
metal::float4 _e160 = texture_array_bounded[0].sample(samp[0], uv);
metal::float4 _e161 = v4_;
v4_ = _e161 + _e160;
metal::float4 _e167 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv);
metal::float4 _e168 = v4_;
v4_ = _e168 + _e167;
metal::float4 _e174 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv);
metal::float4 _e175 = v4_;
v4_ = _e175 + _e174;
metal::float4 _e184 = texture_array_bounded[0].sample(samp[0], uv, metal::bias(0.0));
metal::float4 _e185 = v4_;
v4_ = _e185 + _e184;
metal::float4 _e192 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::bias(0.0));
metal::float4 _e193 = v4_;
v4_ = _e193 + _e192;
metal::float4 _e200 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::bias(0.0));
metal::float4 _e201 = v4_;
v4_ = _e201 + _e200;
float _e210 = texture_array_depth[0].sample_compare(samp_comp[0], uv, 0.0);
float _e211 = v1_;
v1_ = _e211 + _e210;
float _e218 = texture_array_depth[uniform_index].sample_compare(samp_comp[uniform_index], uv, 0.0);
float _e219 = v1_;
v1_ = _e219 + _e218;
float _e226 = texture_array_depth[non_uniform_index].sample_compare(samp_comp[non_uniform_index], uv, 0.0);
float _e227 = v1_;
v1_ = _e227 + _e226;
float _e236 = texture_array_depth[0].sample_compare(samp_comp[0], uv, 0.0);
float _e237 = v1_;
v1_ = _e237 + _e236;
float _e244 = texture_array_depth[uniform_index].sample_compare(samp_comp[uniform_index], uv, 0.0);
float _e245 = v1_;
v1_ = _e245 + _e244;
float _e252 = texture_array_depth[non_uniform_index].sample_compare(samp_comp[non_uniform_index], uv, 0.0);
float _e253 = v1_;
v1_ = _e253 + _e252;
metal::float4 _e261 = texture_array_bounded[0].sample(samp[0], uv, metal::gradient2d(uv, uv));
metal::float4 _e262 = v4_;
v4_ = _e262 + _e261;
metal::float4 _e268 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::gradient2d(uv, uv));
metal::float4 _e269 = v4_;
v4_ = _e269 + _e268;
metal::float4 _e275 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::gradient2d(uv, uv));
metal::float4 _e276 = v4_;
v4_ = _e276 + _e275;
metal::float4 _e285 = texture_array_bounded[0].sample(samp[0], uv, metal::level(0.0));
metal::float4 _e286 = v4_;
v4_ = _e286 + _e285;
metal::float4 _e293 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::level(0.0));
metal::float4 _e294 = v4_;
v4_ = _e294 + _e293;
metal::float4 _e301 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::level(0.0));
metal::float4 _e302 = v4_;
v4_ = _e302 + _e301;
metal::float4 _e307 = v4_;
if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[uniform_index].get_width(), texture_array_storage[uniform_index].get_height()))) {
texture_array_storage[uniform_index].write(_e307, metal::uint2(pix));
if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[0].get_width(), texture_array_storage[0].get_height()))) {
texture_array_storage[0].write(_e307, metal::uint2(pix));
}
metal::float4 _e310 = v4_;
if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[non_uniform_index].get_width(), texture_array_storage[non_uniform_index].get_height()))) {
texture_array_storage[non_uniform_index].write(_e310, metal::uint2(pix));
if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[uniform_index].get_width(), texture_array_storage[uniform_index].get_height()))) {
texture_array_storage[uniform_index].write(_e310, metal::uint2(pix));
}
metal::uint2 _e311 = u2_;
uint _e312 = u1_;
metal::float2 v2_ = static_cast<metal::float2>(_e311 + metal::uint2(_e312));
metal::float4 _e316 = v4_;
float _e323 = v1_;
return main_Output { (_e316 + metal::float4(v2_.x, v2_.y, v2_.x, v2_.y)) + metal::float4(_e323) };
metal::float4 _e313 = v4_;
if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[non_uniform_index].get_width(), texture_array_storage[non_uniform_index].get_height()))) {
texture_array_storage[non_uniform_index].write(_e313, metal::uint2(pix));
}
metal::uint2 _e314 = u2_;
uint _e315 = u1_;
metal::float2 v2_ = static_cast<metal::float2>(_e314 + metal::uint2(_e315));
metal::float4 _e319 = v4_;
float _e326 = v1_;
return main_Output { (_e319 + metal::float4(v2_.x, v2_.y, v2_.x, v2_.y)) + metal::float4(_e326) };
}

2
tests/out/msl/boids.msl
View File

@@ -9,7 +9,6 @@ struct _mslBufferSizes {
uint size2;
};
constexpr constant unsigned NUM_PARTICLES = 1500u;
struct Particle {
metal::float2 pos;
metal::float2 vel;
@@ -27,6 +26,7 @@ typedef Particle type_3[1];
struct Particles {
type_3 particles;
};
constant uint NUM_PARTICLES = 1500u;
struct main_Input {
};

2
tests/out/msl/do-while.msl
View File

@@ -11,7 +11,7 @@ void fb1_(
bool loop_init = true;
while(true) {
if (!loop_init) {
bool _e6 = cond;
bool _e2 = cond;
if (!(cond)) {
break;
}

4
tests/out/msl/empty-global-name.msl
View File

@@ -11,8 +11,8 @@ struct type_1 {
void function(
device type_1& unnamed
) {
int _e8 = unnamed.member;
unnamed.member = _e8 + 1;
int _e4 = unnamed.member;
unnamed.member = _e4 + 1;
return;
}

2
tests/out/msl/globals.msl
View File

@@ -8,7 +8,6 @@ struct _mslBufferSizes {
uint size3;
};
constexpr constant bool Foo_1 = true;
struct type_2 {
float inner[10];
};
@@ -32,6 +31,7 @@ struct type_14 {
struct type_15 {
type_14 inner[2];
};
constant bool Foo_1 = true;
void test_msl_packed_vec3_as_arg(
metal::float3 arg

23
tests/out/msl/image.msl
View File

@@ -4,7 +4,6 @@
using metal::uint;
constant metal::int2 const_type_9_ = {3, 1};
struct main_Input {
};
@@ -126,46 +125,46 @@ fragment texture_sampleOutput texture_sample(
metal::float4 _e14 = image_2d.sample(sampler_reg, tc);
metal::float4 _e15 = a;
a = _e15 + _e14;
metal::float4 _e19 = image_2d.sample(sampler_reg, tc, const_type_9_);
metal::float4 _e19 = image_2d.sample(sampler_reg, tc, metal::int2(3, 1));
metal::float4 _e20 = a;
a = _e20 + _e19;
metal::float4 _e24 = image_2d.sample(sampler_reg, tc, metal::level(2.3));
metal::float4 _e25 = a;
a = _e25 + _e24;
metal::float4 _e29 = image_2d.sample(sampler_reg, tc, metal::level(2.3), const_type_9_);
metal::float4 _e29 = image_2d.sample(sampler_reg, tc, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e30 = a;
a = _e30 + _e29;
metal::float4 _e35 = image_2d.sample(sampler_reg, tc, metal::bias(2.0), const_type_9_);
metal::float4 _e35 = image_2d.sample(sampler_reg, tc, metal::bias(2.0), metal::int2(3, 1));
metal::float4 _e36 = a;
a = _e36 + _e35;
metal::float4 _e41 = image_2d_array.sample(sampler_reg, tc, 0u);
metal::float4 _e42 = a;
a = _e42 + _e41;
metal::float4 _e47 = image_2d_array.sample(sampler_reg, tc, 0u, const_type_9_);
metal::float4 _e47 = image_2d_array.sample(sampler_reg, tc, 0u, metal::int2(3, 1));
metal::float4 _e48 = a;
a = _e48 + _e47;
metal::float4 _e53 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3));
metal::float4 _e54 = a;
a = _e54 + _e53;
metal::float4 _e59 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3), const_type_9_);
metal::float4 _e59 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e60 = a;
a = _e60 + _e59;
metal::float4 _e66 = image_2d_array.sample(sampler_reg, tc, 0u, metal::bias(2.0), const_type_9_);
metal::float4 _e66 = image_2d_array.sample(sampler_reg, tc, 0u, metal::bias(2.0), metal::int2(3, 1));
metal::float4 _e67 = a;
a = _e67 + _e66;
metal::float4 _e72 = image_2d_array.sample(sampler_reg, tc, 0);
metal::float4 _e73 = a;
a = _e73 + _e72;
metal::float4 _e78 = image_2d_array.sample(sampler_reg, tc, 0, const_type_9_);
metal::float4 _e78 = image_2d_array.sample(sampler_reg, tc, 0, metal::int2(3, 1));
metal::float4 _e79 = a;
a = _e79 + _e78;
metal::float4 _e84 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3));
metal::float4 _e85 = a;
a = _e85 + _e84;
metal::float4 _e90 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3), const_type_9_);
metal::float4 _e90 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e91 = a;
a = _e91 + _e90;
metal::float4 _e97 = image_2d_array.sample(sampler_reg, tc, 0, metal::bias(2.0), const_type_9_);
metal::float4 _e97 = image_2d_array.sample(sampler_reg, tc, 0, metal::bias(2.0), metal::int2(3, 1));
metal::float4 _e98 = a;
a = _e98 + _e97;
metal::float4 _e103 = image_cube_array.sample(sampler_reg, tc3_, 0u);
@@ -245,9 +244,9 @@ fragment gatherOutput gather(
) {
metal::float2 tc_2 = metal::float2(0.5);
metal::float4 s2d = image_2d.gather(sampler_reg, tc_2, metal::int2(0), metal::component::y);
metal::float4 s2d_offset = image_2d.gather(sampler_reg, tc_2, const_type_9_, metal::component::w);
metal::float4 s2d_offset = image_2d.gather(sampler_reg, tc_2, metal::int2(3, 1), metal::component::w);
metal::float4 s2d_depth = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5);
metal::float4 s2d_depth_offset = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5, const_type_9_);
metal::float4 s2d_depth_offset = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5, metal::int2(3, 1));
metal::uint4 u = image_2d_u32_.gather(sampler_reg, tc_2);
metal::int4 i = image_2d_i32_.gather(sampler_reg, tc_2);
metal::float4 f = static_cast<metal::float4>(u) + static_cast<metal::float4>(i);

24
tests/out/msl/operators.msl
View File

@@ -14,10 +14,10 @@ struct type_13 {
struct type_14 {
int inner[4];
};
constant metal::float4 v_f32_one = {1.0, 1.0, 1.0, 1.0};
constant metal::float4 v_f32_zero = {0.0, 0.0, 0.0, 0.0};
constant metal::float4 v_f32_half = {0.5, 0.5, 0.5, 0.5};
constant metal::int4 v_i32_one = {1, 1, 1, 1};
constant metal::float4 v_f32_one = metal::float4(1.0, 1.0, 1.0, 1.0);
constant metal::float4 v_f32_zero = metal::float4(0.0, 0.0, 0.0, 0.0);
constant metal::float4 v_f32_half = metal::float4(0.5, 0.5, 0.5, 0.5);
constant metal::int4 v_i32_one = metal::int4(1, 1, 1, 1);
metal::float4 builtins(
) {
@@ -263,15 +263,15 @@ void assignment(
a_1 = _e27 << 2u;
int _e30 = a_1;
a_1 = _e30 >> 1u;
int _e32 = a_1;
a_1 = _e32 + 1;
int _e35 = a_1;
a_1 = _e35 - 1;
int _e33 = a_1;
a_1 = _e33 + 1;
int _e36 = a_1;
a_1 = _e36 - 1;
vec0_ = metal::int3 {};
int _e42 = vec0_.y;
vec0_.y = _e42 + 1;
int _e47 = vec0_.y;
vec0_.y = _e47 - 1;
int _e43 = vec0_.y;
vec0_.y = _e43 + 1;
int _e48 = vec0_.y;
vec0_.y = _e48 - 1;
return;
}

Some files were not shown because too many files have changed in this diff Show More