feat(const_eval): impl. firstTrailingBit

Author: ErichDonGubler

naga/src/proc/constant_evaluator.rs

@@ -1233,6 +1233,9 @@ impl<'a> ConstantEvaluator<'a> {
             crate::MathFunction::ReverseBits => {
                 component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) })
             }
+            crate::MathFunction::FirstTrailingBit => {
+                component_wise_concrete_int(self, span, [arg], |ci| Ok(first_trailing_bit(ci)))
+            }
             crate::MathFunction::FirstLeadingBit => {
                 component_wise_concrete_int(self, span, [arg], |ci| Ok(first_leading_bit(ci)))
             }
@@ -2101,6 +2104,86 @@ impl<'a> ConstantEvaluator<'a> {
     }
 }
 
+fn first_trailing_bit(concrete_int: ConcreteInt<1>) -> ConcreteInt<1> {
+    // NOTE: Bit indices for this built-in start at 0 at the "right" (or LSB). For example, a value
+    // of 1 means the least significant bit is set. Therefore, an input of `0x[80 00…]` would
+    // return a right-to-left bit index of 0.
+    let trailing_zeros_to_bit_idx = |e: u32| -> u32 {
+        match e {
+            idx @ 0..=31 => idx,
+            32 => u32::MAX,
+            _ => panic!("invalid `trailing_zeros` values"),
+        }
+    };
+    match concrete_int {
+        ConcreteInt::U32([e]) => ConcreteInt::U32([trailing_zeros_to_bit_idx(e.trailing_zeros())]),
+        ConcreteInt::I32([e]) => {
+            ConcreteInt::I32([trailing_zeros_to_bit_idx(e.trailing_zeros()) as i32])
+        }
+    }
+}
+
+#[test]
+fn first_trailing_bit_smoke() {
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::I32([0])),
+        ConcreteInt::I32([-1])
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::I32([1])),
+        ConcreteInt::I32([0])
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::I32([2])),
+        ConcreteInt::I32([1])
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::I32([-1])),
+        ConcreteInt::I32([0]),
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::I32([i32::MIN])),
+        ConcreteInt::I32([31]),
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::I32([i32::MAX])),
+        ConcreteInt::I32([0]),
+    );
+    for idx in 0..32 {
+        assert_eq!(
+            first_trailing_bit(ConcreteInt::I32([1 << idx])),
+            ConcreteInt::I32([idx])
+        )
+    }
+
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::U32([0])),
+        ConcreteInt::U32([u32::MAX])
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::U32([1])),
+        ConcreteInt::U32([0])
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::U32([2])),
+        ConcreteInt::U32([1])
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::U32([1 << 31])),
+        ConcreteInt::U32([31]),
+    );
+    assert_eq!(
+        first_trailing_bit(ConcreteInt::U32([u32::MAX])),
+        ConcreteInt::U32([0]),
+    );
+    for idx in 0..32 {
+        assert_eq!(
+            first_trailing_bit(ConcreteInt::U32([1 << idx])),
+            ConcreteInt::U32([idx])
+        )
+    }
+}
+
 fn first_leading_bit(concrete_int: ConcreteInt<1>) -> ConcreteInt<1> {
     // NOTE: Bit indices for this built-in start at 0 at the "right" (or LSB). For example, 1 means
     // the least significant bit is set. Therefore, an input of 1 would return a right-to-left bit

naga/tests/out/glsl/math-functions.main.Fragment.glsl

@@ -67,10 +67,8 @@ void main() {
     int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y);
     ivec2 flb_b = ivec2(-1, -1);
     uvec2 flb_c = uvec2(0u, 0u);
-    int ftb_a = findLSB(-1);
-    uint ftb_b = uint(findLSB(1u));
-    ivec2 ftb_c = findLSB(ivec2(-1));
-    uvec2 ftb_d = uvec2(findLSB(uvec2(1u)));
+    ivec2 ftb_c = ivec2(0, 0);
+    uvec2 ftb_d = uvec2(0u, 0u);
     uvec2 ctz_e = uvec2(32u, 32u);
     ivec2 ctz_f = ivec2(32, 32);
     uvec2 ctz_g = uvec2(0u, 0u);

naga/tests/out/hlsl/math-functions.hlsl

@@ -81,10 +81,8 @@ void main()
     int const_dot = dot(ZeroValueint2(), ZeroValueint2());
     int2 flb_b = int2(-1, -1);
     uint2 flb_c = uint2(0u, 0u);
-    int ftb_a = asint(firstbitlow(-1));
-    uint ftb_b = firstbitlow(1u);
-    int2 ftb_c = asint(firstbitlow((-1).xx));
-    uint2 ftb_d = firstbitlow((1u).xx);
+    int2 ftb_c = int2(0, 0);
+    uint2 ftb_d = uint2(0u, 0u);
     uint2 ctz_e = uint2(32u, 32u);
     int2 ctz_f = int2(32, 32);
     uint2 ctz_g = uint2(0u, 0u);

naga/tests/out/msl/math-functions.msl

@@ -69,10 +69,8 @@ fragment void main_(
     int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y);
     metal::int2 flb_b = metal::int2(-1, -1);
     metal::uint2 flb_c = metal::uint2(0u, 0u);
-    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);
-    metal::uint2 ftb_d = (((metal::ctz(metal::uint2(1u)) + 1) % 33) - 1);
+    metal::int2 ftb_c = metal::int2(0, 0);
+    metal::uint2 ftb_d = metal::uint2(0u, 0u);
     metal::uint2 ctz_e = metal::uint2(32u, 32u);
     metal::int2 ctz_f = metal::int2(32, 32);
     metal::uint2 ctz_g = metal::uint2(0u, 0u);

