From 462418287c381fca16221f68d992a1c5ca359810 Mon Sep 17 00:00:00 2001 From: sagudev <16504129+sagudev@users.noreply.github.com> Date: Thu, 18 Sep 2025 10:42:40 +0200 Subject: [PATCH 1/3] [naga] Add `LiteralVector` and `match_literal_vector!` Co-authored-by: SelkoSays <70065171+SelkoSays@users.noreply.github.com> Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> --- naga/src/proc/constant_evaluator.rs | 386 ++++++++++++++++++++++++++++ 1 file changed, 386 insertions(+) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 5aa9baaa961..2029fa4a2cc 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -268,6 +268,392 @@ gen_component_wise_extractor! { ], } +/// Vectors with a concrete element type. +#[derive(Debug)] +enum LiteralVector { + F64(ArrayVec), + F32(ArrayVec), + F16(ArrayVec), + U32(ArrayVec), + I32(ArrayVec), + U64(ArrayVec), + I64(ArrayVec), + Bool(ArrayVec), + AbstractInt(ArrayVec), + AbstractFloat(ArrayVec), +} + +impl LiteralVector { + #[allow(clippy::missing_const_for_fn, reason = "MSRV")] + fn len(&self) -> usize { + match *self { + LiteralVector::F64(ref v) => v.len(), + LiteralVector::F32(ref v) => v.len(), + LiteralVector::F16(ref v) => v.len(), + LiteralVector::U32(ref v) => v.len(), + LiteralVector::I32(ref v) => v.len(), + LiteralVector::U64(ref v) => v.len(), + LiteralVector::I64(ref v) => v.len(), + LiteralVector::Bool(ref v) => v.len(), + LiteralVector::AbstractInt(ref v) => v.len(), + LiteralVector::AbstractFloat(ref v) => v.len(), + } + } + + /// Creates [`LiteralVector`] of size 1 from single [`Literal`] + fn from_literal(literal: Literal) -> Self { + match literal { + Literal::F64(e) => Self::F64(ArrayVec::from_iter(iter::once(e))), + Literal::F32(e) => Self::F32(ArrayVec::from_iter(iter::once(e))), + Literal::U32(e) => Self::U32(ArrayVec::from_iter(iter::once(e))), + Literal::I32(e) => Self::I32(ArrayVec::from_iter(iter::once(e))), + Literal::U64(e) => Self::U64(ArrayVec::from_iter(iter::once(e))), + Literal::I64(e) => Self::I64(ArrayVec::from_iter(iter::once(e))), + Literal::Bool(e) => Self::Bool(ArrayVec::from_iter(iter::once(e))), + Literal::AbstractInt(e) => Self::AbstractInt(ArrayVec::from_iter(iter::once(e))), + Literal::AbstractFloat(e) => Self::AbstractFloat(ArrayVec::from_iter(iter::once(e))), + Literal::F16(e) => Self::F16(ArrayVec::from_iter(iter::once(e))), + } + } + + /// Creates [`LiteralVector`] from [`ArrayVec`] of [`Literal`]s. + /// Returns error if components types do not match. + /// # Panics + /// Panics if vector is empty + fn from_literal_vec( + components: ArrayVec, + ) -> Result { + assert!(!components.is_empty()); + Ok(match components[0] { + Literal::I32(_) => Self::I32( + components + .iter() + .map(|l| match l { + &Literal::I32(v) => Ok(v), + // TODO: should we handle abstract int here? + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::U32(_) => Self::U32( + components + .iter() + .map(|l| match l { + &Literal::U32(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::I64(_) => Self::I64( + components + .iter() + .map(|l| match l { + &Literal::I64(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::U64(_) => Self::U64( + components + .iter() + .map(|l| match l { + &Literal::U64(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::F32(_) => Self::F32( + components + .iter() + .map(|l| match l { + &Literal::F32(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::F64(_) => Self::F64( + components + .iter() + .map(|l| match l { + &Literal::F64(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::Bool(_) => Self::Bool( + components + .iter() + .map(|l| match l { + &Literal::Bool(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::AbstractInt(_) => Self::AbstractInt( + components + .iter() + .map(|l| match l { + &Literal::AbstractInt(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::AbstractFloat(_) => Self::AbstractFloat( + components + .iter() + .map(|l| match l { + &Literal::AbstractFloat(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + Literal::F16(_) => Self::F16( + components + .iter() + .map(|l| match l { + &Literal::F16(v) => Ok(v), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?, + ), + }) + } + + #[allow(dead_code)] + /// Returns [`ArrayVec`] of [`Literal`]s + fn to_literal_vec(&self) -> ArrayVec { + match *self { + LiteralVector::F64(ref v) => v.iter().map(|e| (Literal::F64(*e))).collect(), + LiteralVector::F32(ref v) => v.iter().map(|e| (Literal::F32(*e))).collect(), + LiteralVector::F16(ref v) => v.iter().map(|e| (Literal::F16(*e))).collect(), + LiteralVector::U32(ref v) => v.iter().map(|e| (Literal::U32(*e))).collect(), + LiteralVector::I32(ref v) => v.iter().map(|e| (Literal::I32(*e))).collect(), + LiteralVector::U64(ref v) => v.iter().map(|e| (Literal::U64(*e))).collect(), + LiteralVector::I64(ref v) => v.iter().map(|e| (Literal::I64(*e))).collect(), + LiteralVector::Bool(ref v) => v.iter().map(|e| (Literal::Bool(*e))).collect(), + LiteralVector::AbstractInt(ref v) => { + v.iter().map(|e| (Literal::AbstractInt(*e))).collect() + } + LiteralVector::AbstractFloat(ref v) => { + v.iter().map(|e| (Literal::AbstractFloat(*e))).collect() + } + } + } + + #[allow(dead_code)] + /// Puts self into eval's expressions arena and returns handle to it + fn register_as_evaluated_expr( + &self, + eval: &mut ConstantEvaluator<'_>, + span: Span, + ) -> Result, ConstantEvaluatorError> { + let lit_vec = self.to_literal_vec(); + assert!(!lit_vec.is_empty()); + let expr = if lit_vec.len() == 1 { + Expression::Literal(lit_vec[0]) + } else { + Expression::Compose { + ty: eval.types.insert( + Type { + name: None, + inner: TypeInner::Vector { + size: match lit_vec.len() { + 2 => crate::VectorSize::Bi, + 3 => crate::VectorSize::Tri, + 4 => crate::VectorSize::Quad, + _ => unreachable!(), + }, + scalar: lit_vec[0].scalar(), + }, + }, + Span::UNDEFINED, + ), + components: lit_vec + .iter() + .map(|&l| eval.register_evaluated_expr(Expression::Literal(l), span)) + .collect::>()?, + } + }; + eval.register_evaluated_expr(expr, span) + } +} + +/// A macro for matching on [`LiteralVector`] variants. +/// +/// `Float` variant expands to `F16`, `F32`, `F64` and `AbstractFloat`. +/// `Integer` variant expands to `I32`, `I64`, `U32`, `U64` and `AbstractInt`. +/// +/// For output both [`Literal`] (fold) and [`LiteralVector`] (map) are supported. +/// +/// Example usage: +/// +/// ```rust,ignore +/// match_literal_vector!(match v => Literal { +/// F16 => |v| {v.sum()}, +/// Integer => |v| {v.sum()}, +/// U32 => |v| -> I32 {v.sum()}, // optionally override return type +/// }) +/// ``` +/// +/// ```rust,ignore +/// match_literal_vector!(match (e1, e2) => LiteralVector { +/// F16 => |e1, e2| {e1+e2}, +/// Integer => |e1, e2| {e1+e2}, +/// U32 => |e1, e2| -> I32 {e1+e2}, // optionally override return type +/// }) +/// ``` +macro_rules! match_literal_vector { + (match $lit_vec:expr => $out:ident { + $( + $ty:ident => |$($var:ident),+| $(-> $ret:ident)? { $body:expr } + ),+ + $(,)? + }) => { + match_literal_vector!(@inner_start $lit_vec; $out; [$($ty),+]; [$({ $($var),+ ; $($ret)? ; $body }),+]) + }; + + (@inner_start + $lit_vec:expr; + $out:ident; + [$($ty:ident),+]; + [$({ $($var:ident),+ ; $($ret:ident)? ; $body:expr }),+] + ) => { + match_literal_vector!(@inner + $lit_vec; + $out; + [$($ty),+]; + [] <> [$({ $($var),+ ; $($ret)? ; $body }),+] + ) + }; + + (@inner + $lit_vec:expr; + $out:ident; + [$ty:ident $(, $ty1:ident)*]; + [$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <> + [$({ $($var:ident),+ ; $($ret:ident)? ; $body:expr }),+] + ) => { + match_literal_vector!(@inner + $ty; + $lit_vec; + $out; + [$($ty1),*]; + [$({$_ty ; $($_var),+ ; $($_ret)? ; $_body}),*] <> + [$({ $($var),+ ; $($ret)? ; $body }),+] + ) + }; + (@inner + Integer; + $lit_vec:expr; + $out:ident; + [$($ty:ident),*]; + [$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <> + [ + { $($var:ident),+ ; $($ret:ident)? ; $body:expr } + $(,{ $($var1:ident),+ ; $($ret1:ident)? ; $body1:expr })* + ] + ) => { + match_literal_vector!(@inner + $lit_vec; + $out; + [U32, I32, U64, I64, AbstractInt $(, $ty)*]; + [$({$_ty ; $($_var),+ ; $($_ret)? ; $_body}),*] <> + [ + { $($var),+ ; $($ret)? ; $body }, // U32 + { $($var),+ ; $($ret)? ; $body }, // I32 + { $($var),+ ; $($ret)? ; $body }, // U64 + { $($var),+ ; $($ret)? ; $body }, // I64 + { $($var),+ ; $($ret)? ; $body } // AbstractInt + $(,{ $($var1),+ ; $($ret1)? ; $body1 })* + ] + ) + }; + (@inner + Float; + $lit_vec:expr; + $out:ident; + [$($ty:ident),*]; + [$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <> + [ + { $($var:ident),+ ; $($ret:ident)? ; $body:expr } + $(,{ $($var1:ident),+ ; $($ret1:ident)? ; $body1:expr })* + ] + ) => { + match_literal_vector!(@inner + $lit_vec; + $out; + [F16, F32, F64, AbstractFloat $(, $ty)*]; + [$({$_ty ; $($_var),+ ; $($_ret)? ; $_body}),*] <> + [ + { $($var),+ ; $($ret)? ; $body }, // F16 + { $($var),+ ; $($ret)? ; $body }, // F32 + { $($var),+ ; $($ret)? ; $body }, // F64 + { $($var),+ ; $($ret)? ; $body } // AbstractFloat + $(,{ $($var1),+ ; $($ret1)? ; $body1 })* + ] + ) + }; + (@inner + $ty:ident; + $lit_vec:expr; + $out:ident; + [$ty1:ident $(,$ty2:ident)*]; + [$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <> [ + { $($var:ident),+ ; $($ret:ident)? ; $body:expr } + $(, { $($var1:ident),+ ; $($ret1:ident)? ; $body1:expr })* + ] + ) => { + match_literal_vector!(@inner + $ty1; + $lit_vec; + $out; + [$($ty2),*]; + [ + $({$_ty ; $($_var),+ ; $($_ret)? ; $_body},)* + { $ty; $($var),+ ; $($ret)? ; $body } + ] <> + [$({ $($var1),+ ; $($ret1)? ; $body1 }),*] + + ) + }; + (@inner + $ty:ident; + $lit_vec:expr; + $out:ident; + []; + [$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <> + [{ $($var:ident),+ ; $($ret:ident)? ; $body:expr }] + ) => { + match_literal_vector!(@inner_finish + $lit_vec; + $out; + [ + $({ $_ty ; $($_var),+ ; $($_ret)? ; $_body },)* + { $ty; $($var),+ ; $($ret)? ; $body } + ] + ) + }; + (@inner_finish + $lit_vec:expr; + $out:ident; + [$({$ty:ident ; $($var:ident),+ ; $($ret:ident)? ; $body:expr}),+] + ) => { + match $lit_vec { + $( + #[allow(unused_parens)] + ($(LiteralVector::$ty(ref $var)),+) => { Ok(match_literal_vector!(@expand_ret $out; $ty $(; $ret)? ; $body)) } + )+ + _ => Err(ConstantEvaluatorError::InvalidMathArg), + } + }; + (@expand_ret $out:ident; $ty:ident; $body:expr) => { + $out::$ty($body) + }; + (@expand_ret $out:ident; $_ty:ident; $ret:ident; $body:expr) => { + $out::$ret($body) + }; +} + #[derive(Debug)] enum Behavior<'a> { Wgsl(WgslRestrictions<'a>), From dd325cca91f8bcfc9ef616f138f0e4c32e4a2864 Mon Sep 17 00:00:00 2001 From: sagudev <16504129+sagudev@users.noreply.github.com> Date: Thu, 18 Sep 2025 10:43:12 +0200 Subject: [PATCH 2/3] [naga] Implement builtins dot, length, distance, normalize in const using LiteralVector Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> --- naga/src/proc/constant_evaluator.rs | 136 ++++++++++++++++++++++++++-- 1 file changed, 129 insertions(+), 7 deletions(-) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 2029fa4a2cc..08c7708edc5 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1730,17 +1730,106 @@ impl<'a> ConstantEvaluator<'a> { self.packed_dot_product(arg, arg1.unwrap(), span, false) } crate::MathFunction::Cross => self.cross_product(arg, arg1.unwrap(), span), + crate::MathFunction::Dot => { + // https://www.w3.org/TR/WGSL/#dot-builtin + let e1 = self.extract_vec(arg, false)?; + let e2 = self.extract_vec(arg1.unwrap(), false)?; + if e1.len() != e2.len() { + return Err(ConstantEvaluatorError::InvalidMathArg); + } + + fn int_dot

