From 2d3005b74576fa19d894c6c58e682514478d4915 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 12:14:51 -0500 Subject: [PATCH] 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 ab12eb16b..560b29e76 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) - `wgpu-types`'s `trace` and `replay` features have been replaced by the `serde` feature. By @KirmesBude in [#5149](https://github.com/gfx-rs/wgpu/pull/5149) - `wgpu-core`'s `serial-pass` feature has been removed. Use `serde` instead. By @KirmesBude in [#5149](https://github.com/gfx-rs/wgpu/pull/5149) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 585daf1fe..6a2991fa6 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 a68b613a3..7f91571dc 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 b3b938da6..61c59f00c 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 a9d37076e..0e6a5b24d 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 442b0ba0d..6e07c6d7a 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 8b8330738..228248b3c 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);