builtins: Add firstTrailingBit

Currently polyfilled for all backends.
HLSL should be able to map this to 'firstbitlow', but there might need
to be some special case handling for 0 (undocumented behavior). For now
just polyfill.

CTS tests: https://github.com/gpuweb/cts/pull/1003

Bug: tint:1367
Bug: tint:1449
Change-Id: I8125b32687196678906e5a9d056b4f2efd885073
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/81502
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

src/tint/builtins.def

@@ -316,6 +316,8 @@ fn exp<N: num>(vec<N, f32>) -> vec<N, f32>
 fn exp2(f32) -> f32
 fn exp2<N: num>(vec<N, f32>) -> vec<N, f32>
 fn faceForward<N: num>(vec<N, f32>, vec<N, f32>, vec<N, f32>) -> vec<N, f32>
+fn firstTrailingBit<T: iu32>(T) -> T
+fn firstTrailingBit<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
 fn floor(f32) -> f32
 fn floor<N: num>(vec<N, f32>) -> vec<N, f32>
 fn fma(f32, f32, f32) -> f32

src/tint/sem/builtin_type.cc

@@ -117,6 +117,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
   if (name == "faceForward") {
     return BuiltinType::kFaceForward;
   }
+  if (name == "firstTrailingBit") {
+    return BuiltinType::kFirstTrailingBit;
+  }
   if (name == "floor") {
     return BuiltinType::kFloor;
   }
