From 8583dc107938c6a55efc5371241f91b715389fae Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:11:45 -0500 Subject: [PATCH 01/20] feat(const_eval): impl. `ceil` --- CHANGELOG.md | 2 ++ naga/src/proc/constant_evaluator.rs | 3 +++ naga/tests/out/glsl/variations.main.Fragment.glsl | 3 +-- 3 files changed, 6 insertions(+), 2 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 53a43857bc..0f47906153 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -66,6 +66,8 @@ Bottom level categories: - `step` - `tan` - `tanh` + - [#5098](https://github.com/gfx-rs/wgpu/pull/5098) by @ErichDonGubler: + - `ceil` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index a84863066e..d1ee99e02f 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -842,6 +842,9 @@ impl<'a> ConstantEvaluator<'a> { Ok([e1.powf(e2)]) }) } + crate::MathFunction::Ceil => { + component_wise_float!(self, span, [arg], |e| { Ok([e.ceil()]) }) + } crate::MathFunction::Clamp => { component_wise_scalar!( self, diff --git a/naga/tests/out/glsl/variations.main.Fragment.glsl b/naga/tests/out/glsl/variations.main.Fragment.glsl index 5ea3eb03cf..25b258987c 100644 --- a/naga/tests/out/glsl/variations.main.Fragment.glsl +++ b/naga/tests/out/glsl/variations.main.Fragment.glsl @@ -8,9 +8,8 @@ uniform highp samplerCube _group_0_binding_0_fs; void main_1() { ivec2 sizeCube = ivec2(0); - float a = 0.0; + float a = 1.0; sizeCube = ivec2(uvec2(textureSize(_group_0_binding_0_fs, 0).xy)); - a = ceil(1.0); return; } From a21270ae623ec8ee2b9ac4b22098e427d1e8ee67 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:12:55 -0500 Subject: [PATCH 02/20] feat(const_eval): impl. `floor` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0f47906153..0005d65b45 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -68,6 +68,7 @@ Bottom level categories: - `tanh` - [#5098](https://github.com/gfx-rs/wgpu/pull/5098) by @ErichDonGubler: - `ceil` + - `floor` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index d1ee99e02f..4c6ae058a9 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -865,6 +865,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Cosh => { component_wise_float!(self, span, [arg], |e| { Ok([e.cosh()]) }) } + crate::MathFunction::Floor => { + component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) + } 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 cb3428c0fdd65f708f29a6b94de3364024bdfbd0 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 5 Jan 2024 10:56:24 -0500 Subject: [PATCH 03/20] feat(const_eval): impl. `countLeadingZeros` with new `component_wise_concrete_int` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 21 ++ .../glsl/math-functions.main.Fragment.glsl | 7 +- naga/tests/out/hlsl/math-functions.hlsl | 7 +- naga/tests/out/msl/math-functions.msl | 6 +- naga/tests/out/spv/math-functions.spvasm | 241 +++++++++--------- naga/tests/out/wgsl/math-functions.wgsl | 6 +- 7 files changed, 147 insertions(+), 142 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0005d65b45..7e5f9c099b 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -68,6 +68,7 @@ Bottom level categories: - `tanh` - [#5098](https://github.com/gfx-rs/wgpu/pull/5098) by @ErichDonGubler: - `ceil` + - `countLeadingZeros` - `floor` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 4c6ae058a9..88c59809a9 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -222,6 +222,18 @@ gen_component_wise_extractor! { ], } +gen_component_wise_extractor! { + component_wise_concrete_int -> ConcreteInt, + literals: [ + U32 => U32: u32, + I32 => I32: i32, + ], + scalar_kinds: [ + Sint, + Uint, + ], +} + #[derive(Debug)] enum Behavior { Wgsl, @@ -865,6 +877,15 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Cosh => { component_wise_float!(self, span, [arg], |e| { Ok([e.cosh()]) }) } + crate::MathFunction::CountLeadingZeros => { + component_wise_concrete_int!(self, span, [arg], |e| { + #[allow(clippy::useless_conversion)] + Ok([e + .leading_zeros() + .try_into() + .expect("bit count overflowed 32 bits, somehow!?")]) + }) + } crate::MathFunction::Floor => { component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl index bf0561f12e..2892157801 100644 --- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/naga/tests/out/glsl/math-functions.main.Fragment.glsl @@ -83,11 +83,8 @@ void main() { ivec2 ctz_f = ivec2(min(uvec2(findLSB(ivec2(0))), uvec2(32u))); uvec2 ctz_g = min(uvec2(findLSB(uvec2(1u))), uvec2(32u)); 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 _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))); + ivec2 clz_c = ivec2(0, 0); + uvec2 clz_d = uvec2(31u, 31u); float lde_a = ldexp(1.0, 2); vec2 lde_b = ldexp(vec2(1.0, 2.0), ivec2(3, 4)); _modf_result_f32_ modf_a = naga_modf(1.5); diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl index 5da3461dae..d576365fc4 100644 --- a/naga/tests/out/hlsl/math-functions.hlsl +++ b/naga/tests/out/hlsl/math-functions.hlsl @@ -93,11 +93,8 @@ void main() int2 ctz_f = asint(min((32u).xx, firstbitlow((0).xx))); uint2 ctz_g = min((32u).xx, firstbitlow((1u).xx)); 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 _expr67 = (-1).xx; - int2 clz_c = (_expr67 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr67))); - uint2 clz_d = ((31u).xx - firstbithigh((1u).xx)); + int2 clz_c = int2(0, 0); + uint2 clz_d = uint2(31u, 31u); float lde_a = ldexp(1.0, 2); float2 lde_b = ldexp(float2(1.0, 2.0), int2(3, 4)); _modf_result_f32_ modf_a = naga_modf(1.5); diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl index 45fbcd00a1..33fae07f78 100644 --- a/naga/tests/out/msl/math-functions.msl +++ b/naga/tests/out/msl/math-functions.msl @@ -88,10 +88,8 @@ fragment void main_( metal::int2 ctz_f = metal::ctz(metal::int2(0)); metal::uint2 ctz_g = metal::ctz(metal::uint2(1u)); metal::int2 ctz_h = metal::ctz(metal::int2(1)); - int clz_a = metal::clz(-1); - uint clz_b = metal::clz(1u); - metal::int2 clz_c = metal::clz(metal::int2(-1)); - metal::uint2 clz_d = metal::clz(metal::uint2(1u)); + metal::int2 clz_c = metal::int2(0, 0); + metal::uint2 clz_d = metal::uint2(31u, 31u); float lde_a = metal::ldexp(1.0, 2); metal::float2 lde_b = metal::ldexp(metal::float2(1.0, 2.0), metal::int2(3, 4)); _modf_result_f32_ modf_a = naga_modf(1.5); diff --git a/naga/tests/out/spv/math-functions.spvasm b/naga/tests/out/spv/math-functions.spvasm index bbbf370970..8450bdd21e 100644 --- a/naga/tests/out/spv/math-functions.spvasm +++ b/naga/tests/out/spv/math-functions.spvasm @@ -1,145 +1,138 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 125 +; Bound: 118 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %15 "main" -OpExecutionMode %15 OriginUpperLeft -OpMemberDecorate %8 0 Offset 0 -OpMemberDecorate %8 1 Offset 4 -OpMemberDecorate %9 0 Offset 0 -OpMemberDecorate %9 1 Offset 8 +OpEntryPoint Fragment %17 "main" +OpExecutionMode %17 OriginUpperLeft OpMemberDecorate %10 0 Offset 0 -OpMemberDecorate %10 1 Offset 16 +OpMemberDecorate %10 1 Offset 4 OpMemberDecorate %11 0 Offset 0 -OpMemberDecorate %11 1 Offset 4 +OpMemberDecorate %11 1 Offset 8 +OpMemberDecorate %12 0 Offset 0 +OpMemberDecorate %12 1 Offset 16 OpMemberDecorate %13 0 Offset 0 -OpMemberDecorate %13 1 Offset 16 +OpMemberDecorate %13 1 Offset 4 +OpMemberDecorate %15 0 Offset 0 +OpMemberDecorate %15 1 Offset 16 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpTypeVector %4 4 %6 = OpTypeInt 32 1 %5 = OpTypeVector %6 2 -%7 = OpTypeVector %4 2 -%8 = OpTypeStruct %4 %4 -%9 = OpTypeStruct %7 %7 -%10 = OpTypeStruct %3 %3 -%11 = OpTypeStruct %4 %6 -%12 = OpTypeVector %6 4 -%13 = OpTypeStruct %3 %12 -%16 = OpTypeFunction %2 -%17 = OpConstant %4 1.0 -%18 = OpConstant %4 0.0 -%19 = OpConstantComposite %3 %18 %18 %18 %18 -%20 = OpConstant %6 -1 -%21 = OpConstantComposite %12 %20 %20 %20 %20 -%22 = OpConstant %4 -1.0 -%23 = OpConstantComposite %3 %22 %22 %22 %22 -%24 = OpConstantNull %5 -%25 = OpTypeInt 32 0 -%26 = OpConstant %25 0 -%27 = OpConstantComposite %5 %20 %20 -%28 = OpConstant %25 1 -%29 = OpTypeVector %25 2 -%30 = OpConstantComposite %29 %28 %28 +%8 = OpTypeInt 32 0 +%7 = OpTypeVector %8 2 +%9 = OpTypeVector %4 2 +%10 = OpTypeStruct %4 %4 +%11 = OpTypeStruct %9 %9 +%12 = OpTypeStruct %3 %3 +%13 = OpTypeStruct %4 %6 +%14 = OpTypeVector %6 4 +%15 = OpTypeStruct %3 %14 +%18 = OpTypeFunction %2 +%19 = OpConstant %4 1.0 +%20 = OpConstant %4 0.0 +%21 = OpConstantComposite %3 %20 %20 %20 %20 +%22 = OpConstant %6 -1 +%23 = OpConstantComposite %14 %22 %22 %22 %22 +%24 = OpConstant %4 -1.0 +%25 = OpConstantComposite %3 %24 %24 %24 %24 +%26 = OpConstantNull %5 +%27 = OpConstant %8 0 +%28 = OpConstantComposite %5 %22 %22 +%29 = OpConstant %8 1 +%30 = OpConstantComposite %7 %29 %29 %31 = OpConstant %6 0 -%32 = OpConstant %25 4294967295 -%33 = OpConstantComposite %29 %26 %26 +%32 = OpConstant %8 4294967295 +%33 = OpConstantComposite %7 %27 %27 %34 = OpConstantComposite %5 %31 %31 %35 = OpConstant %6 1 %36 = OpConstantComposite %5 %35 %35 -%37 = OpConstant %6 2 -%38 = OpConstant %4 2.0 -%39 = OpConstantComposite %7 %17 %38 -%40 = OpConstant %6 3 -%41 = OpConstant %6 4 -%42 = OpConstantComposite %5 %40 %41 -%43 = OpConstant %4 1.5 -%44 = OpConstantComposite %7 %43 %43 -%45 = OpConstantComposite %3 %43 %43 %43 %43 -%52 = OpConstantComposite %3 %17 %17 %17 %17 -%59 = OpConstantNull %6 -%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 -%46 = OpLabel -%47 = OpExtInst %4 %1 Degrees %17 -%48 = OpExtInst %4 %1 Radians %17 -%49 = OpExtInst %3 %1 Degrees %19 -%50 = OpExtInst %3 %1 Radians %19 -%51 = OpExtInst %3 %1 FClamp %19 %19 %52 -%53 = OpExtInst %3 %1 Refract %19 %19 %17 -%54 = OpExtInst %6 %1 SSign %20 -%55 = OpExtInst %12 %1 SSign %21 -%56 = OpExtInst %4 %1 FSign %22 -%57 = OpExtInst %3 %1 FSign %23 -%60 = OpCompositeExtract %6 %24 0 -%61 = OpCompositeExtract %6 %24 0 -%62 = OpIMul %6 %60 %61 -%63 = OpIAdd %6 %59 %62 -%64 = OpCompositeExtract %6 %24 1 -%65 = OpCompositeExtract %6 %24 1 -%66 = OpIMul %6 %64 %65 -%58 = OpIAdd %6 %63 %66 -%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 = OpCompositeExtract %4 %106 0 -%108 = OpExtInst %8 %1 ModfStruct %43 +%37 = OpConstant %8 31 +%38 = OpConstantComposite %5 %31 %31 +%39 = OpConstantComposite %7 %37 %37 +%40 = OpConstant %6 2 +%41 = OpConstant %4 2.0 +%42 = OpConstantComposite %9 %19 %41 +%43 = OpConstant %6 3 +%44 = OpConstant %6 4 +%45 = OpConstantComposite %5 %43 %44 +%46 = OpConstant %4 1.5 +%47 = OpConstantComposite %9 %46 %46 +%48 = OpConstantComposite %3 %46 %46 %46 %46 +%55 = OpConstantComposite %3 %19 %19 %19 %19 +%62 = OpConstantNull %6 +%79 = OpConstant %8 32 +%88 = OpConstantComposite %7 %79 %79 +%17 = OpFunction %2 None %18 +%16 = OpLabel +OpBranch %49 +%49 = OpLabel +%50 = OpExtInst %4 %1 Degrees %19 +%51 = OpExtInst %4 %1 Radians %19 +%52 = OpExtInst %3 %1 Degrees %21 +%53 = OpExtInst %3 %1 Radians %21 +%54 = OpExtInst %3 %1 FClamp %21 %21 %55 +%56 = OpExtInst %3 %1 Refract %21 %21 %19 +%57 = OpExtInst %6 %1 SSign %22 +%58 = OpExtInst %14 %1 SSign %23 +%59 = OpExtInst %4 %1 FSign %24 +%60 = OpExtInst %3 %1 FSign %25 +%63 = OpCompositeExtract %6 %26 0 +%64 = OpCompositeExtract %6 %26 0 +%65 = OpIMul %6 %63 %64 +%66 = OpIAdd %6 %62 %65 +%67 = OpCompositeExtract %6 %26 1 +%68 = OpCompositeExtract %6 %26 1 +%69 = OpIMul %6 %67 %68 +%61 = OpIAdd %6 %66 %69 +%70 = OpExtInst %8 %1 FindUMsb %27 +%71 = OpExtInst %6 %1 FindSMsb %22 +%72 = OpExtInst %5 %1 FindSMsb %28 +%73 = OpExtInst %7 %1 FindUMsb %30 +%74 = OpExtInst %6 %1 FindILsb %22 +%75 = OpExtInst %8 %1 FindILsb %29 +%76 = OpExtInst %5 %1 FindILsb %28 +%77 = OpExtInst %7 %1 FindILsb %30 +%80 = OpExtInst %8 %1 FindILsb %27 +%78 = OpExtInst %8 %1 UMin %79 %80 +%82 = OpExtInst %6 %1 FindILsb %31 +%81 = OpExtInst %6 %1 UMin %79 %82 +%84 = OpExtInst %8 %1 FindILsb %32 +%83 = OpExtInst %8 %1 UMin %79 %84 +%86 = OpExtInst %6 %1 FindILsb %22 +%85 = OpExtInst %6 %1 UMin %79 %86 +%89 = OpExtInst %7 %1 FindILsb %33 +%87 = OpExtInst %7 %1 UMin %88 %89 +%91 = OpExtInst %5 %1 FindILsb %34 +%90 = OpExtInst %5 %1 UMin %88 %91 +%93 = OpExtInst %7 %1 FindILsb %30 +%92 = OpExtInst %7 %1 UMin %88 %93 +%95 = OpExtInst %5 %1 FindILsb %36 +%94 = OpExtInst %5 %1 UMin %88 %95 +%96 = OpExtInst %4 %1 Ldexp %19 %40 +%97 = OpExtInst %9 %1 Ldexp %42 %45 +%98 = OpExtInst %10 %1 ModfStruct %46 +%99 = OpExtInst %10 %1 ModfStruct %46 +%100 = OpCompositeExtract %4 %99 0 +%101 = OpExtInst %10 %1 ModfStruct %46 +%102 = OpCompositeExtract %4 %101 1 +%103 = OpExtInst %11 %1 ModfStruct %47 +%104 = OpExtInst %12 %1 ModfStruct %48 +%105 = OpCompositeExtract %3 %104 1 +%106 = OpCompositeExtract %4 %105 0 +%107 = OpExtInst %11 %1 ModfStruct %47 +%108 = OpCompositeExtract %9 %107 0 %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 = 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 +%110 = OpExtInst %13 %1 FrexpStruct %46 +%111 = OpExtInst %13 %1 FrexpStruct %46 +%112 = OpCompositeExtract %4 %111 0 +%113 = OpExtInst %13 %1 FrexpStruct %46 +%114 = OpCompositeExtract %6 %113 1 +%115 = OpExtInst %15 %1 FrexpStruct %48 +%116 = OpCompositeExtract %14 %115 1 +%117 = OpCompositeExtract %6 %116 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 ce38fee986..68c2ba329f 100644 --- a/naga/tests/out/wgsl/math-functions.wgsl +++ b/naga/tests/out/wgsl/math-functions.wgsl @@ -28,10 +28,8 @@ fn main() { let ctz_f = countTrailingZeros(vec2(0i)); let ctz_g = countTrailingZeros(vec2(1u)); let ctz_h = countTrailingZeros(vec2(1i)); - let clz_a = countLeadingZeros(-1i); - let clz_b = countLeadingZeros(1u); - let clz_c = countLeadingZeros(vec2(-1i)); - let clz_d = countLeadingZeros(vec2(1u)); + let clz_c = vec2(0i, 0i); + let clz_d = vec2(31u, 31u); let lde_a = ldexp(1f, 2i); let lde_b = ldexp(vec2(1f, 2f), vec2(3i, 4i)); let modf_a = modf(1.5f); From d6704a8652e3025b16ff341fbad4ec159c992d71 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 5 Jan 2024 10:56:24 -0500 Subject: [PATCH 04/20] feat(const_eval): impl. `countOneBits` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 9 +++++++++ 2 files changed, 10 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 7e5f9c099b..7fea5b8aa8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -69,6 +69,7 @@ Bottom level categories: - [#5098](https://github.com/gfx-rs/wgpu/pull/5098) by @ErichDonGubler: - `ceil` - `countLeadingZeros` + - `countOneBits` - `floor` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 88c59809a9..29fced23a8 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -886,6 +886,15 @@ impl<'a> ConstantEvaluator<'a> { .expect("bit count overflowed 32 bits, somehow!?")]) }) } + crate::MathFunction::CountOneBits => { + component_wise_concrete_int!(self, span, [arg], |e| { + #[allow(clippy::useless_conversion)] + Ok([e + .count_ones() + .try_into() + .expect("bit count overflowed 32 bits, somehow!?")]) + }) + } crate::MathFunction::Floor => { component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } From e975f3f82d876da29bde46286cd7ae1f72cffa38 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 5 Jan 2024 10:56:24 -0500 Subject: [PATCH 05/20] feat(const_eval): impl. `countTrailingZeros` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 9 ++ .../glsl/math-functions.main.Fragment.glsl | 12 +-- naga/tests/out/hlsl/math-functions.hlsl | 12 +-- naga/tests/out/msl/math-functions.msl | 12 +-- naga/tests/out/spv/math-functions.spvasm | 82 ++++++++----------- naga/tests/out/wgsl/math-functions.wgsl | 12 +-- 7 files changed, 58 insertions(+), 82 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 7fea5b8aa8..78102f4191 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -70,6 +70,7 @@ Bottom level categories: - `ceil` - `countLeadingZeros` - `countOneBits` + - `countTrailingZeros` - `floor` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 29fced23a8..ea8807d7b3 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -895,6 +895,15 @@ impl<'a> ConstantEvaluator<'a> { .expect("bit count overflowed 32 bits, somehow!?")]) }) } + crate::MathFunction::CountTrailingZeros => { + component_wise_concrete_int!(self, span, [arg], |e| { + #[allow(clippy::useless_conversion)] + Ok([e + .trailing_zeros() + .try_into() + .expect("bit count overflowed 32 bits, somehow!?")]) + }) + } crate::MathFunction::Floor => { component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl index 2892157801..a68b613a3b 100644 --- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/naga/tests/out/glsl/math-functions.main.Fragment.glsl @@ -75,14 +75,10 @@ void main() { uint ftb_b = uint(findLSB(1u)); ivec2 ftb_c = findLSB(ivec2(-1)); uvec2 ftb_d = uvec2(findLSB(uvec2(1u))); - uint ctz_a = min(uint(findLSB(0u)), 32u); - int ctz_b = int(min(uint(findLSB(0)), 32u)); - uint ctz_c = min(uint(findLSB(4294967295u)), 32u); - int ctz_d = int(min(uint(findLSB(-1)), 32u)); - uvec2 ctz_e = min(uvec2(findLSB(uvec2(0u))), uvec2(32u)); - ivec2 ctz_f = ivec2(min(uvec2(findLSB(ivec2(0))), uvec2(32u))); - uvec2 ctz_g = min(uvec2(findLSB(uvec2(1u))), uvec2(32u)); - ivec2 ctz_h = ivec2(min(uvec2(findLSB(ivec2(1))), uvec2(32u))); + uvec2 ctz_e = uvec2(32u, 32u); + ivec2 ctz_f = ivec2(32, 32); + uvec2 ctz_g = uvec2(0u, 0u); + ivec2 ctz_h = ivec2(0, 0); ivec2 clz_c = ivec2(0, 0); uvec2 clz_d = uvec2(31u, 31u); float lde_a = ldexp(1.0, 2); diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl index d576365fc4..b3b938da6b 100644 --- a/naga/tests/out/hlsl/math-functions.hlsl +++ b/naga/tests/out/hlsl/math-functions.hlsl @@ -85,14 +85,10 @@ void main() uint ftb_b = firstbitlow(1u); int2 ftb_c = asint(firstbitlow((-1).xx)); uint2 ftb_d = firstbitlow((1u).xx); - uint ctz_a = min(32u, firstbitlow(0u)); - int ctz_b = asint(min(32u, firstbitlow(0))); - uint ctz_c = min(32u, firstbitlow(4294967295u)); - int ctz_d = asint(min(32u, firstbitlow(-1))); - uint2 ctz_e = min((32u).xx, firstbitlow((0u).xx)); - int2 ctz_f = asint(min((32u).xx, firstbitlow((0).xx))); - uint2 ctz_g = min((32u).xx, firstbitlow((1u).xx)); - int2 ctz_h = asint(min((32u).xx, firstbitlow((1).xx))); + uint2 ctz_e = uint2(32u, 32u); + int2 ctz_f = int2(32, 32); + uint2 ctz_g = uint2(0u, 0u); + int2 ctz_h = int2(0, 0); int2 clz_c = int2(0, 0); uint2 clz_d = uint2(31u, 31u); float lde_a = ldexp(1.0, 2); diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl index 33fae07f78..a9d37076e7 100644 --- a/naga/tests/out/msl/math-functions.msl +++ b/naga/tests/out/msl/math-functions.msl @@ -80,14 +80,10 @@ fragment void main_( uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1); metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1); metal::uint2 ftb_d = (((metal::ctz(metal::uint2(1u)) + 1) % 33) - 1); - uint ctz_a = metal::ctz(0u); - int ctz_b = metal::ctz(0); - uint ctz_c = metal::ctz(4294967295u); - int ctz_d = metal::ctz(-1); - metal::uint2 ctz_e = metal::ctz(metal::uint2(0u)); - metal::int2 ctz_f = metal::ctz(metal::int2(0)); - metal::uint2 ctz_g = metal::ctz(metal::uint2(1u)); - metal::int2 ctz_h = metal::ctz(metal::int2(1)); + metal::uint2 ctz_e = metal::uint2(32u, 32u); + metal::int2 ctz_f = metal::int2(32, 32); + metal::uint2 ctz_g = metal::uint2(0u, 0u); + metal::int2 ctz_h = metal::int2(0, 0); metal::int2 clz_c = metal::int2(0, 0); metal::uint2 clz_d = metal::uint2(31u, 31u); float lde_a = metal::ldexp(1.0, 2); diff --git a/naga/tests/out/spv/math-functions.spvasm b/naga/tests/out/spv/math-functions.spvasm index 8450bdd21e..442b0ba0d0 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: 118 +; Bound: 100 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -44,15 +44,15 @@ OpMemberDecorate %15 1 Offset 16 %28 = OpConstantComposite %5 %22 %22 %29 = OpConstant %8 1 %30 = OpConstantComposite %7 %29 %29 -%31 = OpConstant %6 0 -%32 = OpConstant %8 4294967295 -%33 = OpConstantComposite %7 %27 %27 -%34 = OpConstantComposite %5 %31 %31 -%35 = OpConstant %6 1 -%36 = OpConstantComposite %5 %35 %35 -%37 = OpConstant %8 31 -%38 = OpConstantComposite %5 %31 %31 -%39 = OpConstantComposite %7 %37 %37 +%31 = OpConstant %8 32 +%32 = OpConstant %6 32 +%33 = OpConstant %6 0 +%34 = OpConstantComposite %7 %31 %31 +%35 = OpConstantComposite %5 %32 %32 +%36 = OpConstantComposite %7 %27 %27 +%37 = OpConstantComposite %5 %33 %33 +%38 = OpConstant %8 31 +%39 = OpConstantComposite %7 %38 %38 %40 = OpConstant %6 2 %41 = OpConstant %4 2.0 %42 = OpConstantComposite %9 %19 %41 @@ -64,8 +64,6 @@ OpMemberDecorate %15 1 Offset 16 %48 = OpConstantComposite %3 %46 %46 %46 %46 %55 = OpConstantComposite %3 %19 %19 %19 %19 %62 = OpConstantNull %6 -%79 = OpConstant %8 32 -%88 = OpConstantComposite %7 %79 %79 %17 = OpFunction %2 None %18 %16 = OpLabel OpBranch %49 @@ -96,43 +94,27 @@ OpBranch %49 %75 = OpExtInst %8 %1 FindILsb %29 %76 = OpExtInst %5 %1 FindILsb %28 %77 = OpExtInst %7 %1 FindILsb %30 -%80 = OpExtInst %8 %1 FindILsb %27 -%78 = OpExtInst %8 %1 UMin %79 %80 -%82 = OpExtInst %6 %1 FindILsb %31 -%81 = OpExtInst %6 %1 UMin %79 %82 -%84 = OpExtInst %8 %1 FindILsb %32 -%83 = OpExtInst %8 %1 UMin %79 %84 -%86 = OpExtInst %6 %1 FindILsb %22 -%85 = OpExtInst %6 %1 UMin %79 %86 -%89 = OpExtInst %7 %1 FindILsb %33 -%87 = OpExtInst %7 %1 UMin %88 %89 -%91 = OpExtInst %5 %1 FindILsb %34 -%90 = OpExtInst %5 %1 UMin %88 %91 -%93 = OpExtInst %7 %1 FindILsb %30 -%92 = OpExtInst %7 %1 UMin %88 %93 -%95 = OpExtInst %5 %1 FindILsb %36 -%94 = OpExtInst %5 %1 UMin %88 %95 -%96 = OpExtInst %4 %1 Ldexp %19 %40 -%97 = OpExtInst %9 %1 Ldexp %42 %45 -%98 = OpExtInst %10 %1 ModfStruct %46 -%99 = OpExtInst %10 %1 ModfStruct %46 -%100 = OpCompositeExtract %4 %99 0 -%101 = OpExtInst %10 %1 ModfStruct %46 -%102 = OpCompositeExtract %4 %101 1 -%103 = OpExtInst %11 %1 ModfStruct %47 -%104 = OpExtInst %12 %1 ModfStruct %48 -%105 = OpCompositeExtract %3 %104 1 -%106 = OpCompositeExtract %4 %105 0 -%107 = OpExtInst %11 %1 ModfStruct %47 -%108 = OpCompositeExtract %9 %107 0 -%109 = OpCompositeExtract %4 %108 1 -%110 = OpExtInst %13 %1 FrexpStruct %46 -%111 = OpExtInst %13 %1 FrexpStruct %46 -%112 = OpCompositeExtract %4 %111 0 -%113 = OpExtInst %13 %1 FrexpStruct %46 -%114 = OpCompositeExtract %6 %113 1 -%115 = OpExtInst %15 %1 FrexpStruct %48 -%116 = OpCompositeExtract %14 %115 1 -%117 = OpCompositeExtract %6 %116 0 +%78 = OpExtInst %4 %1 Ldexp %19 %40 +%79 = OpExtInst %9 %1 Ldexp %42 %45 +%80 = OpExtInst %10 %1 ModfStruct %46 +%81 = OpExtInst %10 %1 ModfStruct %46 +%82 = OpCompositeExtract %4 %81 0 +%83 = OpExtInst %10 %1 ModfStruct %46 +%84 = OpCompositeExtract %4 %83 1 +%85 = OpExtInst %11 %1 ModfStruct %47 +%86 = OpExtInst %12 %1 ModfStruct %48 +%87 = OpCompositeExtract %3 %86 1 +%88 = OpCompositeExtract %4 %87 0 +%89 = OpExtInst %11 %1 ModfStruct %47 +%90 = OpCompositeExtract %9 %89 0 +%91 = OpCompositeExtract %4 %90 1 +%92 = OpExtInst %13 %1 FrexpStruct %46 +%93 = OpExtInst %13 %1 FrexpStruct %46 +%94 = OpCompositeExtract %4 %93 0 +%95 = OpExtInst %13 %1 FrexpStruct %46 +%96 = OpCompositeExtract %6 %95 1 +%97 = OpExtInst %15 %1 FrexpStruct %48 +%98 = OpCompositeExtract %14 %97 1 +%99 = OpCompositeExtract %6 %98 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 68c2ba329f..8b8330738b 100644 --- a/naga/tests/out/wgsl/math-functions.wgsl +++ b/naga/tests/out/wgsl/math-functions.wgsl @@ -20,14 +20,10 @@ fn main() { let ftb_b = firstTrailingBit(1u); let ftb_c = firstTrailingBit(vec2(-1i)); let ftb_d = firstTrailingBit(vec2(1u)); - let ctz_a = countTrailingZeros(0u); - let ctz_b = countTrailingZeros(0i); - let ctz_c = countTrailingZeros(4294967295u); - let ctz_d = countTrailingZeros(-1i); - let ctz_e = countTrailingZeros(vec2(0u)); - let ctz_f = countTrailingZeros(vec2(0i)); - let ctz_g = countTrailingZeros(vec2(1u)); - let ctz_h = countTrailingZeros(vec2(1i)); + let ctz_e = vec2(32u, 32u); + let ctz_f = vec2(32i, 32i); + let ctz_g = vec2(0u, 0u); + let ctz_h = vec2(0i, 0i); let clz_c = vec2(0i, 0i); let clz_d = vec2(31u, 31u); let lde_a = ldexp(1f, 2i); From d5c163727ba27367e108b7597e955520dfd2a8fe Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 5 Jan 2024 10:56:24 -0500 Subject: [PATCH 06/20] feat(const_eval): impl. `exp` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 78102f4191..049d3522de 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -71,6 +71,7 @@ Bottom level categories: - `countLeadingZeros` - `countOneBits` - `countTrailingZeros` + - `exp` - `floor` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index ea8807d7b3..a97a6e6a53 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -904,6 +904,9 @@ impl<'a> ConstantEvaluator<'a> { .expect("bit count overflowed 32 bits, somehow!?")]) }) } + crate::MathFunction::Exp => { + component_wise_float!(self, span, [arg], |e| { Ok([e.exp()]) }) + } crate::MathFunction::Floor => { component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } From 438b85eb526742ec0b81f826cad56d6868095549 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 18 Jan 2024 17:39:21 -0500 Subject: [PATCH 07/20] feat(const_eval): impl. `degrees` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 049d3522de..e7f3a55f26 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -71,6 +71,7 @@ Bottom level categories: - `countLeadingZeros` - `countOneBits` - `countTrailingZeros` + - `degrees` - `exp` - `floor` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index a97a6e6a53..0a8e4c0971 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -904,6 +904,9 @@ impl<'a> ConstantEvaluator<'a> { .expect("bit count overflowed 32 bits, somehow!?")]) }) } + crate::MathFunction::Degrees => { + component_wise_float!(self, span, [arg], |e| { Ok([e.to_degrees()]) }) + } crate::MathFunction::Exp => { component_wise_float!(self, span, [arg], |e| { Ok([e.exp()]) }) } From 85478b1045fe26adfdac80d258ca22aba1dcaff0 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 5 Jan 2024 10:56:24 -0500 Subject: [PATCH 08/20] feat(const_eval): impl. `exp2` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index e7f3a55f26..52d0c9ce78 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -73,6 +73,7 @@ Bottom level categories: - `countTrailingZeros` - `degrees` - `exp` + - `exp2` - `floor` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 0a8e4c0971..66cccc947a 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -910,6 +910,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Exp => { component_wise_float!(self, span, [arg], |e| { Ok([e.exp()]) }) } + crate::MathFunction::Exp2 => { + component_wise_float!(self, span, [arg], |e| { Ok([e.exp2()]) }) + } crate::MathFunction::Floor => { component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } From feca852bad6eb1d9e16b3537f53d60f7eb4a3831 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:13:43 -0500 Subject: [PATCH 09/20] feat(const_eval): impl. `fma` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 8 ++++++++ 2 files changed, 9 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 52d0c9ce78..714b0dfa01 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -75,6 +75,7 @@ Bottom level categories: - `exp` - `exp2` - `floor` + - `fma` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 66cccc947a..622fd46896 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -916,6 +916,14 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Floor => { component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } + crate::MathFunction::Fma => { + component_wise_float!( + self, + span, + [arg, arg1.unwrap(), arg2.unwrap()], + |e1, e2, e3| { Ok([e1.mul_add(e2, e3)]) } + ) + } 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 264d437983fd645ffa3a95f9a70cd4c1ff02f9f2 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:13:55 -0500 Subject: [PATCH 10/20] feat(const_eval): impl. `inverseSqrt` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 714b0dfa01..db9f741f58 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -76,6 +76,7 @@ Bottom level categories: - `exp2` - `floor` - `fma` + - `inverseSqrt` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 622fd46896..cf83e64264 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -924,6 +924,9 @@ impl<'a> ConstantEvaluator<'a> { |e1, e2, e3| { Ok([e1.mul_add(e2, e3)]) } ) } + crate::MathFunction::InverseSqrt => { + component_wise_float!(self, span, [arg], |e| { Ok([1. / e.sqrt()]) }) + } 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 a4fa846f14391d354010fce3c584c08bddc16a6e Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Thu, 18 Jan 2024 17:10:08 -0500 Subject: [PATCH 11/20] feat(const_eval): impl. `fract` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 7 +++++++ 2 files changed, 8 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index db9f741f58..0a288a0654 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -75,6 +75,7 @@ Bottom level categories: - `exp` - `exp2` - `floor` + - `fract` - `fma` - `inverseSqrt` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index cf83e64264..2398068691 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -916,6 +916,13 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Floor => { component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } + crate::MathFunction::Fract => { + component_wise_float!(self, span, [arg], |e| { + // N.B., Rust's definition of `fract` is `e - e.trunc()`, so we can't use that + // here. + Ok([e - e.floor()]) + }) + } crate::MathFunction::Fma => { component_wise_float!( self, From fdf7ba882a4bff05d0b43c01cc29fb5d35a30f16 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:14:08 -0500 Subject: [PATCH 12/20] feat(const_eval): impl. `log` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0a288a0654..79b2deabe3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -78,6 +78,7 @@ Bottom level categories: - `fract` - `fma` - `inverseSqrt` + - `log` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 2398068691..96192ec827 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -934,6 +934,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::InverseSqrt => { component_wise_float!(self, span, [arg], |e| { Ok([1. / e.sqrt()]) }) } + crate::MathFunction::Log => { + component_wise_float!(self, span, [arg], |e| { Ok([e.ln()]) }) + } 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 72153ea50217c378fe67d19ef01d8f973838ae8a Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:14:18 -0500 Subject: [PATCH 13/20] feat(const_eval): impl. `log2` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 79b2deabe3..261671de6e 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -79,6 +79,7 @@ Bottom level categories: - `fma` - `inverseSqrt` - `log` + - `log2` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 96192ec827..1e132f66d4 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -937,6 +937,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Log => { component_wise_float!(self, span, [arg], |e| { Ok([e.ln()]) }) } + crate::MathFunction::Log2 => { + component_wise_float!(self, span, [arg], |e| { Ok([e.log2()]) }) + } 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 457de810c377a578bffe454a052f4dd4b06b06d9 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:14:30 -0500 Subject: [PATCH 14/20] feat(const_eval): impl. `radians` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 261671de6e..b739e147c3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -80,6 +80,7 @@ Bottom level categories: - `inverseSqrt` - `log` - `log2` + - `radians` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 1e132f66d4..63b2ea0656 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -940,6 +940,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Log2 => { component_wise_float!(self, span, [arg], |e| { Ok([e.log2()]) }) } + crate::MathFunction::Radians => { + component_wise_float!(self, span, [arg], |e1| { Ok([e1.to_radians()]) }) + } 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 53f134ca2d19a31e04d191bf5c1bfda2be511d07 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:14:40 -0500 Subject: [PATCH 15/20] feat(const_eval): impl. `reverseBits` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index b739e147c3..42281bc964 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -81,6 +81,7 @@ Bottom level categories: - `log` - `log2` - `radians` + - `reverseBits` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 63b2ea0656..585daf1fe3 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -943,6 +943,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Radians => { component_wise_float!(self, span, [arg], |e1| { Ok([e1.to_radians()]) }) } + crate::MathFunction::ReverseBits => { + component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) }) + } 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 a1bb6965955a15e9fc0ca4490ef481cdc92e9ecf Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:14:51 -0500 Subject: [PATCH 16/20] feat(const_eval): impl. `sign` with new `component_wise_signed` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 19 +++ .../glsl/math-functions.main.Fragment.glsl | 6 +- naga/tests/out/hlsl/math-functions.hlsl | 6 +- naga/tests/out/msl/math-functions.msl | 15 +- naga/tests/out/spv/math-functions.spvasm | 144 +++++++++--------- naga/tests/out/wgsl/math-functions.wgsl | 6 +- 7 files changed, 102 insertions(+), 95 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 42281bc964..b11712140c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -82,6 +82,7 @@ Bottom level categories: - `log2` - `radians` - `reverseBits` + - `sign` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 585daf1fe3..6a2991fa66 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -234,6 +234,22 @@ gen_component_wise_extractor! { ], } +gen_component_wise_extractor! { + component_wise_signed -> Signed, + literals: [ + AbstractFloat => AbstractFloat: f64, + AbstractInt => AbstractInt: i64, + F32 => F32: f32, + I32 => I32: i32, + ], + scalar_kinds: [ + Sint, + AbstractInt, + Float, + AbstractFloat, + ], +} + #[derive(Debug)] enum Behavior { Wgsl, @@ -975,6 +991,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Saturate => { component_wise_float!(self, span, [arg], |e| { Ok([e.clamp(0., 1.)]) }) } + crate::MathFunction::Sign => { + component_wise_signed!(self, span, [arg], |e| { Ok([e.signum()]) }) + } crate::MathFunction::Sin => { component_wise_float!(self, span, [arg], |e| { Ok([e.sin()]) }) } diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl index a68b613a3b..7f91571dcc 100644 --- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/naga/tests/out/glsl/math-functions.main.Fragment.glsl @@ -62,10 +62,8 @@ void main() { vec4 d = radians(v); vec4 e = clamp(v, vec4(0.0), vec4(1.0)); vec4 g = refract(v, v, 1.0); - int sign_a = sign(-1); - ivec4 sign_b = sign(ivec4(-1)); - float sign_c = sign(-1.0); - vec4 sign_d = sign(vec4(-1.0)); + ivec4 sign_b = ivec4(-1, -1, -1, -1); + vec4 sign_d = vec4(-1.0, -1.0, -1.0, -1.0); int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y); uint first_leading_bit_abs = uint(findMSB(0u)); int flb_a = findMSB(-1); diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl index b3b938da6b..61c59f00c1 100644 --- a/naga/tests/out/hlsl/math-functions.hlsl +++ b/naga/tests/out/hlsl/math-functions.hlsl @@ -72,10 +72,8 @@ void main() float4 d = radians(v); float4 e = saturate(v); float4 g = refract(v, v, 1.0); - int sign_a = sign(-1); - int4 sign_b = sign((-1).xxxx); - float sign_c = sign(-1.0); - float4 sign_d = sign((-1.0).xxxx); + int4 sign_b = int4(-1, -1, -1, -1); + float4 sign_d = float4(-1.0, -1.0, -1.0, -1.0); int const_dot = dot((int2)0, (int2)0); uint first_leading_bit_abs = firstbithigh(0u); int flb_a = asint(firstbithigh(-1)); diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl index a9d37076e7..0e6a5b24dc 100644 --- a/naga/tests/out/msl/math-functions.msl +++ b/naga/tests/out/msl/math-functions.msl @@ -64,18 +64,15 @@ fragment void main_( metal::float4 d = ((v) * 0.017453292519943295474); metal::float4 e = metal::saturate(v); metal::float4 g = metal::refract(v, v, 1.0); - int sign_a = metal::select(metal::select(-1, 1, (-1 > 0)), 0, (-1 == 0)); - metal::int4 _e12 = metal::int4(-1); - metal::int4 sign_b = metal::select(metal::select(int4(-1), int4(1), (_e12 > 0)), 0, (_e12 == 0)); - float sign_c = metal::sign(-1.0); - metal::float4 sign_d = metal::sign(metal::float4(-1.0)); + metal::int4 sign_b = metal::int4(-1, -1, -1, -1); + metal::float4 sign_d = metal::float4(-1.0, -1.0, -1.0, -1.0); int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); 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 _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); + metal::int2 _e29 = metal::int2(-1); + metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e29, ~_e29, _e29 < 0)), int2(-1), _e29 == 0 || _e29 == -1); + metal::uint2 _e32 = metal::uint2(1u); + metal::uint2 flb_c = metal::select(31 - metal::clz(_e32), uint2(-1), _e32 == 0 || _e32 == -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 442b0ba0d0..6e07c6d7a6 100644 --- a/naga/tests/out/spv/math-functions.spvasm +++ b/naga/tests/out/spv/math-functions.spvasm @@ -1,69 +1,69 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 100 +; Bound: 96 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Fragment %17 "main" OpExecutionMode %17 OriginUpperLeft -OpMemberDecorate %10 0 Offset 0 -OpMemberDecorate %10 1 Offset 4 OpMemberDecorate %11 0 Offset 0 -OpMemberDecorate %11 1 Offset 8 +OpMemberDecorate %11 1 Offset 4 OpMemberDecorate %12 0 Offset 0 -OpMemberDecorate %12 1 Offset 16 +OpMemberDecorate %12 1 Offset 8 OpMemberDecorate %13 0 Offset 0 -OpMemberDecorate %13 1 Offset 4 +OpMemberDecorate %13 1 Offset 16 +OpMemberDecorate %14 0 Offset 0 +OpMemberDecorate %14 1 Offset 4 OpMemberDecorate %15 0 Offset 0 OpMemberDecorate %15 1 Offset 16 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpTypeVector %4 4 %6 = OpTypeInt 32 1 -%5 = OpTypeVector %6 2 -%8 = OpTypeInt 32 0 -%7 = OpTypeVector %8 2 -%9 = OpTypeVector %4 2 -%10 = OpTypeStruct %4 %4 -%11 = OpTypeStruct %9 %9 -%12 = OpTypeStruct %3 %3 -%13 = OpTypeStruct %4 %6 -%14 = OpTypeVector %6 4 -%15 = OpTypeStruct %3 %14 +%5 = OpTypeVector %6 4 +%7 = OpTypeVector %6 2 +%9 = OpTypeInt 32 0 +%8 = OpTypeVector %9 2 +%10 = OpTypeVector %4 2 +%11 = OpTypeStruct %4 %4 +%12 = OpTypeStruct %10 %10 +%13 = OpTypeStruct %3 %3 +%14 = OpTypeStruct %4 %6 +%15 = OpTypeStruct %3 %5 %18 = OpTypeFunction %2 %19 = OpConstant %4 1.0 %20 = OpConstant %4 0.0 %21 = OpConstantComposite %3 %20 %20 %20 %20 %22 = OpConstant %6 -1 -%23 = OpConstantComposite %14 %22 %22 %22 %22 +%23 = OpConstantComposite %5 %22 %22 %22 %22 %24 = OpConstant %4 -1.0 %25 = OpConstantComposite %3 %24 %24 %24 %24 -%26 = OpConstantNull %5 -%27 = OpConstant %8 0 -%28 = OpConstantComposite %5 %22 %22 -%29 = OpConstant %8 1 -%30 = OpConstantComposite %7 %29 %29 -%31 = OpConstant %8 32 +%26 = OpConstantNull %7 +%27 = OpConstant %9 0 +%28 = OpConstantComposite %7 %22 %22 +%29 = OpConstant %9 1 +%30 = OpConstantComposite %8 %29 %29 +%31 = OpConstant %9 32 %32 = OpConstant %6 32 %33 = OpConstant %6 0 -%34 = OpConstantComposite %7 %31 %31 -%35 = OpConstantComposite %5 %32 %32 -%36 = OpConstantComposite %7 %27 %27 -%37 = OpConstantComposite %5 %33 %33 -%38 = OpConstant %8 31 -%39 = OpConstantComposite %7 %38 %38 +%34 = OpConstantComposite %8 %31 %31 +%35 = OpConstantComposite %7 %32 %32 +%36 = OpConstantComposite %8 %27 %27 +%37 = OpConstantComposite %7 %33 %33 +%38 = OpConstant %9 31 +%39 = OpConstantComposite %8 %38 %38 %40 = OpConstant %6 2 %41 = OpConstant %4 2.0 -%42 = OpConstantComposite %9 %19 %41 +%42 = OpConstantComposite %10 %19 %41 %43 = OpConstant %6 3 %44 = OpConstant %6 4 -%45 = OpConstantComposite %5 %43 %44 +%45 = OpConstantComposite %7 %43 %44 %46 = OpConstant %4 1.5 -%47 = OpConstantComposite %9 %46 %46 +%47 = OpConstantComposite %10 %46 %46 %48 = OpConstantComposite %3 %46 %46 %46 %46 %55 = OpConstantComposite %3 %19 %19 %19 %19 -%62 = OpConstantNull %6 +%58 = OpConstantNull %6 %17 = OpFunction %2 None %18 %16 = OpLabel OpBranch %49 @@ -74,47 +74,43 @@ OpBranch %49 %53 = OpExtInst %3 %1 Radians %21 %54 = OpExtInst %3 %1 FClamp %21 %21 %55 %56 = OpExtInst %3 %1 Refract %21 %21 %19 -%57 = OpExtInst %6 %1 SSign %22 -%58 = OpExtInst %14 %1 SSign %23 -%59 = OpExtInst %4 %1 FSign %24 -%60 = OpExtInst %3 %1 FSign %25 -%63 = OpCompositeExtract %6 %26 0 -%64 = OpCompositeExtract %6 %26 0 +%59 = OpCompositeExtract %6 %26 0 +%60 = OpCompositeExtract %6 %26 0 +%61 = OpIMul %6 %59 %60 +%62 = OpIAdd %6 %58 %61 +%63 = OpCompositeExtract %6 %26 1 +%64 = OpCompositeExtract %6 %26 1 %65 = OpIMul %6 %63 %64 -%66 = OpIAdd %6 %62 %65 -%67 = OpCompositeExtract %6 %26 1 -%68 = OpCompositeExtract %6 %26 1 -%69 = OpIMul %6 %67 %68 -%61 = OpIAdd %6 %66 %69 -%70 = OpExtInst %8 %1 FindUMsb %27 -%71 = OpExtInst %6 %1 FindSMsb %22 -%72 = OpExtInst %5 %1 FindSMsb %28 -%73 = OpExtInst %7 %1 FindUMsb %30 -%74 = OpExtInst %6 %1 FindILsb %22 -%75 = OpExtInst %8 %1 FindILsb %29 -%76 = OpExtInst %5 %1 FindILsb %28 -%77 = OpExtInst %7 %1 FindILsb %30 -%78 = OpExtInst %4 %1 Ldexp %19 %40 -%79 = OpExtInst %9 %1 Ldexp %42 %45 -%80 = OpExtInst %10 %1 ModfStruct %46 -%81 = OpExtInst %10 %1 ModfStruct %46 -%82 = OpCompositeExtract %4 %81 0 -%83 = OpExtInst %10 %1 ModfStruct %46 -%84 = OpCompositeExtract %4 %83 1 -%85 = OpExtInst %11 %1 ModfStruct %47 -%86 = OpExtInst %12 %1 ModfStruct %48 -%87 = OpCompositeExtract %3 %86 1 -%88 = OpCompositeExtract %4 %87 0 -%89 = OpExtInst %11 %1 ModfStruct %47 -%90 = OpCompositeExtract %9 %89 0 -%91 = OpCompositeExtract %4 %90 1 -%92 = OpExtInst %13 %1 FrexpStruct %46 -%93 = OpExtInst %13 %1 FrexpStruct %46 -%94 = OpCompositeExtract %4 %93 0 -%95 = OpExtInst %13 %1 FrexpStruct %46 -%96 = OpCompositeExtract %6 %95 1 -%97 = OpExtInst %15 %1 FrexpStruct %48 -%98 = OpCompositeExtract %14 %97 1 -%99 = OpCompositeExtract %6 %98 0 +%57 = OpIAdd %6 %62 %65 +%66 = OpExtInst %9 %1 FindUMsb %27 +%67 = OpExtInst %6 %1 FindSMsb %22 +%68 = OpExtInst %7 %1 FindSMsb %28 +%69 = OpExtInst %8 %1 FindUMsb %30 +%70 = OpExtInst %6 %1 FindILsb %22 +%71 = OpExtInst %9 %1 FindILsb %29 +%72 = OpExtInst %7 %1 FindILsb %28 +%73 = OpExtInst %8 %1 FindILsb %30 +%74 = OpExtInst %4 %1 Ldexp %19 %40 +%75 = OpExtInst %10 %1 Ldexp %42 %45 +%76 = OpExtInst %11 %1 ModfStruct %46 +%77 = OpExtInst %11 %1 ModfStruct %46 +%78 = OpCompositeExtract %4 %77 0 +%79 = OpExtInst %11 %1 ModfStruct %46 +%80 = OpCompositeExtract %4 %79 1 +%81 = OpExtInst %12 %1 ModfStruct %47 +%82 = OpExtInst %13 %1 ModfStruct %48 +%83 = OpCompositeExtract %3 %82 1 +%84 = OpCompositeExtract %4 %83 0 +%85 = OpExtInst %12 %1 ModfStruct %47 +%86 = OpCompositeExtract %10 %85 0 +%87 = OpCompositeExtract %4 %86 1 +%88 = OpExtInst %14 %1 FrexpStruct %46 +%89 = OpExtInst %14 %1 FrexpStruct %46 +%90 = OpCompositeExtract %4 %89 0 +%91 = OpExtInst %14 %1 FrexpStruct %46 +%92 = OpCompositeExtract %6 %91 1 +%93 = OpExtInst %15 %1 FrexpStruct %48 +%94 = OpCompositeExtract %5 %93 1 +%95 = OpCompositeExtract %6 %94 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 8b8330738b..228248b3ce 100644 --- a/naga/tests/out/wgsl/math-functions.wgsl +++ b/naga/tests/out/wgsl/math-functions.wgsl @@ -7,10 +7,8 @@ fn main() { let d = radians(v); let e = saturate(v); let g = refract(v, v, 1f); - let sign_a = sign(-1i); - let sign_b = sign(vec4(-1i)); - let sign_c = sign(-1f); - let sign_d = sign(vec4(-1f)); + let sign_b = vec4(-1i, -1i, -1i, -1i); + let sign_d = vec4(-1f, -1f, -1f, -1f); let const_dot = dot(vec2(), vec2()); let first_leading_bit_abs = firstLeadingBit(0u); let flb_a = firstLeadingBit(-1i); From a6fd518b93a57dc9311f4237dcad22e4def1ef50 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:14:59 -0500 Subject: [PATCH 17/20] feat(const_eval): impl. `trunc` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 3 +++ 2 files changed, 4 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index b11712140c..bf304336ed 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -83,6 +83,7 @@ Bottom level categories: - `radians` - `reverseBits` - `sign` + - `trunc` - Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 6a2991fa66..8a69acc820 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1014,6 +1014,9 @@ impl<'a> ConstantEvaluator<'a> { Ok([if edge <= x { 1.0 } else { 0.0 }]) }) } + crate::MathFunction::Trunc => { + component_wise_float!(self, span, [arg], |e| { Ok([e.trunc()]) }) + } fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" ))), From 7d41f29cf2cd4b9986aff3dcf3d95efe8fdc3b99 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:17:11 -0500 Subject: [PATCH 18/20] feat(const_eval): impl. `max` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 5 +++++ 2 files changed, 6 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index bf304336ed..8f0e07ed02 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -80,6 +80,7 @@ Bottom level categories: - `inverseSqrt` - `log` - `log2` + - `max` - `radians` - `reverseBits` - `sign` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 8a69acc820..062209a4e6 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -956,6 +956,11 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::Log2 => { component_wise_float!(self, span, [arg], |e| { Ok([e.log2()]) }) } + crate::MathFunction::Max => { + component_wise_scalar!(self, span, [arg, arg1.unwrap()], |e1, e2| { + Ok([e1.max(e2)]) + }) + } crate::MathFunction::Radians => { component_wise_float!(self, span, [arg], |e1| { Ok([e1.to_radians()]) }) } From 27f2d93c3e7d3fa7fdb457bfdc7a783b30ef8f21 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:17:31 -0500 Subject: [PATCH 19/20] feat(const_eval): impl. `min` --- CHANGELOG.md | 1 + naga/src/proc/constant_evaluator.rs | 5 +++++ 2 files changed, 6 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 8f0e07ed02..7a8b419045 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -81,6 +81,7 @@ Bottom level categories: - `log` - `log2` - `max` + - `min` - `radians` - `reverseBits` - `sign` diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 062209a4e6..95d8484392 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -961,6 +961,11 @@ impl<'a> ConstantEvaluator<'a> { Ok([e1.max(e2)]) }) } + crate::MathFunction::Min => { + component_wise_scalar!(self, span, [arg, arg1.unwrap()], |e1, e2| { + Ok([e1.min(e2)]) + }) + } crate::MathFunction::Radians => { component_wise_float!(self, span, [arg], |e1| { Ok([e1.to_radians()]) }) } From 97534787a1cea59521afa115c34bbdd073388fd8 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:43:30 -0500 Subject: [PATCH 20/20] style(const_eval): match variant decl. order of current `MathFunction` impls. --- naga/src/proc/constant_evaluator.rs | 221 +++++++++++++++------------- 1 file changed, 117 insertions(+), 104 deletions(-) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 95d8484392..b3884b04b1 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -837,7 +837,9 @@ impl<'a> ConstantEvaluator<'a> { )); } + // NOTE: We try to match the declaration order of `MathFunction` here. match fun { + // comparison crate::MathFunction::Abs => { component_wise_scalar(self, span, [arg], |args| match args { Scalar::AbstractFloat([e]) => Ok(Scalar::AbstractFloat([e.abs()])), @@ -847,31 +849,15 @@ 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::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::Min => { + component_wise_scalar!(self, span, [arg, arg1.unwrap()], |e1, e2| { + Ok([e1.min(e2)]) }) } - crate::MathFunction::Ceil => { - component_wise_float!(self, span, [arg], |e| { Ok([e.ceil()]) }) + crate::MathFunction::Max => { + component_wise_scalar!(self, span, [arg, arg1.unwrap()], |e1, e2| { + Ok([e1.max(e2)]) + }) } crate::MathFunction::Clamp => { component_wise_scalar!( @@ -887,90 +873,60 @@ impl<'a> ConstantEvaluator<'a> { } ) } + crate::MathFunction::Saturate => { + component_wise_float!(self, span, [arg], |e| { Ok([e.clamp(0., 1.)]) }) + } + + // trigonometry 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::CountLeadingZeros => { - component_wise_concrete_int!(self, span, [arg], |e| { - #[allow(clippy::useless_conversion)] - Ok([e - .leading_zeros() - .try_into() - .expect("bit count overflowed 32 bits, somehow!?")]) - }) - } - crate::MathFunction::CountOneBits => { - component_wise_concrete_int!(self, span, [arg], |e| { - #[allow(clippy::useless_conversion)] - Ok([e - .count_ones() - .try_into() - .expect("bit count overflowed 32 bits, somehow!?")]) - }) - } - crate::MathFunction::CountTrailingZeros => { - component_wise_concrete_int!(self, span, [arg], |e| { - #[allow(clippy::useless_conversion)] - Ok([e - .trailing_zeros() - .try_into() - .expect("bit count overflowed 32 bits, somehow!?")]) - }) - } - crate::MathFunction::Degrees => { - component_wise_float!(self, span, [arg], |e| { Ok([e.to_degrees()]) }) - } - crate::MathFunction::Exp => { - component_wise_float!(self, span, [arg], |e| { Ok([e.exp()]) }) + crate::MathFunction::Sin => { + component_wise_float!(self, span, [arg], |e| { Ok([e.sin()]) }) } - crate::MathFunction::Exp2 => { - component_wise_float!(self, span, [arg], |e| { Ok([e.exp2()]) }) + crate::MathFunction::Sinh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.sinh()]) }) } - crate::MathFunction::Floor => { - component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) + crate::MathFunction::Tan => { + component_wise_float!(self, span, [arg], |e| { Ok([e.tan()]) }) } - crate::MathFunction::Fract => { - component_wise_float!(self, span, [arg], |e| { - // N.B., Rust's definition of `fract` is `e - e.trunc()`, so we can't use that - // here. - Ok([e - e.floor()]) - }) + crate::MathFunction::Tanh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.tanh()]) }) } - crate::MathFunction::Fma => { - component_wise_float!( - self, - span, - [arg, arg1.unwrap(), arg2.unwrap()], - |e1, e2, e3| { Ok([e1.mul_add(e2, e3)]) } - ) + crate::MathFunction::Acos => { + component_wise_float!(self, span, [arg], |e| { Ok([e.acos()]) }) } - crate::MathFunction::InverseSqrt => { - component_wise_float!(self, span, [arg], |e| { Ok([1. / e.sqrt()]) }) + crate::MathFunction::Asin => { + component_wise_float!(self, span, [arg], |e| { Ok([e.asin()]) }) } - crate::MathFunction::Log => { - component_wise_float!(self, span, [arg], |e| { Ok([e.ln()]) }) + crate::MathFunction::Atan => { + component_wise_float!(self, span, [arg], |e| { Ok([e.atan()]) }) } - crate::MathFunction::Log2 => { - component_wise_float!(self, span, [arg], |e| { Ok([e.log2()]) }) + crate::MathFunction::Asinh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.asinh()]) }) } - crate::MathFunction::Max => { - component_wise_scalar!(self, span, [arg, arg1.unwrap()], |e1, e2| { - Ok([e1.max(e2)]) - }) + crate::MathFunction::Acosh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.acosh()]) }) } - crate::MathFunction::Min => { - component_wise_scalar!(self, span, [arg, arg1.unwrap()], |e1, e2| { - Ok([e1.min(e2)]) - }) + crate::MathFunction::Atanh => { + component_wise_float!(self, span, [arg], |e| { Ok([e.atanh()]) }) } crate::MathFunction::Radians => { component_wise_float!(self, span, [arg], |e1| { Ok([e1.to_radians()]) }) } - crate::MathFunction::ReverseBits => { - component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) }) + crate::MathFunction::Degrees => { + component_wise_float!(self, span, [arg], |e| { Ok([e.to_degrees()]) }) + } + + // decomposition + crate::MathFunction::Ceil => { + component_wise_float!(self, span, [arg], |e| { Ok([e.ceil()]) }) + } + crate::MathFunction::Floor => { + component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) }) } crate::MathFunction::Round => { // TODO: Use `f{32,64}.round_ties_even()` when available on stable. This polyfill @@ -998,35 +954,92 @@ 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.)]) }) + crate::MathFunction::Fract => { + component_wise_float!(self, span, [arg], |e| { + // N.B., Rust's definition of `fract` is `e - e.trunc()`, so we can't use that + // here. + Ok([e - e.floor()]) + }) } - crate::MathFunction::Sign => { - component_wise_signed!(self, span, [arg], |e| { Ok([e.signum()]) }) + crate::MathFunction::Trunc => { + component_wise_float!(self, span, [arg], |e| { Ok([e.trunc()]) }) } - crate::MathFunction::Sin => { - component_wise_float!(self, span, [arg], |e| { Ok([e.sin()]) }) + + // exponent + crate::MathFunction::Exp => { + component_wise_float!(self, span, [arg], |e| { Ok([e.exp()]) }) } - crate::MathFunction::Sinh => { - component_wise_float!(self, span, [arg], |e| { Ok([e.sinh()]) }) + crate::MathFunction::Exp2 => { + component_wise_float!(self, span, [arg], |e| { Ok([e.exp2()]) }) } - crate::MathFunction::Tan => { - component_wise_float!(self, span, [arg], |e| { Ok([e.tan()]) }) + crate::MathFunction::Log => { + component_wise_float!(self, span, [arg], |e| { Ok([e.ln()]) }) } - crate::MathFunction::Tanh => { - component_wise_float!(self, span, [arg], |e| { Ok([e.tanh()]) }) + crate::MathFunction::Log2 => { + component_wise_float!(self, span, [arg], |e| { Ok([e.log2()]) }) } - crate::MathFunction::Sqrt => { - component_wise_float!(self, span, [arg], |e| { Ok([e.sqrt()]) }) + crate::MathFunction::Pow => { + component_wise_float!(self, span, [arg, arg1.unwrap()], |e1, e2| { + Ok([e1.powf(e2)]) + }) + } + + // computational + crate::MathFunction::Sign => { + component_wise_signed!(self, span, [arg], |e| { Ok([e.signum()]) }) + } + crate::MathFunction::Fma => { + component_wise_float!( + self, + span, + [arg, arg1.unwrap(), arg2.unwrap()], + |e1, e2, e3| { Ok([e1.mul_add(e2, e3)]) } + ) } crate::MathFunction::Step => { component_wise_float!(self, span, [arg, arg1.unwrap()], |edge, x| { Ok([if edge <= x { 1.0 } else { 0.0 }]) }) } - crate::MathFunction::Trunc => { - component_wise_float!(self, span, [arg], |e| { Ok([e.trunc()]) }) + crate::MathFunction::Sqrt => { + component_wise_float!(self, span, [arg], |e| { Ok([e.sqrt()]) }) + } + crate::MathFunction::InverseSqrt => { + component_wise_float!(self, span, [arg], |e| { Ok([1. / e.sqrt()]) }) } + + // bits + crate::MathFunction::CountTrailingZeros => { + component_wise_concrete_int!(self, span, [arg], |e| { + #[allow(clippy::useless_conversion)] + Ok([e + .trailing_zeros() + .try_into() + .expect("bit count overflowed 32 bits, somehow!?")]) + }) + } + crate::MathFunction::CountLeadingZeros => { + component_wise_concrete_int!(self, span, [arg], |e| { + #[allow(clippy::useless_conversion)] + Ok([e + .leading_zeros() + .try_into() + .expect("bit count overflowed 32 bits, somehow!?")]) + }) + } + crate::MathFunction::CountOneBits => { + component_wise_concrete_int!(self, span, [arg], |e| { + #[allow(clippy::useless_conversion)] + Ok([e + .count_ones() + .try_into() + .expect("bit count overflowed 32 bits, somehow!?")]) + }) + } + crate::MathFunction::ReverseBits => { + component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) }) + } + fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" ))),