(a: &[P], b: &[P]) -> Result + where + P: num_traits::PrimInt + num_traits::CheckedAdd + num_traits::CheckedMul, + { + a.iter() + .zip(b.iter()) + .map(|(&aa, bb)| aa.checked_mul(bb)) + .try_fold(P::zero(), |acc, x| { + if let Some(x) = x { + acc.checked_add(&x) + } else { + None + } + }) + .ok_or(ConstantEvaluatorError::Overflow( + "in dot built-in".to_string(), + )) + } + + let result = match_literal_vector!(match (e1, e2) => Literal { + Float => |e1, e2| { e1.iter().zip(e2.iter()).map(|(&aa, &bb)| aa * bb).sum() }, + Integer => |e1, e2| { int_dot(e1, e2)? }, + })?; + self.register_evaluated_expr(Expression::Literal(result), span) + } + crate::MathFunction::Length => { + // https://www.w3.org/TR/WGSL/#length-builtin + let e1 = self.extract_vec(arg, true)?; + + fn float_length(e: &[F]) -> F + where + F: core::ops::Mul, + F: num_traits::Float + iter::Sum, + { + e.iter().map(|&ei| ei * ei).sum::().sqrt() + } + + let result = match_literal_vector!(match e1 => Literal { + Float => |e1| { float_length(e1) }, + })?; + self.register_evaluated_expr(Expression::Literal(result), span) + } + crate::MathFunction::Distance => { + // https://www.w3.org/TR/WGSL/#distance-builtin + let e1 = self.extract_vec(arg, true)?; + let e2 = self.extract_vec(arg1.unwrap(), true)?; + if e1.len() != e2.len() { + return Err(ConstantEvaluatorError::InvalidMathArg); + } + + fn float_distance(a: &[F], b: &[F]) -> F + where + F: core::ops::Mul, + F: num_traits::Float + iter::Sum + core::ops::Sub, + { + a.iter() + .zip(b.iter()) + .map(|(&aa, &bb)| aa - bb) + .map(|ei| ei * ei) + .sum::() + .sqrt() + } + let result = match_literal_vector!(match (e1, e2) => Literal { + Float => |e1, e2| { float_distance(e1, e2) }, + })?; + self.register_evaluated_expr(Expression::Literal(result), span) + } + crate::MathFunction::Normalize => { + // https://www.w3.org/TR/WGSL/#normalize-builtin + let e1 = self.extract_vec(arg, true)?; + + fn float_normalize(e: &[F]) -> ArrayVec + where + F: core::ops::Mul, + F: num_traits::Float + iter::Sum, + { + let len = e.iter().map(|&ei| ei * ei).sum::().sqrt(); + e.iter().map(|&ei| ei / len).collect() + } + + let result = match_literal_vector!(match e1 => LiteralVector { + Float => |e1| { float_normalize(e1) }, + })?; + result.register_as_evaluated_expr(self, span) + } // unimplemented crate::MathFunction::Atan2 | crate::MathFunction::Modf | crate::MathFunction::Frexp | crate::MathFunction::Ldexp - | crate::MathFunction::Dot | crate::MathFunction::Outer - | crate::MathFunction::Distance - | crate::MathFunction::Length - | crate::MathFunction::Normalize | crate::MathFunction::FaceForward | crate::MathFunction::Reflect | crate::MathFunction::Refract @@ -1816,8 +1905,8 @@ impl<'a> ConstantEvaluator<'a> { ) -> Result, ConstantEvaluatorError> { use Literal as Li; - let (a, ty) = self.extract_vec::<3>(a)?; - let (b, _) = self.extract_vec::<3>(b)?; + let (a, ty) = self.extract_vec_with_size::<3>(a)?; + let (b, _) = self.extract_vec_with_size::<3>(b)?; let product = match (a, b) { ( @@ -1884,7 +1973,7 @@ impl<'a> ConstantEvaluator<'a> { /// values. /// /// Also return the type handle from the `Compose` expression. - fn extract_vec( + fn extract_vec_with_size( &mut self, expr: Handle, ) -> Result<([Literal; N], Handle), ConstantEvaluatorError> { @@ -1908,6 +1997,39 @@ impl<'a> ConstantEvaluator<'a> { Ok((value, ty)) } + /// Extract the values of a `vecN` from `expr`. + /// + /// Return the value of `expr`, whose type is `vecN` for some + /// vector size `N` and scalar `S`, as an array of `N` [`Literal`] + /// values. + /// + /// Also return the type handle from the `Compose` expression. + fn extract_vec( + &mut self, + expr: Handle, + allow_single: bool, + ) -> Result { + let span = self.expressions.get_span(expr); + let expr = self.eval_zero_value_and_splat(expr, span)?; + + match self.expressions[expr] { + Expression::Literal(literal) if allow_single => { + Ok(LiteralVector::from_literal(literal)) + } + Expression::Compose { ty, ref components } => { + let components: ArrayVec = + crate::proc::flatten_compose(ty, components, self.expressions, self.types) + .map(|expr| match self.expressions[expr] { + Expression::Literal(l) => Ok(l), + _ => Err(ConstantEvaluatorError::InvalidMathArg), + }) + .collect::>()?; + LiteralVector::from_literal_vec(components) + } + _ => Err(ConstantEvaluatorError::InvalidMathArg), + } + } + fn array_length( &mut self, array: Handle, From fc966e8a514ac559a44479140270f4d2d0fb8974 Mon Sep 17 00:00:00 2001 From: sagudev <16504129+sagudev@users.noreply.github.com> Date: Thu, 18 Sep 2025 10:43:20 +0200 Subject: [PATCH 3/3] Update snapshots Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> --- .../out/glsl/wgsl-functions.main.Compute.glsl | 5 +- .../wgsl-math-functions.main.Fragment.glsl | 13 +- naga/tests/out/hlsl/wgsl-functions.hlsl | 3 +- naga/tests/out/hlsl/wgsl-math-functions.hlsl | 5 - naga/tests/out/msl/wgsl-functions.msl | 5 +- naga/tests/out/msl/wgsl-math-functions.msl | 1 - naga/tests/out/spv/wgsl-functions.spvasm | 128 ++++++++---------- naga/tests/out/spv/wgsl-math-functions.spvasm | 122 ++++++++--------- naga/tests/out/wgsl/wgsl-functions.wgsl | 3 +- naga/tests/out/wgsl/wgsl-math-functions.wgsl | 1 - 10 files changed, 120 insertions(+), 166 deletions(-) diff --git a/naga/tests/out/glsl/wgsl-functions.main.Compute.glsl b/naga/tests/out/glsl/wgsl-functions.main.Compute.glsl index ffe08719249..4de519fe485 100644 --- a/naga/tests/out/glsl/wgsl-functions.main.Compute.glsl +++ b/naga/tests/out/glsl/wgsl-functions.main.Compute.glsl @@ -20,10 +20,7 @@ int test_integer_dot_product() { uvec3 a_3_ = uvec3(1u); uvec3 b_3_ = uvec3(1u); uint c_3_ = ( + a_3_.x * b_3_.x + a_3_.y * b_3_.y + a_3_.z * b_3_.z); - ivec4 _e11 = ivec4(4); - ivec4 _e13 = ivec4(2); - int c_4_ = ( + _e11.x * _e13.x + _e11.y * _e13.y + _e11.z * _e13.z + _e11.w * _e13.w); - return c_4_; + return 32; } uint test_packed_integer_dot_product() { diff --git a/naga/tests/out/glsl/wgsl-math-functions.main.Fragment.glsl b/naga/tests/out/glsl/wgsl-math-functions.main.Fragment.glsl index 1877fb3d91e..a82716603e1 100644 --- a/naga/tests/out/glsl/wgsl-math-functions.main.Fragment.glsl +++ b/naga/tests/out/glsl/wgsl-math-functions.main.Fragment.glsl @@ -64,7 +64,6 @@ void main() { vec4 g = refract(v, v, 1.0); ivec4 sign_b = ivec4(-1, -1, -1, -1); vec4 sign_d = vec4(-1.0, -1.0, -1.0, -1.0); - int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y); ivec2 flb_b = ivec2(-1, -1); uvec2 flb_c = uvec2(0u, 0u); ivec2 ftb_c = ivec2(0, 0); @@ -88,12 +87,12 @@ void main() { int frexp_c = naga_frexp(1.5).exp_; int frexp_d = naga_frexp(vec4(1.5, 1.5, 1.5, 1.5)).exp_.x; float quantizeToF16_a = unpackHalf2x16(packHalf2x16(vec2(1.0))).x; - vec2 _e120 = vec2(1.0, 1.0); - vec2 quantizeToF16_b = unpackHalf2x16(packHalf2x16(_e120)); - vec3 _e125 = vec3(1.0, 1.0, 1.0); - vec3 quantizeToF16_c = vec3(unpackHalf2x16(packHalf2x16(_e125.xy)), unpackHalf2x16(packHalf2x16(_e125.zz)).x); - vec4 _e131 = vec4(1.0, 1.0, 1.0, 1.0); - vec4 quantizeToF16_d = vec4(unpackHalf2x16(packHalf2x16(_e131.xy)), unpackHalf2x16(packHalf2x16(_e131.zw))); + vec2 _e118 = vec2(1.0, 1.0); + vec2 quantizeToF16_b = unpackHalf2x16(packHalf2x16(_e118)); + vec3 _e123 = vec3(1.0, 1.0, 1.0); + vec3 quantizeToF16_c = vec3(unpackHalf2x16(packHalf2x16(_e123.xy)), unpackHalf2x16(packHalf2x16(_e123.zz)).x); + vec4 _e129 = vec4(1.0, 1.0, 1.0, 1.0); + vec4 quantizeToF16_d = vec4(unpackHalf2x16(packHalf2x16(_e129.xy)), unpackHalf2x16(packHalf2x16(_e129.zw))); return; } diff --git a/naga/tests/out/hlsl/wgsl-functions.hlsl b/naga/tests/out/hlsl/wgsl-functions.hlsl index bb72985638b..a09d22d4d9c 100644 --- a/naga/tests/out/hlsl/wgsl-functions.hlsl +++ b/naga/tests/out/hlsl/wgsl-functions.hlsl @@ -14,8 +14,7 @@ int test_integer_dot_product() uint3 a_3_ = (1u).xxx; uint3 b_3_ = (1u).xxx; uint c_3_ = dot(a_3_, b_3_); - int c_4_ = dot((int(4)).xxxx, (int(2)).xxxx); - return c_4_; + return int(32); } uint test_packed_integer_dot_product() diff --git a/naga/tests/out/hlsl/wgsl-math-functions.hlsl b/naga/tests/out/hlsl/wgsl-math-functions.hlsl index b31d804a7b2..8da03ba5368 100644 --- a/naga/tests/out/hlsl/wgsl-math-functions.hlsl +++ b/naga/tests/out/hlsl/wgsl-math-functions.hlsl @@ -63,10 +63,6 @@ _frexp_result_vec4_f32_ naga_frexp(float4 arg) { return result; } -int2 ZeroValueint2() { - return (int2)0; -} - void main() { float4 v = (0.0).xxxx; @@ -78,7 +74,6 @@ void main() float4 g = refract(v, v, 1.0); int4 sign_b = int4(int(-1), int(-1), int(-1), int(-1)); float4 sign_d = float4(-1.0, -1.0, -1.0, -1.0); - int const_dot = dot(ZeroValueint2(), ZeroValueint2()); int2 flb_b = int2(int(-1), int(-1)); uint2 flb_c = uint2(0u, 0u); int2 ftb_c = int2(int(0), int(0)); diff --git a/naga/tests/out/msl/wgsl-functions.msl b/naga/tests/out/msl/wgsl-functions.msl index 6e5e1490be7..fd3bdc249ed 100644 --- a/naga/tests/out/msl/wgsl-functions.msl +++ b/naga/tests/out/msl/wgsl-functions.msl @@ -21,10 +21,7 @@ int test_integer_dot_product( metal::uint3 a_3_ = metal::uint3(1u); metal::uint3 b_3_ = metal::uint3(1u); uint c_3_ = ( + a_3_.x * b_3_.x + a_3_.y * b_3_.y + a_3_.z * b_3_.z); - metal::int4 _e11 = metal::int4(4); - metal::int4 _e13 = metal::int4(2); - int c_4_ = ( + _e11.x * _e13.x + _e11.y * _e13.y + _e11.z * _e13.z + _e11.w * _e13.w); - return c_4_; + return 32; } uint test_packed_integer_dot_product( diff --git a/naga/tests/out/msl/wgsl-math-functions.msl b/naga/tests/out/msl/wgsl-math-functions.msl index 97f3ea73174..50a52d56f68 100644 --- a/naga/tests/out/msl/wgsl-math-functions.msl +++ b/naga/tests/out/msl/wgsl-math-functions.msl @@ -66,7 +66,6 @@ fragment void main_( metal::float4 g = metal::refract(v, v, 1.0); metal::int4 sign_b = metal::int4(-1, -1, -1, -1); metal::float4 sign_d = metal::float4(-1.0, -1.0, -1.0, -1.0); - int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); metal::int2 flb_b = metal::int2(-1, -1); metal::uint2 flb_c = metal::uint2(0u, 0u); metal::int2 ftb_c = metal::int2(0, 0); diff --git a/naga/tests/out/spv/wgsl-functions.spvasm b/naga/tests/out/spv/wgsl-functions.spvasm index a84e322ae74..8ddc9953016 100644 --- a/naga/tests/out/spv/wgsl-functions.spvasm +++ b/naga/tests/out/spv/wgsl-functions.spvasm @@ -1,15 +1,15 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 95 +; Bound: 75 OpCapability Shader OpCapability DotProduct OpCapability DotProductInput4x8BitPacked OpExtension "SPV_KHR_integer_dot_product" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %89 "main" -OpExecutionMode %89 LocalSize 1 1 1 +OpEntryPoint GLCompute %69 "main" +OpExecutionMode %69 LocalSize 1 1 1 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpTypeVector %4 2 @@ -27,22 +27,18 @@ OpExecutionMode %89 LocalSize 1 1 1 %22 = OpConstant %6 1 %23 = OpTypeVector %6 3 %24 = OpConstantComposite %23 %22 %22 %22 -%25 = OpConstant %5 4 -%26 = OpTypeVector %5 4 -%27 = OpConstantComposite %26 %25 %25 %25 %25 -%28 = OpConstant %5 2 -%29 = OpConstantComposite %26 %28 %28 %28 %28 -%32 = OpConstantNull %5 -%41 = OpConstantNull %6 -%71 = OpTypeFunction %6 -%72 = OpConstant %6 2 -%73 = OpConstant %6 3 -%74 = OpConstant %6 4 -%75 = OpConstant %6 5 -%76 = OpConstant %6 6 -%77 = OpConstant %6 7 -%78 = OpConstant %6 8 -%90 = OpTypeFunction %2 +%25 = OpConstant %5 32 +%28 = OpConstantNull %5 +%37 = OpConstantNull %6 +%51 = OpTypeFunction %6 +%52 = OpConstant %6 2 +%53 = OpConstant %6 3 +%54 = OpConstant %6 4 +%55 = OpConstant %6 5 +%56 = OpConstant %6 6 +%57 = OpConstant %6 7 +%58 = OpConstant %6 8 +%70 = OpTypeFunction %2 %8 = OpFunction %3 None %9 %7 = OpLabel OpBranch %14 @@ -52,66 +48,50 @@ OpReturnValue %15 OpFunctionEnd %17 = OpFunction %5 None %18 %16 = OpLabel -OpBranch %30 -%30 = OpLabel -%33 = OpCompositeExtract %5 %21 0 -%34 = OpCompositeExtract %5 %21 0 +OpBranch %26 +%26 = OpLabel +%29 = OpCompositeExtract %5 %21 0 +%30 = OpCompositeExtract %5 %21 0 +%31 = OpIMul %5 %29 %30 +%32 = OpIAdd %5 %28 %31 +%33 = OpCompositeExtract %5 %21 1 +%34 = OpCompositeExtract %5 %21 1 %35 = OpIMul %5 %33 %34 -%36 = OpIAdd %5 %32 %35 -%37 = OpCompositeExtract %5 %21 1 -%38 = OpCompositeExtract %5 %21 1 -%39 = OpIMul %5 %37 %38 -%31 = OpIAdd %5 %36 %39 -%42 = OpCompositeExtract %6 %24 0 -%43 = OpCompositeExtract %6 %24 0 +%27 = OpIAdd %5 %32 %35 +%38 = OpCompositeExtract %6 %24 0 +%39 = OpCompositeExtract %6 %24 0 +%40 = OpIMul %6 %38 %39 +%41 = OpIAdd %6 %37 %40 +%42 = OpCompositeExtract %6 %24 1 +%43 = OpCompositeExtract %6 %24 1 %44 = OpIMul %6 %42 %43 %45 = OpIAdd %6 %41 %44 -%46 = OpCompositeExtract %6 %24 1 -%47 = OpCompositeExtract %6 %24 1 +%46 = OpCompositeExtract %6 %24 2 +%47 = OpCompositeExtract %6 %24 2 %48 = OpIMul %6 %46 %47 -%49 = OpIAdd %6 %45 %48 -%50 = OpCompositeExtract %6 %24 2 -%51 = OpCompositeExtract %6 %24 2 -%52 = OpIMul %6 %50 %51 -%40 = OpIAdd %6 %49 %52 -%54 = OpCompositeExtract %5 %27 0 -%55 = OpCompositeExtract %5 %29 0 -%56 = OpIMul %5 %54 %55 -%57 = OpIAdd %5 %32 %56 -%58 = OpCompositeExtract %5 %27 1 -%59 = OpCompositeExtract %5 %29 1 -%60 = OpIMul %5 %58 %59 -%61 = OpIAdd %5 %57 %60 -%62 = OpCompositeExtract %5 %27 2 -%63 = OpCompositeExtract %5 %29 2 -%64 = OpIMul %5 %62 %63 -%65 = OpIAdd %5 %61 %64 -%66 = OpCompositeExtract %5 %27 3 -%67 = OpCompositeExtract %5 %29 3 -%68 = OpIMul %5 %66 %67 -%53 = OpIAdd %5 %65 %68 -OpReturnValue %53 +%36 = OpIAdd %6 %45 %48 +OpReturnValue %25 OpFunctionEnd -%70 = OpFunction %6 None %71 -%69 = OpLabel -OpBranch %79 -%79 = OpLabel -%80 = OpSDot %5 %22 %72 PackedVectorFormat4x8Bit -%81 = OpUDot %6 %73 %74 PackedVectorFormat4x8Bit -%82 = OpIAdd %6 %75 %81 -%83 = OpIAdd %6 %76 %81 -%84 = OpSDot %5 %82 %83 PackedVectorFormat4x8Bit -%85 = OpIAdd %6 %77 %81 -%86 = OpIAdd %6 %78 %81 -%87 = OpUDot %6 %85 %86 PackedVectorFormat4x8Bit -OpReturnValue %87 +%50 = OpFunction %6 None %51 +%49 = OpLabel +OpBranch %59 +%59 = OpLabel +%60 = OpSDot %5 %22 %52 PackedVectorFormat4x8Bit +%61 = OpUDot %6 %53 %54 PackedVectorFormat4x8Bit +%62 = OpIAdd %6 %55 %61 +%63 = OpIAdd %6 %56 %61 +%64 = OpSDot %5 %62 %63 PackedVectorFormat4x8Bit +%65 = OpIAdd %6 %57 %61 +%66 = OpIAdd %6 %58 %61 +%67 = OpUDot %6 %65 %66 PackedVectorFormat4x8Bit +OpReturnValue %67 OpFunctionEnd -%89 = OpFunction %2 None %90 -%88 = OpLabel -OpBranch %91 -%91 = OpLabel -%92 = OpFunctionCall %3 %8 -%93 = OpFunctionCall %5 %17 -%94 = OpFunctionCall %6 %70 +%69 = OpFunction %2 None %70 +%68 = OpLabel +OpBranch %71 +%71 = OpLabel +%72 = OpFunctionCall %3 %8 +%73 = OpFunctionCall %5 %17 +%74 = OpFunctionCall %6 %50 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/wgsl-math-functions.spvasm b/naga/tests/out/spv/wgsl-math-functions.spvasm index 9aba53bc248..aaef540cf1c 100644 --- a/naga/tests/out/spv/wgsl-math-functions.spvasm +++ b/naga/tests/out/spv/wgsl-math-functions.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 95 +; Bound: 85 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -40,76 +40,66 @@ OpMemberDecorate %15 1 Offset 16 %24 = OpConstantComposite %6 %23 %23 %23 %23 %25 = OpConstant %3 -1 %26 = OpConstantComposite %4 %25 %25 %25 %25 -%27 = OpConstantNull %7 +%27 = OpConstant %5 0 %28 = OpConstant %9 4294967295 %29 = OpConstantComposite %7 %23 %23 %30 = OpConstant %9 0 %31 = OpConstantComposite %8 %30 %30 -%32 = OpConstant %5 0 -%33 = OpConstantComposite %7 %32 %32 -%34 = OpConstant %9 32 -%35 = OpConstant %5 32 -%36 = OpConstantComposite %8 %34 %34 -%37 = OpConstantComposite %7 %35 %35 -%38 = OpConstant %9 31 -%39 = OpConstantComposite %8 %38 %38 -%40 = OpConstant %5 2 -%41 = OpConstant %3 2 -%42 = OpConstantComposite %10 %20 %41 -%43 = OpConstant %5 3 -%44 = OpConstant %5 4 -%45 = OpConstantComposite %7 %43 %44 -%46 = OpConstant %3 1.5 -%47 = OpConstantComposite %10 %46 %46 -%48 = OpConstantComposite %4 %46 %46 %46 %46 -%49 = OpConstantComposite %10 %20 %20 -%50 = OpConstantComposite %16 %20 %20 %20 -%51 = OpConstantComposite %4 %20 %20 %20 %20 -%58 = OpConstantComposite %4 %20 %20 %20 %20 -%61 = OpConstantNull %5 +%32 = OpConstantComposite %7 %27 %27 +%33 = OpConstant %9 32 +%34 = OpConstant %5 32 +%35 = OpConstantComposite %8 %33 %33 +%36 = OpConstantComposite %7 %34 %34 +%37 = OpConstant %9 31 +%38 = OpConstantComposite %8 %37 %37 +%39 = OpConstant %5 2 +%40 = OpConstant %3 2 +%41 = OpConstantComposite %10 %20 %40 +%42 = OpConstant %5 3 +%43 = OpConstant %5 4 +%44 = OpConstantComposite %7 %42 %43 +%45 = OpConstant %3 1.5 +%46 = OpConstantComposite %10 %45 %45 +%47 = OpConstantComposite %4 %45 %45 %45 %45 +%48 = OpConstantComposite %10 %20 %20 +%49 = OpConstantComposite %16 %20 %20 %20 +%50 = OpConstantComposite %4 %20 %20 %20 %20 +%57 = OpConstantComposite %4 %20 %20 %20 %20 %18 = OpFunction %2 None %19 %17 = OpLabel -OpBranch %52 -%52 = OpLabel -%53 = OpExtInst %3 %1 Degrees %20 -%54 = OpExtInst %3 %1 Radians %20 -%55 = OpExtInst %4 %1 Degrees %22 -%56 = OpExtInst %4 %1 Radians %22 -%57 = OpExtInst %4 %1 FClamp %22 %22 %58 -%59 = OpExtInst %4 %1 Refract %22 %22 %20 -%62 = OpCompositeExtract %5 %27 0 -%63 = OpCompositeExtract %5 %27 0 -%64 = OpIMul %5 %62 %63 -%65 = OpIAdd %5 %61 %64 -%66 = OpCompositeExtract %5 %27 1 -%67 = OpCompositeExtract %5 %27 1 -%68 = OpIMul %5 %66 %67 -%60 = OpIAdd %5 %65 %68 -%69 = OpExtInst %3 %1 Ldexp %20 %40 -%70 = OpExtInst %10 %1 Ldexp %42 %45 -%71 = OpExtInst %11 %1 ModfStruct %46 -%72 = OpExtInst %11 %1 ModfStruct %46 -%73 = OpCompositeExtract %3 %72 0 -%74 = OpExtInst %11 %1 ModfStruct %46 -%75 = OpCompositeExtract %3 %74 1 -%76 = OpExtInst %12 %1 ModfStruct %47 -%77 = OpExtInst %13 %1 ModfStruct %48 -%78 = OpCompositeExtract %4 %77 1 -%79 = OpCompositeExtract %3 %78 0 -%80 = OpExtInst %12 %1 ModfStruct %47 -%81 = OpCompositeExtract %10 %80 0 -%82 = OpCompositeExtract %3 %81 1 -%83 = OpExtInst %14 %1 FrexpStruct %46 -%84 = OpExtInst %14 %1 FrexpStruct %46 -%85 = OpCompositeExtract %3 %84 0 -%86 = OpExtInst %14 %1 FrexpStruct %46 -%87 = OpCompositeExtract %5 %86 1 -%88 = OpExtInst %15 %1 FrexpStruct %48 -%89 = OpCompositeExtract %6 %88 1 -%90 = OpCompositeExtract %5 %89 0 -%91 = OpQuantizeToF16 %3 %20 -%92 = OpQuantizeToF16 %10 %49 -%93 = OpQuantizeToF16 %16 %50 -%94 = OpQuantizeToF16 %4 %51 +OpBranch %51 +%51 = OpLabel +%52 = OpExtInst %3 %1 Degrees %20 +%53 = OpExtInst %3 %1 Radians %20 +%54 = OpExtInst %4 %1 Degrees %22 +%55 = OpExtInst %4 %1 Radians %22 +%56 = OpExtInst %4 %1 FClamp %22 %22 %57 +%58 = OpExtInst %4 %1 Refract %22 %22 %20 +%59 = OpExtInst %3 %1 Ldexp %20 %39 +%60 = OpExtInst %10 %1 Ldexp %41 %44 +%61 = OpExtInst %11 %1 ModfStruct %45 +%62 = OpExtInst %11 %1 ModfStruct %45 +%63 = OpCompositeExtract %3 %62 0 +%64 = OpExtInst %11 %1 ModfStruct %45 +%65 = OpCompositeExtract %3 %64 1 +%66 = OpExtInst %12 %1 ModfStruct %46 +%67 = OpExtInst %13 %1 ModfStruct %47 +%68 = OpCompositeExtract %4 %67 1 +%69 = OpCompositeExtract %3 %68 0 +%70 = OpExtInst %12 %1 ModfStruct %46 +%71 = OpCompositeExtract %10 %70 0 +%72 = OpCompositeExtract %3 %71 1 +%73 = OpExtInst %14 %1 FrexpStruct %45 +%74 = OpExtInst %14 %1 FrexpStruct %45 +%75 = OpCompositeExtract %3 %74 0 +%76 = OpExtInst %14 %1 FrexpStruct %45 +%77 = OpCompositeExtract %5 %76 1 +%78 = OpExtInst %15 %1 FrexpStruct %47 +%79 = OpCompositeExtract %6 %78 1 +%80 = OpCompositeExtract %5 %79 0 +%81 = OpQuantizeToF16 %3 %20 +%82 = OpQuantizeToF16 %10 %48 +%83 = OpQuantizeToF16 %16 %49 +%84 = OpQuantizeToF16 %4 %50 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/wgsl-functions.wgsl b/naga/tests/out/wgsl/wgsl-functions.wgsl index 69ca01cf1f8..511e4ab83f9 100644 --- a/naga/tests/out/wgsl/wgsl-functions.wgsl +++ b/naga/tests/out/wgsl/wgsl-functions.wgsl @@ -12,8 +12,7 @@ fn test_integer_dot_product() -> i32 { let a_3_ = vec3(1u); let b_3_ = vec3(1u); let c_3_ = dot(a_3_, b_3_); - let c_4_ = dot(vec4(4i), vec4(2i)); - return c_4_; + return 32i; } fn test_packed_integer_dot_product() -> u32 { diff --git a/naga/tests/out/wgsl/wgsl-math-functions.wgsl b/naga/tests/out/wgsl/wgsl-math-functions.wgsl index 082afb933f3..20570ab5b2a 100644 --- a/naga/tests/out/wgsl/wgsl-math-functions.wgsl +++ b/naga/tests/out/wgsl/wgsl-math-functions.wgsl @@ -9,7 +9,6 @@ fn main() { let g = refract(v, v, 1f); let sign_b = vec4(-1i, -1i, -1i, -1i); let sign_d = vec4(-1f, -1f, -1f, -1f); - let const_dot = dot(vec2(), vec2()); let flb_b = vec2(-1i, -1i); let flb_c = vec2(0u, 0u); let ftb_c = vec2(0i, 0i);