Repository: gfx-rs/wgpu
Pull request #5101 (2024-01-19 21:03:03)

Impl. const. eval. for "first bit" numeric built-ins
**Connections**
_Link to the issues addressed by this PR, or dependent PRs in other repositories_

Soft dependency on https://github.com/gfx-rs/wgpu/pull/5098.

**Description**
_Describe what problem this is solving, and how it's solved._

I think it's self-evident, but I'm happy to write something longer if it's not.

**Testing**
_Explain how this change is tested._

Some testing is evident from changes to Naga's snapshot tests, but there's not full coverage of these. I would accept a request to add tests before merging.

<!-- 
Thanks for filing! The codeowners file will automatically request reviews from the appropriate teams.

After you get a review and have addressed any comments, please explicitly re-request a review from the
person(s) who reviewed your changes. This will make sure it gets re-added to their review queue - you're no bothering us!
-->

**Checklist**

- [x] Run `cargo fmt`.
- [x] Run `cargo clippy`. If applicable, add:
  - ~~`--target wasm32-unknown-unknown`~~
  - ~~`--target wasm32-unknown-emscripten`~~
- [x] Run `cargo xtask test` to run tests.
- [x] Add change to `CHANGELOG.md`. See simple instructions inside file.


cwfitzgerald: Any chance of this being revived?
ErichDonGubler: @cwfitzgerald: It's not dead, just sleeping! 😉 The only thing this PR needs is prioritization from the Firefox side. I'll tag @jimblandy for comment.
ErichDonGubler: Rebased.
teoxoy: Just rebased this via GH's UI since it was running into timeouts that got fixed in https://github.com/gfx-rs/wgpu/pull/6030.

Patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index c9eccafcd..9567c74f0 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -39,6 +39,12 @@ Bottom level categories:
 
 ## Unreleased
 