@@ -407,6 +410,8 @@ const char* str(BuiltinType i) {
       return "exp2";
     case BuiltinType::kFaceForward:
       return "faceForward";
+    case BuiltinType::kFirstTrailingBit:
+      return "firstTrailingBit";
     case BuiltinType::kFloor:
       return "floor";
     case BuiltinType::kFma:

src/tint/sem/builtin_type.h

@@ -63,6 +63,7 @@ enum class BuiltinType {
   kExp,
   kExp2,
   kFaceForward,
+  kFirstTrailingBit,
   kFloor,
   kFma,
   kFract,

src/tint/transform/builtin_polyfill.cc

@@ -45,10 +45,7 @@ struct BuiltinPolyfill::State {
   /// @return the polyfill function name
   Symbol countLeadingZeros(const sem::Type* ty) {
     auto name = b.Symbols().New("tint_count_leading_zeros");
-    uint32_t width = 1;
+    uint32_t width = WidthOf(ty);
-    if (auto* v = ty->As<sem::Vector>()) {
-      width = v->Width();
-    }
 
     // Returns either u32 or vecN<u32>
     auto U = [&]() -> const ast::Type* {
@@ -113,10 +110,7 @@ struct BuiltinPolyfill::State {
   /// @return the polyfill function name
   Symbol countTrailingZeros(const sem::Type* ty) {
     auto name = b.Symbols().New("tint_count_trailing_zeros");
-    uint32_t width = 1;
+    uint32_t width = WidthOf(ty);
-    if (auto* v = ty->As<sem::Vector>()) {
-      width = v->Width();
-    }
 
     // Returns either u32 or vecN<u32>
     auto U = [&]() -> const ast::Type* {
@@ -182,8 +176,87 @@ struct BuiltinPolyfill::State {
     return name;
   }
 
+  /// Builds the polyfill function for the `firstTrailingBit` builtin
+  /// @param ty the parameter and return type for the function
+  /// @return the polyfill function name
+  Symbol firstTrailingBit(const sem::Type* ty) {
+    auto name = b.Symbols().New("tint_first_trailing_bit");
+    uint32_t width = WidthOf(ty);
+
+    // Returns either u32 or vecN<u32>
+    auto U = [&]() -> const ast::Type* {
+      if (width == 1) {
+        return b.ty.u32();
+      }
+      return b.ty.vec<ProgramBuilder::u32>(width);
+    };
+    auto V = [&](uint32_t value) -> const ast::Expression* {
+      if (width == 1) {
+        return b.Expr(value);
+      }
+      return b.Construct(b.ty.vec<ProgramBuilder::u32>(width), value);
+    };
+    auto B = [&](const ast::Expression* value) -> const ast::Expression* {
+      if (width == 1) {
+        return b.Construct<bool>(value);
+      }
+      return b.Construct(b.ty.vec<bool>(width), value);
+    };
+    b.Func(name, {b.Param("v", T(ty))}, T(ty),
+           {
+               // var x = U(v);
+               b.Decl(b.Var("x", nullptr, b.Construct(U(), b.Expr("v")))),
+               // let b16 = select(16, 0, bool(x & 0x0000ffff));
+               b.Decl(b.Const("b16", nullptr,
+                              b.Call("select", V(16), V(0),
+                                     B(b.And("x", V(0x0000ffff)))))),
+               // x = x >> b16;
+               b.Assign("x", b.Shr("x", "b16")),
+               // let b8  = select(8,  0, bool(x & 0x000000ff));
+               b.Decl(b.Const(
+                   "b8", nullptr,
+                   b.Call("select", V(8), V(0), B(b.And("x", V(0x000000ff)))))),
+               // x = x >> b8;
+               b.Assign("x", b.Shr("x", "b8")),
+               // let b4  = select(4,  0, bool(x & 0x0000000f));
+               b.Decl(b.Const(
+                   "b4", nullptr,
+                   b.Call("select", V(4), V(0), B(b.And("x", V(0x0000000f)))))),
+               // x = x >> b4;
+               b.Assign("x", b.Shr("x", "b4")),
+               // let b2  = select(2,  0, bool(x & 0x00000003));
+               b.Decl(b.Const(
+                   "b2", nullptr,
+                   b.Call("select", V(2), V(0), B(b.And("x", V(0x00000003)))))),
+               // x = x >> b2;
+               b.Assign("x", b.Shr("x", "b2")),
+               // let b1  = select(1,  0, bool(x & 0x00000001));
+               b.Decl(b.Const(
+                   "b1", nullptr,
+                   b.Call("select", V(1), V(0), B(b.And("x", V(0x00000001)))))),
+               // let is_zero  = select(0, 0xffffffff, x == 0);
+               b.Decl(b.Const("is_zero", nullptr,
+                              b.Call("select", V(0), V(0xffffffff),
+                                     b.Equal("x", V(0))))),
+               // return R(b16 | b8 | b4 | b2 | b1 | is_zero);
+               b.Return(b.Construct(
+                   T(ty),
+                   b.Or(b.Or(b.Or(b.Or(b.Or("b16", "b8"), "b4"), "b2"), "b1"),
+                        "is_zero"))),
+           });
+    return name;
+  }
+
  private:
-  const ast::Type* T(const sem::Type* ty) { return CreateASTTypeFor(ctx, ty); }
+  const ast::Type* T(const sem::Type* ty) const {
+    return CreateASTTypeFor(ctx, ty);
+  }
+  uint32_t WidthOf(const sem::Type* ty) const {
+    if (auto* v = ty->As<sem::Vector>()) {
+      return v->Width();
+    }
+    return 1;
+  }
 };
 
 BuiltinPolyfill::BuiltinPolyfill() = default;
@@ -209,6 +282,11 @@ bool BuiltinPolyfill::ShouldRun(const Program* program,
                 return true;
               }
               break;
+            case sem::BuiltinType::kFirstTrailingBit:
+              if (builtins.first_trailing_bit) {
+                return true;
+              }
+              break;
             default:
               break;
           }
@@ -252,6 +330,13 @@ void BuiltinPolyfill::Run(CloneContext& ctx,
                   });
                 }
                 break;
+              case sem::BuiltinType::kFirstTrailingBit:
+                if (builtins.first_trailing_bit) {
+                  polyfill = utils::GetOrCreate(polyfills, builtin, [&] {
+                    return s.firstTrailingBit(builtin->ReturnType());
+                  });
+                }
+                break;
               default:
                 break;
             }

src/tint/transform/builtin_polyfill.h

@@ -34,6 +34,8 @@ class BuiltinPolyfill : public Castable<BuiltinPolyfill, Transform> {
     bool count_leading_zeros = false;
     /// Should `countTrailingZeros()` be polyfilled?
     bool count_trailing_zeros = false;
+    /// Should `firstTrailingBit()` be polyfilled?
+    bool first_trailing_bit = false;
   };
 
   /// Config is consumed by the BuiltinPolyfill transform.

src/tint/transform/builtin_polyfill_test.cc

@@ -348,6 +348,160 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// firstTrailingBit
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillFirstTrailingBit() {
+  BuiltinPolyfill::Builtins builtins;
+  builtins.first_trailing_bit = true;
+  DataMap data;
+  data.Add<BuiltinPolyfill::Config>(builtins);
+  return data;
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunFirstTrailingBit) {
+  auto* src = R"(
+fn f() {
+  firstTrailingBit(0xf);
+}
+)";
+
+  EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+  EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(src, polyfillFirstTrailingBit()));
+}
+
+TEST_F(BuiltinPolyfillTest, FirstTrailingBit_i32) {
+  auto* src = R"(
+fn f() {
+  let r : i32 = firstTrailingBit(15);
+}
+)";
+
+  auto* expect = R"(
+fn tint_first_trailing_bit(v : i32) -> i32 {
+  var x = u32(v);
+  let b16 = select(16u, 0u, bool((x & 65535u)));
+  x = (x >> b16);
+  let b8 = select(8u, 0u, bool((x & 255u)));
+  x = (x >> b8);
+  let b4 = select(4u, 0u, bool((x & 15u)));
+  x = (x >> b4);
+  let b2 = select(2u, 0u, bool((x & 3u)));
+  x = (x >> b2);
+  let b1 = select(1u, 0u, bool((x & 1u)));
+  let is_zero = select(0u, 4294967295u, (x == 0u));
+  return i32((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+fn f() {
+  let r : i32 = tint_first_trailing_bit(15);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillFirstTrailingBit());
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, FirstTrailingBit_u32) {
+  auto* src = R"(
+fn f() {
+  let r : u32 = firstTrailingBit(15u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_first_trailing_bit(v : u32) -> u32 {
+  var x = u32(v);
+  let b16 = select(16u, 0u, bool((x & 65535u)));
+  x = (x >> b16);
+  let b8 = select(8u, 0u, bool((x & 255u)));
+  x = (x >> b8);
+  let b4 = select(4u, 0u, bool((x & 15u)));
+  x = (x >> b4);
+  let b2 = select(2u, 0u, bool((x & 3u)));
+  x = (x >> b2);
+  let b1 = select(1u, 0u, bool((x & 1u)));
+  let is_zero = select(0u, 4294967295u, (x == 0u));
+  return u32((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+fn f() {
+  let r : u32 = tint_first_trailing_bit(15u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillFirstTrailingBit());
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, FirstTrailingBit_vec3_i32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<i32> = firstTrailingBit(vec3<i32>(15));
+}
+)";
+
+  auto* expect = R"(
+fn tint_first_trailing_bit(v : vec3<i32>) -> vec3<i32> {
+  var x = vec3<u32>(v);
+  let b16 = select(vec3<u32>(16u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(65535u))));
+  x = (x >> b16);
+  let b8 = select(vec3<u32>(8u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(255u))));
+  x = (x >> b8);
+  let b4 = select(vec3<u32>(4u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(15u))));
+  x = (x >> b4);
+  let b2 = select(vec3<u32>(2u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(3u))));
+  x = (x >> b2);
+  let b1 = select(vec3<u32>(1u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(1u))));
+  let is_zero = select(vec3<u32>(0u), vec3<u32>(4294967295u), (x == vec3<u32>(0u)));
+  return vec3<i32>((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+fn f() {
+  let r : vec3<i32> = tint_first_trailing_bit(vec3<i32>(15));
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillFirstTrailingBit());
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, FirstTrailingBit_vec3_u32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<u32> = firstTrailingBit(vec3<u32>(15u));
+}
+)";
+
+  auto* expect = R"(
+fn tint_first_trailing_bit(v : vec3<u32>) -> vec3<u32> {
+  var x = vec3<u32>(v);
+  let b16 = select(vec3<u32>(16u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(65535u))));
+  x = (x >> b16);
+  let b8 = select(vec3<u32>(8u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(255u))));
+  x = (x >> b8);
+  let b4 = select(vec3<u32>(4u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(15u))));
+  x = (x >> b4);
+  let b2 = select(vec3<u32>(2u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(3u))));
+  x = (x >> b2);
+  let b1 = select(vec3<u32>(1u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(1u))));
+  let is_zero = select(vec3<u32>(0u), vec3<u32>(4294967295u), (x == vec3<u32>(0u)));
+  return vec3<u32>((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+fn f() {
+  let r : vec3<u32> = tint_first_trailing_bit(vec3<u32>(15u));
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillFirstTrailingBit());
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/tint/transform/glsl.cc

@@ -56,6 +56,7 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const {
     BuiltinPolyfill::Builtins polyfills;
     polyfills.count_leading_zeros = true;
     polyfills.count_trailing_zeros = true;
+    polyfills.first_trailing_bit = true;
     data.Add<BuiltinPolyfill::Config>(polyfills);
     manager.Add<BuiltinPolyfill>();
   }

src/tint/writer/hlsl/generator_impl.cc

@@ -143,8 +143,11 @@ SanitizedResult Sanitize(
 
   {  // Builtin polyfills
     transform::BuiltinPolyfill::Builtins polyfills;
+    // TODO(crbug.com/tint/1449): Some of these can map to HLSL's `firstbitlow`
+    // and `firstbithigh`.
     polyfills.count_leading_zeros = true;
     polyfills.count_trailing_zeros = true;
+    polyfills.first_trailing_bit = true;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);
     manager.Add<transform::BuiltinPolyfill>();
   }

src/tint/writer/msl/generator_impl.cc

@@ -59,6 +59,7 @@
 #include "src/tint/sem/vector_type.h"
 #include "src/tint/sem/void_type.h"
 #include "src/tint/transform/array_length_from_uniform.h"
+#include "src/tint/transform/builtin_polyfill.h"
 #include "src/tint/transform/canonicalize_entry_point_io.h"
 #include "src/tint/transform/external_texture_transform.h"
 #include "src/tint/transform/manager.h"
@@ -127,6 +128,13 @@ SanitizedResult Sanitize(
   transform::Manager manager;
   transform::DataMap data;
 
+  {  // Builtin polyfills
+    transform::BuiltinPolyfill::Builtins polyfills;
+    polyfills.first_trailing_bit = true;
+    data.Add<transform::BuiltinPolyfill::Config>(polyfills);
+    manager.Add<transform::BuiltinPolyfill>();
+  }
+
   // Build the config for the internal ArrayLengthFromUniform transform.
   transform::ArrayLengthFromUniform::Config array_length_from_uniform_cfg(
       array_length_from_uniform.ubo_binding);

src/tint/writer/spirv/builder.cc

@@ -263,6 +263,7 @@ SanitizedResult Sanitize(const Program* in,
     transform::BuiltinPolyfill::Builtins polyfills;
     polyfills.count_leading_zeros = true;
     polyfills.count_trailing_zeros = true;
+    polyfills.first_trailing_bit = true;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);
     manager.Add<transform::BuiltinPolyfill>();
   }

test/tint/builtins/gen/firstTrailingBit/110f2c.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(vec<4, u32>) -> vec<4, u32>
+fn firstTrailingBit_110f2c() {
+  var res: vec4<u32> = firstTrailingBit(vec4<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_110f2c();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_110f2c();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_110f2c();
+}

test/tint/builtins/gen/firstTrailingBit/110f2c.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+uvec4 tint_first_trailing_bit(uvec4 v) {
+  uvec4 x = uvec4(v);
+  uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
+  x = (x >> b16);
+  uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
+  x = (x >> b8);
+  uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
+  x = (x >> b4);
+  uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
+  x = (x >> b2);
+  uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
+  uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
+  return uvec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_110f2c() {
+  uvec4 res = tint_first_trailing_bit(uvec4(0u, 0u, 0u, 0u));
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_110f2c();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+uvec4 tint_first_trailing_bit(uvec4 v) {
+  uvec4 x = uvec4(v);
+  uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
+  x = (x >> b16);
+  uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
+  x = (x >> b8);
+  uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
+  x = (x >> b4);
+  uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
+  x = (x >> b2);
+  uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
+  uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
+  return uvec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_110f2c() {
+  uvec4 res = tint_first_trailing_bit(uvec4(0u, 0u, 0u, 0u));
+}
+
+void fragment_main() {
+  firstTrailingBit_110f2c();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uvec4 tint_first_trailing_bit(uvec4 v) {
+  uvec4 x = uvec4(v);
+  uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
+  x = (x >> b16);
+  uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
+  x = (x >> b8);
+  uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
+  x = (x >> b4);
+  uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
+  x = (x >> b2);
+  uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
+  uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
+  return uvec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_110f2c() {
+  uvec4 res = tint_first_trailing_bit(uvec4(0u, 0u, 0u, 0u));
+}
+
+void compute_main() {
+  firstTrailingBit_110f2c();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/110f2c.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+uint4 tint_first_trailing_bit(uint4 v) {
+  uint4 x = uint4(v);
+  const uint4 b16 = (bool4((x & uint4((65535u).xxxx))) ? uint4((0u).xxxx) : uint4((16u).xxxx));
+  x = (x >> b16);
+  const uint4 b8 = (bool4((x & uint4((255u).xxxx))) ? uint4((0u).xxxx) : uint4((8u).xxxx));
+  x = (x >> b8);
+  const uint4 b4 = (bool4((x & uint4((15u).xxxx))) ? uint4((0u).xxxx) : uint4((4u).xxxx));
+  x = (x >> b4);
+  const uint4 b2 = (bool4((x & uint4((3u).xxxx))) ? uint4((0u).xxxx) : uint4((2u).xxxx));
+  x = (x >> b2);
+  const uint4 b1 = (bool4((x & uint4((1u).xxxx))) ? uint4((0u).xxxx) : uint4((1u).xxxx));
+  const uint4 is_zero = ((x == uint4((0u).xxxx)) ? uint4((4294967295u).xxxx) : uint4((0u).xxxx));
+  return uint4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_110f2c() {
+  uint4 res = tint_first_trailing_bit(uint4(0u, 0u, 0u, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_110f2c();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_110f2c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_110f2c();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/110f2c.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint4 tint_first_trailing_bit(uint4 v) {
+  uint4 x = uint4(v);
+  uint4 const b16 = select(uint4(16u), uint4(0u), bool4((x & uint4(65535u))));
+  x = (x >> b16);
+  uint4 const b8 = select(uint4(8u), uint4(0u), bool4((x & uint4(255u))));
+  x = (x >> b8);
+  uint4 const b4 = select(uint4(4u), uint4(0u), bool4((x & uint4(15u))));
+  x = (x >> b4);
+  uint4 const b2 = select(uint4(2u), uint4(0u), bool4((x & uint4(3u))));
+  x = (x >> b2);
+  uint4 const b1 = select(uint4(1u), uint4(0u), bool4((x & uint4(1u))));
+  uint4 const is_zero = select(uint4(0u), uint4(4294967295u), (x == uint4(0u)));
+  return uint4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_110f2c() {
+  uint4 res = tint_first_trailing_bit(uint4());
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_110f2c();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_110f2c();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_110f2c();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/110f2c.wgsl.expected.spvasm

@@ -0,0 +1,142 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 100
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_110f2c "firstTrailingBit_110f2c"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+          %9 = OpTypeFunction %v4uint %v4uint
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %18 = OpConstantNull %v4uint
+       %bool = OpTypeBool
+     %v4bool = OpTypeVector %bool 4
+ %uint_65535 = OpConstant %uint 65535
+         %25 = OpConstantComposite %v4uint %uint_65535 %uint_65535 %uint_65535 %uint_65535
+     %uint_0 = OpConstant %uint 0
+         %28 = OpConstantComposite %v4uint %uint_0 %uint_0 %uint_0 %uint_0
+    %uint_16 = OpConstant %uint 16
+         %30 = OpConstantComposite %v4uint %uint_16 %uint_16 %uint_16 %uint_16
+   %uint_255 = OpConstant %uint 255
+         %37 = OpConstantComposite %v4uint %uint_255 %uint_255 %uint_255 %uint_255
+     %uint_8 = OpConstant %uint 8
+         %40 = OpConstantComposite %v4uint %uint_8 %uint_8 %uint_8 %uint_8
+    %uint_15 = OpConstant %uint 15
+         %47 = OpConstantComposite %v4uint %uint_15 %uint_15 %uint_15 %uint_15
+     %uint_4 = OpConstant %uint 4
+         %50 = OpConstantComposite %v4uint %uint_4 %uint_4 %uint_4 %uint_4
+     %uint_3 = OpConstant %uint 3
+         %57 = OpConstantComposite %v4uint %uint_3 %uint_3 %uint_3 %uint_3
+     %uint_2 = OpConstant %uint 2
+         %60 = OpConstantComposite %v4uint %uint_2 %uint_2 %uint_2 %uint_2
+     %uint_1 = OpConstant %uint 1
+         %67 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
+%uint_4294967295 = OpConstant %uint 4294967295
+         %73 = OpConstantComposite %v4uint %uint_4294967295 %uint_4294967295 %uint_4294967295 %uint_4294967295
+       %void = OpTypeVoid
+         %80 = OpTypeFunction %void
+         %86 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %v4uint None %9
+          %v = OpFunctionParameter %v4uint
+         %14 = OpLabel
+          %x = OpVariable %_ptr_Function_v4uint Function %18
+               OpStore %x %v
+         %23 = OpLoad %v4uint %x
+         %26 = OpBitwiseAnd %v4uint %23 %25
+         %20 = OpINotEqual %v4bool %26 %18
+         %19 = OpSelect %v4uint %20 %28 %30
+         %31 = OpLoad %v4uint %x
+         %32 = OpShiftRightLogical %v4uint %31 %19
+               OpStore %x %32
+         %35 = OpLoad %v4uint %x
+         %38 = OpBitwiseAnd %v4uint %35 %37
+         %34 = OpINotEqual %v4bool %38 %18
+         %33 = OpSelect %v4uint %34 %28 %40
+         %41 = OpLoad %v4uint %x
+         %42 = OpShiftRightLogical %v4uint %41 %33
+               OpStore %x %42
+         %45 = OpLoad %v4uint %x
+         %48 = OpBitwiseAnd %v4uint %45 %47
+         %44 = OpINotEqual %v4bool %48 %18
+         %43 = OpSelect %v4uint %44 %28 %50
+         %51 = OpLoad %v4uint %x
+         %52 = OpShiftRightLogical %v4uint %51 %43
+               OpStore %x %52
+         %55 = OpLoad %v4uint %x
+         %58 = OpBitwiseAnd %v4uint %55 %57
+         %54 = OpINotEqual %v4bool %58 %18
+         %53 = OpSelect %v4uint %54 %28 %60
+         %61 = OpLoad %v4uint %x
+         %62 = OpShiftRightLogical %v4uint %61 %53
+               OpStore %x %62
+         %65 = OpLoad %v4uint %x
+         %68 = OpBitwiseAnd %v4uint %65 %67
+         %64 = OpINotEqual %v4bool %68 %18
+         %63 = OpSelect %v4uint %64 %28 %67
+         %70 = OpLoad %v4uint %x
+         %71 = OpIEqual %v4bool %70 %28
+         %69 = OpSelect %v4uint %71 %73 %28
+         %75 = OpBitwiseOr %v4uint %19 %33
+         %76 = OpBitwiseOr %v4uint %75 %43
+         %77 = OpBitwiseOr %v4uint %76 %53
+         %78 = OpBitwiseOr %v4uint %77 %63
+         %79 = OpBitwiseOr %v4uint %78 %69
+               OpReturnValue %79
+               OpFunctionEnd
+%firstTrailingBit_110f2c = OpFunction %void None %80
+         %83 = OpLabel
+        %res = OpVariable %_ptr_Function_v4uint Function %18
+         %84 = OpFunctionCall %v4uint %tint_first_trailing_bit %18
+               OpStore %res %84
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %86
+         %88 = OpLabel
+         %89 = OpFunctionCall %void %firstTrailingBit_110f2c
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %80
+         %91 = OpLabel
+         %92 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %92
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %80
+         %95 = OpLabel
+         %96 = OpFunctionCall %void %firstTrailingBit_110f2c
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %80
+         %98 = OpLabel
+         %99 = OpFunctionCall %void %firstTrailingBit_110f2c
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/110f2c.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_110f2c() {
+  var res : vec4<u32> = firstTrailingBit(vec4<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_110f2c();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_110f2c();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_110f2c();
+}

test/tint/builtins/gen/firstTrailingBit/3a2acc.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(i32) -> i32
+fn firstTrailingBit_3a2acc() {
+  var res: i32 = firstTrailingBit(1);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_3a2acc();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_3a2acc();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_3a2acc();
+}

test/tint/builtins/gen/firstTrailingBit/3a2acc.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+int tint_first_trailing_bit(int v) {
+  uint x = uint(v);
+  uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_3a2acc() {
+  int res = tint_first_trailing_bit(1);
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_3a2acc();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+int tint_first_trailing_bit(int v) {
+  uint x = uint(v);
+  uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_3a2acc() {
+  int res = tint_first_trailing_bit(1);
+}
+
+void fragment_main() {
+  firstTrailingBit_3a2acc();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+int tint_first_trailing_bit(int v) {
+  uint x = uint(v);
+  uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_3a2acc() {
+  int res = tint_first_trailing_bit(1);
+}
+
+void compute_main() {
+  firstTrailingBit_3a2acc();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/3a2acc.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+int tint_first_trailing_bit(int v) {
+  uint x = uint(v);
+  const uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  const uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  const uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  const uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  const uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  const uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_3a2acc() {
+  int res = tint_first_trailing_bit(1);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_3a2acc();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_3a2acc();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_3a2acc();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/3a2acc.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_first_trailing_bit(int v) {
+  uint x = uint(v);
+  uint const b16 = select(16u, 0u, bool((x & 65535u)));
+  x = (x >> b16);
+  uint const b8 = select(8u, 0u, bool((x & 255u)));
+  x = (x >> b8);
+  uint const b4 = select(4u, 0u, bool((x & 15u)));
+  x = (x >> b4);
+  uint const b2 = select(2u, 0u, bool((x & 3u)));
+  x = (x >> b2);
+  uint const b1 = select(1u, 0u, bool((x & 1u)));
+  uint const is_zero = select(0u, 4294967295u, (x == 0u));
+  return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_3a2acc() {
+  int res = tint_first_trailing_bit(1);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_3a2acc();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_3a2acc();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_3a2acc();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/3a2acc.wgsl.expected.spvasm

@@ -0,0 +1,135 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 91
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_3a2acc "firstTrailingBit_3a2acc"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+          %9 = OpTypeFunction %int %int
+       %uint = OpTypeInt 32 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %18 = OpConstantNull %uint
+       %bool = OpTypeBool
+ %uint_65535 = OpConstant %uint 65535
+     %uint_0 = OpConstant %uint 0
+    %uint_16 = OpConstant %uint 16
+   %uint_255 = OpConstant %uint 255
+     %uint_8 = OpConstant %uint 8
+    %uint_15 = OpConstant %uint 15
+     %uint_4 = OpConstant %uint 4
+     %uint_3 = OpConstant %uint 3
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%uint_4294967295 = OpConstant %uint 4294967295
+       %void = OpTypeVoid
+         %68 = OpTypeFunction %void
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %76 = OpConstantNull %int
+         %77 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %int None %9
+          %v = OpFunctionParameter %int
+         %13 = OpLabel
+          %x = OpVariable %_ptr_Function_uint Function %18
+         %14 = OpBitcast %uint %v
+               OpStore %x %14
+         %22 = OpLoad %uint %x
+         %24 = OpBitwiseAnd %uint %22 %uint_65535
+         %20 = OpINotEqual %bool %24 %18
+         %19 = OpSelect %uint %20 %uint_0 %uint_16
+         %27 = OpLoad %uint %x
+         %28 = OpShiftRightLogical %uint %27 %19
+               OpStore %x %28
+         %31 = OpLoad %uint %x
+         %33 = OpBitwiseAnd %uint %31 %uint_255
+         %30 = OpINotEqual %bool %33 %18
+         %29 = OpSelect %uint %30 %uint_0 %uint_8
+         %35 = OpLoad %uint %x
+         %36 = OpShiftRightLogical %uint %35 %29
+               OpStore %x %36
+         %39 = OpLoad %uint %x
+         %41 = OpBitwiseAnd %uint %39 %uint_15
+         %38 = OpINotEqual %bool %41 %18
+         %37 = OpSelect %uint %38 %uint_0 %uint_4
+         %43 = OpLoad %uint %x
+         %44 = OpShiftRightLogical %uint %43 %37
+               OpStore %x %44
+         %47 = OpLoad %uint %x
+         %49 = OpBitwiseAnd %uint %47 %uint_3
+         %46 = OpINotEqual %bool %49 %18
+         %45 = OpSelect %uint %46 %uint_0 %uint_2
+         %51 = OpLoad %uint %x
+         %52 = OpShiftRightLogical %uint %51 %45
+               OpStore %x %52
+         %55 = OpLoad %uint %x
+         %57 = OpBitwiseAnd %uint %55 %uint_1
+         %54 = OpINotEqual %bool %57 %18
+         %53 = OpSelect %uint %54 %uint_0 %uint_1
+         %59 = OpLoad %uint %x
+         %60 = OpIEqual %bool %59 %uint_0
+         %58 = OpSelect %uint %60 %uint_4294967295 %uint_0
+         %63 = OpBitwiseOr %uint %19 %29
+         %64 = OpBitwiseOr %uint %63 %37
+         %65 = OpBitwiseOr %uint %64 %45
+         %66 = OpBitwiseOr %uint %65 %53
+         %67 = OpBitwiseOr %uint %66 %58
+         %62 = OpBitcast %int %67
+               OpReturnValue %62
+               OpFunctionEnd
+%firstTrailingBit_3a2acc = OpFunction %void None %68
+         %71 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %76
+         %72 = OpFunctionCall %int %tint_first_trailing_bit %int_1
+               OpStore %res %72
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %77
+         %79 = OpLabel
+         %80 = OpFunctionCall %void %firstTrailingBit_3a2acc
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %68
+         %82 = OpLabel
+         %83 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %83
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %68
+         %86 = OpLabel
+         %87 = OpFunctionCall %void %firstTrailingBit_3a2acc
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %68
+         %89 = OpLabel
+         %90 = OpFunctionCall %void %firstTrailingBit_3a2acc
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/3a2acc.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_3a2acc() {
+  var res : i32 = firstTrailingBit(1);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_3a2acc();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_3a2acc();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_3a2acc();
+}

test/tint/builtins/gen/firstTrailingBit/45eb10.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(vec<2, u32>) -> vec<2, u32>
+fn firstTrailingBit_45eb10() {
+  var res: vec2<u32> = firstTrailingBit(vec2<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_45eb10();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_45eb10();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_45eb10();
+}

test/tint/builtins/gen/firstTrailingBit/45eb10.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+uvec2 tint_first_trailing_bit(uvec2 v) {
+  uvec2 x = uvec2(v);
+  uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
+  x = (x >> b16);
+  uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
+  x = (x >> b8);
+  uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
+  x = (x >> b4);
+  uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
+  x = (x >> b2);
+  uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
+  uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
+  return uvec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_45eb10() {
+  uvec2 res = tint_first_trailing_bit(uvec2(0u, 0u));
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_45eb10();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+uvec2 tint_first_trailing_bit(uvec2 v) {
+  uvec2 x = uvec2(v);
+  uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
+  x = (x >> b16);
+  uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
+  x = (x >> b8);
+  uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
+  x = (x >> b4);
+  uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
+  x = (x >> b2);
+  uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
+  uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
+  return uvec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_45eb10() {
+  uvec2 res = tint_first_trailing_bit(uvec2(0u, 0u));
+}
+
+void fragment_main() {
+  firstTrailingBit_45eb10();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uvec2 tint_first_trailing_bit(uvec2 v) {
+  uvec2 x = uvec2(v);
+  uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
+  x = (x >> b16);
+  uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
+  x = (x >> b8);
+  uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
+  x = (x >> b4);
+  uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
+  x = (x >> b2);
+  uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
+  uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
+  return uvec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_45eb10() {
+  uvec2 res = tint_first_trailing_bit(uvec2(0u, 0u));
+}
+
+void compute_main() {
+  firstTrailingBit_45eb10();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/45eb10.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+uint2 tint_first_trailing_bit(uint2 v) {
+  uint2 x = uint2(v);
+  const uint2 b16 = (bool2((x & uint2((65535u).xx))) ? uint2((0u).xx) : uint2((16u).xx));
+  x = (x >> b16);
+  const uint2 b8 = (bool2((x & uint2((255u).xx))) ? uint2((0u).xx) : uint2((8u).xx));
+  x = (x >> b8);
+  const uint2 b4 = (bool2((x & uint2((15u).xx))) ? uint2((0u).xx) : uint2((4u).xx));
+  x = (x >> b4);
+  const uint2 b2 = (bool2((x & uint2((3u).xx))) ? uint2((0u).xx) : uint2((2u).xx));
+  x = (x >> b2);
+  const uint2 b1 = (bool2((x & uint2((1u).xx))) ? uint2((0u).xx) : uint2((1u).xx));
+  const uint2 is_zero = ((x == uint2((0u).xx)) ? uint2((4294967295u).xx) : uint2((0u).xx));
+  return uint2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_45eb10() {
+  uint2 res = tint_first_trailing_bit(uint2(0u, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_45eb10();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_45eb10();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_45eb10();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/45eb10.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint2 tint_first_trailing_bit(uint2 v) {
+  uint2 x = uint2(v);
+  uint2 const b16 = select(uint2(16u), uint2(0u), bool2((x & uint2(65535u))));
+  x = (x >> b16);
+  uint2 const b8 = select(uint2(8u), uint2(0u), bool2((x & uint2(255u))));
+  x = (x >> b8);
+  uint2 const b4 = select(uint2(4u), uint2(0u), bool2((x & uint2(15u))));
+  x = (x >> b4);
+  uint2 const b2 = select(uint2(2u), uint2(0u), bool2((x & uint2(3u))));
+  x = (x >> b2);
+  uint2 const b1 = select(uint2(1u), uint2(0u), bool2((x & uint2(1u))));
+  uint2 const is_zero = select(uint2(0u), uint2(4294967295u), (x == uint2(0u)));
+  return uint2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_45eb10() {
+  uint2 res = tint_first_trailing_bit(uint2());
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_45eb10();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_45eb10();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_45eb10();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/45eb10.wgsl.expected.spvasm

@@ -0,0 +1,142 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 100
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_45eb10 "firstTrailingBit_45eb10"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %uint = OpTypeInt 32 0
+     %v2uint = OpTypeVector %uint 2
+          %9 = OpTypeFunction %v2uint %v2uint
+%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %18 = OpConstantNull %v2uint
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+ %uint_65535 = OpConstant %uint 65535
+         %25 = OpConstantComposite %v2uint %uint_65535 %uint_65535
+     %uint_0 = OpConstant %uint 0
+         %28 = OpConstantComposite %v2uint %uint_0 %uint_0
+    %uint_16 = OpConstant %uint 16
+         %30 = OpConstantComposite %v2uint %uint_16 %uint_16
+   %uint_255 = OpConstant %uint 255
+         %37 = OpConstantComposite %v2uint %uint_255 %uint_255
+     %uint_8 = OpConstant %uint 8
+         %40 = OpConstantComposite %v2uint %uint_8 %uint_8
+    %uint_15 = OpConstant %uint 15
+         %47 = OpConstantComposite %v2uint %uint_15 %uint_15
+     %uint_4 = OpConstant %uint 4
+         %50 = OpConstantComposite %v2uint %uint_4 %uint_4
+     %uint_3 = OpConstant %uint 3
+         %57 = OpConstantComposite %v2uint %uint_3 %uint_3
+     %uint_2 = OpConstant %uint 2
+         %60 = OpConstantComposite %v2uint %uint_2 %uint_2
+     %uint_1 = OpConstant %uint 1
+         %67 = OpConstantComposite %v2uint %uint_1 %uint_1
+%uint_4294967295 = OpConstant %uint 4294967295
+         %73 = OpConstantComposite %v2uint %uint_4294967295 %uint_4294967295
+       %void = OpTypeVoid
+         %80 = OpTypeFunction %void
+         %86 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %v2uint None %9
+          %v = OpFunctionParameter %v2uint
+         %14 = OpLabel
+          %x = OpVariable %_ptr_Function_v2uint Function %18
+               OpStore %x %v
+         %23 = OpLoad %v2uint %x
+         %26 = OpBitwiseAnd %v2uint %23 %25
+         %20 = OpINotEqual %v2bool %26 %18
+         %19 = OpSelect %v2uint %20 %28 %30
+         %31 = OpLoad %v2uint %x
+         %32 = OpShiftRightLogical %v2uint %31 %19
+               OpStore %x %32
+         %35 = OpLoad %v2uint %x
+         %38 = OpBitwiseAnd %v2uint %35 %37
+         %34 = OpINotEqual %v2bool %38 %18
+         %33 = OpSelect %v2uint %34 %28 %40
+         %41 = OpLoad %v2uint %x
+         %42 = OpShiftRightLogical %v2uint %41 %33
+               OpStore %x %42
+         %45 = OpLoad %v2uint %x
+         %48 = OpBitwiseAnd %v2uint %45 %47
+         %44 = OpINotEqual %v2bool %48 %18
+         %43 = OpSelect %v2uint %44 %28 %50
+         %51 = OpLoad %v2uint %x
+         %52 = OpShiftRightLogical %v2uint %51 %43
+               OpStore %x %52
+         %55 = OpLoad %v2uint %x
+         %58 = OpBitwiseAnd %v2uint %55 %57
+         %54 = OpINotEqual %v2bool %58 %18
+         %53 = OpSelect %v2uint %54 %28 %60
+         %61 = OpLoad %v2uint %x
+         %62 = OpShiftRightLogical %v2uint %61 %53
+               OpStore %x %62
+         %65 = OpLoad %v2uint %x
+         %68 = OpBitwiseAnd %v2uint %65 %67
+         %64 = OpINotEqual %v2bool %68 %18
+         %63 = OpSelect %v2uint %64 %28 %67
+         %70 = OpLoad %v2uint %x
+         %71 = OpIEqual %v2bool %70 %28
+         %69 = OpSelect %v2uint %71 %73 %28
+         %75 = OpBitwiseOr %v2uint %19 %33
+         %76 = OpBitwiseOr %v2uint %75 %43
+         %77 = OpBitwiseOr %v2uint %76 %53
+         %78 = OpBitwiseOr %v2uint %77 %63
+         %79 = OpBitwiseOr %v2uint %78 %69
+               OpReturnValue %79
+               OpFunctionEnd
+%firstTrailingBit_45eb10 = OpFunction %void None %80
+         %83 = OpLabel
+        %res = OpVariable %_ptr_Function_v2uint Function %18
+         %84 = OpFunctionCall %v2uint %tint_first_trailing_bit %18
+               OpStore %res %84
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %86
+         %88 = OpLabel
+         %89 = OpFunctionCall %void %firstTrailingBit_45eb10
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %80
+         %91 = OpLabel
+         %92 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %92
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %80
+         %95 = OpLabel
+         %96 = OpFunctionCall %void %firstTrailingBit_45eb10
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %80
+         %98 = OpLabel
+         %99 = OpFunctionCall %void %firstTrailingBit_45eb10
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/45eb10.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_45eb10() {
+  var res : vec2<u32> = firstTrailingBit(vec2<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_45eb10();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_45eb10();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_45eb10();
+}

test/tint/builtins/gen/firstTrailingBit/47d475.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(u32) -> u32
+fn firstTrailingBit_47d475() {
+  var res: u32 = firstTrailingBit(1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_47d475();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_47d475();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_47d475();
+}

test/tint/builtins/gen/firstTrailingBit/47d475.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+uint tint_first_trailing_bit(uint v) {
+  uint x = uint(v);
+  uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_47d475() {
+  uint res = tint_first_trailing_bit(1u);
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_47d475();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+uint tint_first_trailing_bit(uint v) {
+  uint x = uint(v);
+  uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_47d475() {
+  uint res = tint_first_trailing_bit(1u);
+}
+
+void fragment_main() {
+  firstTrailingBit_47d475();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uint tint_first_trailing_bit(uint v) {
+  uint x = uint(v);
+  uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_47d475() {
+  uint res = tint_first_trailing_bit(1u);
+}
+
+void compute_main() {
+  firstTrailingBit_47d475();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/47d475.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+uint tint_first_trailing_bit(uint v) {
+  uint x = uint(v);
+  const uint b16 = (bool((x & 65535u)) ? 0u : 16u);
+  x = (x >> b16);
+  const uint b8 = (bool((x & 255u)) ? 0u : 8u);
+  x = (x >> b8);
+  const uint b4 = (bool((x & 15u)) ? 0u : 4u);
+  x = (x >> b4);
+  const uint b2 = (bool((x & 3u)) ? 0u : 2u);
+  x = (x >> b2);
+  const uint b1 = (bool((x & 1u)) ? 0u : 1u);
+  const uint is_zero = ((x == 0u) ? 4294967295u : 0u);
+  return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_47d475() {
+  uint res = tint_first_trailing_bit(1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_47d475();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_47d475();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_47d475();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/47d475.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint tint_first_trailing_bit(uint v) {
+  uint x = uint(v);
+  uint const b16 = select(16u, 0u, bool((x & 65535u)));
+  x = (x >> b16);
+  uint const b8 = select(8u, 0u, bool((x & 255u)));
+  x = (x >> b8);
+  uint const b4 = select(4u, 0u, bool((x & 15u)));
+  x = (x >> b4);
+  uint const b2 = select(2u, 0u, bool((x & 3u)));
+  x = (x >> b2);
+  uint const b1 = select(1u, 0u, bool((x & 1u)));
+  uint const is_zero = select(0u, 4294967295u, (x == 0u));
+  return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_47d475() {
+  uint res = tint_first_trailing_bit(1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_47d475();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_47d475();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_47d475();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/47d475.wgsl.expected.spvasm

@@ -0,0 +1,129 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 87
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_47d475 "firstTrailingBit_47d475"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %uint = OpTypeInt 32 0
+          %9 = OpTypeFunction %uint %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+       %bool = OpTypeBool
+ %uint_65535 = OpConstant %uint 65535
+     %uint_0 = OpConstant %uint 0
+    %uint_16 = OpConstant %uint 16
+   %uint_255 = OpConstant %uint 255
+     %uint_8 = OpConstant %uint 8
+    %uint_15 = OpConstant %uint 15
+     %uint_4 = OpConstant %uint 4
+     %uint_3 = OpConstant %uint 3
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%uint_4294967295 = OpConstant %uint 4294967295
+       %void = OpTypeVoid
+         %67 = OpTypeFunction %void
+         %73 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %uint None %9
+          %v = OpFunctionParameter %uint
+         %13 = OpLabel
+          %x = OpVariable %_ptr_Function_uint Function %17
+               OpStore %x %v
+         %21 = OpLoad %uint %x
+         %23 = OpBitwiseAnd %uint %21 %uint_65535
+         %19 = OpINotEqual %bool %23 %17
+         %18 = OpSelect %uint %19 %uint_0 %uint_16
+         %26 = OpLoad %uint %x
+         %27 = OpShiftRightLogical %uint %26 %18
+               OpStore %x %27
+         %30 = OpLoad %uint %x
+         %32 = OpBitwiseAnd %uint %30 %uint_255
+         %29 = OpINotEqual %bool %32 %17
+         %28 = OpSelect %uint %29 %uint_0 %uint_8
+         %34 = OpLoad %uint %x
+         %35 = OpShiftRightLogical %uint %34 %28
+               OpStore %x %35
+         %38 = OpLoad %uint %x
+         %40 = OpBitwiseAnd %uint %38 %uint_15
+         %37 = OpINotEqual %bool %40 %17
+         %36 = OpSelect %uint %37 %uint_0 %uint_4
+         %42 = OpLoad %uint %x
+         %43 = OpShiftRightLogical %uint %42 %36
+               OpStore %x %43
+         %46 = OpLoad %uint %x
+         %48 = OpBitwiseAnd %uint %46 %uint_3
+         %45 = OpINotEqual %bool %48 %17
+         %44 = OpSelect %uint %45 %uint_0 %uint_2
+         %50 = OpLoad %uint %x
+         %51 = OpShiftRightLogical %uint %50 %44
+               OpStore %x %51
+         %54 = OpLoad %uint %x
+         %56 = OpBitwiseAnd %uint %54 %uint_1
+         %53 = OpINotEqual %bool %56 %17
+         %52 = OpSelect %uint %53 %uint_0 %uint_1
+         %58 = OpLoad %uint %x
+         %59 = OpIEqual %bool %58 %uint_0
+         %57 = OpSelect %uint %59 %uint_4294967295 %uint_0
+         %62 = OpBitwiseOr %uint %18 %28
+         %63 = OpBitwiseOr %uint %62 %36
+         %64 = OpBitwiseOr %uint %63 %44
+         %65 = OpBitwiseOr %uint %64 %52
+         %66 = OpBitwiseOr %uint %65 %57
+               OpReturnValue %66
+               OpFunctionEnd
+%firstTrailingBit_47d475 = OpFunction %void None %67
+         %70 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %17
+         %71 = OpFunctionCall %uint %tint_first_trailing_bit %uint_1
+               OpStore %res %71
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %73
+         %75 = OpLabel
+         %76 = OpFunctionCall %void %firstTrailingBit_47d475
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %67
+         %78 = OpLabel
+         %79 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %79
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %67
+         %82 = OpLabel
+         %83 = OpFunctionCall %void %firstTrailingBit_47d475
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %67
+         %85 = OpLabel
+         %86 = OpFunctionCall %void %firstTrailingBit_47d475
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/47d475.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_47d475() {
+  var res : u32 = firstTrailingBit(1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_47d475();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_47d475();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_47d475();
+}

test/tint/builtins/gen/firstTrailingBit/50c072.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(vec<2, i32>) -> vec<2, i32>
+fn firstTrailingBit_50c072() {
+  var res: vec2<i32> = firstTrailingBit(vec2<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_50c072();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_50c072();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_50c072();
+}

test/tint/builtins/gen/firstTrailingBit/50c072.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+ivec2 tint_first_trailing_bit(ivec2 v) {
+  uvec2 x = uvec2(v);
+  uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
+  x = (x >> b16);
+  uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
+  x = (x >> b8);
+  uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
+  x = (x >> b4);
+  uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
+  x = (x >> b2);
+  uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
+  uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
+  return ivec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_50c072() {
+  ivec2 res = tint_first_trailing_bit(ivec2(0, 0));
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_50c072();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+ivec2 tint_first_trailing_bit(ivec2 v) {
+  uvec2 x = uvec2(v);
+  uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
+  x = (x >> b16);
+  uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
+  x = (x >> b8);
+  uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
+  x = (x >> b4);
+  uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
+  x = (x >> b2);
+  uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
+  uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
+  return ivec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_50c072() {
+  ivec2 res = tint_first_trailing_bit(ivec2(0, 0));
+}
+
+void fragment_main() {
+  firstTrailingBit_50c072();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+ivec2 tint_first_trailing_bit(ivec2 v) {
+  uvec2 x = uvec2(v);
+  uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
+  x = (x >> b16);
+  uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
+  x = (x >> b8);
+  uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
+  x = (x >> b4);
+  uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
+  x = (x >> b2);
+  uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
+  uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
+  return ivec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_50c072() {
+  ivec2 res = tint_first_trailing_bit(ivec2(0, 0));
+}
+
+void compute_main() {
+  firstTrailingBit_50c072();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/50c072.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+int2 tint_first_trailing_bit(int2 v) {
+  uint2 x = uint2(v);
+  const uint2 b16 = (bool2((x & uint2((65535u).xx))) ? uint2((0u).xx) : uint2((16u).xx));
+  x = (x >> b16);
+  const uint2 b8 = (bool2((x & uint2((255u).xx))) ? uint2((0u).xx) : uint2((8u).xx));
+  x = (x >> b8);
+  const uint2 b4 = (bool2((x & uint2((15u).xx))) ? uint2((0u).xx) : uint2((4u).xx));
+  x = (x >> b4);
+  const uint2 b2 = (bool2((x & uint2((3u).xx))) ? uint2((0u).xx) : uint2((2u).xx));
+  x = (x >> b2);
+  const uint2 b1 = (bool2((x & uint2((1u).xx))) ? uint2((0u).xx) : uint2((1u).xx));
+  const uint2 is_zero = ((x == uint2((0u).xx)) ? uint2((4294967295u).xx) : uint2((0u).xx));
+  return int2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_50c072() {
+  int2 res = tint_first_trailing_bit(int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_50c072();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_50c072();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_50c072();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/50c072.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int2 tint_first_trailing_bit(int2 v) {
+  uint2 x = uint2(v);
+  uint2 const b16 = select(uint2(16u), uint2(0u), bool2((x & uint2(65535u))));
+  x = (x >> b16);
+  uint2 const b8 = select(uint2(8u), uint2(0u), bool2((x & uint2(255u))));
+  x = (x >> b8);
+  uint2 const b4 = select(uint2(4u), uint2(0u), bool2((x & uint2(15u))));
+  x = (x >> b4);
+  uint2 const b2 = select(uint2(2u), uint2(0u), bool2((x & uint2(3u))));
+  x = (x >> b2);
+  uint2 const b1 = select(uint2(1u), uint2(0u), bool2((x & uint2(1u))));
+  uint2 const is_zero = select(uint2(0u), uint2(4294967295u), (x == uint2(0u)));
+  return int2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_50c072() {
+  int2 res = tint_first_trailing_bit(int2());
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_50c072();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_50c072();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_50c072();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/50c072.wgsl.expected.spvasm

@@ -0,0 +1,148 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 104
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_50c072 "firstTrailingBit_50c072"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+          %9 = OpTypeFunction %v2int %v2int
+       %uint = OpTypeInt 32 0
+     %v2uint = OpTypeVector %uint 2
+%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %20 = OpConstantNull %v2uint
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+ %uint_65535 = OpConstant %uint 65535
+         %27 = OpConstantComposite %v2uint %uint_65535 %uint_65535
+     %uint_0 = OpConstant %uint 0
+         %30 = OpConstantComposite %v2uint %uint_0 %uint_0
+    %uint_16 = OpConstant %uint 16
+         %32 = OpConstantComposite %v2uint %uint_16 %uint_16
+   %uint_255 = OpConstant %uint 255
+         %39 = OpConstantComposite %v2uint %uint_255 %uint_255
+     %uint_8 = OpConstant %uint 8
+         %42 = OpConstantComposite %v2uint %uint_8 %uint_8
+    %uint_15 = OpConstant %uint 15
+         %49 = OpConstantComposite %v2uint %uint_15 %uint_15
+     %uint_4 = OpConstant %uint 4
+         %52 = OpConstantComposite %v2uint %uint_4 %uint_4
+     %uint_3 = OpConstant %uint 3
+         %59 = OpConstantComposite %v2uint %uint_3 %uint_3
+     %uint_2 = OpConstant %uint 2
+         %62 = OpConstantComposite %v2uint %uint_2 %uint_2
+     %uint_1 = OpConstant %uint 1
+         %69 = OpConstantComposite %v2uint %uint_1 %uint_1
+%uint_4294967295 = OpConstant %uint 4294967295
+         %75 = OpConstantComposite %v2uint %uint_4294967295 %uint_4294967295
+       %void = OpTypeVoid
+         %82 = OpTypeFunction %void
+         %87 = OpConstantNull %v2int
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %90 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %v2int None %9
+          %v = OpFunctionParameter %v2int
+         %14 = OpLabel
+          %x = OpVariable %_ptr_Function_v2uint Function %20
+         %15 = OpBitcast %v2uint %v
+               OpStore %x %15
+         %25 = OpLoad %v2uint %x
+         %28 = OpBitwiseAnd %v2uint %25 %27
+         %22 = OpINotEqual %v2bool %28 %20
+         %21 = OpSelect %v2uint %22 %30 %32
+         %33 = OpLoad %v2uint %x
+         %34 = OpShiftRightLogical %v2uint %33 %21
+               OpStore %x %34
+         %37 = OpLoad %v2uint %x
+         %40 = OpBitwiseAnd %v2uint %37 %39
+         %36 = OpINotEqual %v2bool %40 %20
+         %35 = OpSelect %v2uint %36 %30 %42
+         %43 = OpLoad %v2uint %x
+         %44 = OpShiftRightLogical %v2uint %43 %35
+               OpStore %x %44
+         %47 = OpLoad %v2uint %x
+         %50 = OpBitwiseAnd %v2uint %47 %49
+         %46 = OpINotEqual %v2bool %50 %20
+         %45 = OpSelect %v2uint %46 %30 %52
+         %53 = OpLoad %v2uint %x
+         %54 = OpShiftRightLogical %v2uint %53 %45
+               OpStore %x %54
+         %57 = OpLoad %v2uint %x
+         %60 = OpBitwiseAnd %v2uint %57 %59
+         %56 = OpINotEqual %v2bool %60 %20
+         %55 = OpSelect %v2uint %56 %30 %62
+         %63 = OpLoad %v2uint %x
+         %64 = OpShiftRightLogical %v2uint %63 %55
+               OpStore %x %64
+         %67 = OpLoad %v2uint %x
+         %70 = OpBitwiseAnd %v2uint %67 %69
+         %66 = OpINotEqual %v2bool %70 %20
+         %65 = OpSelect %v2uint %66 %30 %69
+         %72 = OpLoad %v2uint %x
+         %73 = OpIEqual %v2bool %72 %30
+         %71 = OpSelect %v2uint %73 %75 %30
+         %77 = OpBitwiseOr %v2uint %21 %35
+         %78 = OpBitwiseOr %v2uint %77 %45
+         %79 = OpBitwiseOr %v2uint %78 %55
+         %80 = OpBitwiseOr %v2uint %79 %65
+         %81 = OpBitwiseOr %v2uint %80 %71
+         %76 = OpBitcast %v2int %81
+               OpReturnValue %76
+               OpFunctionEnd
+%firstTrailingBit_50c072 = OpFunction %void None %82
+         %85 = OpLabel
+        %res = OpVariable %_ptr_Function_v2int Function %87
+         %86 = OpFunctionCall %v2int %tint_first_trailing_bit %87
+               OpStore %res %86
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %90
+         %92 = OpLabel
+         %93 = OpFunctionCall %void %firstTrailingBit_50c072
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %82
+         %95 = OpLabel
+         %96 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %96
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %82
+         %99 = OpLabel
+        %100 = OpFunctionCall %void %firstTrailingBit_50c072
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %82
+        %102 = OpLabel
+        %103 = OpFunctionCall %void %firstTrailingBit_50c072
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/50c072.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_50c072() {
+  var res : vec2<i32> = firstTrailingBit(vec2<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_50c072();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_50c072();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_50c072();
+}

test/tint/builtins/gen/firstTrailingBit/7496d6.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(vec<3, i32>) -> vec<3, i32>
+fn firstTrailingBit_7496d6() {
+  var res: vec3<i32> = firstTrailingBit(vec3<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_7496d6();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_7496d6();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_7496d6();
+}

test/tint/builtins/gen/firstTrailingBit/7496d6.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+ivec3 tint_first_trailing_bit(ivec3 v) {
+  uvec3 x = uvec3(v);
+  uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
+  x = (x >> b16);
+  uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
+  x = (x >> b8);
+  uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
+  x = (x >> b4);
+  uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
+  x = (x >> b2);
+  uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
+  uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
+  return ivec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_7496d6() {
+  ivec3 res = tint_first_trailing_bit(ivec3(0, 0, 0));
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_7496d6();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+ivec3 tint_first_trailing_bit(ivec3 v) {
+  uvec3 x = uvec3(v);
+  uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
+  x = (x >> b16);
+  uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
+  x = (x >> b8);
+  uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
+  x = (x >> b4);
+  uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
+  x = (x >> b2);
+  uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
+  uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
+  return ivec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_7496d6() {
+  ivec3 res = tint_first_trailing_bit(ivec3(0, 0, 0));
+}
+
+void fragment_main() {
+  firstTrailingBit_7496d6();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+ivec3 tint_first_trailing_bit(ivec3 v) {
+  uvec3 x = uvec3(v);
+  uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
+  x = (x >> b16);
+  uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
+  x = (x >> b8);
+  uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
+  x = (x >> b4);
+  uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
+  x = (x >> b2);
+  uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
+  uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
+  return ivec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_7496d6() {
+  ivec3 res = tint_first_trailing_bit(ivec3(0, 0, 0));
+}
+
+void compute_main() {
+  firstTrailingBit_7496d6();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/7496d6.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+int3 tint_first_trailing_bit(int3 v) {
+  uint3 x = uint3(v);
+  const uint3 b16 = (bool3((x & uint3((65535u).xxx))) ? uint3((0u).xxx) : uint3((16u).xxx));
+  x = (x >> b16);
+  const uint3 b8 = (bool3((x & uint3((255u).xxx))) ? uint3((0u).xxx) : uint3((8u).xxx));
+  x = (x >> b8);
+  const uint3 b4 = (bool3((x & uint3((15u).xxx))) ? uint3((0u).xxx) : uint3((4u).xxx));
+  x = (x >> b4);
+  const uint3 b2 = (bool3((x & uint3((3u).xxx))) ? uint3((0u).xxx) : uint3((2u).xxx));
+  x = (x >> b2);
+  const uint3 b1 = (bool3((x & uint3((1u).xxx))) ? uint3((0u).xxx) : uint3((1u).xxx));
+  const uint3 is_zero = ((x == uint3((0u).xxx)) ? uint3((4294967295u).xxx) : uint3((0u).xxx));
+  return int3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_7496d6() {
+  int3 res = tint_first_trailing_bit(int3(0, 0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_7496d6();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_7496d6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_7496d6();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/7496d6.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int3 tint_first_trailing_bit(int3 v) {
+  uint3 x = uint3(v);
+  uint3 const b16 = select(uint3(16u), uint3(0u), bool3((x & uint3(65535u))));
+  x = (x >> b16);
+  uint3 const b8 = select(uint3(8u), uint3(0u), bool3((x & uint3(255u))));
+  x = (x >> b8);
+  uint3 const b4 = select(uint3(4u), uint3(0u), bool3((x & uint3(15u))));
+  x = (x >> b4);
+  uint3 const b2 = select(uint3(2u), uint3(0u), bool3((x & uint3(3u))));
+  x = (x >> b2);
+  uint3 const b1 = select(uint3(1u), uint3(0u), bool3((x & uint3(1u))));
+  uint3 const is_zero = select(uint3(0u), uint3(4294967295u), (x == uint3(0u)));
+  return int3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_7496d6() {
+  int3 res = tint_first_trailing_bit(int3());
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_7496d6();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_7496d6();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_7496d6();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/7496d6.wgsl.expected.spvasm

@@ -0,0 +1,148 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 104
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_7496d6 "firstTrailingBit_7496d6"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+          %9 = OpTypeFunction %v3int %v3int
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %20 = OpConstantNull %v3uint
+       %bool = OpTypeBool
+     %v3bool = OpTypeVector %bool 3
+ %uint_65535 = OpConstant %uint 65535
+         %27 = OpConstantComposite %v3uint %uint_65535 %uint_65535 %uint_65535
+     %uint_0 = OpConstant %uint 0
+         %30 = OpConstantComposite %v3uint %uint_0 %uint_0 %uint_0
+    %uint_16 = OpConstant %uint 16
+         %32 = OpConstantComposite %v3uint %uint_16 %uint_16 %uint_16
+   %uint_255 = OpConstant %uint 255
+         %39 = OpConstantComposite %v3uint %uint_255 %uint_255 %uint_255
+     %uint_8 = OpConstant %uint 8
+         %42 = OpConstantComposite %v3uint %uint_8 %uint_8 %uint_8
+    %uint_15 = OpConstant %uint 15
+         %49 = OpConstantComposite %v3uint %uint_15 %uint_15 %uint_15
+     %uint_4 = OpConstant %uint 4
+         %52 = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_4
+     %uint_3 = OpConstant %uint 3
+         %59 = OpConstantComposite %v3uint %uint_3 %uint_3 %uint_3
+     %uint_2 = OpConstant %uint 2
+         %62 = OpConstantComposite %v3uint %uint_2 %uint_2 %uint_2
+     %uint_1 = OpConstant %uint 1
+         %69 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
+%uint_4294967295 = OpConstant %uint 4294967295
+         %75 = OpConstantComposite %v3uint %uint_4294967295 %uint_4294967295 %uint_4294967295
+       %void = OpTypeVoid
+         %82 = OpTypeFunction %void
+         %87 = OpConstantNull %v3int
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %90 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %v3int None %9
+          %v = OpFunctionParameter %v3int
+         %14 = OpLabel
+          %x = OpVariable %_ptr_Function_v3uint Function %20
+         %15 = OpBitcast %v3uint %v
+               OpStore %x %15
+         %25 = OpLoad %v3uint %x
+         %28 = OpBitwiseAnd %v3uint %25 %27
+         %22 = OpINotEqual %v3bool %28 %20
+         %21 = OpSelect %v3uint %22 %30 %32
+         %33 = OpLoad %v3uint %x
+         %34 = OpShiftRightLogical %v3uint %33 %21
+               OpStore %x %34
+         %37 = OpLoad %v3uint %x
+         %40 = OpBitwiseAnd %v3uint %37 %39
+         %36 = OpINotEqual %v3bool %40 %20
+         %35 = OpSelect %v3uint %36 %30 %42
+         %43 = OpLoad %v3uint %x
+         %44 = OpShiftRightLogical %v3uint %43 %35
+               OpStore %x %44
+         %47 = OpLoad %v3uint %x
+         %50 = OpBitwiseAnd %v3uint %47 %49
+         %46 = OpINotEqual %v3bool %50 %20
+         %45 = OpSelect %v3uint %46 %30 %52
+         %53 = OpLoad %v3uint %x
+         %54 = OpShiftRightLogical %v3uint %53 %45
+               OpStore %x %54
+         %57 = OpLoad %v3uint %x
+         %60 = OpBitwiseAnd %v3uint %57 %59
+         %56 = OpINotEqual %v3bool %60 %20
+         %55 = OpSelect %v3uint %56 %30 %62
+         %63 = OpLoad %v3uint %x
+         %64 = OpShiftRightLogical %v3uint %63 %55
+               OpStore %x %64
+         %67 = OpLoad %v3uint %x
+         %70 = OpBitwiseAnd %v3uint %67 %69
+         %66 = OpINotEqual %v3bool %70 %20
+         %65 = OpSelect %v3uint %66 %30 %69
+         %72 = OpLoad %v3uint %x
+         %73 = OpIEqual %v3bool %72 %30
+         %71 = OpSelect %v3uint %73 %75 %30
+         %77 = OpBitwiseOr %v3uint %21 %35
+         %78 = OpBitwiseOr %v3uint %77 %45
+         %79 = OpBitwiseOr %v3uint %78 %55
+         %80 = OpBitwiseOr %v3uint %79 %65
+         %81 = OpBitwiseOr %v3uint %80 %71
+         %76 = OpBitcast %v3int %81
+               OpReturnValue %76
+               OpFunctionEnd
+%firstTrailingBit_7496d6 = OpFunction %void None %82
+         %85 = OpLabel
+        %res = OpVariable %_ptr_Function_v3int Function %87
+         %86 = OpFunctionCall %v3int %tint_first_trailing_bit %87
+               OpStore %res %86
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %90
+         %92 = OpLabel
+         %93 = OpFunctionCall %void %firstTrailingBit_7496d6
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %82
+         %95 = OpLabel
+         %96 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %96
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %82
+         %99 = OpLabel
+        %100 = OpFunctionCall %void %firstTrailingBit_7496d6
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %82
+        %102 = OpLabel
+        %103 = OpFunctionCall %void %firstTrailingBit_7496d6
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/7496d6.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_7496d6() {
+  var res : vec3<i32> = firstTrailingBit(vec3<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_7496d6();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_7496d6();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_7496d6();
+}

test/tint/builtins/gen/firstTrailingBit/86551b.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(vec<4, i32>) -> vec<4, i32>
+fn firstTrailingBit_86551b() {
+  var res: vec4<i32> = firstTrailingBit(vec4<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_86551b();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_86551b();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_86551b();
+}

test/tint/builtins/gen/firstTrailingBit/86551b.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+ivec4 tint_first_trailing_bit(ivec4 v) {
+  uvec4 x = uvec4(v);
+  uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
+  x = (x >> b16);
+  uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
+  x = (x >> b8);
+  uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
+  x = (x >> b4);
+  uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
+  x = (x >> b2);
+  uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
+  uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
+  return ivec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_86551b() {
+  ivec4 res = tint_first_trailing_bit(ivec4(0, 0, 0, 0));
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_86551b();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+ivec4 tint_first_trailing_bit(ivec4 v) {
+  uvec4 x = uvec4(v);
+  uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
+  x = (x >> b16);
+  uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
+  x = (x >> b8);
+  uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
+  x = (x >> b4);
+  uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
+  x = (x >> b2);
+  uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
+  uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
+  return ivec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_86551b() {
+  ivec4 res = tint_first_trailing_bit(ivec4(0, 0, 0, 0));
+}
+
+void fragment_main() {
+  firstTrailingBit_86551b();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+ivec4 tint_first_trailing_bit(ivec4 v) {
+  uvec4 x = uvec4(v);
+  uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
+  x = (x >> b16);
+  uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
+  x = (x >> b8);
+  uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
+  x = (x >> b4);
+  uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
+  x = (x >> b2);
+  uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
+  uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
+  return ivec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_86551b() {
+  ivec4 res = tint_first_trailing_bit(ivec4(0, 0, 0, 0));
+}
+
+void compute_main() {
+  firstTrailingBit_86551b();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/86551b.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+int4 tint_first_trailing_bit(int4 v) {
+  uint4 x = uint4(v);
+  const uint4 b16 = (bool4((x & uint4((65535u).xxxx))) ? uint4((0u).xxxx) : uint4((16u).xxxx));
+  x = (x >> b16);
+  const uint4 b8 = (bool4((x & uint4((255u).xxxx))) ? uint4((0u).xxxx) : uint4((8u).xxxx));
+  x = (x >> b8);
+  const uint4 b4 = (bool4((x & uint4((15u).xxxx))) ? uint4((0u).xxxx) : uint4((4u).xxxx));
+  x = (x >> b4);
+  const uint4 b2 = (bool4((x & uint4((3u).xxxx))) ? uint4((0u).xxxx) : uint4((2u).xxxx));
+  x = (x >> b2);
+  const uint4 b1 = (bool4((x & uint4((1u).xxxx))) ? uint4((0u).xxxx) : uint4((1u).xxxx));
+  const uint4 is_zero = ((x == uint4((0u).xxxx)) ? uint4((4294967295u).xxxx) : uint4((0u).xxxx));
+  return int4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_86551b() {
+  int4 res = tint_first_trailing_bit(int4(0, 0, 0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_86551b();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_86551b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_86551b();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/86551b.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int4 tint_first_trailing_bit(int4 v) {
+  uint4 x = uint4(v);
+  uint4 const b16 = select(uint4(16u), uint4(0u), bool4((x & uint4(65535u))));
+  x = (x >> b16);
+  uint4 const b8 = select(uint4(8u), uint4(0u), bool4((x & uint4(255u))));
+  x = (x >> b8);
+  uint4 const b4 = select(uint4(4u), uint4(0u), bool4((x & uint4(15u))));
+  x = (x >> b4);
+  uint4 const b2 = select(uint4(2u), uint4(0u), bool4((x & uint4(3u))));
+  x = (x >> b2);
+  uint4 const b1 = select(uint4(1u), uint4(0u), bool4((x & uint4(1u))));
+  uint4 const is_zero = select(uint4(0u), uint4(4294967295u), (x == uint4(0u)));
+  return int4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_86551b() {
+  int4 res = tint_first_trailing_bit(int4());
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_86551b();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_86551b();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_86551b();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/86551b.wgsl.expected.spvasm

@@ -0,0 +1,148 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 104
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_86551b "firstTrailingBit_86551b"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
+          %9 = OpTypeFunction %v4int %v4int
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %20 = OpConstantNull %v4uint
+       %bool = OpTypeBool
+     %v4bool = OpTypeVector %bool 4
+ %uint_65535 = OpConstant %uint 65535
+         %27 = OpConstantComposite %v4uint %uint_65535 %uint_65535 %uint_65535 %uint_65535
+     %uint_0 = OpConstant %uint 0
+         %30 = OpConstantComposite %v4uint %uint_0 %uint_0 %uint_0 %uint_0
+    %uint_16 = OpConstant %uint 16
+         %32 = OpConstantComposite %v4uint %uint_16 %uint_16 %uint_16 %uint_16
+   %uint_255 = OpConstant %uint 255
+         %39 = OpConstantComposite %v4uint %uint_255 %uint_255 %uint_255 %uint_255
+     %uint_8 = OpConstant %uint 8
+         %42 = OpConstantComposite %v4uint %uint_8 %uint_8 %uint_8 %uint_8
+    %uint_15 = OpConstant %uint 15
+         %49 = OpConstantComposite %v4uint %uint_15 %uint_15 %uint_15 %uint_15
+     %uint_4 = OpConstant %uint 4
+         %52 = OpConstantComposite %v4uint %uint_4 %uint_4 %uint_4 %uint_4
+     %uint_3 = OpConstant %uint 3
+         %59 = OpConstantComposite %v4uint %uint_3 %uint_3 %uint_3 %uint_3
+     %uint_2 = OpConstant %uint 2
+         %62 = OpConstantComposite %v4uint %uint_2 %uint_2 %uint_2 %uint_2
+     %uint_1 = OpConstant %uint 1
+         %69 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
+%uint_4294967295 = OpConstant %uint 4294967295
+         %75 = OpConstantComposite %v4uint %uint_4294967295 %uint_4294967295 %uint_4294967295 %uint_4294967295
+       %void = OpTypeVoid
+         %82 = OpTypeFunction %void
+         %87 = OpConstantNull %v4int
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %90 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %v4int None %9
+          %v = OpFunctionParameter %v4int
+         %14 = OpLabel
+          %x = OpVariable %_ptr_Function_v4uint Function %20
+         %15 = OpBitcast %v4uint %v
+               OpStore %x %15
+         %25 = OpLoad %v4uint %x
+         %28 = OpBitwiseAnd %v4uint %25 %27
+         %22 = OpINotEqual %v4bool %28 %20
+         %21 = OpSelect %v4uint %22 %30 %32
+         %33 = OpLoad %v4uint %x
+         %34 = OpShiftRightLogical %v4uint %33 %21
+               OpStore %x %34
+         %37 = OpLoad %v4uint %x
+         %40 = OpBitwiseAnd %v4uint %37 %39
+         %36 = OpINotEqual %v4bool %40 %20
+         %35 = OpSelect %v4uint %36 %30 %42
+         %43 = OpLoad %v4uint %x
+         %44 = OpShiftRightLogical %v4uint %43 %35
+               OpStore %x %44
+         %47 = OpLoad %v4uint %x
+         %50 = OpBitwiseAnd %v4uint %47 %49
+         %46 = OpINotEqual %v4bool %50 %20
+         %45 = OpSelect %v4uint %46 %30 %52
+         %53 = OpLoad %v4uint %x
+         %54 = OpShiftRightLogical %v4uint %53 %45
+               OpStore %x %54
+         %57 = OpLoad %v4uint %x
+         %60 = OpBitwiseAnd %v4uint %57 %59
+         %56 = OpINotEqual %v4bool %60 %20
+         %55 = OpSelect %v4uint %56 %30 %62
+         %63 = OpLoad %v4uint %x
+         %64 = OpShiftRightLogical %v4uint %63 %55
+               OpStore %x %64
+         %67 = OpLoad %v4uint %x
+         %70 = OpBitwiseAnd %v4uint %67 %69
+         %66 = OpINotEqual %v4bool %70 %20
+         %65 = OpSelect %v4uint %66 %30 %69
+         %72 = OpLoad %v4uint %x
+         %73 = OpIEqual %v4bool %72 %30
+         %71 = OpSelect %v4uint %73 %75 %30
+         %77 = OpBitwiseOr %v4uint %21 %35
+         %78 = OpBitwiseOr %v4uint %77 %45
+         %79 = OpBitwiseOr %v4uint %78 %55
+         %80 = OpBitwiseOr %v4uint %79 %65
+         %81 = OpBitwiseOr %v4uint %80 %71
+         %76 = OpBitcast %v4int %81
+               OpReturnValue %76
+               OpFunctionEnd
+%firstTrailingBit_86551b = OpFunction %void None %82
+         %85 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %87
+         %86 = OpFunctionCall %v4int %tint_first_trailing_bit %87
+               OpStore %res %86
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %90
+         %92 = OpLabel
+         %93 = OpFunctionCall %void %firstTrailingBit_86551b
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %82
+         %95 = OpLabel
+         %96 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %96
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %82
+         %99 = OpLabel
+        %100 = OpFunctionCall %void %firstTrailingBit_86551b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %82
+        %102 = OpLabel
+        %103 = OpFunctionCall %void %firstTrailingBit_86551b
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/86551b.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_86551b() {
+  var res : vec4<i32> = firstTrailingBit(vec4<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_86551b();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_86551b();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_86551b();
+}

test/tint/builtins/gen/firstTrailingBit/cb51ce.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/builtin-gen
+// using the template:
+//   test/tint/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/tint/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn firstTrailingBit(vec<3, u32>) -> vec<3, u32>
+fn firstTrailingBit_cb51ce() {
+  var res: vec3<u32> = firstTrailingBit(vec3<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_cb51ce();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_cb51ce();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_cb51ce();
+}

test/tint/builtins/gen/firstTrailingBit/cb51ce.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+uvec3 tint_first_trailing_bit(uvec3 v) {
+  uvec3 x = uvec3(v);
+  uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
+  x = (x >> b16);
+  uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
+  x = (x >> b8);
+  uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
+  x = (x >> b4);
+  uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
+  x = (x >> b2);
+  uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
+  uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
+  return uvec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_cb51ce() {
+  uvec3 res = tint_first_trailing_bit(uvec3(0u, 0u, 0u));
+}
+
+vec4 vertex_main() {
+  firstTrailingBit_cb51ce();
+  return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+uvec3 tint_first_trailing_bit(uvec3 v) {
+  uvec3 x = uvec3(v);
+  uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
+  x = (x >> b16);
+  uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
+  x = (x >> b8);
+  uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
+  x = (x >> b4);
+  uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
+  x = (x >> b2);
+  uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
+  uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
+  return uvec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_cb51ce() {
+  uvec3 res = tint_first_trailing_bit(uvec3(0u, 0u, 0u));
+}
+
+void fragment_main() {
+  firstTrailingBit_cb51ce();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uvec3 tint_first_trailing_bit(uvec3 v) {
+  uvec3 x = uvec3(v);
+  uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
+  x = (x >> b16);
+  uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
+  x = (x >> b8);
+  uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
+  x = (x >> b4);
+  uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
+  x = (x >> b2);
+  uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
+  uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
+  return uvec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_cb51ce() {
+  uvec3 res = tint_first_trailing_bit(uvec3(0u, 0u, 0u));
+}
+
+void compute_main() {
+  firstTrailingBit_cb51ce();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/cb51ce.wgsl.expected.hlsl

@@ -0,0 +1,45 @@
+uint3 tint_first_trailing_bit(uint3 v) {
+  uint3 x = uint3(v);
+  const uint3 b16 = (bool3((x & uint3((65535u).xxx))) ? uint3((0u).xxx) : uint3((16u).xxx));
+  x = (x >> b16);
+  const uint3 b8 = (bool3((x & uint3((255u).xxx))) ? uint3((0u).xxx) : uint3((8u).xxx));
+  x = (x >> b8);
+  const uint3 b4 = (bool3((x & uint3((15u).xxx))) ? uint3((0u).xxx) : uint3((4u).xxx));
+  x = (x >> b4);
+  const uint3 b2 = (bool3((x & uint3((3u).xxx))) ? uint3((0u).xxx) : uint3((2u).xxx));
+  x = (x >> b2);
+  const uint3 b1 = (bool3((x & uint3((1u).xxx))) ? uint3((0u).xxx) : uint3((1u).xxx));
+  const uint3 is_zero = ((x == uint3((0u).xxx)) ? uint3((4294967295u).xxx) : uint3((0u).xxx));
+  return uint3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_cb51ce() {
+  uint3 res = tint_first_trailing_bit(uint3(0u, 0u, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_cb51ce();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  firstTrailingBit_cb51ce();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  firstTrailingBit_cb51ce();
+  return;
+}

test/tint/builtins/gen/firstTrailingBit/cb51ce.wgsl.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint3 tint_first_trailing_bit(uint3 v) {
+  uint3 x = uint3(v);
+  uint3 const b16 = select(uint3(16u), uint3(0u), bool3((x & uint3(65535u))));
+  x = (x >> b16);
+  uint3 const b8 = select(uint3(8u), uint3(0u), bool3((x & uint3(255u))));
+  x = (x >> b8);
+  uint3 const b4 = select(uint3(4u), uint3(0u), bool3((x & uint3(15u))));
+  x = (x >> b4);
+  uint3 const b2 = select(uint3(2u), uint3(0u), bool3((x & uint3(3u))));
+  x = (x >> b2);
+  uint3 const b1 = select(uint3(1u), uint3(0u), bool3((x & uint3(1u))));
+  uint3 const is_zero = select(uint3(0u), uint3(4294967295u), (x == uint3(0u)));
+  return uint3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
+}
+
+void firstTrailingBit_cb51ce() {
+  uint3 res = tint_first_trailing_bit(uint3());
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  firstTrailingBit_cb51ce();
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  firstTrailingBit_cb51ce();
+  return;
+}
+
+kernel void compute_main() {
+  firstTrailingBit_cb51ce();
+  return;
+}
+

test/tint/builtins/gen/firstTrailingBit/cb51ce.wgsl.expected.spvasm

@@ -0,0 +1,142 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 100
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %tint_first_trailing_bit "tint_first_trailing_bit"
+               OpName %v "v"
+               OpName %x "x"
+               OpName %firstTrailingBit_cb51ce "firstTrailingBit_cb51ce"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+          %9 = OpTypeFunction %v3uint %v3uint
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %18 = OpConstantNull %v3uint
+       %bool = OpTypeBool
+     %v3bool = OpTypeVector %bool 3
+ %uint_65535 = OpConstant %uint 65535
+         %25 = OpConstantComposite %v3uint %uint_65535 %uint_65535 %uint_65535
+     %uint_0 = OpConstant %uint 0
+         %28 = OpConstantComposite %v3uint %uint_0 %uint_0 %uint_0
+    %uint_16 = OpConstant %uint 16
+         %30 = OpConstantComposite %v3uint %uint_16 %uint_16 %uint_16
+   %uint_255 = OpConstant %uint 255
+         %37 = OpConstantComposite %v3uint %uint_255 %uint_255 %uint_255
+     %uint_8 = OpConstant %uint 8
+         %40 = OpConstantComposite %v3uint %uint_8 %uint_8 %uint_8
+    %uint_15 = OpConstant %uint 15
+         %47 = OpConstantComposite %v3uint %uint_15 %uint_15 %uint_15
+     %uint_4 = OpConstant %uint 4
+         %50 = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_4
+     %uint_3 = OpConstant %uint 3
+         %57 = OpConstantComposite %v3uint %uint_3 %uint_3 %uint_3
+     %uint_2 = OpConstant %uint 2
+         %60 = OpConstantComposite %v3uint %uint_2 %uint_2 %uint_2
+     %uint_1 = OpConstant %uint 1
+         %67 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
+%uint_4294967295 = OpConstant %uint 4294967295
+         %73 = OpConstantComposite %v3uint %uint_4294967295 %uint_4294967295 %uint_4294967295
+       %void = OpTypeVoid
+         %80 = OpTypeFunction %void
+         %86 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_first_trailing_bit = OpFunction %v3uint None %9
+          %v = OpFunctionParameter %v3uint
+         %14 = OpLabel
+          %x = OpVariable %_ptr_Function_v3uint Function %18
+               OpStore %x %v
+         %23 = OpLoad %v3uint %x
+         %26 = OpBitwiseAnd %v3uint %23 %25
+         %20 = OpINotEqual %v3bool %26 %18
+         %19 = OpSelect %v3uint %20 %28 %30
+         %31 = OpLoad %v3uint %x
+         %32 = OpShiftRightLogical %v3uint %31 %19
+               OpStore %x %32
+         %35 = OpLoad %v3uint %x
+         %38 = OpBitwiseAnd %v3uint %35 %37
+         %34 = OpINotEqual %v3bool %38 %18
+         %33 = OpSelect %v3uint %34 %28 %40
+         %41 = OpLoad %v3uint %x
+         %42 = OpShiftRightLogical %v3uint %41 %33
+               OpStore %x %42
+         %45 = OpLoad %v3uint %x
+         %48 = OpBitwiseAnd %v3uint %45 %47
+         %44 = OpINotEqual %v3bool %48 %18
+         %43 = OpSelect %v3uint %44 %28 %50
+         %51 = OpLoad %v3uint %x
+         %52 = OpShiftRightLogical %v3uint %51 %43
+               OpStore %x %52
+         %55 = OpLoad %v3uint %x
+         %58 = OpBitwiseAnd %v3uint %55 %57
+         %54 = OpINotEqual %v3bool %58 %18
+         %53 = OpSelect %v3uint %54 %28 %60
+         %61 = OpLoad %v3uint %x
+         %62 = OpShiftRightLogical %v3uint %61 %53
+               OpStore %x %62
+         %65 = OpLoad %v3uint %x
+         %68 = OpBitwiseAnd %v3uint %65 %67
+         %64 = OpINotEqual %v3bool %68 %18
+         %63 = OpSelect %v3uint %64 %28 %67
+         %70 = OpLoad %v3uint %x
+         %71 = OpIEqual %v3bool %70 %28
+         %69 = OpSelect %v3uint %71 %73 %28
+         %75 = OpBitwiseOr %v3uint %19 %33
+         %76 = OpBitwiseOr %v3uint %75 %43
+         %77 = OpBitwiseOr %v3uint %76 %53
+         %78 = OpBitwiseOr %v3uint %77 %63
+         %79 = OpBitwiseOr %v3uint %78 %69
+               OpReturnValue %79
+               OpFunctionEnd
+%firstTrailingBit_cb51ce = OpFunction %void None %80
+         %83 = OpLabel
+        %res = OpVariable %_ptr_Function_v3uint Function %18
+         %84 = OpFunctionCall %v3uint %tint_first_trailing_bit %18
+               OpStore %res %84
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %86
+         %88 = OpLabel
+         %89 = OpFunctionCall %void %firstTrailingBit_cb51ce
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %80
+         %91 = OpLabel
+         %92 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %92
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %80
+         %95 = OpLabel
+         %96 = OpFunctionCall %void %firstTrailingBit_cb51ce
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %80
+         %98 = OpLabel
+         %99 = OpFunctionCall %void %firstTrailingBit_cb51ce
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/firstTrailingBit/cb51ce.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn firstTrailingBit_cb51ce() {
+  var res : vec3<u32> = firstTrailingBit(vec3<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  firstTrailingBit_cb51ce();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  firstTrailingBit_cb51ce();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  firstTrailingBit_cb51ce();
+}