From a10d042fa59953b22af4f32e6117e3762ce2ce83 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Sun, 10 Dec 2023 16:51:22 -0800 Subject: [PATCH] [naga wgsl-in] Support abstract operands to binary operators. --- CHANGELOG.md | 5 +- naga/src/front/wgsl/error.rs | 20 ++++ naga/src/front/wgsl/lower/conversion.rs | 80 ++++++++++++-- naga/src/front/wgsl/lower/mod.rs | 52 ++++++++- naga/src/proc/constant_evaluator.rs | 54 ++++++++-- naga/tests/in/abstract-types-operators.wgsl | 45 ++++++++ .../out/msl/abstract-types-operators.msl | 73 +++++++++++++ .../out/spv/abstract-types-operators.spvasm | 101 ++++++++++++++++++ .../tests/out/spv/debug-symbol-terrain.spvasm | 24 +++-- .../out/wgsl/abstract-types-operators.wgsl | 68 ++++++++++++ naga/tests/snapshots.rs | 4 + 11 files changed, 498 insertions(+), 28 deletions(-) create mode 100644 naga/tests/in/abstract-types-operators.wgsl create mode 100644 naga/tests/out/msl/abstract-types-operators.msl create mode 100644 naga/tests/out/spv/abstract-types-operators.spvasm create mode 100644 naga/tests/out/wgsl/abstract-types-operators.wgsl diff --git a/CHANGELOG.md b/CHANGELOG.md index 6aacbbec3e2..09c5485a6bd 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -68,7 +68,10 @@ This feature allowed you to call `global_id` on any wgpu opaque handle to get a #### Naga -- Naga'sn WGSL front and back ends now have experimental support for 64-bit floating-point literals: `1.0lf` denotes an `f64` value. There has been experimental support for an `f64` type for a while, but until now there was no syntax for writing literals with that type. As before, Naga module validation rejects `f64` values unless `naga::valid::Capabilities::FLOAT64` is requested. By @jimblandy in [#4747](https://github.com/gfx-rs/wgpu/pull/4747). +- Naga's WGSL front end now allows binary operators to produce values with abstract types, rather than concretizing thir operands. By @jimblandy in [#4850](https://github.com/gfx-rs/wgpu/pull/4850). + +- Naga's WGSL front and back ends now have experimental support for 64-bit floating-point literals: `1.0lf` denotes an `f64` value. There has been experimental support for an `f64` type for a while, but until now there was no syntax for writing literals with that type. As before, Naga module validation rejects `f64` values unless `naga::valid::Capabilities::FLOAT64` is requested. By @jimblandy in [#4747](https://github.com/gfx-rs/wgpu/pull/4747). + - Naga constant evaluation can now process binary operators whose operands are both vectors. By @jimblandy in [#4861](https://github.com/gfx-rs/wgpu/pull/4861). ### Changes diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs index fffa28817b0..5b3657f1f1b 100644 --- a/naga/src/front/wgsl/error.rs +++ b/naga/src/front/wgsl/error.rs @@ -257,6 +257,12 @@ pub enum Error<'a> { source_span: Span, source_type: String, }, + AutoConversionLeafScalar { + dest_span: Span, + dest_scalar: String, + source_span: Span, + source_type: String, + }, ConcretizationFailed { expr_span: Span, expr_type: String, @@ -738,6 +744,20 @@ impl<'a> Error<'a> { ], notes: vec![], }, + Error::AutoConversionLeafScalar { dest_span, ref dest_scalar, source_span, ref source_type } => ParseError { + message: format!("automatic conversions cannot convert elements of `{source_type}` to `{dest_scalar}`"), + labels: vec![ + ( + dest_span, + format!("a value with elements of type {dest_scalar} is required here").into(), + ), + ( + source_span, + format!("this expression has type {source_type}").into(), + ) + ], + notes: vec![], + }, Error::ConcretizationFailed { expr_span, ref expr_type, ref scalar, ref inner } => ParseError { message: format!("failed to convert expression to a concrete type: {}", inner), labels: vec![ diff --git a/naga/src/front/wgsl/lower/conversion.rs b/naga/src/front/wgsl/lower/conversion.rs index 7378f0f63e5..2a2690f096a 100644 --- a/naga/src/front/wgsl/lower/conversion.rs +++ b/naga/src/front/wgsl/lower/conversion.rs @@ -51,21 +51,80 @@ impl<'source, 'temp, 'out> super::ExpressionContext<'source, 'temp, 'out> { } }; - let converted = if let crate::TypeInner::Array { .. } = *goal_inner { - let span = self.get_expression_span(expr); + self.convert_leaf_scalar(expr, expr_span, goal_scalar) + } + + /// Try to convert `expr`'s leaf scalar to `goal` using automatic conversions. + /// + /// If no conversions are necessary, return `expr` unchanged. + /// + /// If automatic conversions cannot convert `expr` to `goal_scalar`, return + /// an [`AutoConversionLeafScalar`] 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<_>`. + /// + /// [`AutoConversionLeafScalar`]: super::Error::AutoConversionLeafScalar + pub fn try_automatic_conversion_for_leaf_scalar( + &mut self, + expr: Handle, + goal_scalar: crate::Scalar, + goal_span: Span, + ) -> Result, super::Error<'source>> { + let expr_span = self.get_expression_span(expr); + let expr_resolution = super::resolve!(self, expr); + let types = &self.module.types; + let expr_inner = expr_resolution.inner_with(types); + + let make_error = || { + let gctx = &self.module.to_ctx(); + let source_type = expr_resolution.to_wgsl(gctx); + super::Error::AutoConversionLeafScalar { + dest_span: goal_span, + dest_scalar: goal_scalar.to_wgsl(), + source_span: expr_span, + source_type, + } + }; + + let expr_scalar = match expr_inner.scalar() { + Some(scalar) => scalar, + None => return Err(make_error()), + }; + + if expr_scalar == goal_scalar { + return Ok(expr); + } + + if !expr_scalar.automatically_converts_to(goal_scalar) { + return Err(make_error()); + } + + assert!(expr_scalar.is_abstract()); + + self.convert_leaf_scalar(expr, expr_span, goal_scalar) + } + + fn convert_leaf_scalar( + &mut self, + expr: Handle, + expr_span: Span, + goal_scalar: crate::Scalar, + ) -> Result, super::Error<'source>> { + let expr_inner = super::resolve_inner!(self, expr); + if let crate::TypeInner::Array { .. } = *expr_inner { self.as_const_evaluator() - .cast_array(expr, goal_scalar, span) - .map_err(|err| super::Error::ConstantEvaluatorError(err, span))? + .cast_array(expr, goal_scalar, expr_span) + .map_err(|err| super::Error::ConstantEvaluatorError(err, expr_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) + self.append_expression(cast, expr_span) + } } /// Try to convert `exprs` to `goal_ty` using WGSL's automatic conversions. @@ -428,6 +487,11 @@ impl crate::Scalar { } } + /// Return `true` if automatic conversions will covert `self` to `goal`. + pub fn automatically_converts_to(self, goal: Self) -> bool { + self.automatic_conversion_combine(goal) == Some(goal) + } + const fn concretize(self) -> Self { use crate::ScalarKind as Sk; match self.kind { diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 6486e6cf6b4..a4a49da0612 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1602,11 +1602,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { return Ok(Typed::Reference(pointer)); } ast::Expression::Binary { op, left, right } => { - // Load both operands. - let mut left = self.expression(left, ctx)?; - let mut right = self.expression(right, ctx)?; - ctx.binary_op_splat(op, &mut left, &mut right)?; - Typed::Plain(crate::Expression::Binary { op, left, right }) + self.binary(op, left, right, span, ctx)? } ast::Expression::Call { ref function, @@ -1737,6 +1733,52 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { expr.try_map(|handle| ctx.append_expression(handle, span)) } + fn binary( + &mut self, + op: crate::BinaryOperator, + left: Handle>, + right: Handle>, + span: Span, + ctx: &mut ExpressionContext<'source, '_, '_>, + ) -> Result, Error<'source>> { + // Load both operands. + let mut left = self.expression_for_abstract(left, ctx)?; + let mut right = self.expression_for_abstract(right, ctx)?; + + // Convert `scalar op vector` to `vector op vector` by introducing + // `Splat` expressions. + ctx.binary_op_splat(op, &mut left, &mut right)?; + + // Apply automatic conversions. + match op { + // Shift operators require the right operand to be `u32` or + // `vecN`. We can let the validator sort out vector length + // issues, but the right operand must be, or convert to, a u32 leaf + // scalar. + crate::BinaryOperator::ShiftLeft | crate::BinaryOperator::ShiftRight => { + right = + ctx.try_automatic_conversion_for_leaf_scalar(right, crate::Scalar::U32, span)?; + } + + // All other operators follow the same pattern: reconcile the + // scalar leaf types. If there's no reconciliation possible, + // leave the expressions as they are: validation will report the + // problem. + _ => { + ctx.grow_types(left)?; + ctx.grow_types(right)?; + if let Ok(consensus_scalar) = + ctx.automatic_conversion_consensus([left, right].iter()) + { + ctx.convert_to_leaf_scalar(&mut left, consensus_scalar)?; + ctx.convert_to_leaf_scalar(&mut right, consensus_scalar)?; + } + } + } + + Ok(Typed::Plain(crate::Expression::Binary { op, left, right })) + } + /// Generate Naga IR for call expressions and statements, and type /// constructor expressions. /// diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 5c71bc167db..81bcb35e58f 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -141,7 +141,7 @@ pub enum ConstantEvaluatorError { InvalidAccessIndexTy, #[error("Constants don't support array length expressions")] ArrayLength, - #[error("Cannot cast type `{from}` to `{to}`")] + #[error("Cannot cast scalar components of expression `{from}` to type `{to}`")] InvalidCastArg { from: String, to: String }, #[error("Cannot apply the unary op to the argument")] InvalidUnaryOpArg, @@ -989,15 +989,11 @@ impl<'a> ConstantEvaluator<'a> { let expr = self.eval_zero_value(expr, span)?; let make_error = || -> Result<_, ConstantEvaluatorError> { - let ty = self.resolve_type(expr)?; + let from = format!("{:?} {:?}", expr, self.expressions[expr]); - #[cfg(feature = "wgsl-in")] - let from = ty.to_wgsl(&self.to_ctx()); #[cfg(feature = "wgsl-in")] let to = target.to_wgsl(); - #[cfg(not(feature = "wgsl-in"))] - let from = format!("{ty:?}"); #[cfg(not(feature = "wgsl-in"))] let to = format!("{target:?}"); @@ -1325,6 +1321,47 @@ impl<'a> ConstantEvaluator<'a> { BinaryOperator::Modulo => a % b, _ => return Err(ConstantEvaluatorError::InvalidBinaryOpArgs), }), + (Literal::AbstractInt(a), Literal::AbstractInt(b)) => { + Literal::AbstractInt(match op { + BinaryOperator::Add => a.checked_add(b).ok_or_else(|| { + ConstantEvaluatorError::Overflow("addition".into()) + })?, + BinaryOperator::Subtract => a.checked_sub(b).ok_or_else(|| { + ConstantEvaluatorError::Overflow("subtraction".into()) + })?, + BinaryOperator::Multiply => a.checked_mul(b).ok_or_else(|| { + ConstantEvaluatorError::Overflow("multiplication".into()) + })?, + BinaryOperator::Divide => a.checked_div(b).ok_or_else(|| { + if b == 0 { + ConstantEvaluatorError::DivisionByZero + } else { + ConstantEvaluatorError::Overflow("division".into()) + } + })?, + BinaryOperator::Modulo => a.checked_rem(b).ok_or_else(|| { + if b == 0 { + ConstantEvaluatorError::RemainderByZero + } else { + ConstantEvaluatorError::Overflow("remainder".into()) + } + })?, + BinaryOperator::And => a & b, + BinaryOperator::ExclusiveOr => a ^ b, + BinaryOperator::InclusiveOr => a | b, + _ => return Err(ConstantEvaluatorError::InvalidBinaryOpArgs), + }) + } + (Literal::AbstractFloat(a), Literal::AbstractFloat(b)) => { + Literal::AbstractFloat(match op { + BinaryOperator::Add => a + b, + BinaryOperator::Subtract => a - b, + BinaryOperator::Multiply => a * b, + BinaryOperator::Divide => a / b, + BinaryOperator::Modulo => a % b, + _ => return Err(ConstantEvaluatorError::InvalidBinaryOpArgs), + }) + } (Literal::Bool(a), Literal::Bool(b)) => Literal::Bool(match op { BinaryOperator::LogicalAnd => a && b, BinaryOperator::LogicalOr => a || b, @@ -1550,7 +1587,10 @@ impl<'a> ConstantEvaluator<'a> { }; Tr::Value(TypeInner::Vector { scalar, size }) } - _ => return Err(ConstantEvaluatorError::SubexpressionsAreNotConstant), + _ => { + log::debug!("resolve_type: SubexpressionsAreNotConstant"); + return Err(ConstantEvaluatorError::SubexpressionsAreNotConstant); + } }; Ok(resolution) diff --git a/naga/tests/in/abstract-types-operators.wgsl b/naga/tests/in/abstract-types-operators.wgsl new file mode 100644 index 00000000000..2c06ef43e49 --- /dev/null +++ b/naga/tests/in/abstract-types-operators.wgsl @@ -0,0 +1,45 @@ +const plus_fafaf: f32 = 1.0 + 2.0; +const plus_fafai: f32 = 1.0 + 2; +const plus_faf_f: f32 = 1.0 + 2f; +const plus_faiaf: f32 = 1 + 2.0; +const plus_faiai: f32 = 1 + 2; +const plus_fai_f: f32 = 1 + 2f; +const plus_f_faf: f32 = 1f + 2.0; +const plus_f_fai: f32 = 1f + 2; +const plus_f_f_f: f32 = 1f + 2f; + +const plus_iaiai: i32 = 1 + 2; +const plus_iai_i: i32 = 1 + 2i; +const plus_i_iai: i32 = 1i + 2; +const plus_i_i_i: i32 = 1i + 2i; + +const plus_uaiai: u32 = 1 + 2; +const plus_uai_u: u32 = 1 + 2u; +const plus_u_uai: u32 = 1u + 2; +const plus_u_u_u: u32 = 1u + 2u; + +fn runtime_values() { + var f: f32 = 42; + var i: i32 = 43; + var u: u32 = 44; + + var plus_fafaf: f32 = 1.0 + 2.0; + var plus_fafai: f32 = 1.0 + 2; + var plus_faf_f: f32 = 1.0 + f; + var plus_faiaf: f32 = 1 + 2.0; + var plus_faiai: f32 = 1 + 2; + var plus_fai_f: f32 = 1 + f; + var plus_f_faf: f32 = f + 2.0; + var plus_f_fai: f32 = f + 2; + var plus_f_f_f: f32 = f + f; + + var plus_iaiai: i32 = 1 + 2; + var plus_iai_i: i32 = 1 + i; + var plus_i_iai: i32 = i + 2; + var plus_i_i_i: i32 = i + i; + + var plus_uaiai: u32 = 1 + 2; + var plus_uai_u: u32 = 1 + u; + var plus_u_uai: u32 = u + 2; + var plus_u_u_u: u32 = u + u; +} diff --git a/naga/tests/out/msl/abstract-types-operators.msl b/naga/tests/out/msl/abstract-types-operators.msl new file mode 100644 index 00000000000..b0185731523 --- /dev/null +++ b/naga/tests/out/msl/abstract-types-operators.msl @@ -0,0 +1,73 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +constant float plus_fafaf_1 = 3.0; +constant float plus_fafai_1 = 3.0; +constant float plus_faf_f_1 = 3.0; +constant float plus_faiaf_1 = 3.0; +constant float plus_faiai_1 = 3.0; +constant float plus_fai_f_1 = 3.0; +constant float plus_f_faf_1 = 3.0; +constant float plus_f_fai_1 = 3.0; +constant float plus_f_f_f_1 = 3.0; +constant int plus_iaiai_1 = 3; +constant int plus_iai_i_1 = 3; +constant int plus_i_iai_1 = 3; +constant int plus_i_i_i_1 = 3; +constant uint plus_uaiai_1 = 3u; +constant uint plus_uai_u_1 = 3u; +constant uint plus_u_uai_1 = 3u; +constant uint plus_u_u_u_1 = 3u; + +void runtime_values( +) { + float f = 42.0; + int i = 43; + uint u = 44u; + float plus_fafaf = 3.0; + float plus_fafai = 3.0; + float plus_faf_f = {}; + float plus_faiaf = 3.0; + float plus_faiai = 3.0; + float plus_fai_f = {}; + float plus_f_faf = {}; + float plus_f_fai = {}; + float plus_f_f_f = {}; + int plus_iaiai = 3; + int plus_iai_i = {}; + int plus_i_iai = {}; + int plus_i_i_i = {}; + uint plus_uaiai = 3u; + uint plus_uai_u = {}; + uint plus_u_uai = {}; + uint plus_u_u_u = {}; + float _e8 = f; + plus_faf_f = 1.0 + _e8; + float _e14 = f; + plus_fai_f = 1.0 + _e14; + float _e18 = f; + plus_f_faf = _e18 + 2.0; + float _e22 = f; + plus_f_fai = _e22 + 2.0; + float _e26 = f; + float _e27 = f; + plus_f_f_f = _e26 + _e27; + int _e31 = i; + plus_iai_i = 1 + _e31; + int _e35 = i; + plus_i_iai = _e35 + 2; + int _e39 = i; + int _e40 = i; + plus_i_i_i = _e39 + _e40; + uint _e44 = u; + plus_uai_u = 1u + _e44; + uint _e48 = u; + plus_u_uai = _e48 + 2u; + uint _e52 = u; + uint _e53 = u; + plus_u_u_u = _e52 + _e53; + return; +} diff --git a/naga/tests/out/spv/abstract-types-operators.spvasm b/naga/tests/out/spv/abstract-types-operators.spvasm new file mode 100644 index 00000000000..ea9f8582951 --- /dev/null +++ b/naga/tests/out/spv/abstract-types-operators.spvasm @@ -0,0 +1,101 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 81 +OpCapability Shader +OpCapability Linkage +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +%2 = OpTypeVoid +%3 = OpTypeFloat 32 +%4 = OpTypeInt 32 1 +%5 = OpTypeInt 32 0 +%6 = OpConstant %3 3.0 +%7 = OpConstant %4 3 +%8 = OpConstant %5 3 +%11 = OpTypeFunction %2 +%12 = OpConstant %3 42.0 +%13 = OpConstant %4 43 +%14 = OpConstant %5 44 +%15 = OpConstant %3 1.0 +%16 = OpConstant %3 2.0 +%17 = OpConstant %4 1 +%18 = OpConstant %4 2 +%19 = OpConstant %5 1 +%20 = OpConstant %5 2 +%22 = OpTypePointer Function %3 +%24 = OpTypePointer Function %4 +%26 = OpTypePointer Function %5 +%30 = OpConstantNull %3 +%34 = OpConstantNull %3 +%36 = OpConstantNull %3 +%38 = OpConstantNull %3 +%40 = OpConstantNull %3 +%43 = OpConstantNull %4 +%45 = OpConstantNull %4 +%47 = OpConstantNull %4 +%50 = OpConstantNull %5 +%52 = OpConstantNull %5 +%54 = OpConstantNull %5 +%10 = OpFunction %2 None %11 +%9 = OpLabel +%53 = OpVariable %26 Function %54 +%48 = OpVariable %26 Function %8 +%42 = OpVariable %24 Function %43 +%37 = OpVariable %22 Function %38 +%32 = OpVariable %22 Function %6 +%28 = OpVariable %22 Function %6 +%23 = OpVariable %24 Function %13 +%51 = OpVariable %26 Function %52 +%46 = OpVariable %24 Function %47 +%41 = OpVariable %24 Function %7 +%35 = OpVariable %22 Function %36 +%31 = OpVariable %22 Function %6 +%27 = OpVariable %22 Function %6 +%21 = OpVariable %22 Function %12 +%49 = OpVariable %26 Function %50 +%44 = OpVariable %24 Function %45 +%39 = OpVariable %22 Function %40 +%33 = OpVariable %22 Function %34 +%29 = OpVariable %22 Function %30 +%25 = OpVariable %26 Function %14 +OpBranch %55 +%55 = OpLabel +%56 = OpLoad %3 %21 +%57 = OpFAdd %3 %15 %56 +OpStore %29 %57 +%58 = OpLoad %3 %21 +%59 = OpFAdd %3 %15 %58 +OpStore %33 %59 +%60 = OpLoad %3 %21 +%61 = OpFAdd %3 %60 %16 +OpStore %35 %61 +%62 = OpLoad %3 %21 +%63 = OpFAdd %3 %62 %16 +OpStore %37 %63 +%64 = OpLoad %3 %21 +%65 = OpLoad %3 %21 +%66 = OpFAdd %3 %64 %65 +OpStore %39 %66 +%67 = OpLoad %4 %23 +%68 = OpIAdd %4 %17 %67 +OpStore %42 %68 +%69 = OpLoad %4 %23 +%70 = OpIAdd %4 %69 %18 +OpStore %44 %70 +%71 = OpLoad %4 %23 +%72 = OpLoad %4 %23 +%73 = OpIAdd %4 %71 %72 +OpStore %46 %73 +%74 = OpLoad %5 %25 +%75 = OpIAdd %5 %19 %74 +OpStore %49 %75 +%76 = OpLoad %5 %25 +%77 = OpIAdd %5 %76 %20 +OpStore %51 %77 +%78 = OpLoad %5 %25 +%79 = OpLoad %5 %25 +%80 = OpIAdd %5 %78 %79 +OpStore %53 %80 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/debug-symbol-terrain.spvasm b/naga/tests/out/spv/debug-symbol-terrain.spvasm index 36fae9d60a8..623b8dc2c1b 100644 --- a/naga/tests/out/spv/debug-symbol-terrain.spvasm +++ b/naga/tests/out/spv/debug-symbol-terrain.spvasm @@ -555,9 +555,9 @@ OpDecorate %582 Location 0 %79 = OpConstantComposite %4 %78 %78 %78 %80 = OpConstantComposite %4 %74 %74 %74 %81 = OpConstant %5 2.0 -%82 = OpConstant %5 1.7928429 -%83 = OpConstant %5 0.85373473 -%84 = OpConstantComposite %4 %82 %82 %82 +%82 = OpConstant %5 0.85373473 +%83 = OpConstant %5 1.7928429 +%84 = OpConstantComposite %4 %83 %83 %83 %85 = OpConstant %5 130.0 %87 = OpTypePointer Function %6 %88 = OpConstantNull %6 @@ -659,6 +659,7 @@ OpBranch %60 %60 = OpLabel OpLine %3 10 52 %61 = OpVectorTimesScalar %4 %52 %55 +OpLine %3 10 63 OpLine %3 10 50 %62 = OpFAdd %4 %61 %57 %63 = OpFMul %4 %62 %52 @@ -755,6 +756,7 @@ OpLine %3 21 28 %154 = OpVectorShuffle %6 %153 %153 2 3 %155 = OpDot %5 %152 %154 %156 = OpCompositeConstruct %4 %145 %150 %155 +OpLine %3 21 28 %157 = OpFSub %4 %79 %156 OpLine %3 21 24 %158 = OpExtInst %4 %1 FMax %157 %80 @@ -772,17 +774,21 @@ OpLine %3 23 9 %164 = OpFMul %4 %162 %163 OpLine %3 23 5 OpStore %94 %164 -OpLine %3 24 13 +OpLine %3 24 18 %165 = OpVectorShuffle %4 %73 %73 3 3 3 %166 = OpFMul %4 %144 %165 %167 = OpExtInst %4 %1 Fract %166 +OpLine %3 24 13 %168 = OpVectorTimesScalar %4 %167 %81 +OpLine %3 24 37 OpLine %3 24 13 %169 = OpFSub %4 %168 %57 OpLine %3 25 13 %170 = OpExtInst %4 %1 FAbs %169 +OpLine %3 25 22 OpLine %3 25 13 %171 = OpFSub %4 %170 %79 +OpLine %3 26 24 OpLine %3 26 14 %172 = OpFAdd %4 %169 %79 %173 = OpExtInst %4 %1 Floor %172 @@ -790,11 +796,13 @@ OpLine %3 27 14 %174 = OpFSub %4 %169 %173 OpLine %3 1 1 %175 = OpLoad %4 %94 -OpLine %3 28 9 +OpLine %3 28 53 %176 = OpFMul %4 %174 %174 %177 = OpFMul %4 %171 %171 %178 = OpFAdd %4 %176 %177 -%179 = OpVectorTimesScalar %4 %178 %83 +OpLine %3 28 14 +%179 = OpVectorTimesScalar %4 %178 %82 +OpLine %3 28 9 %180 = OpFSub %4 %84 %179 %181 = OpFMul %4 %175 %180 OpLine %3 28 5 @@ -817,9 +825,10 @@ OpLine %3 29 13 %196 = OpFMul %6 %193 %195 %197 = OpFAdd %6 %192 %196 %198 = OpCompositeConstruct %4 %188 %197 -OpLine %3 30 12 +OpLine %3 30 19 %199 = OpLoad %4 %94 %200 = OpDot %5 %199 %198 +OpLine %3 30 12 %201 = OpFMul %5 %85 %200 OpReturnValue %201 OpFunctionEnd @@ -1149,6 +1158,7 @@ OpLine %3 163 14 %429 = OpCompositeConstruct %6 %424 %428 OpLine %3 165 30 %430 = OpVectorTimesScalar %6 %429 %81 +OpLine %3 165 30 %431 = OpFAdd %6 %419 %430 OpLine %3 165 20 %432 = OpCompositeConstruct %7 %431 %74 %56 diff --git a/naga/tests/out/wgsl/abstract-types-operators.wgsl b/naga/tests/out/wgsl/abstract-types-operators.wgsl new file mode 100644 index 00000000000..00e858f9e9a --- /dev/null +++ b/naga/tests/out/wgsl/abstract-types-operators.wgsl @@ -0,0 +1,68 @@ +const plus_fafaf_1: f32 = 3.0; +const plus_fafai_1: f32 = 3.0; +const plus_faf_f_1: f32 = 3.0; +const plus_faiaf_1: f32 = 3.0; +const plus_faiai_1: f32 = 3.0; +const plus_fai_f_1: f32 = 3.0; +const plus_f_faf_1: f32 = 3.0; +const plus_f_fai_1: f32 = 3.0; +const plus_f_f_f_1: f32 = 3.0; +const plus_iaiai_1: i32 = 3i; +const plus_iai_i_1: i32 = 3i; +const plus_i_iai_1: i32 = 3i; +const plus_i_i_i_1: i32 = 3i; +const plus_uaiai_1: u32 = 3u; +const plus_uai_u_1: u32 = 3u; +const plus_u_uai_1: u32 = 3u; +const plus_u_u_u_1: u32 = 3u; + +fn runtime_values() { + var f: f32 = 42.0; + var i: i32 = 43i; + var u: u32 = 44u; + var plus_fafaf: f32 = 3.0; + var plus_fafai: f32 = 3.0; + var plus_faf_f: f32; + var plus_faiaf: f32 = 3.0; + var plus_faiai: f32 = 3.0; + var plus_fai_f: f32; + var plus_f_faf: f32; + var plus_f_fai: f32; + var plus_f_f_f: f32; + var plus_iaiai: i32 = 3i; + var plus_iai_i: i32; + var plus_i_iai: i32; + var plus_i_i_i: i32; + var plus_uaiai: u32 = 3u; + var plus_uai_u: u32; + var plus_u_uai: u32; + var plus_u_u_u: u32; + + let _e8 = f; + plus_faf_f = (1.0 + _e8); + let _e14 = f; + plus_fai_f = (1.0 + _e14); + let _e18 = f; + plus_f_faf = (_e18 + 2.0); + let _e22 = f; + plus_f_fai = (_e22 + 2.0); + let _e26 = f; + let _e27 = f; + plus_f_f_f = (_e26 + _e27); + let _e31 = i; + plus_iai_i = (1i + _e31); + let _e35 = i; + plus_i_iai = (_e35 + 2i); + let _e39 = i; + let _e40 = i; + plus_i_i_i = (_e39 + _e40); + let _e44 = u; + plus_uai_u = (1u + _e44); + let _e48 = u; + plus_u_uai = (_e48 + 2u); + let _e52 = u; + let _e53 = u; + plus_u_u_u = (_e52 + _e53); + return; +} + diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index c201b206a2e..6b934de55ba 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -798,6 +798,10 @@ fn convert_wgsl() { "abstract-types-var", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, ), + ( + "abstract-types-operators", + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, + ), ]; for &(name, targets) in inputs.iter() {