naga/tests/out/spv/math-functions.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 94
+; Bound: 87
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
@@ -44,71 +44,64 @@ OpMemberDecorate %15 1 Offset 16
 %28 = OpConstantComposite  %7  %22 %22
 %29 = OpConstant  %9  0
 %30 = OpConstantComposite  %8  %29 %29
-%31 = OpConstant  %9  1
-%32 = OpConstantComposite  %7  %22 %22
-%33 = OpConstantComposite  %8  %31 %31
-%34 = OpConstant  %9  32
-%35 = OpConstant  %6  32
-%36 = OpConstant  %6  0
-%37 = OpConstantComposite  %8  %34 %34
-%38 = OpConstantComposite  %7  %35 %35
-%39 = OpConstantComposite  %7  %36 %36
-%40 = OpConstant  %9  31
-%41 = OpConstantComposite  %8  %40 %40
-%42 = OpConstant  %6  2
-%43 = OpConstant  %4  2.0
-%44 = OpConstantComposite  %10  %19 %43
-%45 = OpConstant  %6  3
-%46 = OpConstant  %6  4
-%47 = OpConstantComposite  %7  %45 %46
-%48 = OpConstant  %4  1.5
-%49 = OpConstantComposite  %10  %48 %48
-%50 = OpConstantComposite  %3  %48 %48 %48 %48
-%57 = OpConstantComposite  %3  %19 %19 %19 %19
-%60 = OpConstantNull  %6
+%31 = OpConstant  %6  0
+%32 = OpConstantComposite  %7  %31 %31
+%33 = OpConstant  %9  32
+%34 = OpConstant  %6  32
+%35 = OpConstantComposite  %8  %33 %33
+%36 = OpConstantComposite  %7  %34 %34
+%37 = OpConstant  %9  31
+%38 = OpConstantComposite  %8  %37 %37
+%39 = OpConstant  %6  2
+%40 = OpConstant  %4  2.0
+%41 = OpConstantComposite  %10  %19 %40
+%42 = OpConstant  %6  3
+%43 = OpConstant  %6  4
+%44 = OpConstantComposite  %7  %42 %43
+%45 = OpConstant  %4  1.5
+%46 = OpConstantComposite  %10  %45 %45
+%47 = OpConstantComposite  %3  %45 %45 %45 %45
+%54 = OpConstantComposite  %3  %19 %19 %19 %19
+%57 = OpConstantNull  %6
 %17 = OpFunction  %2  None %18
 %16 = OpLabel
