diff --git a/CHANGELOG.md b/CHANGELOG.md index 61da7136c6..41e68c7357 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) - `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) 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);