diff --git a/CHANGELOG.md b/CHANGELOG.md index 41cd22eef..a6ceb886a 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) - `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 29fced23a..ea8807d7b 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 289215780..a68b613a3 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 d576365fc..b3b938da6 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 33fae07f7..a9d37076e 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 8450bdd21..442b0ba0d 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 68c2ba329..8b8330738 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);