-OpBranch %51
-%51 = OpLabel
-%52 = OpExtInst  %4  %1 Degrees %19
-%53 = OpExtInst  %4  %1 Radians %19
-%54 = OpExtInst  %3  %1 Degrees %21
-%55 = OpExtInst  %3  %1 Radians %21
-%56 = OpExtInst  %3  %1 FClamp %21 %21 %57
-%58 = OpExtInst  %3  %1 Refract %21 %21 %19
-%61 = OpCompositeExtract  %6  %26 0
-%62 = OpCompositeExtract  %6  %26 0
-%63 = OpIMul  %6  %61 %62
-%64 = OpIAdd  %6  %60 %63
-%65 = OpCompositeExtract  %6  %26 1
-%66 = OpCompositeExtract  %6  %26 1
-%67 = OpIMul  %6  %65 %66
-%59 = OpIAdd  %6  %64 %67
-%68 = OpExtInst  %6  %1 FindILsb %22
-%69 = OpExtInst  %9  %1 FindILsb %31
-%70 = OpExtInst  %7  %1 FindILsb %32
-%71 = OpExtInst  %8  %1 FindILsb %33
-%72 = OpExtInst  %4  %1 Ldexp %19 %42
-%73 = OpExtInst  %10  %1 Ldexp %44 %47
-%74 = OpExtInst  %11  %1 ModfStruct %48
-%75 = OpExtInst  %11  %1 ModfStruct %48
-%76 = OpCompositeExtract  %4  %75 0
-%77 = OpExtInst  %11  %1 ModfStruct %48
+OpBranch %48
+%48 = OpLabel
+%49 = OpExtInst  %4  %1 Degrees %19
+%50 = OpExtInst  %4  %1 Radians %19
+%51 = OpExtInst  %3  %1 Degrees %21
+%52 = OpExtInst  %3  %1 Radians %21
+%53 = OpExtInst  %3  %1 FClamp %21 %21 %54
+%55 = OpExtInst  %3  %1 Refract %21 %21 %19
+%58 = OpCompositeExtract  %6  %26 0
+%59 = OpCompositeExtract  %6  %26 0
+%60 = OpIMul  %6  %58 %59
+%61 = OpIAdd  %6  %57 %60
+%62 = OpCompositeExtract  %6  %26 1
+%63 = OpCompositeExtract  %6  %26 1
+%64 = OpIMul  %6  %62 %63
+%56 = OpIAdd  %6  %61 %64
+%65 = OpExtInst  %4  %1 Ldexp %19 %39
+%66 = OpExtInst  %10  %1 Ldexp %41 %44
+%67 = OpExtInst  %11  %1 ModfStruct %45
+%68 = OpExtInst  %11  %1 ModfStruct %45
+%69 = OpCompositeExtract  %4  %68 0
+%70 = OpExtInst  %11  %1 ModfStruct %45
+%71 = OpCompositeExtract  %4  %70 1
+%72 = OpExtInst  %12  %1 ModfStruct %46
+%73 = OpExtInst  %13  %1 ModfStruct %47
+%74 = OpCompositeExtract  %3  %73 1
+%75 = OpCompositeExtract  %4  %74 0
+%76 = OpExtInst  %12  %1 ModfStruct %46
+%77 = OpCompositeExtract  %10  %76 0
 %78 = OpCompositeExtract  %4  %77 1
-%79 = OpExtInst  %12  %1 ModfStruct %49
-%80 = OpExtInst  %13  %1 ModfStruct %50
-%81 = OpCompositeExtract  %3  %80 1
-%82 = OpCompositeExtract  %4  %81 0
-%83 = OpExtInst  %12  %1 ModfStruct %49
-%84 = OpCompositeExtract  %10  %83 0
-%85 = OpCompositeExtract  %4  %84 1
-%86 = OpExtInst  %14  %1 FrexpStruct %48
-%87 = OpExtInst  %14  %1 FrexpStruct %48
-%88 = OpCompositeExtract  %4  %87 0
-%89 = OpExtInst  %14  %1 FrexpStruct %48
-%90 = OpCompositeExtract  %6  %89 1
-%91 = OpExtInst  %15  %1 FrexpStruct %50
-%92 = OpCompositeExtract  %5  %91 1
-%93 = OpCompositeExtract  %6  %92 0
+%79 = OpExtInst  %14  %1 FrexpStruct %45
+%80 = OpExtInst  %14  %1 FrexpStruct %45
+%81 = OpCompositeExtract  %4  %80 0
+%82 = OpExtInst  %14  %1 FrexpStruct %45
+%83 = OpCompositeExtract  %6  %82 1
+%84 = OpExtInst  %15  %1 FrexpStruct %47
+%85 = OpCompositeExtract  %5  %84 1
+%86 = OpCompositeExtract  %6  %85 0
 OpReturn
 OpFunctionEnd

naga/tests/out/wgsl/math-functions.wgsl

@@ -12,10 +12,8 @@ fn main() {
     let const_dot = dot(vec2<i32>(), vec2<i32>());
     let flb_b = vec2<i32>(-1i, -1i);
     let flb_c = vec2<u32>(0u, 0u);
-    let ftb_a = firstTrailingBit(-1i);
-    let ftb_b = firstTrailingBit(1u);
-    let ftb_c = firstTrailingBit(vec2(-1i));
-    let ftb_d = firstTrailingBit(vec2(1u));
+    let ftb_c = vec2<i32>(0i, 0i);
+    let ftb_d = vec2<u32>(0u, 0u);
     let ctz_e = vec2<u32>(32u, 32u);
     let ctz_f = vec2<i32>(32i, 32i);
     let ctz_g = vec2<u32>(0u, 0u);