nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-25 08:13:27 +00:00
Code Issues Packages Projects Releases Wiki Activity

[naga wgsl-in] Implement abstract types for consts, constructors.

Browse Source
This commit is contained in:
Jim Blandy 2023-11-22 13:50:47 -08:00 committed by Teodor Tanasoaia
parent 0b79599a8e
commit 601f235b34
22 changed files with 1595 additions and 793 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

40
CHANGELOG.md
View File

@ -101,8 +101,44 @@ Passing an owned value `window` to `Surface` will return a `Surface<'static>`. S
- Introduce a new `Scalar` struct type for use in Naga's IR, and update all frontend, middle, and backend code appropriately. By @jimblandy in [#4673](https://github.com/gfx-rs/wgpu/pull/4673).
- Add more metal keywords. By @fornwall in [#4707](https://github.com/gfx-rs/wgpu/pull/4707).
- Implement WGSL abstract types (by @jimblandy):
- Add a new `naga::Literal` variant, `I64`, for signed 64-bit literals. [#4711](https://github.com/gfx-rs/wgpu/pull/4711)
- Add partial support for WGSL abstract types (@jimblandy in [#4743](https://github.com/gfx-rs/wgpu/pull/4743)).
Abstract types make numeric literals easier to use, by
automatically converting literals and other constant expressions
from abstract numeric types to concrete types when safe and
necessary. For example, to build a vector of floating-point
numbers, Naga previously made you write:
vec3<f32>(1.0, 2.0, 3.0)
With this change, you can now simply write:
vec3<f32>(1, 2, 3)
Even though the literals are abstract integers, Naga recognizes
that it is safe and necessary to convert them to `f32` values in
order to build the vector. You can also use abstract values as
initializers for global constants, like this:
const unit_x: vec2<f32> = vec2(1, 0);
The literals `1` and `0` are abstract integers, and the expression
`vec2(1, 0)` is an abstract vector. However, Naga recognizes that
it can convert that to the concrete type `vec2<f32>` to satisfy
the given type of `unit_x`.
The WGSL specification permits abstract integers and
floating-point values in almost all contexts, but Naga's support
for this is still incomplete. Many WGSL operators and builtin
functions are specified to produce abstract results when applied
to abstract inputs, but for now Naga simply concretizes them all
before applying the operation. We will expand Naga's abstract type
support in subsequent pull requests.
As part of this work, the public types `naga::ScalarKind` and
`naga::Literal` now have new variants, `AbstractInt` and `AbstractFloat`.
- Add a new `naga::Literal` variant, `I64`, for signed 64-bit literals. [#4711](https://github.com/gfx-rs/wgpu/pull/4711)
- Emit and init `struct` member padding always. By @ErichDonGubler in [#4701](https://github.com/gfx-rs/wgpu/pull/4701).

2
naga/Cargo.toml
View File

@ -31,7 +31,7 @@ deserialize = ["serde", "bitflags/serde", "indexmap/serde"]
arbitrary = ["dep:arbitrary", "bitflags/arbitrary", "indexmap/arbitrary"]
spv-in = ["petgraph", "spirv"]
spv-out = ["spirv"]
wgsl-in = ["hexf-parse", "unicode-xid"]
wgsl-in = ["hexf-parse", "unicode-xid", "compact"]
wgsl-out = []
hlsl-out = []
compact = []

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

@ -251,6 +251,12 @@ pub enum Error<'a> {
ExpectedPositiveArrayLength(Span),
MissingWorkgroupSize(Span),
ConstantEvaluatorError(ConstantEvaluatorError, Span),
AutoConversion {
dest_span: Span,
dest_type: String,
source_span: Span,
source_type: String,
},
}
impl<'a> Error<'a> {
@ -712,6 +718,20 @@ impl<'a> Error<'a> {
)],
notes: vec![],
},
Error::AutoConversion { dest_span, ref dest_type, source_span, ref source_type } => ParseError {
message: format!("automatic conversions cannot convert `{source_type}` to `{dest_type}`"),
labels: vec![
(
dest_span,
format!("a value of type {dest_type} is required here").into(),
),
(
source_span,
format!("this expression has type {source_type}").into(),
)
],
notes: vec![],
}
}
}
}

213
naga/src/front/wgsl/lower/construction.rs
View File

@ -116,13 +116,15 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
components: &[Handle<ast::Expression<'source>>],
ctx: &mut ExpressionContext<'source, '_, '_>,
) -> Result<Handle<crate::Expression>, Error<'source>> {
use crate::proc::TypeResolution as Tr;
let constructor_h = self.constructor(constructor, ctx)?;
let components = match *components {
[] => Components::None,
[component] => {
let span = ctx.ast_expressions.get_span(component);
let component = self.expression(component, ctx)?;
let component = self.expression_for_abstract(component, ctx)?;
let ty_inner = super::resolve_inner!(ctx, component);
Components::One {
@ -134,13 +136,17 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
ref ast_components @ [_, _, ..] => {
let components = ast_components
.iter()
.map(|&expr| self.expression(expr, ctx))
.map(|&expr| self.expression_for_abstract(expr, ctx))
.collect::<Result<_, _>>()?;
let spans = ast_components
.iter()
.map(|&expr| ctx.ast_expressions.get_span(expr))
.collect();
for &component in &components {
ctx.grow_types(component)?;
}
Components::Many { components, spans }
}
};
@ -288,18 +294,13 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
// Vector constructor (splat)
(
Components::One {
component,
ty_inner: &crate::TypeInner::Scalar(src_scalar),
mut component,
ty_inner: &crate::TypeInner::Scalar(_),
..
},
Constructor::Type((
_,
&crate::TypeInner::Vector {
size,
scalar: dst_scalar,
},
)),
) if dst_scalar == src_scalar => {
Constructor::Type((_, &crate::TypeInner::Vector { size, scalar })),
) => {
ctx.convert_slice_to_common_scalar(std::slice::from_mut(&mut component), scalar)?;
expr = crate::Expression::Splat {
size,
value: component,
@ -307,37 +308,82 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
}
// Vector constructor (by elements), partial
(Components::Many { components, spans }, Constructor::PartialVector { size }) => {
let scalar =
component_scalar_from_constructor_args(&components, ctx).map_err(|index| {
(
Components::Many {
mut components,
spans,
},
Constructor::PartialVector { size },
) => {
let consensus_scalar =
automatic_conversion_consensus(&components, ctx).map_err(|index| {
Error::InvalidConstructorComponentType(spans[index], index as i32)
})?;
let inner = scalar.to_inner_vector(size);
ctx.convert_slice_to_common_scalar(&mut components, consensus_scalar)?;
let inner = consensus_scalar.to_inner_vector(size);
let ty = ctx.ensure_type_exists(inner);
expr = crate::Expression::Compose { ty, components };
}
// Vector constructor (by elements), full type given
(
Components::Many { components, .. },
Constructor::Type((ty, &crate::TypeInner::Vector { .. })),
Components::Many { mut components, .. },
Constructor::Type((ty, &crate::TypeInner::Vector { scalar, .. })),
) => {
ctx.try_automatic_conversions_for_vector(&mut components, scalar, ty_span)?;
expr = crate::Expression::Compose { ty, components };
}
// Matrix constructor (by elements)
// Matrix constructor (by elements), partial
(
Components::Many { components, spans },
Components::Many {
mut components,
spans,
},
Constructor::PartialMatrix { columns, rows },
)
| (
Components::Many { components, spans },
Constructor::Type((_, &crate::TypeInner::Matrix { columns, rows, .. })),
) if components.len() == columns as usize * rows as usize => {
let scalar =
component_scalar_from_constructor_args(&components, ctx).map_err(|index| {
let consensus_scalar =
automatic_conversion_consensus(&components, ctx).map_err(|index| {
Error::InvalidConstructorComponentType(spans[index], index as i32)
})?;
ctx.convert_slice_to_common_scalar(&mut components, consensus_scalar)?;
let vec_ty = ctx.ensure_type_exists(consensus_scalar.to_inner_vector(rows));
let components = components
.chunks(rows as usize)
.map(|vec_components| {
ctx.append_expression(
crate::Expression::Compose {
ty: vec_ty,
components: Vec::from(vec_components),
},
Default::default(),
)
})
.collect::<Result<Vec<_>, _>>()?;
let ty = ctx.ensure_type_exists(crate::TypeInner::Matrix {
columns,
rows,
scalar: consensus_scalar,
});
expr = crate::Expression::Compose { ty, components };
}
// Matrix constructor (by elements), type given
(
Components::Many { mut components, .. },
Constructor::Type((
_,
&crate::TypeInner::Matrix {
columns,
rows,
scalar,
},
)),
) if components.len() == columns as usize * rows as usize => {
let element = Tr::Value(crate::TypeInner::Scalar(scalar));
ctx.try_automatic_conversions_slice(&mut components, &element, ty_span)?;
let vec_ty = ctx.ensure_type_exists(scalar.to_inner_vector(rows));
let components = components
@ -363,28 +409,55 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
// Matrix constructor (by columns)
(
Components::Many { components, spans },
Components::Many {
mut components,
spans,
},
Constructor::PartialMatrix { columns, rows },
)
| (
Components::Many { components, spans },
Components::Many {
mut components,
spans,
},
Constructor::Type((_, &crate::TypeInner::Matrix { columns, rows, .. })),
) => {
let scalar =
component_scalar_from_constructor_args(&components, ctx).map_err(|index| {
let consensus_scalar =
automatic_conversion_consensus(&components, ctx).map_err(|index| {
Error::InvalidConstructorComponentType(spans[index], index as i32)
})?;
ctx.convert_slice_to_common_scalar(&mut components, consensus_scalar)?;
let ty = ctx.ensure_type_exists(crate::TypeInner::Matrix {
columns,
rows,
scalar,
scalar: consensus_scalar,
});
expr = crate::Expression::Compose { ty, components };
}
// Array constructor - infer type
(components, Constructor::PartialArray) => {
let components = components.into_components_vec();
let mut components = components.into_components_vec();
if let Ok(consensus_scalar) = automatic_conversion_consensus(&components, ctx) {
// Note that this will *not* necessarily convert all the
// components to the same type! The `automatic_conversion_consensus`
// function only considers the parameters' leaf scalar
// types; the parameters themselves could be any mix of
// vectors, matrices, and scalars.
//
// But *if* it is possible for this array construction
// expression to be well-typed at all, then all the
// parameters must have the same type constructors (vec,
// matrix, scalar) applied to their leaf scalars, so
// reconciling their scalars is always the right thing to
// do. And if this array construction is not well-typed,
// these conversions will not make it so, and we can let
// validation catch the error.
ctx.convert_slice_to_common_scalar(&mut components, consensus_scalar)?;
} else {
// There's no consensus scalar. Emit the `Compose`
// expression anyway, and let validation catch the problem.
}
let base = ctx.register_type(components[0])?;
@ -403,15 +476,30 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
expr = crate::Expression::Compose { ty, components };
}
// Array or Struct constructor
// Array constructor, explicit type
(components, Constructor::Type((ty, &crate::TypeInner::Array { base, .. }))) => {
let mut components = components.into_components_vec();
ctx.try_automatic_conversions_slice(&mut components, &Tr::Handle(base), span)?;
expr = crate::Expression::Compose { ty, components };
}
// Struct constructor
(
components,
Constructor::Type((
ty,
&crate::TypeInner::Array { .. } | &crate::TypeInner::Struct { .. },
)),
Constructor::Type((ty, &crate::TypeInner::Struct { ref members, .. })),
) => {
let components = components.into_components_vec();
let mut components = components.into_components_vec();
let struct_ty_span = ctx.module.types.get_span(ty);
// Make a vector of the members' type handles in advance, to
// avoid borrowing `members` from `ctx` while we generate
// new code.
let members: Vec<Handle<crate::Type>> = members.iter().map(|m| m.ty).collect();
for (component, &ty) in components.iter_mut().zip(&members) {
*component =
ctx.try_automatic_conversions(*component, &Tr::Handle(ty), struct_ty_span)?;
}
expr = crate::Expression::Compose { ty, components };
}
@ -504,12 +592,11 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
}
}
/// Compute a vector or matrix's scalar type from those of its
/// constructor arguments.
/// Find the consensus scalar of `components` under WGSL's automatic
/// conversions.
///
/// Given `components`, the arguments given to a vector or matrix
/// constructor, return the scalar type of the vector or matrix's
/// elements.
/// If `components` can all be converted to any common scalar via
/// WGSL's automatic conversions, return the best such scalar.
///
/// The `components` slice must not be empty. All elements' types must
/// have been resolved.
@ -518,20 +605,36 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
/// constructors, return `Err(i)`, where `i` is the index in
/// `components` of some problematic argument.
///
/// This function doesn't fully type-check the arguments, so it may
/// return `Ok` even when the Naga validator will reject the resulting
/// This function doesn't fully type-check the arguments - it only
/// considers their leaf scalar types. This means it may return `Ok`
/// even when the Naga validator will reject the resulting
/// construction expression later.
fn component_scalar_from_constructor_args(
fn automatic_conversion_consensus(
components: &[Handle<crate::Expression>],
ctx: &mut ExpressionContext<'_, '_, '_>,
ctx: &ExpressionContext<'_, '_, '_>,
) -> Result<Scalar, usize> {
// Since we don't yet implement abstract types, we can settle for
// just inspecting the first element.
let first = components[0];
ctx.grow_types(first).map_err(|_| 0_usize)?;
let inner = ctx.typifier()[first].inner_with(&ctx.module.types);
match inner.scalar() {
Some(scalar) => Ok(scalar),
None => Err(0),
let types = &ctx.module.types;
let mut inners = components
.iter()
.map(|&c| ctx.typifier()[c].inner_with(types));
log::debug!(
"wgsl automatic_conversion_consensus: {:?}",
inners
.clone()
.map(|inner| inner.to_wgsl(&ctx.module.to_ctx()))
.collect::<Vec<String>>()
);
let mut best = inners.next().unwrap().scalar().ok_or(0_usize)?;
for (inner, i) in inners.zip(1..) {
let scalar = inner.scalar().ok_or(i)?;
match best.automatic_conversion_combine(scalar) {
Some(new_best) => {
best = new_best;
}
None => return Err(i),
}
}
log::debug!(" consensus: {:?}", best.to_wgsl());
Ok(best)
}

375
naga/src/front/wgsl/lower/conversion.rs Normal file
View File

@ -0,0 +1,375 @@
//! WGSL's automatic conversions for abstract types.
use crate::{Handle, Span};
impl<'source, 'temp, 'out> super::ExpressionContext<'source, 'temp, 'out> {
/// Try to use WGSL's automatic conversions to convert `expr` to `goal_ty`.
///
/// If no conversions are necessary, return `expr` unchanged.
///
/// If automatic conversions cannot convert `expr` to `goal_ty`, return an
/// [`AutoConversion`] error.
///
/// Although the Load Rule is one of the automatic conversions, this
/// function assumes it has already been applied if appropriate, as
/// indicated by the fact that the Rust type of `expr` is not `Typed<_>`.
///
/// [`AutoConversion`]: super::Error::AutoConversion
pub fn try_automatic_conversions(
&mut self,
expr: Handle<crate::Expression>,
goal_ty: &crate::proc::TypeResolution,
goal_span: Span,
) -> Result<Handle<crate::Expression>, super::Error<'source>> {
let expr_span = self.get_expression_span(expr);
// Keep the TypeResolution so we can get type names for
// structs in error messages.
let expr_resolution = super::resolve!(self, expr);
let types = &self.module.types;
let expr_inner = expr_resolution.inner_with(types);
let goal_inner = goal_ty.inner_with(types);
// If `expr` already has the requested type, we're done.
if expr_inner.equivalent(goal_inner, types) {
return Ok(expr);
}
let (_expr_scalar, goal_scalar) =
match expr_inner.automatically_converts_to(goal_inner, types) {
Some(scalars) => scalars,
None => {
let gctx = &self.module.to_ctx();
let source_type = expr_resolution.to_wgsl(gctx);
let dest_type = goal_ty.to_wgsl(gctx);
return Err(super::Error::AutoConversion {
dest_span: goal_span,
dest_type,
source_span: expr_span,
source_type,
});
}
};
let converted = if let crate::TypeInner::Array { .. } = *goal_inner {
let span = self.get_expression_span(expr);
self.as_const_evaluator()
.cast_array(expr, goal_scalar, span)
.map_err(|err| super::Error::ConstantEvaluatorError(err, span))?
} else {
let cast = crate::Expression::As {
expr,
kind: goal_scalar.kind,
convert: Some(goal_scalar.width),
};
self.append_expression(cast, expr_span)?
};
Ok(converted)
}
/// Try to convert `exprs` to `goal_ty` using WGSL's automatic conversions.
pub fn try_automatic_conversions_slice(
&mut self,
exprs: &mut [Handle<crate::Expression>],
goal_ty: &crate::proc::TypeResolution,
goal_span: Span,
) -> Result<(), super::Error<'source>> {
for expr in exprs.iter_mut() {
*expr = self.try_automatic_conversions(*expr, goal_ty, goal_span)?;
}
Ok(())
}
/// Apply WGSL's automatic conversions to a vector constructor's arguments.
///
/// When calling a vector constructor like `vec3<f32>(...)`, the parameters
/// can be a mix of scalars and vectors, with the latter being spread out to
/// contribute each of their components as a component of the new value.
/// When the element type is explicit, as with `<f32>` in the example above,
/// WGSL's automatic conversions should convert abstract scalar and vector
/// parameters to the constructor's required scalar type.
pub fn try_automatic_conversions_for_vector(
&mut self,
exprs: &mut [Handle<crate::Expression>],
goal_scalar: crate::Scalar,
goal_span: Span,
) -> Result<(), super::Error<'source>> {
use crate::proc::TypeResolution as Tr;
use crate::TypeInner as Ti;
let goal_scalar_res = Tr::Value(Ti::Scalar(goal_scalar));
for (i, expr) in exprs.iter_mut().enumerate() {
// Keep the TypeResolution so we can get full type names
// in error messages.
let expr_resolution = super::resolve!(self, *expr);
let types = &self.module.types;
let expr_inner = expr_resolution.inner_with(types);
match *expr_inner {
Ti::Scalar(_) => {
*expr = self.try_automatic_conversions(*expr, &goal_scalar_res, goal_span)?;
}
Ti::Vector { size, scalar: _ } => {
let goal_vector_res = Tr::Value(Ti::Vector {
size,
scalar: goal_scalar,
});
*expr = self.try_automatic_conversions(*expr, &goal_vector_res, goal_span)?;
}
_ => {
let span = self.get_expression_span(*expr);
return Err(super::Error::InvalidConstructorComponentType(
span, i as i32,
));
}
}
}
Ok(())
}
/// Convert all expressions in `exprs` to a common scalar type.
///
/// Note that the caller is responsible for making sure these
/// conversions are actually justified. This function simply
/// generates `As` expressions, regardless of whether they are
/// permitted WGSL automatic conversions. Callers intending to
/// implement automatic conversions need to determine for
/// themselves whether the casts we we generate are justified,
/// perhaps by calling `TypeInner::automatically_converts_to` or
/// `Scalar::automatic_conversion_combine`.
pub fn convert_slice_to_common_scalar(
&mut self,
exprs: &mut [Handle<crate::Expression>],
goal: crate::Scalar,
) -> Result<(), super::Error<'source>> {
for expr in exprs.iter_mut() {
let inner = super::resolve_inner!(self, *expr);
// Do nothing if `inner` doesn't even have leaf scalars;
// it's a type error that validation will catch.
if inner.scalar() != Some(goal) {
let cast = crate::Expression::As {
expr: *expr,
kind: goal.kind,
convert: Some(goal.width),
};
let expr_span = self.get_expression_span(*expr);
*expr = self.append_expression(cast, expr_span)?;
}
}
Ok(())
}
/// Return an expression for the concretized value of `expr`.
///
/// If `expr` is already concrete, return it unchanged.
pub fn concretize(
&mut self,
mut expr: Handle<crate::Expression>,
) -> Result<Handle<crate::Expression>, super::Error<'source>> {
let inner = super::resolve_inner!(self, expr);
if let Some(scalar) = inner.automatically_convertible_scalar(&self.module.types) {
let concretized = scalar.concretize();
if concretized != scalar {
let span = self.get_expression_span(expr);
expr = self
.as_const_evaluator()
.cast_array(expr, concretized, span)
.map_err(|err| super::Error::ConstantEvaluatorError(err, span))?;
}
}
Ok(expr)
}
}
impl crate::TypeInner {
/// Determine whether `self` automatically converts to `goal`.
///
/// If WGSL's automatic conversions (excluding the Load Rule) will
/// convert `self` to `goal`, then return a pair `(from, to)`,
/// where `from` and `to` are the scalar types of the leaf values
/// of `self` and `goal`.
///
/// This function assumes that `self` and `goal` are different
/// types. Callers should first check whether any conversion is
/// needed at all.
///
/// If the automatic conversions cannot convert `self` to `goal`,
/// return `None`.
fn automatically_converts_to(
&self,
goal: &Self,
types: &crate::UniqueArena<crate::Type>,
) -> Option<(crate::Scalar, crate::Scalar)> {
use crate::ScalarKind as Sk;
use crate::TypeInner as Ti;
// Automatic conversions only change the scalar type of a value's leaves
// (e.g., `vec4<AbstractFloat>` to `vec4<f32>`), never the type
// constructors applied to those scalar types (e.g., never scalar to
// `vec4`, or `vec2` to `vec3`). So first we check that the type
// constructors match, extracting the leaf scalar types in the process.
let expr_scalar;
let goal_scalar;
match (self, goal) {
(&Ti::Scalar(expr), &Ti::Scalar(goal)) => {
expr_scalar = expr;
goal_scalar = goal;
}
(
&Ti::Vector {
size: expr_size,
scalar: expr,
},
&Ti::Vector {
size: goal_size,
scalar: goal,
},
) if expr_size == goal_size => {
expr_scalar = expr;
goal_scalar = goal;
}
(
&Ti::Matrix {
rows: expr_rows,
columns: expr_columns,
scalar: expr,
},
&Ti::Matrix {
rows: goal_rows,
columns: goal_columns,
scalar: goal,
},
) if expr_rows == goal_rows && expr_columns == goal_columns => {
expr_scalar = expr;
goal_scalar = goal;
}
(
&Ti::Array {
base: expr_base,
size: expr_size,
stride: _,
},
&Ti::Array {
base: goal_base,
size: goal_size,
stride: _,
},
) if expr_size == goal_size => {
return types[expr_base]
.inner
.automatically_converts_to(&types[goal_base].inner, types);
}
_ => return None,
}
match (expr_scalar.kind, goal_scalar.kind) {
(Sk::AbstractFloat, Sk::Float) => {}
(Sk::AbstractInt, Sk::Sint | Sk::Uint | Sk::AbstractFloat | Sk::Float) => {}
_ => return None,
}
log::trace!(" okay: expr {expr_scalar:?}, goal {goal_scalar:?}");
Some((expr_scalar, goal_scalar))
}
fn automatically_convertible_scalar(
&self,
types: &crate::UniqueArena<crate::Type>,
) -> Option<crate::Scalar> {
use crate::TypeInner as Ti;
match *self {
Ti::Scalar(scalar) | Ti::Vector { scalar, .. } | Ti::Matrix { scalar, .. } => {
Some(scalar)
}
Ti::Array { base, .. } => types[base].inner.automatically_convertible_scalar(types),
Ti::Atomic(_)
| Ti::Pointer { .. }
| Ti::ValuePointer { .. }
| Ti::Struct { .. }
| Ti::Image { .. }
| Ti::Sampler { .. }
| Ti::AccelerationStructure
| Ti::RayQuery
| Ti::BindingArray { .. } => None,
}
}
}
impl crate::Scalar {
/// Find the common type of `self` and `other` under WGSL's
/// automatic conversions.
///
/// If there are any scalars to which WGSL's automatic conversions
/// will convert both `self` and `other`, return the best such
/// scalar. Otherwise, return `None`.
pub const fn automatic_conversion_combine(self, other: Self) -> Option<crate::Scalar> {
use crate::ScalarKind as Sk;
match (self.kind, other.kind) {
// When the kinds match...
(Sk::AbstractFloat, Sk::AbstractFloat)
| (Sk::AbstractInt, Sk::AbstractInt)
| (Sk::Sint, Sk::Sint)
| (Sk::Uint, Sk::Uint)
| (Sk::Float, Sk::Float)
| (Sk::Bool, Sk::Bool) => {
if self.width == other.width {
// ... either no conversion is necessary ...
Some(self)
} else {
// ... or no conversion is possible.
// We never convert concrete to concrete, and
// abstract types should have only one size.
None
}
}
// AbstractInt converts to AbstractFloat.
(Sk::AbstractFloat | Sk::AbstractInt, Sk::AbstractFloat | Sk::AbstractInt) => {
Some(Self {
kind: Sk::AbstractFloat,
width: crate::ABSTRACT_WIDTH,
})
}
// AbstractFloat converts to Float.
(Sk::AbstractFloat, Sk::Float) => Some(Self::float(other.width)),
(Sk::Float, Sk::AbstractFloat) => Some(Self::float(self.width)),
// AbstractInt converts to concrete integer or float.
(Sk::AbstractInt, kind @ (Sk::Uint | Sk::Sint | Sk::Float)) => Some(Self {
kind,
width: other.width,
}),
(kind @ (Sk::Uint | Sk::Sint | Sk::Float), Sk::AbstractInt) => Some(Self {
kind,
width: self.width,
}),
// AbstractFloat can't be reconciled with concrete integer types.
(Sk::AbstractFloat, Sk::Uint | Sk::Sint) | (Sk::Uint | Sk::Sint, Sk::AbstractFloat) => {
None
}
// Nothing can be reconciled with `bool`.
(Sk::Bool, _) | (_, Sk::Bool) => None,
// Different concrete types cannot be reconciled.
(Sk::Sint | Sk::Uint | Sk::Float, Sk::Sint | Sk::Uint | Sk::Float) => None,
}
}
const fn concretize(self) -> Self {
use crate::ScalarKind as Sk;
match self.kind {
Sk::Sint | Sk::Uint | Sk::Float | Sk::Bool => self,
Sk::AbstractInt => Self::I32,
Sk::AbstractFloat => Self::F32,
}
}
}

69
naga/src/front/wgsl/lower/mod.rs
View File

@ -11,6 +11,7 @@ use crate::proc::{
use crate::{Arena, FastHashMap, FastIndexMap, Handle, Span};
mod construction;
mod conversion;
/// Resolves the inner type of a given expression.
///
@ -66,6 +67,7 @@ macro_rules! resolve {
&$ctx.typifier()[$expr]
}};
}
pub(super) use resolve;
/// State for constructing a `crate::Module`.
pub struct GlobalContext<'source, 'temp, 'out> {
@ -903,29 +905,39 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
}
ast::GlobalDeclKind::Const(ref c) => {
let mut ectx = ctx.as_const();
let init = self.expression(c.init, &mut ectx)?;
let inferred_type = ectx.register_type(init)?;
let mut init = self.expression_for_abstract(c.init, &mut ectx)?;
let explicit_ty =
c.ty.map(|ty| self.resolve_ast_type(ty, &mut ctx))
.transpose()?;
if let Some(explicit) = explicit_ty {
if explicit != inferred_type {
let gctx = ctx.module.to_ctx();
return Err(Error::InitializationTypeMismatch {
name: c.name.span,
expected: explicit.to_wgsl(&gctx),
got: inferred_type.to_wgsl(&gctx),
});
}
let ty;
if let Some(explicit_ty) = c.ty {
let explicit_ty =
self.resolve_ast_type(explicit_ty, &mut ectx.as_global())?;
let explicit_ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
init = ectx
.try_automatic_conversions(init, &explicit_ty_res, c.name.span)
.map_err(|error| match error {
Error::AutoConversion {
dest_span: _,
dest_type,
source_span: _,
source_type,
} => Error::InitializationTypeMismatch {
name: c.name.span,
expected: dest_type,
got: source_type,
},
other => other,
})?;
ty = explicit_ty;
} else {
init = ectx.concretize(init)?;
ty = ectx.register_type(init)?;
}
let handle = ctx.module.constants.append(
crate::Constant {
name: Some(c.name.name.to_string()),
r#override: crate::Override::None,
ty: inferred_type,
ty,
init,
},
span,
@ -951,6 +963,11 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
}
}
// Constant evaluation may leave abstract-typed literals and
// compositions in expression arenas, so we need to compact the module
// to remove unused expressions and types.
crate::compact::compact(&mut module);
Ok(module)
}
@ -1449,10 +1466,25 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
}
/// Lower `expr` and apply the Load Rule if possible.
///
/// For the time being, this concretizes abstract values, to support
/// consumers that haven't been adapted to consume them yet. Consumers
/// prepared for abstract values can call [`expression_for_abstract`].
///
/// [`expression_for_abstract`]: Lowerer::expression_for_abstract
fn expression(
&mut self,
expr: Handle<ast::Expression<'source>>,
ctx: &mut ExpressionContext<'source, '_, '_>,
) -> Result<Handle<crate::Expression>, Error<'source>> {
let expr = self.expression_for_abstract(expr, ctx)?;
ctx.concretize(expr)
}
fn expression_for_abstract(
&mut self,
expr: Handle<ast::Expression<'source>>,
ctx: &mut ExpressionContext<'source, '_, '_>,
) -> Result<Handle<crate::Expression>, Error<'source>> {
let expr = self.expression_for_reference(expr, ctx)?;
ctx.apply_load_rule(expr)
@ -1473,8 +1505,9 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
ast::Literal::Number(Number::I32(i)) => crate::Literal::I32(i),
ast::Literal::Number(Number::U32(u)) => crate::Literal::U32(u),
ast::Literal::Number(Number::F64(f)) => crate::Literal::F64(f),
ast::Literal::Number(_) => {
unreachable!("got abstract numeric type when not expected");
ast::Literal::Number(Number::AbstractInt(i)) => crate::Literal::AbstractInt(i),
ast::Literal::Number(Number::AbstractFloat(f)) => {
crate::Literal::AbstractFloat(f)
}
ast::Literal::Bool(b) => crate::Literal::Bool(b),
};

45
naga/src/front/wgsl/parse/lexer.rs
View File

@ -465,13 +465,13 @@ fn test_numbers() {
sub_test(
"0x123 0X123u 1u 123 0 0i 0x3f",
&[
Token::Number(Ok(Number::I32(291))),
Token::Number(Ok(Number::AbstractInt(291))),
Token::Number(Ok(Number::U32(291))),
Token::Number(Ok(Number::U32(1))),
Token::Number(Ok(Number::I32(123))),
Token::Number(Ok(Number::AbstractInt(123))),
Token::Number(Ok(Number::AbstractInt(0))),
Token::Number(Ok(Number::I32(0))),
Token::Number(Ok(Number::I32(0))),
Token::Number(Ok(Number::I32(63))),
Token::Number(Ok(Number::AbstractInt(63))),
],
);
// decimal floating point
@ -479,17 +479,17 @@ fn test_numbers() {
"0.e+4f 01. .01 12.34 .0f 0h 1e-3 0xa.fp+2 0x1P+4f 0X.3 0x3p+2h 0X1.fp-4 0x3.2p+2h",
&[
Token::Number(Ok(Number::F32(0.))),
Token::Number(Ok(Number::F32(1.))),
Token::Number(Ok(Number::F32(0.01))),
Token::Number(Ok(Number::F32(12.34))),
Token::Number(Ok(Number::AbstractFloat(1.))),
Token::Number(Ok(Number::AbstractFloat(0.01))),
Token::Number(Ok(Number::AbstractFloat(12.34))),
Token::Number(Ok(Number::F32(0.))),
Token::Number(Err(NumberError::UnimplementedF16)),
Token::Number(Ok(Number::F32(0.001))),
Token::Number(Ok(Number::F32(43.75))),
Token::Number(Ok(Number::AbstractFloat(0.001))),
Token::Number(Ok(Number::AbstractFloat(43.75))),
Token::Number(Ok(Number::F32(16.))),
Token::Number(Ok(Number::F32(0.1875))),
Token::Number(Ok(Number::AbstractFloat(0.1875))),
Token::Number(Err(NumberError::UnimplementedF16)),
Token::Number(Ok(Number::F32(0.12109375))),
Token::Number(Ok(Number::AbstractFloat(0.12109375))),
Token::Number(Err(NumberError::UnimplementedF16)),
],
);
@ -635,7 +635,7 @@ fn double_floats() {
Token::Number(Ok(Number::F64(0.0625))),
Token::Number(Ok(Number::F64(0.0625))),
Token::Number(Ok(Number::F64(10.0))),
Token::Number(Ok(Number::I32(10))),
Token::Number(Ok(Number::AbstractInt(10))),
Token::Word("l"),
],
)
@ -646,13 +646,16 @@ fn test_tokens() {
sub_test("id123_OK", &[Token::Word("id123_OK")]);
sub_test(
"92No",
&[Token::Number(Ok(Number::I32(92))), Token::Word("No")],
&[
Token::Number(Ok(Number::AbstractInt(92))),
Token::Word("No"),
],
);
sub_test(
"2u3o",
&[
Token::Number(Ok(Number::U32(2))),
Token::Number(Ok(Number::I32(3))),
Token::Number(Ok(Number::AbstractInt(3))),
Token::Word("o"),
],
);
@ -660,7 +663,7 @@ fn test_tokens() {
"2.4f44po",
&[
Token::Number(Ok(Number::F32(2.4))),
Token::Number(Ok(Number::I32(44))),
Token::Number(Ok(Number::AbstractInt(44))),
Token::Word("po"),
],
);
@ -699,13 +702,13 @@ fn test_tokens() {
&[
// The 'f' suffixes are taken as a hex digit:
// the fractional part is 0x2f / 256.
Token::Number(Ok(Number::F32(1.0 + 0x2f as f32 / 256.0))),
Token::Number(Ok(Number::F32(1.0 + 0x2f as f32 / 256.0))),
Token::Number(Ok(Number::F32(1.125))),
Token::Number(Ok(Number::AbstractFloat(1.0 + 0x2f as f64 / 256.0))),
Token::Number(Ok(Number::AbstractFloat(1.0 + 0x2f as f64 / 256.0))),
Token::Number(Ok(Number::AbstractFloat(1.125))),
Token::Word("h"),
Token::Number(Ok(Number::F32(1.125))),
Token::Number(Ok(Number::AbstractFloat(1.125))),
Token::Word("H"),
Token::Number(Ok(Number::F32(1.125))),
Token::Number(Ok(Number::AbstractFloat(1.125))),
Token::Word("lf"),
],
)
@ -719,7 +722,7 @@ fn test_variable_decl() {
Token::Attribute,
Token::Word("group"),
Token::Paren('('),
Token::Number(Ok(Number::I32(0))),
Token::Number(Ok(Number::AbstractInt(0))),
Token::Paren(')'),
Token::Word("var"),
Token::Paren('<'),

28
naga/src/front/wgsl/parse/number.rs
View File

@ -20,37 +20,11 @@ pub enum Number {
F64(f64),
}
impl Number {
/// Convert abstract numbers to a plausible concrete counterpart.
///
/// Return concrete numbers unchanged. If the conversion would be
/// lossy, return an error.
fn abstract_to_concrete(self) -> Result<Number, NumberError> {
match self {
Number::AbstractInt(num) => i32::try_from(num)
.map(Number::I32)
.map_err(|_| NumberError::NotRepresentable),
Number::AbstractFloat(num) => {
let num = num as f32;
if num.is_finite() {
Ok(Number::F32(num))
} else {
Err(NumberError::NotRepresentable)
}
}
num => Ok(num),
}
}
}
// TODO: when implementing Creation-Time Expressions, remove the ability to match the minus sign
pub(in crate::front::wgsl) fn consume_number(input: &str) -> (Token<'_>, &str) {
let (result, rest) = parse(input);
(
Token::Number(result.and_then(Number::abstract_to_concrete)),
rest,
)
(Token::Number(result), rest)
}
enum Kind {

4
naga/src/front/wgsl/tests.rs
View File

@ -76,7 +76,7 @@ fn parse_type_cast() {
assert!(parse_str(
"
fn main() {
let x: vec2<f32> = vec2<f32>(0);
let x: vec2<f32> = vec2<f32>(0i, 0i);
}
",
)
@ -313,7 +313,7 @@ fn parse_texture_load() {
"
var t: texture_3d<u32>;
fn foo() {
let r: vec4<u32> = textureLoad(t, vec3<u32>(0.0, 1.0, 2.0), 1);
let r: vec4<u32> = textureLoad(t, vec3<u32>(0u, 1u, 2u), 1);
}
",
)

109
naga/src/proc/constant_evaluator.rs
View File

@ -277,6 +277,17 @@ impl<'a> ConstantEvaluator<'a> {
}
}
pub fn to_ctx(&self) -> crate::proc::GlobalCtx {
crate::proc::GlobalCtx {
types: self.types,
constants: self.constants,
const_expressions: match self.function_local_data {
Some(ref data) => data.const_expressions,
None => self.expressions,
},
}
}
fn check(&self, expr: Handle<Expression>) -> Result<(), ConstantEvaluatorError> {
if let Some(ref function_local_data) = self.function_local_data {
if !function_local_data.expression_constness.is_const(expr) {
@ -1035,7 +1046,21 @@ impl<'a> ConstantEvaluator<'a> {
return Err(ConstantEvaluatorError::InvalidCastArg)
}
}),
_ => return Err(ConstantEvaluatorError::InvalidCastArg),
Sc::ABSTRACT_FLOAT => Literal::AbstractFloat(match literal {
Literal::AbstractInt(v) => {
// Overflow is forbidden, but inexact conversions
// are fine. The range of f64 is far larger than
// that of i64, so we don't have to check anything
// here.
v as f64
}
Literal::AbstractFloat(v) => v,
_ => return Err(ConstantEvaluatorError::InvalidCastArg),
}),
_ => {
log::debug!("Constant evaluator refused to convert value to {target:?}");
return Err(ConstantEvaluatorError::InvalidCastArg);
}
};
Expression::Literal(literal)
}
@ -1085,6 +1110,66 @@ impl<'a> ConstantEvaluator<'a> {
self.register_evaluated_expr(expr, span)
}
/// Convert the scalar leaves of `expr` to `target`, handling arrays.
///
/// `expr` must be a `Compose` expression whose type is a scalar, vector,
/// matrix, or nested arrays of such.
///
/// This is basically the same as the [`cast`] method, except that that
/// should only handle Naga [`As`] expressions, which cannot convert arrays.
///
/// Treat `span` as the location of the resulting expression.
///
/// [`cast`]: ConstantEvaluator::cast
/// [`As`]: crate::Expression::As
pub fn cast_array(
&mut self,
expr: Handle<Expression>,
target: crate::Scalar,
span: Span,
) -> Result<Handle<Expression>, ConstantEvaluatorError> {
let Expression::Compose { ty, ref components } = self.expressions[expr] else {
return self.cast(expr, target, span);
};
let crate::TypeInner::Array { base: _, size, stride: _ } = self.types[ty].inner else {
return self.cast(expr, target, span);
};
let mut components = components.clone();
for component in &mut components {
*component = self.cast_array(*component, target, span)?;
}
let first = components
.first()
.ok_or(ConstantEvaluatorError::InvalidCastArg)?;
let new_base = match self.resolve_type(*first)? {
crate::proc::TypeResolution::Handle(ty) => ty,
crate::proc::TypeResolution::Value(inner) => {
self.types.insert(Type { name: None, inner }, span)
}
};
let new_base_stride = self.types[new_base].inner.size(self.to_ctx());
let new_array_ty = self.types.insert(
Type {
name: None,
inner: TypeInner::Array {
base: new_base,
size,
stride: new_base_stride,
},
},
span,
);
let compose = Expression::Compose {
ty: new_array_ty,
components,
};
self.register_evaluated_expr(compose, span)
}
fn unary_op(
&mut self,
op: UnaryOperator,
@ -1339,6 +1424,28 @@ impl<'a> ConstantEvaluator<'a> {
Ok(self.expressions.append(expr, span))
}
}
fn resolve_type(
&self,
expr: Handle<Expression>,
) -> Result<crate::proc::TypeResolution, ConstantEvaluatorError> {
use crate::proc::TypeResolution as Tr;
use crate::Expression as Ex;
let resolution = match self.expressions[expr] {
Ex::Literal(ref literal) => Tr::Value(literal.ty_inner()),
Ex::Constant(c) => Tr::Handle(self.constants[c].ty),
Ex::ZeroValue(ty) | Ex::Compose { ty, .. } => Tr::Handle(ty),
Ex::Splat { size, value } => {
let Tr::Value(TypeInner::Scalar(scalar)) = self.resolve_type(value)? else {
return Err(ConstantEvaluatorError::SplatScalarOnly);
};
Tr::Value(TypeInner::Vector { scalar, size })
}
_ => return Err(ConstantEvaluatorError::SubexpressionsAreNotConstant),
};
Ok(resolution)
}
}
#[cfg(test)]

50
naga/src/valid/expression.rs
View File

@ -670,37 +670,23 @@ impl super::Validator {
let good = match op {
Bo::Add | Bo::Subtract => match *left_inner {
Ti::Scalar(scalar) | Ti::Vector { scalar, .. } => match scalar.kind {
Sk::Uint
| Sk::Sint
| Sk::Float
| Sk::AbstractInt
| Sk::AbstractFloat => left_inner == right_inner,
Sk::Bool => false,
Sk::Uint | Sk::Sint | Sk::Float => left_inner == right_inner,
Sk::Bool | Sk::AbstractInt | Sk::AbstractFloat => false,
},
Ti::Matrix { .. } => left_inner == right_inner,
_ => false,
},
Bo::Divide | Bo::Modulo => match *left_inner {
Ti::Scalar(scalar) | Ti::Vector { scalar, .. } => match scalar.kind {
Sk::Uint
| Sk::Sint
| Sk::Float
| Sk::AbstractInt
| Sk::AbstractFloat => left_inner == right_inner,
Sk::Bool => false,
Sk::Uint | Sk::Sint | Sk::Float => left_inner == right_inner,
Sk::Bool | Sk::AbstractInt | Sk::AbstractFloat => false,
},
_ => false,
},
Bo::Multiply => {
let kind_allowed = match left_inner.scalar_kind() {
Some(
Sk::Uint
| Sk::Sint
| Sk::Float
| Sk::AbstractInt
| Sk::AbstractFloat,
) => true,
Some(Sk::Bool) | None => false,
Some(Sk::Uint | Sk::Sint | Sk::Float) => true,
Some(Sk::Bool | Sk::AbstractInt | Sk::AbstractFloat) | None => false,
};
let types_match = match (left_inner, right_inner) {
// Straight scalar and mixed scalar/vector.
@ -776,12 +762,8 @@ impl super::Validator {
Bo::Less | Bo::LessEqual | Bo::Greater | Bo::GreaterEqual => {
match *left_inner {
Ti::Scalar(scalar) | Ti::Vector { scalar, .. } => match scalar.kind {
Sk::Uint
| Sk::Sint
| Sk::Float
| Sk::AbstractInt
| Sk::AbstractFloat => left_inner == right_inner,
Sk::Bool => false,
Sk::Uint | Sk::Sint | Sk::Float => left_inner == right_inner,
Sk::Bool | Sk::AbstractInt | Sk::AbstractFloat => false,
},
ref other => {
log::error!("Op {:?} left type {:?}", op, other);
@ -802,10 +784,8 @@ impl super::Validator {
},
Bo::And | Bo::InclusiveOr => match *left_inner {
Ti::Scalar(scalar) | Ti::Vector { scalar, .. } => match scalar.kind {
Sk::Bool | Sk::Sint | Sk::Uint | Sk::AbstractInt => {
left_inner == right_inner
}
Sk::Float | Sk::AbstractFloat => false,
Sk::Bool | Sk::Sint | Sk::Uint => left_inner == right_inner,
Sk::Float | Sk::AbstractInt | Sk::AbstractFloat => false,
},
ref other => {
log::error!("Op {:?} left type {:?}", op, other);
@ -814,8 +794,8 @@ impl super::Validator {
},
Bo::ExclusiveOr => match *left_inner {
Ti::Scalar(scalar) | Ti::Vector { scalar, .. } => match scalar.kind {
Sk::Sint | Sk::Uint | Sk::AbstractInt => left_inner == right_inner,
Sk::Bool | Sk::Float | Sk::AbstractFloat => false,
Sk::Sint | Sk::Uint => left_inner == right_inner,
Sk::Bool | Sk::Float | Sk::AbstractInt | Sk::AbstractFloat => false,
},
ref other => {
log::error!("Op {:?} left type {:?}", op, other);
@ -843,10 +823,8 @@ impl super::Validator {
}
};
match base_scalar.kind {
Sk::Sint | Sk::Uint | Sk::AbstractInt => {
base_size.is_ok() && base_size == shift_size
}
Sk::Float | Sk::AbstractFloat | Sk::Bool => false,
Sk::Sint | Sk::Uint => base_size.is_ok() && base_size == shift_size,
Sk::Float | Sk::AbstractInt | Sk::AbstractFloat | Sk::Bool => false,
}
}
};

2
naga/src/valid/function.rs
View File

@ -112,7 +112,7 @@ pub enum FunctionError {
InvalidStorePointer(Handle<crate::Expression>),
#[error("The value {0:?} can not be stored")]
InvalidStoreValue(Handle<crate::Expression>),
#[error("Store of {value:?} into {pointer:?} doesn't have matching types")]
#[error("The type of {value:?} doesn't match the type stored in {pointer:?}")]
InvalidStoreTypes {
pointer: Handle<crate::Expression>,
value: Handle<crate::Expression>,

77
naga/tests/in/abstract-types.wgsl Normal file
View File

@ -0,0 +1,77 @@
// i/x: type inferred / explicit
// vX/mX/aX: vector / matrix / array of X
// where X: u/i/f: u32 / i32 / f32
// s: vector splat
// r: vector spread (vector arg to vector constructor)
// p: "partial" constructor (type parameter inferred)
// u/i/f/ai/af: u32 / i32 / f32 / abstract float / abstract integer as parameter
// _: just for alignment
// Ensure that:
// - the inferred type is correct.
// - all parameters' types are considered.
// - all parameters are converted to the consensus type.
const xvupaiai: vec2<u32> = vec2(42, 43);
const xvfpaiai: vec2<f32> = vec2(44, 45);
const xvupuai: vec2<u32> = vec2(42u, 43);
const xvupaiu: vec2<u32> = vec2(42, 43u);
const xvuuai: vec2<u32> = vec2<u32>(42u, 43);
const xvuaiu: vec2<u32> = vec2<u32>(42, 43u);
const xmfpaiaiaiai: mat2x2<f32> = mat2x2(1, 2, 3, 4);
const xmfpafaiaiai: mat2x2<f32> = mat2x2(1.0, 2, 3, 4);
const xmfpaiafaiai: mat2x2<f32> = mat2x2(1, 2.0, 3, 4);
const xmfpaiaiafai: mat2x2<f32> = mat2x2(1, 2, 3.0, 4);
const xmfpaiaiaiaf: mat2x2<f32> = mat2x2(1, 2, 3, 4.0);
const ivispai = vec2(1);
const ivfspaf = vec2(1.0);
const ivis_ai = vec2<i32>(1);
const ivus_ai = vec2<u32>(1);
const ivfs_ai = vec2<f32>(1);
const ivfs_af = vec2<f32>(1.0);
const iafafaf = array<f32, 2>(1.0, 2.0);
const iafaiai = array<f32, 2>(1, 2);
const iafpafaf = array(1.0, 2.0);
const iafpaiaf = array(1, 2.0);
const iafpafai = array(1.0, 2);
const xafpafaf: array<f32, 2> = array(1.0, 2.0);
struct S {
f: f32,
i: i32,
u: u32,
}
const s_f_i_u: S = S(1.0f, 1i, 1u);
const s_f_iai: S = S(1.0f, 1i, 1);
const s_fai_u: S = S(1.0f, 1, 1u);
const s_faiai: S = S(1.0f, 1, 1);
const saf_i_u: S = S(1.0, 1i, 1u);
const saf_iai: S = S(1.0, 1i, 1);
const safai_u: S = S(1.0, 1, 1u);
const safaiai: S = S(1.0, 1, 1);
// Vector construction with spreads
const ivfr_f__f = vec3<f32>(vec2<f32>(1.0f, 2.0f), 3.0f);
const ivfr_f_af = vec3<f32>(vec2<f32>(1.0f, 2.0f), 3.0 );
const ivfraf__f = vec3<f32>(vec2 (1.0 , 2.0 ), 3.0f);
const ivfraf_af = vec3<f32>(vec2 (1.0 , 2.0 ), 3.0 );
const ivf__fr_f = vec3<f32>(1.0f, vec2<f32>(2.0f, 3.0f));
const ivf__fraf = vec3<f32>(1.0f, vec2 (2.0 , 3.0 ));
const ivf_afr_f = vec3<f32>(1.0 , vec2<f32>(2.0f, 3.0f));
const ivf_afraf = vec3<f32>(1.0 , vec2 (2.0 , 3.0 ));
const ivfr_f_ai = vec3<f32>(vec2<f32>(1.0f, 2.0f), 3 );
const ivfrai__f = vec3<f32>(vec2 (1 , 2 ), 3.0f);
const ivfrai_ai = vec3<f32>(vec2 (1 , 2 ), 3 );
const ivf__frai = vec3<f32>(1.0f, vec2 (2 , 3 ));
const ivf_air_f = vec3<f32>(1 , vec2<f32>(2.0f, 3.0f));
const ivf_airai = vec3<f32>(1 , vec2 (2 , 3 ));

16
naga/tests/out/ir/access.compact.ron
View File

@ -1629,6 +1629,14 @@
1: "foo",
},
body: [
Emit((
start: 0,
end: 0,
)),
Emit((
start: 0,
end: 0,
)),
Emit((
start: 2,
end: 3,
@ -2192,6 +2200,14 @@
],
result: None,
),
Emit((
start: 0,
end: 0,
)),
Emit((
start: 0,
end: 0,
)),
Emit((
start: 3,
end: 4,

1138
naga/tests/out/ir/access.ron
View File

File diff suppressed because it is too large Load Diff

6
naga/tests/out/ir/collatz.ron
View File

@ -60,11 +60,7 @@
init: None,
),
],
const_expressions: [
Literal(I32(0)),
Literal(I32(0)),
Literal(I32(1)),
],
const_expressions: [],
functions: [
(
name: Some("collatz_iterations"),

59
naga/tests/out/msl/abstract-types.msl Normal file
View File

@ -0,0 +1,59 @@
// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
struct type_5 {
float inner[2];
};
struct S {
float f;
int i;
uint u;
};
constant metal::uint2 xvupaiai = metal::uint2(42u, 43u);
constant metal::float2 xvfpaiai = metal::float2(44.0, 45.0);
constant metal::uint2 xvupuai = metal::uint2(42u, 43u);
constant metal::uint2 xvupaiu = metal::uint2(42u, 43u);
constant metal::uint2 xvuuai = metal::uint2(42u, 43u);
constant metal::uint2 xvuaiu = metal::uint2(42u, 43u);
constant metal::float2x2 xmfpaiaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 xmfpafaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 xmfpaiafaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 xmfpaiaiafai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 xmfpaiaiaiaf = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::int2 ivispai = metal::int2(1);
constant metal::float2 ivfspaf = metal::float2(1.0);
constant metal::int2 ivis_ai = metal::int2(1);
constant metal::uint2 ivus_ai = metal::uint2(1u);
constant metal::float2 ivfs_ai = metal::float2(1.0);
constant metal::float2 ivfs_af = metal::float2(1.0);
constant type_5 iafafaf = type_5 {1.0, 2.0};
constant type_5 iafaiai = type_5 {1.0, 2.0};
constant type_5 iafpafaf = type_5 {1.0, 2.0};
constant type_5 iafpaiaf = type_5 {1.0, 2.0};
constant type_5 iafpafai = type_5 {1.0, 2.0};
constant type_5 xafpafaf = type_5 {1.0, 2.0};
constant S s_f_i_u = S {1.0, 1, 1u};
constant S s_f_iai = S {1.0, 1, 1u};
constant S s_fai_u = S {1.0, 1, 1u};
constant S s_faiai = S {1.0, 1, 1u};
constant S saf_i_u = S {1.0, 1, 1u};
constant S saf_iai = S {1.0, 1, 1u};
constant S safai_u = S {1.0, 1, 1u};
constant S safaiai = S {1.0, 1, 1u};
constant metal::float3 ivfr_f_f = metal::float3(metal::float2(1.0, 2.0), 3.0);
constant metal::float3 ivfr_f_af = metal::float3(metal::float2(1.0, 2.0), 3.0);
constant metal::float3 ivfraf_f = metal::float3(metal::float2(1.0, 2.0), 3.0);
constant metal::float3 ivfraf_af = metal::float3(metal::float2(1.0, 2.0), 3.0);
constant metal::float3 ivf_fr_f = metal::float3(1.0, metal::float2(2.0, 3.0));
constant metal::float3 ivf_fraf = metal::float3(1.0, metal::float2(2.0, 3.0));
constant metal::float3 ivf_afr_f = metal::float3(1.0, metal::float2(2.0, 3.0));
constant metal::float3 ivf_afraf = metal::float3(1.0, metal::float2(2.0, 3.0));
constant metal::float3 ivfr_f_ai = metal::float3(metal::float2(1.0, 2.0), 3.0);
constant metal::float3 ivfrai_f = metal::float3(metal::float2(1.0, 2.0), 3.0);
constant metal::float3 ivfrai_ai = metal::float3(metal::float2(1.0, 2.0), 3.0);
constant metal::float3 ivf_frai = metal::float3(1.0, metal::float2(2.0, 3.0));
constant metal::float3 ivf_air_f = metal::float3(1.0, metal::float2(2.0, 3.0));
constant metal::float3 ivf_airai = metal::float3(1.0, metal::float2(2.0, 3.0));

46
naga/tests/out/spv/abstract-types.spvasm Normal file
View File

@ -0,0 +1,46 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 36
OpCapability Shader
OpCapability Linkage
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpDecorate %10 ArrayStride 4
OpMemberDecorate %12 0 Offset 0
OpMemberDecorate %12 1 Offset 4
OpMemberDecorate %12 2 Offset 8
%2 = OpTypeVoid
%4 = OpTypeInt 32 0
%3 = OpTypeVector %4 2
%6 = OpTypeFloat 32
%5 = OpTypeVector %6 2
%7 = OpTypeMatrix %5 2
%9 = OpTypeInt 32 1
%8 = OpTypeVector %9 2
%11 = OpConstant %4 2
%10 = OpTypeArray %6 %11
%12 = OpTypeStruct %6 %9 %4
%13 = OpTypeVector %6 3
%14 = OpConstant %4 42
%15 = OpConstant %4 43
%16 = OpConstantComposite %3 %14 %15
%17 = OpConstant %6 44.0
%18 = OpConstant %6 45.0
%19 = OpConstantComposite %5 %17 %18
%20 = OpConstant %6 1.0
%21 = OpConstant %6 2.0
%22 = OpConstantComposite %5 %20 %21
%23 = OpConstant %6 3.0
%24 = OpConstant %6 4.0
%25 = OpConstantComposite %5 %23 %24
%26 = OpConstantComposite %7 %22 %25
%27 = OpConstant %9 1
%28 = OpConstantComposite %8 %27 %27
%29 = OpConstantComposite %5 %20 %20
%30 = OpConstant %4 1
%31 = OpConstantComposite %3 %30 %30
%32 = OpConstantComposite %10 %20 %21
%33 = OpConstantComposite %12 %20 %27 %30
%34 = OpConstantComposite %13 %20 %21 %23
%35 = OpConstantComposite %5 %21 %23

8
naga/tests/out/spv/ray-query.spvasm
View File

@ -66,10 +66,10 @@ OpMemberDecorate %18 0 Offset 0
%47 = OpConstantComposite %5 %27 %25 %27
%48 = OpConstant %4 4
%49 = OpConstant %4 255
%50 = OpConstant %6 0.1
%51 = OpConstant %6 100.0
%52 = OpConstantComposite %5 %27 %27 %27
%53 = OpConstantComposite %14 %48 %49 %50 %51 %52 %47
%50 = OpConstantComposite %5 %27 %27 %27
%51 = OpConstant %6 0.1
%52 = OpConstant %6 100.0
%53 = OpConstantComposite %14 %48 %49 %51 %52 %50 %47
%55 = OpTypePointer Function %13
%72 = OpConstant %4 1
%85 = OpTypePointer StorageBuffer %4

52
naga/tests/out/wgsl/abstract-types.wgsl Normal file
View File

@ -0,0 +1,52 @@
struct S {
f: f32,
i: i32,
u: u32,
}
const xvupaiai: vec2<u32> = vec2<u32>(42u, 43u);
const xvfpaiai: vec2<f32> = vec2<f32>(44.0, 45.0);
const xvupuai: vec2<u32> = vec2<u32>(42u, 43u);
const xvupaiu: vec2<u32> = vec2<u32>(42u, 43u);
const xvuuai: vec2<u32> = vec2<u32>(42u, 43u);
const xvuaiu: vec2<u32> = vec2<u32>(42u, 43u);
const xmfpaiaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const xmfpafaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const xmfpaiafaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const xmfpaiaiafai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const xmfpaiaiaiaf: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const ivispai: vec2<i32> = vec2(1);
const ivfspaf: vec2<f32> = vec2(1.0);
const ivis_ai: vec2<i32> = vec2(1);
const ivus_ai: vec2<u32> = vec2(1u);
const ivfs_ai: vec2<f32> = vec2(1.0);
const ivfs_af: vec2<f32> = vec2(1.0);
const iafafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const iafaiai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const iafpafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const iafpaiaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const iafpafai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const xafpafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const s_f_i_u: S = S(1.0, 1, 1u);
const s_f_iai: S = S(1.0, 1, 1u);
const s_fai_u: S = S(1.0, 1, 1u);
const s_faiai: S = S(1.0, 1, 1u);
const saf_i_u: S = S(1.0, 1, 1u);
const saf_iai: S = S(1.0, 1, 1u);
const safai_u: S = S(1.0, 1, 1u);
const safaiai: S = S(1.0, 1, 1u);
const ivfr_f_f: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivfr_f_af: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivfraf_f: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivfraf_af: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivf_fr_f: vec3<f32> = vec3<f32>(1.0, vec2<f32>(2.0, 3.0));
const ivf_fraf: vec3<f32> = vec3<f32>(1.0, vec2<f32>(2.0, 3.0));
const ivf_afr_f: vec3<f32> = vec3<f32>(1.0, vec2<f32>(2.0, 3.0));
const ivf_afraf: vec3<f32> = vec3<f32>(1.0, vec2<f32>(2.0, 3.0));
const ivfr_f_ai: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivfrai_f: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivfrai_ai: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivf_frai: vec3<f32> = vec3<f32>(1.0, vec2<f32>(2.0, 3.0));
const ivf_air_f: vec3<f32> = vec3<f32>(1.0, vec2<f32>(2.0, 3.0));
const ivf_airai: vec3<f32> = vec3<f32>(1.0, vec2<f32>(2.0, 3.0));

4
naga/tests/snapshots.rs
View File

@ -783,6 +783,10 @@ fn convert_wgsl() {
"f64",
Targets::SPIRV | Targets::GLSL | Targets::HLSL | Targets::WGSL,
),
(
"abstract-types",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
),
];
for &(name, targets) in inputs.iter() {

25
naga/tests/wgsl_errors.rs
View File

@ -502,7 +502,7 @@ fn let_type_mismatch() {
r#"
const x: i32 = 1.0;
"#,
r#"error: the type of `x` is expected to be `i32`, but got `f32`
r#"error: the type of `x` is expected to be `i32`, but got `{AbstractFloat}`
┌─ wgsl:2:19
│
2 │ const x: i32 = 1.0;
@ -1996,9 +1996,12 @@ fn binding_array_non_struct() {
#[test]
fn compaction_preserves_spans() {
let source = r#"
const a: i32 = -(-(-(-42i)));
const b: vec2<u32> = vec2<u32>(42u, 43i);
"#; // ^^^^^^^^^^^^^^^^^^^ correct error span: 68..87
fn f() {
var a: i32 = -(-(-(-42i)));
var x: i32;
x = 42u;
}
"#; // ^^^ correct error span: 95..98
let mut module = naga::front::wgsl::parse_str(source).expect("source ought to parse");
naga::compact::compact(&mut module);
let err = naga::valid::Validator::new(
@ -2011,10 +2014,18 @@ fn compaction_preserves_spans() {
// Ideally this would all just be a `matches!` with a big pattern,
// but the `Span` API is full of opaque structs.
let mut spans = err.spans();
let first_span = spans.next().expect("error should have at least one span").0;
// The first span is the whole function.
let _ = spans.next().expect("error should have at least one span");
// The second span is the assignment destination.
let dest_span = spans
.next()
.expect("error should have at least two spans")
.0;
if !matches!(
first_span.to_range(),
Some(std::ops::Range { start: 68, end: 87 })
dest_span.to_range(),
Some(std::ops::Range { start: 95, end: 98 })
) {
panic!("Error message has wrong span:\n\n{err:#?}");
}