+### New Features
+
+#### Naga
+
+* Support constant evaluation for `firstLeadingBit` and `firstTrailingBit` numeric built-ins in WGSL. Front-ends that translate to these built-ins also benefit from constant evaluation. By @ErichDonGubler in [#5101](https://github.com/gfx-rs/wgpu/pull/5101).
+
 ### Bug Fixes
 
 #### General
diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs
index 7ad1f3c59..fe7048054 100644
--- a/naga/src/back/glsl/mod.rs
+++ b/naga/src/back/glsl/mod.rs
@@ -3647,8 +3647,8 @@ impl<'a, W: Write> Writer<'a, W> {
 
                         return Ok(());
                     }
-                    Mf::FindLsb => "findLSB",
-                    Mf::FindMsb => "findMSB",
+                    Mf::FirstTrailingBit => "findLSB",
+                    Mf::FirstLeadingBit => "findMSB",
                     // data packing
                     Mf::Pack4x8snorm => "packSnorm4x8",
                     Mf::Pack4x8unorm => "packUnorm4x8",
@@ -3722,8 +3722,10 @@ impl<'a, W: Write> Writer<'a, W> {
 
                 // Some GLSL functions always return signed integers (like findMSB),
                 // so they need to be cast to uint if the argument is also an uint.
-                let ret_might_need_int_to_uint =
-                    matches!(fun, Mf::FindLsb | Mf::FindMsb | Mf::CountOneBits | Mf::Abs);
+                let ret_might_need_int_to_uint = matches!(
+                    fun,
+                    Mf::FirstTrailingBit | Mf::FirstLeadingBit | Mf::CountOneBits | Mf::Abs
+                );
 
                 // Some GLSL functions only accept signed integers (like abs),
                 // so they need their argument cast from uint to int.
diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs
index 982bf0cfe..85d943e85 100644
--- a/naga/src/back/hlsl/writer.rs
+++ b/naga/src/back/hlsl/writer.rs
@@ -3063,8 +3063,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     Mf::CountLeadingZeros => Function::CountLeadingZeros,
                     Mf::CountOneBits => Function::MissingIntOverload("countbits"),
                     Mf::ReverseBits => Function::MissingIntOverload("reversebits"),
-                    Mf::FindLsb => Function::MissingIntReturnType("firstbitlow"),
-                    Mf::FindMsb => Function::MissingIntReturnType("firstbithigh"),
+                    Mf::FirstTrailingBit => Function::MissingIntReturnType("firstbitlow"),
+                    Mf::FirstLeadingBit => Function::MissingIntReturnType("firstbithigh"),
                     Mf::ExtractBits => Function::Regular(EXTRACT_BITS_FUNCTION),
                     Mf::InsertBits => Function::Regular(INSERT_BITS_FUNCTION),
                     // Data Packing
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 7ec22009b..b112bb369 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -1875,8 +1875,8 @@ impl<W: Write> Writer<W> {
                     Mf::ReverseBits => "reverse_bits",
                     Mf::ExtractBits => "",
                     Mf::InsertBits => "",
-                    Mf::FindLsb => "",
-                    Mf::FindMsb => "",
+                    Mf::FirstTrailingBit => "",
+                    Mf::FirstLeadingBit => "",
                     // data packing
                     Mf::Pack4x8snorm => "pack_float_to_snorm4x8",
                     Mf::Pack4x8unorm => "pack_float_to_unorm4x8",
@@ -1920,7 +1920,7 @@ impl<W: Write> Writer<W> {
                         self.put_expression(arg1.unwrap(), context, false)?;
                         write!(self.out, ")")?;
                     }
-                    Mf::FindLsb => {
+                    Mf::FirstTrailingBit => {
                         let scalar = context.resolve_type(arg).scalar().unwrap();
                         let constant = scalar.width * 8 + 1;
 
@@ -1928,7 +1928,7 @@ impl<W: Write> Writer<W> {
                         self.put_expression(arg, context, true)?;
                         write!(self.out, ") + 1) % {constant}) - 1)")?;
                     }
-                    Mf::FindMsb => {
+                    Mf::FirstLeadingBit => {
                         let inner = context.resolve_type(arg);
                         let scalar = inner.scalar().unwrap();
                         let constant = scalar.width * 8 - 1;
@@ -2702,7 +2702,7 @@ impl<W: Write> Writer<W> {
                             }
                         }
                     }
-                    crate::MathFunction::FindMsb
+                    crate::MathFunction::FirstLeadingBit
                     | crate::MathFunction::Pack4xI8
                     | crate::MathFunction::Pack4xU8
                     | crate::MathFunction::Unpack4xI8
diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs
index 33f892aa4..9fb948586 100644
--- a/naga/src/back/spv/block.rs
+++ b/naga/src/back/spv/block.rs
@@ -1183,13 +1183,13 @@ impl<'w> BlockContext<'w> {
                             count_id,
                         ))
                     }
-                    Mf::FindLsb => MathOp::Ext(spirv::GLOp::FindILsb),
-                    Mf::FindMsb => {
+                    Mf::FirstTrailingBit => MathOp::Ext(spirv::GLOp::FindILsb),
+                    Mf::FirstLeadingBit => {
                         if arg_ty.scalar_width() == Some(4) {
                             let thing = match arg_scalar_kind {
                                 Some(crate::ScalarKind::Uint) => spirv::GLOp::FindUMsb,
                                 Some(crate::ScalarKind::Sint) => spirv::GLOp::FindSMsb,
-                                other => unimplemented!("Unexpected findMSB({:?})", other),
+                                other => unimplemented!("Unexpected firstLeadingBit({:?})", other),
                             };
                             MathOp::Ext(thing)
                         } else {
diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs
index 8cd37830e..6a069113e 100644
--- a/naga/src/back/wgsl/writer.rs
+++ b/naga/src/back/wgsl/writer.rs
@@ -1710,8 +1710,8 @@ impl<W: Write> Writer<W> {
                     Mf::ReverseBits => Function::Regular("reverseBits"),
                     Mf::ExtractBits => Function::Regular("extractBits"),
                     Mf::InsertBits => Function::Regular("insertBits"),
-                    Mf::FindLsb => Function::Regular("firstTrailingBit"),
-                    Mf::FindMsb => Function::Regular("firstLeadingBit"),
+                    Mf::FirstTrailingBit => Function::Regular("firstTrailingBit"),
+                    Mf::FirstLeadingBit => Function::Regular("firstLeadingBit"),
                     // data packing
                     Mf::Pack4x8snorm => Function::Regular("pack4x8snorm"),
                     Mf::Pack4x8unorm => Function::Regular("pack4x8unorm"),
diff --git a/naga/src/front/glsl/builtins.rs b/naga/src/front/glsl/builtins.rs
index cbb9b9938..f76ce7754 100644
--- a/naga/src/front/glsl/builtins.rs
+++ b/naga/src/front/glsl/builtins.rs
@@ -646,8 +646,8 @@ fn inject_standard_builtins(
                 "bitfieldReverse" => MathFunction::ReverseBits,
                 "bitfieldExtract" => MathFunction::ExtractBits,
                 "bitfieldInsert" => MathFunction::InsertBits,
-                "findLSB" => MathFunction::FindLsb,
-                "findMSB" => MathFunction::FindMsb,
+                "findLSB" => MathFunction::FirstTrailingBit,
+                "findMSB" => MathFunction::FirstLeadingBit,
                 _ => unreachable!(),
             };
 
@@ -695,8 +695,12 @@ fn inject_standard_builtins(
                 // we need to cast the return type of findLsb / findMsb
                 let mc = if scalar.kind == Sk::Uint {
                     match mc {
-                        MacroCall::MathFunction(MathFunction::FindLsb) => MacroCall::FindLsbUint,
-                        MacroCall::MathFunction(MathFunction::FindMsb) => MacroCall::FindMsbUint,
+                        MacroCall::MathFunction(MathFunction::FirstTrailingBit) => {
+                            MacroCall::FindLsbUint
+                        }
+                        MacroCall::MathFunction(MathFunction::FirstLeadingBit) => {
+                            MacroCall::FindMsbUint
+                        }
                         mc => mc,
                     }
                 } else {
@@ -1787,8 +1791,8 @@ impl MacroCall {
             )?,
             mc @ (MacroCall::FindLsbUint | MacroCall::FindMsbUint) => {
                 let fun = match mc {
-                    MacroCall::FindLsbUint => MathFunction::FindLsb,
-                    MacroCall::FindMsbUint => MathFunction::FindMsb,
+                    MacroCall::FindLsbUint => MathFunction::FirstTrailingBit,
+                    MacroCall::FindMsbUint => MathFunction::FirstLeadingBit,
                     _ => unreachable!(),
                 };
                 let res = ctx.add_expression(
diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs
index d154712b2..809aff767 100644
--- a/naga/src/front/spv/mod.rs
+++ b/naga/src/front/spv/mod.rs
@@ -3026,8 +3026,8 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                         Glo::UnpackHalf2x16 => Mf::Unpack2x16float,
                         Glo::UnpackUnorm2x16 => Mf::Unpack2x16unorm,
                         Glo::UnpackSnorm2x16 => Mf::Unpack2x16snorm,
-                        Glo::FindILsb => Mf::FindLsb,
-                        Glo::FindUMsb | Glo::FindSMsb => Mf::FindMsb,
+                        Glo::FindILsb => Mf::FirstTrailingBit,
+                        Glo::FindUMsb | Glo::FindSMsb => Mf::FirstLeadingBit,
                         // TODO: https://github.com/gfx-rs/naga/issues/2526
                         Glo::Modf | Glo::Frexp => return Err(Error::UnsupportedExtInst(inst_id)),
                         Glo::IMix
diff --git a/naga/src/front/wgsl/parse/conv.rs b/naga/src/front/wgsl/parse/conv.rs
index 49b15dfa8..80f05db59 100644
--- a/naga/src/front/wgsl/parse/conv.rs
+++ b/naga/src/front/wgsl/parse/conv.rs
@@ -235,8 +235,8 @@ pub fn map_standard_fun(word: &str) -> Option<crate::MathFunction> {
         "reverseBits" => Mf::ReverseBits,
         "extractBits" => Mf::ExtractBits,
         "insertBits" => Mf::InsertBits,
-        "firstTrailingBit" => Mf::FindLsb,
-        "firstLeadingBit" => Mf::FindMsb,
+        "firstTrailingBit" => Mf::FirstTrailingBit,
+        "firstLeadingBit" => Mf::FirstLeadingBit,
         // data packing
         "pack4x8snorm" => Mf::Pack4x8snorm,
         "pack4x8unorm" => Mf::Pack4x8unorm,
diff --git a/naga/src/lib.rs b/naga/src/lib.rs
index 8ed752792..94edec915 100644
--- a/naga/src/lib.rs
+++ b/naga/src/lib.rs
@@ -1198,8 +1198,8 @@ pub enum MathFunction {
     ReverseBits,
     ExtractBits,
     InsertBits,
-    FindLsb,
-    FindMsb,
+    FirstTrailingBit,
+    FirstLeadingBit,
     // data packing
     Pack4x8snorm,
     Pack4x8unorm,
diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs
index b5c821f41..deaa9c93c 100644
--- a/naga/src/proc/constant_evaluator.rs
+++ b/naga/src/proc/constant_evaluator.rs
@@ -27,6 +27,8 @@ macro_rules! gen_component_wise_extractor {
         scalar_kinds: [$( $scalar_kind:ident ),* $(,)?],
     ) => {
         /// A subset of [`Literal`]s intended to be used for implementing numeric built-ins.
+        #[derive(Debug)]
+        #[cfg_attr(test, derive(PartialEq))]
         enum $target<const N: usize> {
             $(
                 #[doc = concat!(
@@ -1231,6 +1233,12 @@ 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)))
+            }
 
             fun => Err(ConstantEvaluatorError::NotImplemented(format!(
                 "{fun:?} built-in function"
@@ -2096,6 +2104,174 @@ 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,
+            _ => unreachable!(),
+        }
+    };
+    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
+    // index of 0.
+    let rtl_to_ltr_bit_idx = |e: u32| -> u32 {
+        match e {
+            idx @ 0..=31 => 31 - idx,
+            32 => u32::MAX,
+            _ => unreachable!(),
+        }
+    };
+    match concrete_int {
+        ConcreteInt::I32([e]) => ConcreteInt::I32([{
+            let rtl_bit_index = if e.is_negative() {
+                e.leading_ones()
+            } else {
+                e.leading_zeros()
+            };
+            rtl_to_ltr_bit_idx(rtl_bit_index) as i32
+        }]),
+        ConcreteInt::U32([e]) => ConcreteInt::U32([rtl_to_ltr_bit_idx(e.leading_zeros())]),
+    }
+}
+
+#[test]
+fn first_leading_bit_smoke() {
+    assert_eq!(
+        first_leading_bit(ConcreteInt::I32([-1])),
+        ConcreteInt::I32([-1])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::I32([0])),
+        ConcreteInt::I32([-1])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::I32([1])),
+        ConcreteInt::I32([0])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::I32([-2])),
+        ConcreteInt::I32([0])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::I32([1234 + 4567])),
+        ConcreteInt::I32([12])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::I32([i32::MAX])),
+        ConcreteInt::I32([30])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::I32([i32::MIN])),
+        ConcreteInt::I32([30])
+    );
+    // NOTE: Ignore the sign bit, which is a separate (above) case.
+    for idx in 0..(32 - 1) {
+        assert_eq!(
+            first_leading_bit(ConcreteInt::I32([1 << idx])),
+            ConcreteInt::I32([idx])
+        );
+    }
+    for idx in 1..(32 - 1) {
+        assert_eq!(
+            first_leading_bit(ConcreteInt::I32([-(1 << idx)])),
+            ConcreteInt::I32([idx - 1])
+        );
+    }
+
+    assert_eq!(
+        first_leading_bit(ConcreteInt::U32([0])),
+        ConcreteInt::U32([u32::MAX])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::U32([1])),
+        ConcreteInt::U32([0])
+    );
+    assert_eq!(
+        first_leading_bit(ConcreteInt::U32([u32::MAX])),
+        ConcreteInt::U32([31])
+    );
+    for idx in 0..32 {
+        assert_eq!(
+            first_leading_bit(ConcreteInt::U32([1 << idx])),
+            ConcreteInt::U32([idx])
+        )
+    }
+}
+
 /// Trait for conversions of abstract values to concrete types.
 trait TryFromAbstract<T>: Sized {
     /// Convert an abstract literal `value` to `Self`.
diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs
index 86d2b11f2..41273c5c7 100644
--- a/naga/src/proc/mod.rs
+++ b/naga/src/proc/mod.rs
@@ -484,8 +484,8 @@ impl super::MathFunction {
             Self::ReverseBits => 1,
             Self::ExtractBits => 3,
             Self::InsertBits => 4,
-            Self::FindLsb => 1,
-            Self::FindMsb => 1,
+            Self::FirstTrailingBit => 1,
+            Self::FirstLeadingBit => 1,
             // data packing
             Self::Pack4x8snorm => 1,
             Self::Pack4x8unorm => 1,
diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs
index 0a02900c4..d8af0cd23 100644
--- a/naga/src/proc/typifier.rs
+++ b/naga/src/proc/typifier.rs
@@ -788,8 +788,8 @@ impl<'a> ResolveContext<'a> {
                     Mf::ReverseBits |
                     Mf::ExtractBits |
                     Mf::InsertBits |
-                    Mf::FindLsb |
-                    Mf::FindMsb => match *res_arg.inner_with(types)  {
+                    Mf::FirstTrailingBit |
+                    Mf::FirstLeadingBit => match *res_arg.inner_with(types)  {
                         Ti::Scalar(scalar @ crate::Scalar {
                             kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                             ..
diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs
index 89bceae06..116560bb6 100644
--- a/naga/src/valid/expression.rs
+++ b/naga/src/valid/expression.rs
@@ -1350,8 +1350,8 @@ impl super::Validator {
                     | Mf::CountTrailingZeros
                     | Mf::CountOneBits
                     | Mf::ReverseBits
-                    | Mf::FindMsb
-                    | Mf::FindLsb => {
+                    | Mf::FirstLeadingBit
+                    | Mf::FirstTrailingBit => {
                         if arg1_ty.is_some() || arg2_ty.is_some() || arg3_ty.is_some() {
                             return Err(ExpressionError::WrongArgumentCount(fun));
                         }


Test patch:
diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl
index 7f91571dc..4ab85269e 100644
--- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl
+++ b/naga/tests/out/glsl/math-functions.main.Fragment.glsl
@@ -65,14 +65,10 @@ void main() {
     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);
-    ivec2 flb_b = findMSB(ivec2(-1));
-    uvec2 flb_c = uvec2(findMSB(uvec2(1u)));
-    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 flb_b = ivec2(-1, -1);
+    uvec2 flb_c = uvec2(0u, 0u);
+    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);
diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl
index c1a771c25..a02b2b128 100644
--- a/naga/tests/out/hlsl/math-functions.hlsl
+++ b/naga/tests/out/hlsl/math-functions.hlsl
@@ -79,14 +79,10 @@ void main()
     int4 sign_b = int4(-1, -1, -1, -1);
     float4 sign_d = float4(-1.0, -1.0, -1.0, -1.0);
     int const_dot = dot(ZeroValueint2(), ZeroValueint2());
-    uint first_leading_bit_abs = firstbithigh(0u);
-    int flb_a = asint(firstbithigh(-1));
-    int2 flb_b = asint(firstbithigh((-1).xx));
-    uint2 flb_c = firstbithigh((1u).xx);
-    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 flb_b = int2(-1, -1);
+    uint2 flb_c = uint2(0u, 0u);
+    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);
diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl
index 0e6a5b24d..559002c39 100644
--- a/naga/tests/out/msl/math-functions.msl
+++ b/naga/tests/out/msl/math-functions.msl
@@ -67,16 +67,10 @@ fragment void main_(
     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 _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);
-    metal::uint2 ftb_d = (((metal::ctz(metal::uint2(1u)) + 1) % 33) - 1);
+    metal::int2 flb_b = metal::int2(-1, -1);
+    metal::uint2 flb_c = metal::uint2(0u, 0u);
+    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);
diff --git a/naga/tests/out/spv/math-functions.spvasm b/naga/tests/out/spv/math-functions.spvasm
index 6e07c6d7a..366857f91 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: 96
+; Bound: 87
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
@@ -40,77 +40,68 @@ OpMemberDecorate %15 1 Offset 16
 %24 = OpConstant  %4  -1.0
 %25 = OpConstantComposite  %3  %24 %24 %24 %24
 %26 = OpConstantNull  %7
-%27 = OpConstant  %9  0
+%27 = OpConstant  %9  4294967295
 %28 = OpConstantComposite  %7  %22 %22
-%29 = OpConstant  %9  1
+%29 = OpConstant  %9  0
 %30 = OpConstantComposite  %8  %29 %29
-%31 = OpConstant  %9  32
-%32 = OpConstant  %6  32
-%33 = OpConstant  %6  0
-%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  %10  %19 %41
-%43 = OpConstant  %6  3
-%44 = OpConstant  %6  4
-%45 = OpConstantComposite  %7  %43 %44
-%46 = OpConstant  %4  1.5
-%47 = OpConstantComposite  %10  %46 %46
-%48 = OpConstantComposite  %3  %46 %46 %46 %46
-%55 = OpConstantComposite  %3  %19 %19 %19 %19
-%58 = 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 %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
+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 = OpCompositeExtract  %6  %26 0
-%61 = OpIMul  %6  %59 %60
-%62 = OpIAdd  %6  %58 %61
+%60 = OpIMul  %6  %58 %59
+%61 = OpIAdd  %6  %57 %60
+%62 = OpCompositeExtract  %6  %26 1
 %63 = OpCompositeExtract  %6  %26 1
-%64 = OpCompositeExtract  %6  %26 1
-%65 = OpIMul  %6  %63 %64
-%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
+%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  %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
\ 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 228248b3c..2271bb9cb 100644
--- a/naga/tests/out/wgsl/math-functions.wgsl
+++ b/naga/tests/out/wgsl/math-functions.wgsl
@@ -10,14 +10,10 @@ fn main() {
     let sign_b = vec4<i32>(-1i, -1i, -1i, -1i);
     let sign_d = vec4<f32>(-1f, -1f, -1f, -1f);
     let const_dot = dot(vec2<i32>(), vec2<i32>());
-    let first_leading_bit_abs = firstLeadingBit(0u);
-    let flb_a = firstLeadingBit(-1i);
-    let flb_b = firstLeadingBit(vec2(-1i));
-    let flb_c = firstLeadingBit(vec2(1u));
-    let ftb_a = firstTrailingBit(-1i);
-    let ftb_b = firstTrailingBit(1u);
-    let ftb_c = firstTrailingBit(vec2(-1i));
-    let ftb_d = firstTrailingBit(vec2(1u));
+    let flb_b = vec2<i32>(-1i, -1i);
+    let flb_c = vec2<u32>(0u, 0u);
+    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);