From 3aa944975c4766919a2e58c94115fa38142fa492 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 4 Jan 2024 16:58:38 -0500 Subject: [PATCH 01/21] chore: run `rustfmt` in `naga` --- naga/src/proc/constant_evaluator.rs | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 82efeece4e..45dffbd4a0 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1144,7 +1144,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); }; From b077d6ca546b89ae2d3cf8cbeb192bb485d8fdde Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 10:37:43 -0500 Subject: [PATCH 02/21] refactor(const_eval): add `component_wise_float` helper, reimpl. `math_pow` --- Cargo.lock | 1 + naga/Cargo.toml | 1 + naga/src/lib.rs | 4 + naga/src/proc/constant_evaluator.rs | 259 +++++++++++++++++++++------- 4 files changed, 207 insertions(+), 58 deletions(-) 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 45dffbd4a0..e020c1692d 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1,9 +1,209 @@ +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_float -> Float, + literals: [ + AbstractFloat => Abstract: f64, + F32 => F32: f32, + ], + scalar_kinds: [ + Float, + AbstractFloat, + ], +} + #[derive(Debug)] enum Behavior { Wgsl, @@ -606,64 +806,7 @@ impl<'a> ConstantEvaluator<'a> { 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))) - } - ( - &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, - } - } - _ => return Err(ConstantEvaluatorError::InvalidMathArg), - }; - - self.register_evaluated_expr(expr, span) + component_wise_float!(self, span, [e1, e2], |e1, e2| { Ok([e1.powf(e2)]) }) } fn math_clamp( From 85b8457d63e08b0b368b187425d33831c9c086c1 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Wed, 10 Jan 2024 17:30:24 -0500 Subject: [PATCH 03/21] feat(const_eval): impl. `abs` with new `component_wise_scalar` --- CHANGELOG.md | 8 ++ naga/src/proc/constant_evaluator.rs | 27 ++++ .../glsl/math-functions.main.Fragment.glsl | 6 +- naga/tests/out/hlsl/math-functions.hlsl | 6 +- naga/tests/out/msl/math-functions.msl | 11 +- naga/tests/out/spv/math-functions.spvasm | 115 +++++++++--------- naga/tests/out/wgsl/math-functions.wgsl | 2 +- 7 files changed, 104 insertions(+), 71 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index c4361ea48a..604399bdcb 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -37,6 +37,14 @@ 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` + ## v0.19.0 (2024-01-17) This release includes: diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index e020c1692d..13b4dd7cef 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -192,6 +192,24 @@ macro_rules! gen_component_wise_extractor { }; } +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: [ @@ -792,6 +810,15 @@ impl<'a> ConstantEvaluator<'a> { } match fun { + 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 + }) + } 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!( 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/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)); From 3064566d9ac7dd20b159f3fffe1d42941789a31b Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 10:41:22 -0500 Subject: [PATCH 04/21] feat(const_eval): impl. `round` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 26 ++++++++++++++++++++++++++ 2 files changed, 27 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 604399bdcb..0528f88f7f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -44,6 +44,7 @@ Bottom level categories: - 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` + - `round` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 13b4dd7cef..b472937d09 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -821,6 +821,32 @@ impl<'a> ConstantEvaluator<'a> { } crate::MathFunction::Pow => self.math_pow(arg, arg1.unwrap(), span), crate::MathFunction::Clamp => self.math_clamp(arg, arg1.unwrap(), arg2.unwrap(), span), + 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 { + (x.abs() - 0.5).copysign(x) + } + } else { + x.round() + } + } + 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)])), + }) + } fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" ))), From abf4947d0b8e61ff5808b03f7d908def1385b698 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 10:41:29 -0500 Subject: [PATCH 05/21] feat(const_eval): impl. `saturate` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0528f88f7f..c033e572f0 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -45,6 +45,7 @@ Bottom level categories: - [#4879](https://github.com/gfx-rs/wgpu/pull/4879) by @ErichDonGubler: - `abs` - `round` + - `saturate` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index b472937d09..e89c1dd2c7 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -847,6 +847,9 @@ impl<'a> ConstantEvaluator<'a> { Float::F32([e]) => Ok(Float::F32([(round_ties_even(e as f64) as f32)])), }) } + crate::MathFunction::Saturate => { + component_wise_float!(self, span, [arg], |e| { Ok([e.clamp(0., 1.)]) }) + } fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" ))), From 1f265d36ba5c87720a0e353d655aeca1539226fb Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 10:41:35 -0500 Subject: [PATCH 06/21] feat(const_eval): impl. `sin` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 + .../tests/out/spv/debug-symbol-terrain.spvasm | 64 +++++++++---------- 3 files changed, 36 insertions(+), 32 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index c033e572f0..587c7c9e6b 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -46,6 +46,7 @@ Bottom level categories: - `abs` - `round` - `saturate` + - `sin` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index e89c1dd2c7..33002d5c47 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -850,6 +850,9 @@ impl<'a> ConstantEvaluator<'a> { 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()]) }) + } fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" ))), diff --git a/naga/tests/out/spv/debug-symbol-terrain.spvasm b/naga/tests/out/spv/debug-symbol-terrain.spvasm index 623b8dc2c1..a818ee6fea 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 %210 "x" +OpName %212 "v" +OpName %214 "a" +OpName %215 "i" OpName %255 "p" OpName %256 "min_max_height" OpName %257 "terrain_point" @@ -575,9 +575,10 @@ 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.47942555 +%211 = OpConstantNull %6 +%213 = OpTypePointer Function %5 +%216 = OpTypePointer Function %8 %258 = OpTypeFunction %4 %6 %6 %271 = OpTypeFunction %14 %6 %6 %272 = OpConstant %5 0.1 @@ -835,22 +836,21 @@ 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 +%212 = OpVariable %213 Function %74 +%215 = OpVariable %216 Function %135 +%210 = OpVariable %87 Function %211 +%214 = OpVariable %213 Function %78 +OpBranch %217 +%217 = OpLabel OpLine %3 36 13 -%217 = OpVectorTimesScalar %6 %203 %206 +%218 = OpVectorTimesScalar %6 %203 %206 OpLine %3 36 5 -OpStore %209 %217 +OpStore %210 %218 OpLine %3 39 17 OpLine %3 40 24 -%218 = OpExtInst %5 %1 Cos %78 +%219 = OpExtInst %5 %1 Cos %78 OpLine %3 40 14 -%219 = OpExtInst %5 %1 Sin %78 -%220 = OpCompositeConstruct %6 %218 %219 +%220 = OpCompositeConstruct %6 %219 %209 OpLine %3 41 15 %221 = OpCompositeExtract %5 %220 0 %222 = OpCompositeExtract %5 %220 1 @@ -867,7 +867,7 @@ OpLoopMerge %230 %232 None OpBranch %231 %231 = OpLabel OpLine %3 43 22 -%233 = OpLoad %8 %214 +%233 = OpLoad %8 %215 %234 = OpULessThan %112 %233 %205 OpLine %3 43 21 OpSelectionMerge %235 None @@ -878,44 +878,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 %212 +%240 = OpLoad %5 %214 +%241 = OpLoad %6 %210 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 %212 %244 OpLine %3 45 13 -%245 = OpLoad %6 %209 +%245 = OpLoad %6 %210 %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 %210 %248 OpLine %3 1 1 -%249 = OpLoad %5 %213 +%249 = OpLoad %5 %214 OpLine %3 46 13 %250 = OpFMul %5 %249 %78 OpLine %3 46 9 -OpStore %213 %250 +OpStore %214 %250 OpBranch %238 %238 = OpLabel OpBranch %232 %232 = OpLabel OpLine %3 1 1 -%251 = OpLoad %8 %214 +%251 = OpLoad %8 %215 OpLine %3 43 43 %252 = OpIAdd %8 %251 %126 OpLine %3 43 39 -OpStore %214 %252 +OpStore %215 %252 OpBranch %229 %230 = OpLabel OpLine %3 1 1 -%253 = OpLoad %5 %211 +%253 = OpLoad %5 %212 OpReturnValue %253 OpFunctionEnd %257 = OpFunction %4 None %258 @@ -1193,8 +1193,8 @@ OpReturn OpFunctionEnd %465 = OpFunction %2 None %346 %453 = OpLabel -%468 = OpVariable %212 Function %74 -%469 = OpVariable %215 Function %135 +%468 = OpVariable %213 Function %74 +%469 = OpVariable %216 Function %135 %456 = OpLoad %8 %455 %459 = OpLoad %7 %457 %462 = OpLoad %6 %460 From 1e00b307f16ea5b2d8bbc715594439edb2128e5c Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 10:41:39 -0500 Subject: [PATCH 07/21] feat(const_eval): impl. `sinh` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 587c7c9e6b..914ed448d3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -47,6 +47,7 @@ Bottom level categories: - `round` - `saturate` - `sin` + - `sinh` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 33002d5c47..edff331e18 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -853,6 +853,9 @@ impl<'a> ConstantEvaluator<'a> { 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()]) }) + } fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" ))), From 4480270f83a84e8cac0395f5707864faff02d0ae Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 10:41:44 -0500 Subject: [PATCH 08/21] feat(const_eval): impl. `sqrt` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 914ed448d3..14776b65ac 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -48,6 +48,7 @@ Bottom level categories: - `saturate` - `sin` - `sinh` + - `sqrt` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index edff331e18..1a36a59c43 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -856,6 +856,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Sinh => { component_wise_float!(self, span, [arg], |e| { Ok([e.sinh()]) }) } + crate::MathFunction::Sqrt => { + component_wise_float!(self, span, [arg], |e| { Ok([e.sqrt()]) }) + } fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" ))), From 5f4766cc8016e38e56f42ef1aebc8b6e2c5f5ba4 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 10:41:51 -0500 Subject: [PATCH 09/21] feat(const_eval): impl. `step` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 5 +++++ 2 files changed, 6 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 14776b65ac..d6b1ccc534 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -49,6 +49,7 @@ Bottom level categories: - `sin` - `sinh` - `sqrt` + - `step` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 1a36a59c43..eefb470b3a 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -859,6 +859,11 @@ impl<'a> ConstantEvaluator<'a> { 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" ))), From 295537a5bab4faa0149f2161cd3f83c77f74ea1e Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 14 Dec 2023 11:20:45 -0500 Subject: [PATCH 10/21] feat(const_eval)!: reimpl. `clamp` --- naga/src/proc/constant_evaluator.rs | 123 ++++------------------------ 1 file changed, 14 insertions(+), 109 deletions(-) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index eefb470b3a..bb5a624109 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -820,7 +820,20 @@ impl<'a> ConstantEvaluator<'a> { }) } crate::MathFunction::Pow => self.math_pow(arg, arg1.unwrap(), span), - crate::MathFunction::Clamp => self.math_clamp(arg, arg1.unwrap(), arg2.unwrap(), span), + crate::MathFunction::Clamp => { + component_wise_scalar!( + self, + span, + [arg, arg1.unwrap(), arg2.unwrap()], + |e, low, high| { + if low > high { + Err(ConstantEvaluatorError::InvalidClamp) + } else { + Ok([e.clamp(low, high)]) + } + } + ) + } 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], @@ -879,114 +892,6 @@ impl<'a> ConstantEvaluator<'a> { component_wise_float!(self, span, [e1, e2], |e1, e2| { Ok([e1.powf(e2)]) }) } - 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)) => { - if low > high { - return Err(ConstantEvaluatorError::InvalidClamp); - } else { - Literal::U32(e.clamp(low, high)) - } - } - (Literal::F32(e), Literal::F32(low), Literal::F32(high)) => { - if low > high { - return Err(ConstantEvaluatorError::InvalidClamp); - } else { - Literal::F32(e.clamp(low, high)) - } - } - _ => 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, - } - } - _ => return Err(ConstantEvaluatorError::InvalidMathArg), - }; - - self.register_evaluated_expr(expr, span) - } - fn array_length( &mut self, array: Handle, From 58c86f057f98af0afa3882d2be111e9f54f91504 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 5 Jan 2024 10:56:24 -0500 Subject: [PATCH 11/21] feat(const_eval): impl. `tan` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index d6b1ccc534..586c943f31 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -50,6 +50,7 @@ Bottom level categories: - `sinh` - `sqrt` - `step` + - `tan` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index bb5a624109..bdf9eec202 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -869,6 +869,9 @@ impl<'a> ConstantEvaluator<'a> { 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::Sqrt => { component_wise_float!(self, span, [arg], |e| { Ok([e.sqrt()]) }) } From 9e0cb4bce3f167c8aadae9432ce4537c571a1d29 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 5 Jan 2024 10:56:24 -0500 Subject: [PATCH 12/21] feat(const_eval): impl. `tanh` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 586c943f31..a06e271c5c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -51,6 +51,7 @@ Bottom level categories: - `sqrt` - `step` - `tan` + - `tanh` ## v0.19.0 (2024-01-17) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index bdf9eec202..029671db8a 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -872,6 +872,9 @@ impl<'a> ConstantEvaluator<'a> { 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()]) }) } From 845b3a96b23790a8601bf6dc97d7ab477eb0b0ce Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 07:47:37 -0500 Subject: [PATCH 13/21] feat(const_eval): impl. `cos` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 + .../tests/out/spv/debug-symbol-terrain.spvasm | 73 +++++++++---------- 3 files changed, 40 insertions(+), 37 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index a06e271c5c..1b474ec20c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -44,6 +44,7 @@ Bottom level categories: - 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` + - `cos` - `round` - `saturate` - `sin` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 029671db8a..cfd57a1767 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -834,6 +834,9 @@ impl<'a> ConstantEvaluator<'a> { } ) } + crate::MathFunction::Cos => { + component_wise_float!(self, span, [arg], |e| { Ok([e.cos()]) }) + } 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], diff --git a/naga/tests/out/spv/debug-symbol-terrain.spvasm b/naga/tests/out/spv/debug-symbol-terrain.spvasm index a818ee6fea..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 %210 "x" -OpName %212 "v" -OpName %214 "a" -OpName %215 "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,10 +575,12 @@ OpDecorate %582 Location 0 %206 = OpConstant %5 0.01 %207 = OpConstant %5 100.0 %208 = OpConstantComposite %6 %207 %207 -%209 = OpConstant %5 0.47942555 -%211 = OpConstantNull %6 -%213 = OpTypePointer Function %5 -%216 = 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 @@ -836,27 +838,24 @@ OpFunctionEnd %204 = OpFunction %5 None %68 %203 = OpFunctionParameter %6 %202 = OpLabel -%212 = OpVariable %213 Function %74 -%215 = OpVariable %216 Function %135 -%210 = OpVariable %87 Function %211 -%214 = OpVariable %213 Function %78 -OpBranch %217 -%217 = 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 -%218 = OpVectorTimesScalar %6 %203 %206 +%220 = OpVectorTimesScalar %6 %203 %206 OpLine %3 36 5 -OpStore %210 %218 +OpStore %212 %220 OpLine %3 39 17 -OpLine %3 40 24 -%219 = OpExtInst %5 %1 Cos %78 OpLine %3 40 14 -%220 = OpCompositeConstruct %6 %219 %209 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 %215 +%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 %212 -%240 = OpLoad %5 %214 -%241 = OpLoad %6 %210 +%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 %212 %244 +OpStore %214 %244 OpLine %3 45 13 -%245 = OpLoad %6 %210 +%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 %210 %248 +OpStore %212 %248 OpLine %3 1 1 -%249 = OpLoad %5 %214 +%249 = OpLoad %5 %216 OpLine %3 46 13 %250 = OpFMul %5 %249 %78 OpLine %3 46 9 -OpStore %214 %250 +OpStore %216 %250 OpBranch %238 %238 = OpLabel OpBranch %232 %232 = OpLabel OpLine %3 1 1 -%251 = OpLoad %8 %215 +%251 = OpLoad %8 %217 OpLine %3 43 43 %252 = OpIAdd %8 %251 %126 OpLine %3 43 39 -OpStore %215 %252 +OpStore %217 %252 OpBranch %229 %230 = OpLabel OpLine %3 1 1 -%253 = OpLoad %5 %212 +%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 %213 Function %74 -%469 = OpVariable %216 Function %135 +%468 = OpVariable %215 Function %74 +%469 = OpVariable %218 Function %135 %456 = OpLoad %8 %455 %459 = OpLoad %7 %457 %462 = OpLoad %6 %460 From 4b31966b8637b68266cc276f9096ead18211f871 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 07:47:37 -0500 Subject: [PATCH 14/21] feat(const_eval): impl. `cosh` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 1b474ec20c..80704dc27f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -45,6 +45,7 @@ Bottom level categories: - [#4879](https://github.com/gfx-rs/wgpu/pull/4879) by @ErichDonGubler: - `abs` - `cos` + - `cosh` - `round` - `saturate` - `sin` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index cfd57a1767..4692355a56 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -837,6 +837,9 @@ impl<'a> ConstantEvaluator<'a> { 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], From f3a6b56c88df6edcd10739b7493d757d730cf91d Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 08:01:52 -0500 Subject: [PATCH 15/21] feat(const_eval): impl. `acos` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 80704dc27f..ff937bbdd7 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -44,6 +44,7 @@ Bottom level categories: - 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` - `cos` - `cosh` - `round` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 4692355a56..b56d4c6999 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -819,6 +819,9 @@ impl<'a> ConstantEvaluator<'a> { Scalar::U32([e]) => Ok(Scalar::U32([e])), // TODO: just re-use the expression, ezpz }) } + crate::MathFunction::Acos => { + component_wise_float!(self, span, [arg], |e| { Ok([e.acos()]) }) + } crate::MathFunction::Pow => self.math_pow(arg, arg1.unwrap(), span), crate::MathFunction::Clamp => { component_wise_scalar!( From 4c93d2098e36bc84477bdf31a6676fc2fac8a6a0 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 08:01:52 -0500 Subject: [PATCH 16/21] feat(const_eval): impl. `acosh` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index ff937bbdd7..685b0755f8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -45,6 +45,7 @@ Bottom level categories: - [#4879](https://github.com/gfx-rs/wgpu/pull/4879) by @ErichDonGubler: - `abs` - `acos` + - `acosh` - `cos` - `cosh` - `round` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index b56d4c6999..598e0a92bd 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -822,6 +822,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Acos => { component_wise_float!(self, span, [arg], |e| { Ok([e.acos()]) }) } + crate::MathFunction::Acosh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.acosh()]) }) + } crate::MathFunction::Pow => self.math_pow(arg, arg1.unwrap(), span), crate::MathFunction::Clamp => { component_wise_scalar!( From acad5c67464b72d4692a6e7852a2eca976263c5e Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 08:02:02 -0500 Subject: [PATCH 17/21] feat(const_eval): impl. `asin` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 685b0755f8..0a5a137283 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -46,6 +46,7 @@ Bottom level categories: - `abs` - `acos` - `acosh` + - `asin` - `cos` - `cosh` - `round` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 598e0a92bd..ca0a596040 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -825,6 +825,9 @@ impl<'a> ConstantEvaluator<'a> { 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::Pow => self.math_pow(arg, arg1.unwrap(), span), crate::MathFunction::Clamp => { component_wise_scalar!( From 2c101dd349e888958219166f228365aa269dfc0d Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 08:02:02 -0500 Subject: [PATCH 18/21] feat(const_eval): impl. `asinh` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0a5a137283..5ef25b7e3a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -47,6 +47,7 @@ Bottom level categories: - `acos` - `acosh` - `asin` + - `asinh` - `cos` - `cosh` - `round` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index ca0a596040..c6f9922ddc 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -828,6 +828,9 @@ impl<'a> ConstantEvaluator<'a> { 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::Pow => self.math_pow(arg, arg1.unwrap(), span), crate::MathFunction::Clamp => { component_wise_scalar!( From 47accec3991d223b46cda020eb07329216c3a407 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 08:02:12 -0500 Subject: [PATCH 19/21] feat(const_eval): impl. `atan` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 5ef25b7e3a..a5540a0582 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -48,6 +48,7 @@ Bottom level categories: - `acosh` - `asin` - `asinh` + - `atan` - `cos` - `cosh` - `round` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index c6f9922ddc..4bfdb8112a 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -831,6 +831,9 @@ impl<'a> ConstantEvaluator<'a> { 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::Pow => self.math_pow(arg, arg1.unwrap(), span), crate::MathFunction::Clamp => { component_wise_scalar!( From a0c9e2661fba7e42a5648a31c4754f6f2610ebd2 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 08:02:12 -0500 Subject: [PATCH 20/21] feat(const_eval): impl. `atanh` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index a5540a0582..9bfea9f358 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -49,6 +49,7 @@ Bottom level categories: - `asin` - `asinh` - `atan` + - `atanh` - `cos` - `cosh` - `round` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 4bfdb8112a..3b3ba56209 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -834,6 +834,9 @@ impl<'a> ConstantEvaluator<'a> { 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 => self.math_pow(arg, arg1.unwrap(), span), crate::MathFunction::Clamp => { component_wise_scalar!( From af20cb3e0aced29a31e2292b4f28a859a0035080 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 22 Dec 2023 08:05:15 -0500 Subject: [PATCH 21/21] refactor(const_eval): inline `math_pow` --- naga/src/proc/constant_evaluator.rs | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 3b3ba56209..5cffdba86e 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -837,7 +837,11 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Atanh => { component_wise_float!(self, span, [arg], |e| { Ok([e.atanh()]) }) } - crate::MathFunction::Pow => self.math_pow(arg, arg1.unwrap(), span), + crate::MathFunction::Pow => { + component_wise_float!(self, span, [arg, arg1.unwrap()], |e1, e2| { + Ok([e1.powf(e2)]) + }) + } crate::MathFunction::Clamp => { component_wise_scalar!( self, @@ -913,15 +917,6 @@ impl<'a> ConstantEvaluator<'a> { } } - fn math_pow( - &mut self, - e1: Handle, - e2: Handle, - span: Span, - ) -> Result, ConstantEvaluatorError> { - component_wise_float!(self, span, [e1, e2], |e1, e2| { Ok([e1.powf(e2)]) }) - } - fn array_length( &mut self, array: Handle,