diff --git a/CHANGELOG.md b/CHANGELOG.md index c4361ea48a..9bfea9f358 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -37,6 +37,30 @@ Bottom level categories: - Hal --> +## Unreleased + +### New features + +- Many numeric built-ins have had a constant evaluation implementation added for them, which allows them to be used in a `const` context: + - [#4879](https://github.com/gfx-rs/wgpu/pull/4879) by @ErichDonGubler: + - `abs` + - `acos` + - `acosh` + - `asin` + - `asinh` + - `atan` + - `atanh` + - `cos` + - `cosh` + - `round` + - `saturate` + - `sin` + - `sinh` + - `sqrt` + - `step` + - `tan` + - `tanh` + ## v0.19.0 (2024-01-17) This release includes: diff --git a/Cargo.lock b/Cargo.lock index c8d235b287..240dec3584 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -2055,6 +2055,7 @@ name = "naga" version = "0.19.0" dependencies = [ "arbitrary", + "arrayvec 0.7.4", "bincode", "bit-set", "bitflags 2.4.1", diff --git a/naga/Cargo.toml b/naga/Cargo.toml index a13e4f9196..8e53d457f2 100644 --- a/naga/Cargo.toml +++ b/naga/Cargo.toml @@ -60,6 +60,7 @@ petgraph = { version = "0.6", optional = true } pp-rs = { version = "0.2.1", optional = true } hexf-parse = { version = "0.2.1", optional = true } unicode-xid = { version = "0.2.3", optional = true } +arrayvec.workspace = true [target.'cfg(not(target_arch = "wasm32"))'.dev-dependencies] criterion = { version = "0.5", features = [] } diff --git a/naga/src/lib.rs b/naga/src/lib.rs index b27ebc6764..bfd8359d88 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -458,6 +458,10 @@ pub enum VectorSize { Quad = 4, } +impl VectorSize { + const MAX: usize = Self::Quad as u8 as usize; +} + /// Primitive type for a scalar. #[repr(u8)] #[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)] diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 82efeece4e..5cffdba86e 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1,9 +1,227 @@ +use std::iter; + +use arrayvec::ArrayVec; + use crate::{ arena::{Arena, Handle, UniqueArena}, ArraySize, BinaryOperator, Constant, Expression, Literal, ScalarKind, Span, Type, TypeInner, UnaryOperator, }; +/// A macro that allows dollar signs (`$`) to be emitted by other macros. Useful for generating +/// `macro_rules!` items that, in turn, emit their own `macro_rules!` items. +/// +/// Technique stolen directly from +/// . +macro_rules! with_dollar_sign { + ($($body:tt)*) => { + macro_rules! __with_dollar_sign { $($body)* } + __with_dollar_sign!($); + } +} + +macro_rules! gen_component_wise_extractor { + ( + $ident:ident -> $target:ident, + literals: [$( $literal:ident => $mapping:ident: $ty:ident ),+ $(,)?], + scalar_kinds: [$( $scalar_kind:ident ),* $(,)?], + ) => { + /// A subset of [`Literal`]s intended to be used for implementing numeric built-ins. + enum $target { + $( + #[doc = concat!( + "Maps to [`Literal::", + stringify!($mapping), + "`]", + )] + $mapping([$ty; N]), + )+ + } + + impl From<$target<1>> for Expression { + fn from(value: $target<1>) -> Self { + match value { + $( + $target::$mapping([value]) => { + Expression::Literal(Literal::$literal(value)) + } + )+ + } + } + } + + #[doc = concat!( + "Attempts to evaluate multiple `exprs` as a combined [`", + stringify!($target), + "`] to pass to `handler`. ", + )] + /// If `exprs` are vectors of the same length, `handler` is called for each corresponding + /// component of each vector. + /// + /// `handler`'s output is registered as a new expression. If `exprs` are vectors of the + /// same length, a new vector expression is registered, composed of each component emitted + /// by `handler`. + fn $ident( + eval: &mut ConstantEvaluator<'_>, + span: Span, + exprs: [Handle; N], + mut handler: F, + ) -> Result, ConstantEvaluatorError> + where + $target: Into, + F: FnMut($target) -> Result<$target, ConstantEvaluatorError> + Clone, + { + assert!(N > 0); + let err = ConstantEvaluatorError::InvalidMathArg; + let mut exprs = exprs.into_iter(); + + macro_rules! sanitize { + ($expr:expr) => { + eval.eval_zero_value_and_splat($expr, span) + .map(|expr| &eval.expressions[expr]) + }; + } + + let new_expr = match sanitize!(exprs.next().unwrap())? { + $( + &Expression::Literal(Literal::$literal(x)) => iter::once(Ok(x)) + .chain(exprs.map(|expr| { + sanitize!(expr).and_then(|expr| match expr { + &Expression::Literal(Literal::$literal(x)) => Ok(x), + _ => Err(err.clone()), + }) + })) + .collect::, _>>() + .map(|a| a.into_inner().unwrap()) + .map($target::$mapping) + .and_then(|comps| Ok(handler(comps)?.into())), + )+ + &Expression::Compose { ty, ref components } => match &eval.types[ty].inner { + &TypeInner::Vector { size: _, scalar } => match scalar.kind { + $(ScalarKind::$scalar_kind)|* => { + let first_ty = ty; + let mut component_groups = + ArrayVec::, N>::new(); + component_groups.push(crate::proc::flatten_compose( + first_ty, + components, + eval.expressions, + eval.types, + ).collect()); + component_groups.extend( + exprs + .map(|expr| { + sanitize!(expr).and_then(|expr| match expr { + &Expression::Compose { ty, ref components } + if &eval.types[ty].inner + == &eval.types[first_ty].inner => + { + Ok(crate::proc::flatten_compose( + ty, + components, + eval.expressions, + eval.types, + ).collect()) + } + _ => Err(err.clone()), + }) + }) + .collect::, _>>( + )?, + ); + let component_groups = component_groups.into_inner().unwrap(); + let mut new_components = + ArrayVec::<_, { crate::VectorSize::MAX }>::new(); + for idx in 0..N { + let group = component_groups + .iter() + .map(|cs| cs[idx]) + .collect::>() + .into_inner() + .unwrap(); + new_components.push($ident( + eval, + span, + group, + handler.clone(), + )?); + } + Ok(Expression::Compose { + ty: first_ty, + components: new_components.into_iter().collect(), + }) + } + _ => return Err(err), + }, + _ => return Err(err), + }, + _ => return Err(err), + }?; + eval.register_evaluated_expr(new_expr, span) + } + + with_dollar_sign! { + ($d:tt) => { + #[allow(unused)] + #[doc = concat!( + "A convenience macro for using the same RHS for each [`", + stringify!($target), + "`] variant in a call to [`", + stringify!($ident), + "`].", + )] + macro_rules! $ident { + ( + $eval:expr, + $span:expr, + [$d ($d expr:expr),+ $d (,)?], + |$d ($d arg:ident),+| $d tt:tt + ) => { + $ident($eval, $span, [$d ($d expr),+], |args| match args { + $( + $target::$mapping([$d ($d arg),+]) => { + let res = $d tt; + Result::map(res, $target::$mapping) + }, + )+ + }) + }; + } + }; + } + }; +} + +gen_component_wise_extractor! { + component_wise_scalar -> Scalar, + literals: [ + AbstractFloat => AbstractFloat: f64, + F32 => F32: f32, + AbstractInt => AbstractInt: i64, + U32 => U32: u32, + I32 => I32: i32, + ], + scalar_kinds: [ + Float, + AbstractFloat, + Sint, + Uint, + AbstractInt, + ], +} + +gen_component_wise_extractor! { + component_wise_float -> Float, + literals: [ + AbstractFloat => Abstract: f64, + F32 => F32: f32, + ], + scalar_kinds: [ + Float, + AbstractFloat, + ], +} + #[derive(Debug)] enum Behavior { Wgsl, @@ -592,186 +810,111 @@ impl<'a> ConstantEvaluator<'a> { } match fun { - crate::MathFunction::Pow => self.math_pow(arg, arg1.unwrap(), span), - crate::MathFunction::Clamp => self.math_clamp(arg, arg1.unwrap(), arg2.unwrap(), span), - fun => Err(ConstantEvaluatorError::NotImplemented(format!( - "{fun:?} built-in function" - ))), - } - } - - fn math_pow( - &mut self, - e1: Handle, - e2: Handle, - span: Span, - ) -> Result, ConstantEvaluatorError> { - let e1 = self.eval_zero_value_and_splat(e1, span)?; - let e2 = self.eval_zero_value_and_splat(e2, span)?; - - let expr = match (&self.expressions[e1], &self.expressions[e2]) { - (&Expression::Literal(Literal::F32(a)), &Expression::Literal(Literal::F32(b))) => { - Expression::Literal(Literal::F32(a.powf(b))) + crate::MathFunction::Abs => { + component_wise_scalar(self, span, [arg], |args| match args { + Scalar::AbstractFloat([e]) => Ok(Scalar::AbstractFloat([e.abs()])), + Scalar::F32([e]) => Ok(Scalar::F32([e.abs()])), + Scalar::AbstractInt([e]) => Ok(Scalar::AbstractInt([e.abs()])), + Scalar::I32([e]) => Ok(Scalar::I32([e.wrapping_abs()])), + Scalar::U32([e]) => Ok(Scalar::U32([e])), // TODO: just re-use the expression, ezpz + }) } - ( - &Expression::Compose { - components: ref src_components0, - ty: ty0, - }, - &Expression::Compose { - components: ref src_components1, - ty: ty1, - }, - ) if ty0 == ty1 - && matches!( - self.types[ty0].inner, - crate::TypeInner::Vector { - scalar: crate::Scalar { - kind: ScalarKind::Float, - .. - }, - .. - } - ) => - { - let mut components: Vec<_> = crate::proc::flatten_compose( - ty0, - src_components0, - self.expressions, - self.types, - ) - .chain(crate::proc::flatten_compose( - ty1, - src_components1, - self.expressions, - self.types, - )) - .collect(); - - let mid = components.len() / 2; - let (first, last) = components.split_at_mut(mid); - for (a, b) in first.iter_mut().zip(&*last) { - *a = self.math_pow(*a, *b, span)?; - } - components.truncate(mid); - - Expression::Compose { - ty: ty0, - components, - } + crate::MathFunction::Acos => { + component_wise_float!(self, span, [arg], |e| { Ok([e.acos()]) }) } - _ => return Err(ConstantEvaluatorError::InvalidMathArg), - }; - - self.register_evaluated_expr(expr, span) - } - - fn math_clamp( - &mut self, - e: Handle, - low: Handle, - high: Handle, - span: Span, - ) -> Result, ConstantEvaluatorError> { - let e = self.eval_zero_value_and_splat(e, span)?; - let low = self.eval_zero_value_and_splat(low, span)?; - let high = self.eval_zero_value_and_splat(high, span)?; - - let expr = match ( - &self.expressions[e], - &self.expressions[low], - &self.expressions[high], - ) { - (&Expression::Literal(e), &Expression::Literal(low), &Expression::Literal(high)) => { - let literal = match (e, low, high) { - (Literal::I32(e), Literal::I32(low), Literal::I32(high)) => { - if low > high { - return Err(ConstantEvaluatorError::InvalidClamp); - } else { - Literal::I32(e.clamp(low, high)) - } - } - (Literal::U32(e), Literal::U32(low), Literal::U32(high)) => { + crate::MathFunction::Acosh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.acosh()]) }) + } + crate::MathFunction::Asin => { + component_wise_float!(self, span, [arg], |e| { Ok([e.asin()]) }) + } + crate::MathFunction::Asinh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.asinh()]) }) + } + crate::MathFunction::Atan => { + component_wise_float!(self, span, [arg], |e| { Ok([e.atan()]) }) + } + crate::MathFunction::Atanh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.atanh()]) }) + } + crate::MathFunction::Pow => { + component_wise_float!(self, span, [arg, arg1.unwrap()], |e1, e2| { + Ok([e1.powf(e2)]) + }) + } + crate::MathFunction::Clamp => { + component_wise_scalar!( + self, + span, + [arg, arg1.unwrap(), arg2.unwrap()], + |e, low, high| { if low > high { - return Err(ConstantEvaluatorError::InvalidClamp); + Err(ConstantEvaluatorError::InvalidClamp) } else { - Literal::U32(e.clamp(low, high)) + Ok([e.clamp(low, high)]) } } - (Literal::F32(e), Literal::F32(low), Literal::F32(high)) => { - if low > high { - return Err(ConstantEvaluatorError::InvalidClamp); + ) + } + crate::MathFunction::Cos => { + component_wise_float!(self, span, [arg], |e| { Ok([e.cos()]) }) + } + crate::MathFunction::Cosh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.cosh()]) }) + } + crate::MathFunction::Round => { + // TODO: Use `f{32,64}.round_ties_even()` when available on stable. This polyfill + // is shamelessly [~~stolen from~~ inspired by `ndarray-image`][polyfill source], + // which has licensing compatible with ours. See also + // . + // + // [polyfill source]: https://github.com/imeka/ndarray-ndimage/blob/8b14b4d6ecfbc96a8a052f802e342a7049c68d8f/src/lib.rs#L98 + fn round_ties_even(x: f64) -> f64 { + let i = x as i64; + let f = (x - i as f64).abs(); + if f == 0.5 { + if i & 1 == 1 { + // -1.5, 1.5, 3.5, ... + (x.abs() + 0.5).copysign(x) } else { - Literal::F32(e.clamp(low, high)) + (x.abs() - 0.5).copysign(x) } + } else { + x.round() } - _ => return Err(ConstantEvaluatorError::InvalidMathArg), - }; - Expression::Literal(literal) - } - ( - &Expression::Compose { - components: ref src_components0, - ty: ty0, - }, - &Expression::Compose { - components: ref src_components1, - ty: ty1, - }, - &Expression::Compose { - components: ref src_components2, - ty: ty2, - }, - ) if ty0 == ty1 - && ty0 == ty2 - && matches!( - self.types[ty0].inner, - crate::TypeInner::Vector { - scalar: crate::Scalar { - kind: ScalarKind::Float, - .. - }, - .. - } - ) => - { - let mut components: Vec<_> = crate::proc::flatten_compose( - ty0, - src_components0, - self.expressions, - self.types, - ) - .chain(crate::proc::flatten_compose( - ty1, - src_components1, - self.expressions, - self.types, - )) - .chain(crate::proc::flatten_compose( - ty2, - src_components2, - self.expressions, - self.types, - )) - .collect(); - - let chunk_size = components.len() / 3; - let (es, rem) = components.split_at_mut(chunk_size); - let (lows, highs) = rem.split_at(chunk_size); - for ((e, low), high) in es.iter_mut().zip(lows).zip(highs) { - *e = self.math_clamp(*e, *low, *high, span)?; - } - components.truncate(chunk_size); - - Expression::Compose { - ty: ty0, - components, } + component_wise_float(self, span, [arg], |e| match e { + Float::Abstract([e]) => Ok(Float::Abstract([round_ties_even(e)])), + Float::F32([e]) => Ok(Float::F32([(round_ties_even(e as f64) as f32)])), + }) } - _ => return Err(ConstantEvaluatorError::InvalidMathArg), - }; - - self.register_evaluated_expr(expr, span) + crate::MathFunction::Saturate => { + component_wise_float!(self, span, [arg], |e| { Ok([e.clamp(0., 1.)]) }) + } + crate::MathFunction::Sin => { + component_wise_float!(self, span, [arg], |e| { Ok([e.sin()]) }) + } + crate::MathFunction::Sinh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.sinh()]) }) + } + crate::MathFunction::Tan => { + component_wise_float!(self, span, [arg], |e| { Ok([e.tan()]) }) + } + crate::MathFunction::Tanh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.tanh()]) }) + } + crate::MathFunction::Sqrt => { + component_wise_float!(self, span, [arg], |e| { Ok([e.sqrt()]) }) + } + crate::MathFunction::Step => { + component_wise_float!(self, span, [arg, arg1.unwrap()], |edge, x| { + Ok([if edge <= x { 1.0 } else { 0.0 }]) + }) + } + fun => Err(ConstantEvaluatorError::NotImplemented(format!( + "{fun:?} built-in function" + ))), + } } fn array_length( @@ -1144,7 +1287,12 @@ impl<'a> ConstantEvaluator<'a> { return self.cast(expr, target, span); }; - let crate::TypeInner::Array { base: _, size, stride: _ } = self.types[ty].inner else { + let crate::TypeInner::Array { + base: _, + size, + stride: _, + } = self.types[ty].inner + else { return self.cast(expr, target, span); }; diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl index ed81535ab5..bf0561f12e 100644 --- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/naga/tests/out/glsl/math-functions.main.Fragment.glsl @@ -67,7 +67,7 @@ void main() { float sign_c = sign(-1.0); vec4 sign_d = sign(vec4(-1.0)); int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y); - uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); + uint first_leading_bit_abs = uint(findMSB(0u)); int flb_a = findMSB(-1); ivec2 flb_b = findMSB(ivec2(-1)); uvec2 flb_c = uvec2(findMSB(uvec2(1u))); @@ -85,8 +85,8 @@ void main() { ivec2 ctz_h = ivec2(min(uvec2(findLSB(ivec2(1))), uvec2(32u))); int clz_a = (-1 < 0 ? 0 : 31 - findMSB(-1)); uint clz_b = uint(31 - findMSB(1u)); - ivec2 _e68 = ivec2(-1); - ivec2 clz_c = mix(ivec2(31) - findMSB(_e68), ivec2(0), lessThan(_e68, ivec2(0))); + ivec2 _e67 = ivec2(-1); + ivec2 clz_c = mix(ivec2(31) - findMSB(_e67), ivec2(0), lessThan(_e67, ivec2(0))); uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); float lde_a = ldexp(1.0, 2); vec2 lde_b = ldexp(vec2(1.0, 2.0), ivec2(3, 4)); diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl index 53d3acf0c1..5da3461dae 100644 --- a/naga/tests/out/hlsl/math-functions.hlsl +++ b/naga/tests/out/hlsl/math-functions.hlsl @@ -77,7 +77,7 @@ void main() float sign_c = sign(-1.0); float4 sign_d = sign((-1.0).xxxx); int const_dot = dot((int2)0, (int2)0); - uint first_leading_bit_abs = firstbithigh(abs(0u)); + uint first_leading_bit_abs = firstbithigh(0u); int flb_a = asint(firstbithigh(-1)); int2 flb_b = asint(firstbithigh((-1).xx)); uint2 flb_c = firstbithigh((1u).xx); @@ -95,8 +95,8 @@ void main() int2 ctz_h = asint(min((32u).xx, firstbitlow((1).xx))); int clz_a = (-1 < 0 ? 0 : 31 - asint(firstbithigh(-1))); uint clz_b = (31u - firstbithigh(1u)); - int2 _expr68 = (-1).xx; - int2 clz_c = (_expr68 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr68))); + int2 _expr67 = (-1).xx; + int2 clz_c = (_expr67 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr67))); uint2 clz_d = ((31u).xx - firstbithigh((1u).xx)); float lde_a = ldexp(1.0, 2); float2 lde_b = ldexp(float2(1.0, 2.0), int2(3, 4)); diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl index d93e502dc6..45fbcd00a1 100644 --- a/naga/tests/out/msl/math-functions.msl +++ b/naga/tests/out/msl/math-functions.msl @@ -70,13 +70,12 @@ fragment void main_( float sign_c = metal::sign(-1.0); metal::float4 sign_d = metal::sign(metal::float4(-1.0)); int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); - uint _e23 = metal::abs(0u); - uint first_leading_bit_abs = metal::select(31 - metal::clz(_e23), uint(-1), _e23 == 0 || _e23 == -1); + uint first_leading_bit_abs = metal::select(31 - metal::clz(0u), uint(-1), 0u == 0 || 0u == -1); int flb_a = metal::select(31 - metal::clz(metal::select(-1, ~-1, -1 < 0)), int(-1), -1 == 0 || -1 == -1); - metal::int2 _e28 = metal::int2(-1); - metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e28, ~_e28, _e28 < 0)), int2(-1), _e28 == 0 || _e28 == -1); - metal::uint2 _e31 = metal::uint2(1u); - metal::uint2 flb_c = metal::select(31 - metal::clz(_e31), uint2(-1), _e31 == 0 || _e31 == -1); + metal::int2 _e27 = metal::int2(-1); + metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e27, ~_e27, _e27 < 0)), int2(-1), _e27 == 0 || _e27 == -1); + metal::uint2 _e30 = metal::uint2(1u); + metal::uint2 flb_c = metal::select(31 - metal::clz(_e30), uint2(-1), _e30 == 0 || _e30 == -1); int ftb_a = (((metal::ctz(-1) + 1) % 33) - 1); uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1); metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1); diff --git a/naga/tests/out/spv/debug-symbol-terrain.spvasm b/naga/tests/out/spv/debug-symbol-terrain.spvasm index 623b8dc2c1..fd8e7f5df3 100644 --- a/naga/tests/out/spv/debug-symbol-terrain.spvasm +++ b/naga/tests/out/spv/debug-symbol-terrain.spvasm @@ -366,10 +366,10 @@ OpName %91 "x12" OpName %94 "m" OpName %203 "p" OpName %204 "fbm" -OpName %209 "x" -OpName %211 "v" -OpName %213 "a" -OpName %214 "i" +OpName %212 "x" +OpName %214 "v" +OpName %216 "a" +OpName %217 "i" OpName %255 "p" OpName %256 "min_max_height" OpName %257 "terrain_point" @@ -575,9 +575,12 @@ OpDecorate %582 Location 0 %206 = OpConstant %5 0.01 %207 = OpConstant %5 100.0 %208 = OpConstantComposite %6 %207 %207 -%210 = OpConstantNull %6 -%212 = OpTypePointer Function %5 -%215 = OpTypePointer Function %8 +%209 = OpConstant %5 0.87758255 +%210 = OpConstant %5 0.47942555 +%211 = OpConstantComposite %6 %209 %210 +%213 = OpConstantNull %6 +%215 = OpTypePointer Function %5 +%218 = OpTypePointer Function %8 %258 = OpTypeFunction %4 %6 %6 %271 = OpTypeFunction %14 %6 %6 %272 = OpConstant %5 0.1 @@ -835,28 +838,24 @@ OpFunctionEnd %204 = OpFunction %5 None %68 %203 = OpFunctionParameter %6 %202 = OpLabel -%211 = OpVariable %212 Function %74 -%214 = OpVariable %215 Function %135 -%209 = OpVariable %87 Function %210 -%213 = OpVariable %212 Function %78 -OpBranch %216 -%216 = OpLabel +%214 = OpVariable %215 Function %74 +%217 = OpVariable %218 Function %135 +%212 = OpVariable %87 Function %213 +%216 = OpVariable %215 Function %78 +OpBranch %219 +%219 = OpLabel OpLine %3 36 13 -%217 = OpVectorTimesScalar %6 %203 %206 +%220 = OpVectorTimesScalar %6 %203 %206 OpLine %3 36 5 -OpStore %209 %217 +OpStore %212 %220 OpLine %3 39 17 -OpLine %3 40 24 -%218 = OpExtInst %5 %1 Cos %78 OpLine %3 40 14 -%219 = OpExtInst %5 %1 Sin %78 -%220 = OpCompositeConstruct %6 %218 %219 OpLine %3 41 15 -%221 = OpCompositeExtract %5 %220 0 -%222 = OpCompositeExtract %5 %220 1 -%223 = OpCompositeExtract %5 %220 1 +%221 = OpCompositeExtract %5 %211 0 +%222 = OpCompositeExtract %5 %211 1 +%223 = OpCompositeExtract %5 %211 1 %224 = OpFNegate %5 %223 -%225 = OpCompositeExtract %5 %220 0 +%225 = OpCompositeExtract %5 %211 0 %226 = OpCompositeConstruct %6 %221 %222 %227 = OpCompositeConstruct %6 %224 %225 %228 = OpCompositeConstruct %9 %226 %227 @@ -867,7 +866,7 @@ OpLoopMerge %230 %232 None OpBranch %231 %231 = OpLabel OpLine %3 43 22 -%233 = OpLoad %8 %214 +%233 = OpLoad %8 %217 %234 = OpULessThan %112 %233 %205 OpLine %3 43 21 OpSelectionMerge %235 None @@ -878,44 +877,44 @@ OpBranch %230 OpBranch %237 %237 = OpLabel OpLine %3 1 1 -%239 = OpLoad %5 %211 -%240 = OpLoad %5 %213 -%241 = OpLoad %6 %209 +%239 = OpLoad %5 %214 +%240 = OpLoad %5 %216 +%241 = OpLoad %6 %212 OpLine %3 44 21 %242 = OpFunctionCall %5 %67 %241 OpLine %3 44 13 %243 = OpFMul %5 %240 %242 %244 = OpFAdd %5 %239 %243 OpLine %3 44 9 -OpStore %211 %244 +OpStore %214 %244 OpLine %3 45 13 -%245 = OpLoad %6 %209 +%245 = OpLoad %6 %212 %246 = OpMatrixTimesVector %6 %228 %245 OpLine %3 45 13 %247 = OpVectorTimesScalar %6 %246 %81 %248 = OpFAdd %6 %247 %208 OpLine %3 45 9 -OpStore %209 %248 +OpStore %212 %248 OpLine %3 1 1 -%249 = OpLoad %5 %213 +%249 = OpLoad %5 %216 OpLine %3 46 13 %250 = OpFMul %5 %249 %78 OpLine %3 46 9 -OpStore %213 %250 +OpStore %216 %250 OpBranch %238 %238 = OpLabel OpBranch %232 %232 = OpLabel OpLine %3 1 1 -%251 = OpLoad %8 %214 +%251 = OpLoad %8 %217 OpLine %3 43 43 %252 = OpIAdd %8 %251 %126 OpLine %3 43 39 -OpStore %214 %252 +OpStore %217 %252 OpBranch %229 %230 = OpLabel OpLine %3 1 1 -%253 = OpLoad %5 %211 +%253 = OpLoad %5 %214 OpReturnValue %253 OpFunctionEnd %257 = OpFunction %4 None %258 @@ -1193,8 +1192,8 @@ OpReturn OpFunctionEnd %465 = OpFunction %2 None %346 %453 = OpLabel -%468 = OpVariable %212 Function %74 -%469 = OpVariable %215 Function %135 +%468 = OpVariable %215 Function %74 +%469 = OpVariable %218 Function %135 %456 = OpLoad %8 %455 %459 = OpLoad %7 %457 %462 = OpLoad %6 %460 diff --git a/naga/tests/out/spv/math-functions.spvasm b/naga/tests/out/spv/math-functions.spvasm index ba3e7cffb9..bbbf370970 100644 --- a/naga/tests/out/spv/math-functions.spvasm +++ b/naga/tests/out/spv/math-functions.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 126 +; Bound: 125 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -61,10 +61,10 @@ OpMemberDecorate %13 1 Offset 16 %45 = OpConstantComposite %3 %43 %43 %43 %43 %52 = OpConstantComposite %3 %17 %17 %17 %17 %59 = OpConstantNull %6 -%77 = OpConstant %25 32 -%86 = OpConstantComposite %29 %77 %77 -%95 = OpConstant %6 31 -%100 = OpConstantComposite %5 %95 %95 +%76 = OpConstant %25 32 +%85 = OpConstantComposite %29 %76 %76 +%94 = OpConstant %6 31 +%99 = OpConstantComposite %5 %94 %94 %15 = OpFunction %2 None %16 %14 = OpLabel OpBranch %46 @@ -87,60 +87,59 @@ OpBranch %46 %65 = OpCompositeExtract %6 %24 1 %66 = OpIMul %6 %64 %65 %58 = OpIAdd %6 %63 %66 -%67 = OpCopyObject %25 %26 -%68 = OpExtInst %25 %1 FindUMsb %67 -%69 = OpExtInst %6 %1 FindSMsb %20 -%70 = OpExtInst %5 %1 FindSMsb %27 -%71 = OpExtInst %29 %1 FindUMsb %30 -%72 = OpExtInst %6 %1 FindILsb %20 -%73 = OpExtInst %25 %1 FindILsb %28 -%74 = OpExtInst %5 %1 FindILsb %27 -%75 = OpExtInst %29 %1 FindILsb %30 -%78 = OpExtInst %25 %1 FindILsb %26 -%76 = OpExtInst %25 %1 UMin %77 %78 -%80 = OpExtInst %6 %1 FindILsb %31 -%79 = OpExtInst %6 %1 UMin %77 %80 -%82 = OpExtInst %25 %1 FindILsb %32 -%81 = OpExtInst %25 %1 UMin %77 %82 -%84 = OpExtInst %6 %1 FindILsb %20 -%83 = OpExtInst %6 %1 UMin %77 %84 -%87 = OpExtInst %29 %1 FindILsb %33 -%85 = OpExtInst %29 %1 UMin %86 %87 -%89 = OpExtInst %5 %1 FindILsb %34 -%88 = OpExtInst %5 %1 UMin %86 %89 -%91 = OpExtInst %29 %1 FindILsb %30 -%90 = OpExtInst %29 %1 UMin %86 %91 -%93 = OpExtInst %5 %1 FindILsb %36 -%92 = OpExtInst %5 %1 UMin %86 %93 -%96 = OpExtInst %6 %1 FindUMsb %20 -%94 = OpISub %6 %95 %96 -%98 = OpExtInst %6 %1 FindUMsb %28 -%97 = OpISub %25 %95 %98 -%101 = OpExtInst %5 %1 FindUMsb %27 -%99 = OpISub %5 %100 %101 -%103 = OpExtInst %5 %1 FindUMsb %30 -%102 = OpISub %29 %100 %103 -%104 = OpExtInst %4 %1 Ldexp %17 %37 -%105 = OpExtInst %7 %1 Ldexp %39 %42 +%67 = OpExtInst %25 %1 FindUMsb %26 +%68 = OpExtInst %6 %1 FindSMsb %20 +%69 = OpExtInst %5 %1 FindSMsb %27 +%70 = OpExtInst %29 %1 FindUMsb %30 +%71 = OpExtInst %6 %1 FindILsb %20 +%72 = OpExtInst %25 %1 FindILsb %28 +%73 = OpExtInst %5 %1 FindILsb %27 +%74 = OpExtInst %29 %1 FindILsb %30 +%77 = OpExtInst %25 %1 FindILsb %26 +%75 = OpExtInst %25 %1 UMin %76 %77 +%79 = OpExtInst %6 %1 FindILsb %31 +%78 = OpExtInst %6 %1 UMin %76 %79 +%81 = OpExtInst %25 %1 FindILsb %32 +%80 = OpExtInst %25 %1 UMin %76 %81 +%83 = OpExtInst %6 %1 FindILsb %20 +%82 = OpExtInst %6 %1 UMin %76 %83 +%86 = OpExtInst %29 %1 FindILsb %33 +%84 = OpExtInst %29 %1 UMin %85 %86 +%88 = OpExtInst %5 %1 FindILsb %34 +%87 = OpExtInst %5 %1 UMin %85 %88 +%90 = OpExtInst %29 %1 FindILsb %30 +%89 = OpExtInst %29 %1 UMin %85 %90 +%92 = OpExtInst %5 %1 FindILsb %36 +%91 = OpExtInst %5 %1 UMin %85 %92 +%95 = OpExtInst %6 %1 FindUMsb %20 +%93 = OpISub %6 %94 %95 +%97 = OpExtInst %6 %1 FindUMsb %28 +%96 = OpISub %25 %94 %97 +%100 = OpExtInst %5 %1 FindUMsb %27 +%98 = OpISub %5 %99 %100 +%102 = OpExtInst %5 %1 FindUMsb %30 +%101 = OpISub %29 %99 %102 +%103 = OpExtInst %4 %1 Ldexp %17 %37 +%104 = OpExtInst %7 %1 Ldexp %39 %42 +%105 = OpExtInst %8 %1 ModfStruct %43 %106 = OpExtInst %8 %1 ModfStruct %43 -%107 = OpExtInst %8 %1 ModfStruct %43 -%108 = OpCompositeExtract %4 %107 0 -%109 = OpExtInst %8 %1 ModfStruct %43 -%110 = OpCompositeExtract %4 %109 1 -%111 = OpExtInst %9 %1 ModfStruct %44 -%112 = OpExtInst %10 %1 ModfStruct %45 -%113 = OpCompositeExtract %3 %112 1 -%114 = OpCompositeExtract %4 %113 0 -%115 = OpExtInst %9 %1 ModfStruct %44 -%116 = OpCompositeExtract %7 %115 0 -%117 = OpCompositeExtract %4 %116 1 +%107 = OpCompositeExtract %4 %106 0 +%108 = OpExtInst %8 %1 ModfStruct %43 +%109 = OpCompositeExtract %4 %108 1 +%110 = OpExtInst %9 %1 ModfStruct %44 +%111 = OpExtInst %10 %1 ModfStruct %45 +%112 = OpCompositeExtract %3 %111 1 +%113 = OpCompositeExtract %4 %112 0 +%114 = OpExtInst %9 %1 ModfStruct %44 +%115 = OpCompositeExtract %7 %114 0 +%116 = OpCompositeExtract %4 %115 1 +%117 = OpExtInst %11 %1 FrexpStruct %43 %118 = OpExtInst %11 %1 FrexpStruct %43 -%119 = OpExtInst %11 %1 FrexpStruct %43 -%120 = OpCompositeExtract %4 %119 0 -%121 = OpExtInst %11 %1 FrexpStruct %43 -%122 = OpCompositeExtract %6 %121 1 -%123 = OpExtInst %13 %1 FrexpStruct %45 -%124 = OpCompositeExtract %12 %123 1 -%125 = OpCompositeExtract %6 %124 0 +%119 = OpCompositeExtract %4 %118 0 +%120 = OpExtInst %11 %1 FrexpStruct %43 +%121 = OpCompositeExtract %6 %120 1 +%122 = OpExtInst %13 %1 FrexpStruct %45 +%123 = OpCompositeExtract %12 %122 1 +%124 = OpCompositeExtract %6 %123 0 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/math-functions.wgsl b/naga/tests/out/wgsl/math-functions.wgsl index 92f391038d..ce38fee986 100644 --- a/naga/tests/out/wgsl/math-functions.wgsl +++ b/naga/tests/out/wgsl/math-functions.wgsl @@ -12,7 +12,7 @@ fn main() { let sign_c = sign(-1f); let sign_d = sign(vec4(-1f)); let const_dot = dot(vec2(), vec2()); - let first_leading_bit_abs = firstLeadingBit(abs(0u)); + let first_leading_bit_abs = firstLeadingBit(0u); let flb_a = firstLeadingBit(-1i); let flb_b = firstLeadingBit(vec2(-1i)); let flb_c = firstLeadingBit(vec2(1u));