builtins: Add countLeadingZeros

Requires polyfilling for all but the MSL backend.

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

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

Doxyfile

@@ -1081,13 +1081,6 @@ VERBATIM_HEADERS       = YES
 
 ALPHABETICAL_INDEX     = YES
 
-# The COLS_IN_ALPHA_INDEX tag can be used to specify the number of columns in
-# which the alphabetical index list will be split.
-# Minimum value: 1, maximum value: 20, default value: 5.
-# This tag requires that the tag ALPHABETICAL_INDEX is set to YES.
-
-COLS_IN_ALPHA_INDEX    = 5
-
 # In case all classes in a project start with a common prefix, all classes will
 # be put under the same header in the alphabetical index. The IGNORE_PREFIX tag
 # can be used to specify a prefix (or a list of prefixes) that should be ignored

src/tint/BUILD.gn

@@ -436,6 +436,8 @@ libtint_source_set("libtint_core_all_src") {
     "transform/array_length_from_uniform.h",
     "transform/binding_remapper.cc",
     "transform/binding_remapper.h",
+    "transform/builtin_polyfill.cc",
+    "transform/builtin_polyfill.h",
     "transform/calculate_array_length.cc",
     "transform/calculate_array_length.h",
     "transform/canonicalize_entry_point_io.cc",

src/tint/CMakeLists.txt

@@ -312,6 +312,8 @@ set(TINT_LIB_SRCS
   transform/array_length_from_uniform.h
   transform/binding_remapper.cc
   transform/binding_remapper.h
+  transform/builtin_polyfill.cc
+  transform/builtin_polyfill.h
   transform/calculate_array_length.cc
   transform/calculate_array_length.h
   transform/combine_samplers.cc
@@ -1006,6 +1008,7 @@ if(TINT_BUILD_TESTS)
       transform/add_spirv_block_attribute_test.cc
       transform/array_length_from_uniform_test.cc
       transform/binding_remapper_test.cc
+      transform/builtin_polyfill_test.cc
       transform/calculate_array_length_test.cc
       transform/canonicalize_entry_point_io_test.cc
       transform/combine_samplers_test.cc

src/tint/builtins.def

@@ -286,6 +286,8 @@ fn cos(f32) -> f32
 fn cos<N: num>(vec<N, f32>) -> vec<N, f32>
 fn cosh(f32) -> f32
 fn cosh<N: num>(vec<N, f32>) -> vec<N, f32>
+fn countLeadingZeros<T: iu32>(T) -> T
+fn countLeadingZeros<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
 fn countOneBits<T: iu32>(T) -> T
 fn countOneBits<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
 fn cross(vec3<f32>, vec3<f32>) -> vec3<f32>

src/tint/program_builder.h

@@ -1751,6 +1751,46 @@ class ProgramBuilder {
                                          Expr(std::forward<RHS>(rhs)));
   }
 
+  /// @param lhs the left hand argument to the greater than operation
+  /// @param rhs the right hand argument to the greater than operation
+  /// @returns a `ast::BinaryExpression` of `lhs` > `rhs`
+  template <typename LHS, typename RHS>
+  const ast::BinaryExpression* GreaterThan(LHS&& lhs, RHS&& rhs) {
+    return create<ast::BinaryExpression>(ast::BinaryOp::kGreaterThan,
+                                         Expr(std::forward<LHS>(lhs)),
+                                         Expr(std::forward<RHS>(rhs)));
+  }
+
+  /// @param lhs the left hand argument to the greater than or equal operation
+  /// @param rhs the right hand argument to the greater than or equal operation
+  /// @returns a `ast::BinaryExpression` of `lhs` >= `rhs`
+  template <typename LHS, typename RHS>
+  const ast::BinaryExpression* GreaterThanEqual(LHS&& lhs, RHS&& rhs) {
+    return create<ast::BinaryExpression>(ast::BinaryOp::kGreaterThanEqual,
+                                         Expr(std::forward<LHS>(lhs)),
+                                         Expr(std::forward<RHS>(rhs)));
+  }
+
+  /// @param lhs the left hand argument to the less than operation
+  /// @param rhs the right hand argument to the less than operation
+  /// @returns a `ast::BinaryExpression` of `lhs` < `rhs`
+  template <typename LHS, typename RHS>
+  const ast::BinaryExpression* LessThan(LHS&& lhs, RHS&& rhs) {
+    return create<ast::BinaryExpression>(ast::BinaryOp::kLessThan,
+                                         Expr(std::forward<LHS>(lhs)),
+                                         Expr(std::forward<RHS>(rhs)));
+  }
+
+  /// @param lhs the left hand argument to the less than or equal operation
+  /// @param rhs the right hand argument to the less than or equal operation
+  /// @returns a `ast::BinaryExpression` of `lhs` <= `rhs`
+  template <typename LHS, typename RHS>
+  const ast::BinaryExpression* LessThanEqual(LHS&& lhs, RHS&& rhs) {
+    return create<ast::BinaryExpression>(ast::BinaryOp::kLessThanEqual,
+                                         Expr(std::forward<LHS>(lhs)),
+                                         Expr(std::forward<RHS>(rhs)));
+  }
+
   /// @param lhs the left hand argument to the equal expression
   /// @param rhs the right hand argument to the equal expression
   /// @returns a `ast::BinaryExpression` comparing `lhs` equal to `rhs`

src/tint/sem/builtin_type.cc

@@ -66,6 +66,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
   if (name == "cosh") {
     return BuiltinType::kCosh;
   }
+  if (name == "countLeadingZeros") {
+    return BuiltinType::kCountLeadingZeros;
+  }
   if (name == "countOneBits") {
     return BuiltinType::kCountOneBits;
   }
@@ -367,6 +370,8 @@ const char* str(BuiltinType i) {
       return "cos";
     case BuiltinType::kCosh:
       return "cosh";
+    case BuiltinType::kCountLeadingZeros:
+      return "countLeadingZeros";
     case BuiltinType::kCountOneBits:
       return "countOneBits";
     case BuiltinType::kCross:

src/tint/sem/builtin_type.h

@@ -46,6 +46,7 @@ enum class BuiltinType {
   kClamp,
   kCos,
   kCosh,
+  kCountLeadingZeros,
   kCountOneBits,
   kCross,
   kDegrees,

src/tint/transform/builtin_polyfill.cc

@@ -0,0 +1,188 @@
+// Copyright 2022 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.
+
+#include "src/tint/transform/builtin_polyfill.h"
+
+#include <unordered_map>
+
+#include "src/tint/program_builder.h"
+#include "src/tint/sem/builtin.h"
+#include "src/tint/sem/call.h"
+#include "src/tint/utils/map.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::BuiltinPolyfill);
+TINT_INSTANTIATE_TYPEINFO(tint::transform::BuiltinPolyfill::Config);
+
+namespace tint {
+namespace transform {
+
+/// The PIMPL state for the BuiltinPolyfill transform
+struct BuiltinPolyfill::State {
+  /// Constructor
+  /// @param c the CloneContext
+  explicit State(CloneContext& c) : ctx(c) {}
+
+  /// The clone context
+  CloneContext& ctx;
+  /// The destination program builder
+  ProgramBuilder& b = *ctx.dst;
+  /// The source clone context
+  const sem::Info& sem = ctx.src->Sem();
+
+  /// Builds the polyfill function for the `countLeadingZeros` builtin
+  /// @param ty the parameter and return type for the function
+  /// @return the polyfill function name
+  Symbol countLeadingZeros(const sem::Type* ty) {
+    auto name = b.Symbols().New("tint_count_leading_zeros");
+    uint32_t width = 1;
+    if (auto* v = ty->As<sem::Vector>()) {
+      width = v->Width();
+    }
+
+    // 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);
+    };
+    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(0, 16, x <= 0x0000ffff);
+            b.Decl(b.Const("b16", nullptr,
+                           b.Call("select", V(0), V(16),
+                                  b.LessThanEqual("x", V(0x0000ffff))))),
+            // x = x << b16;
+            b.Assign("x", b.Shl("x", "b16")),
+            // let b8  = select(0, 8,  x <= 0x00ffffff);
+            b.Decl(b.Const("b8", nullptr,
+                           b.Call("select", V(0), V(8),
+                                  b.LessThanEqual("x", V(0x00ffffff))))),
+            // x = x << b8;
+            b.Assign("x", b.Shl("x", "b8")),
+            // let b4  = select(0, 4,  x <= 0x0fffffff);
+            b.Decl(b.Const("b4", nullptr,
+                           b.Call("select", V(0), V(4),
+                                  b.LessThanEqual("x", V(0x0fffffff))))),
+            // x = x << b4;
+            b.Assign("x", b.Shl("x", "b4")),
+            // let b2  = select(0, 2,  x <= 0x3fffffff);
+            b.Decl(b.Const("b2", nullptr,
+                           b.Call("select", V(0), V(2),
+                                  b.LessThanEqual("x", V(0x3fffffff))))),
+            // x = x << b2;
+            b.Assign("x", b.Shl("x", "b2")),
+            // let b1  = select(0, 1,  x <= 0x7fffffff);
+            b.Decl(b.Const("b1", nullptr,
+                           b.Call("select", V(0), V(1),
+                                  b.LessThanEqual("x", V(0x7fffffff))))),
+            // let is_zero  = select(0, 1, x == 0);
+            b.Decl(b.Const("is_zero", nullptr,
+                           b.Call("select", V(0), V(1), b.Equal("x", V(0))))),
+            // return R((b16 | b8 | b4 | b2 | b1) + zero);
+            b.Return(b.Construct(
+                T(ty),
+                b.Add(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); }
+};
+
+BuiltinPolyfill::BuiltinPolyfill() = default;
+
+BuiltinPolyfill::~BuiltinPolyfill() = default;
+
+bool BuiltinPolyfill::ShouldRun(const Program* program,
+                                const DataMap& data) const {
+  if (auto* cfg = data.Get<Config>()) {
+    auto builtins = cfg->builtins;
+    auto& sem = program->Sem();
+    for (auto* node : program->ASTNodes().Objects()) {
+      if (auto* call = sem.Get<sem::Call>(node)) {
+        if (auto* builtin = call->Target()->As<sem::Builtin>()) {
+          switch (builtin->Type()) {
+            case sem::BuiltinType::kCountLeadingZeros:
+              if (builtins.count_leading_zeros) {
+                return true;
+              }
+              break;
+            default:
+              break;
+          }
+        }
+      }
+    }
+  }
+  return false;
+}
+
+void BuiltinPolyfill::Run(CloneContext& ctx,
+                          const DataMap& data,
+                          DataMap&) const {
+  auto* cfg = data.Get<Config>();
+  if (!cfg) {
+    ctx.Clone();
+    return;
+  }
+
+  std::unordered_map<const sem::Builtin*, Symbol> polyfills;
+
+  ctx.ReplaceAll(
+      [&](const ast::CallExpression* expr) -> const ast::CallExpression* {
+        auto builtins = cfg->builtins;
+        State s{ctx};
+        if (auto* call = s.sem.Get<sem::Call>(expr)) {
+          if (auto* builtin = call->Target()->As<sem::Builtin>()) {
+            Symbol polyfill;
+            switch (builtin->Type()) {
+              case sem::BuiltinType::kCountLeadingZeros:
+                if (builtins.count_leading_zeros) {
+                  polyfill = utils::GetOrCreate(polyfills, builtin, [&] {
+                    return s.countLeadingZeros(builtin->ReturnType());
+                  });
+                }
+                break;
+              default:
+                break;
+            }
+            if (polyfill.IsValid()) {
+              return s.b.Call(polyfill, ctx.Clone(call->Declaration()->args));
+            }
+          }
+        }
+        return nullptr;
+      });
+
+  ctx.Clone();
+}
+
+BuiltinPolyfill::Config::Config(const Builtins& b) : builtins(b) {}
+BuiltinPolyfill::Config::Config(const Config&) = default;
+BuiltinPolyfill::Config::~Config() = default;
+
+}  // namespace transform
+}  // namespace tint

src/tint/transform/builtin_polyfill.h

@@ -0,0 +1,77 @@
+// Copyright 2022 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.
+
+#ifndef SRC_TINT_TRANSFORM_BUILTIN_POLYFILL_H_
+#define SRC_TINT_TRANSFORM_BUILTIN_POLYFILL_H_
+
+#include "src/tint/transform/transform.h"
+
+namespace tint {
+namespace transform {
+
+/// Implements builtins for backends that do not have a native implementation.
+class BuiltinPolyfill : public Castable<BuiltinPolyfill, Transform> {
+ public:
+  /// Constructor
+  BuiltinPolyfill();
+  /// Destructor
+  ~BuiltinPolyfill() override;
+
+  /// Specifies the builtins that should be polyfilled by the transform.
+  struct Builtins {
+    /// Should `countLeadingZeros()` be polyfilled?
+    bool count_leading_zeros = false;
+  };
+
+  /// Config is consumed by the BuiltinPolyfill transform.
+  /// Config specifies the builtins that should be polyfilled.
+  struct Config : public Castable<Data, transform::Data> {
+    /// Constructor
+    /// @param b the list of builtins to polyfill
+    explicit Config(const Builtins& b);
+
+    /// Copy constructor
+    Config(const Config&);
+
+    /// Destructor
+    ~Config() override;
+
+    /// The builtins to polyfill
+    const Builtins builtins;
+  };
+
+  /// @param program the program to inspect
+  /// @param data optional extra transform-specific input data
+  /// @returns true if this transform should be run for the given program
+  bool ShouldRun(const Program* program,
+                 const DataMap& data = {}) const override;
+
+ protected:
+  struct State;
+
+  /// Runs the transform using the CloneContext built for transforming a
+  /// program. Run() is responsible for calling Clone() on the CloneContext.
+  /// @param ctx the CloneContext primed with the input program and
+  /// ProgramBuilder
+  /// @param inputs optional extra transform-specific input data
+  /// @param outputs optional extra transform-specific output data
+  void Run(CloneContext& ctx,
+           const DataMap& inputs,
+           DataMap& outputs) const override;
+};
+
+}  // namespace transform
+}  // namespace tint
+
+#endif  // SRC_TINT_TRANSFORM_BUILTIN_POLYFILL_H_

src/tint/transform/builtin_polyfill_test.cc

@@ -0,0 +1,199 @@
+// Copyright 2022 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.
+
+#include "src/tint/transform/builtin_polyfill.h"
+
+#include <utility>
+
+#include "src/tint/transform/test_helper.h"
+
+namespace tint {
+namespace transform {
+namespace {
+
+using BuiltinPolyfillTest = TransformTest;
+
+TEST_F(BuiltinPolyfillTest, ShouldRunEmptyModule) {
+  auto* src = R"()";
+
+  EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+}
+
+TEST_F(BuiltinPolyfillTest, EmptyModule) {
+  auto* src = R"()";
+
+  auto* expect = src;
+
+  auto got = Run<BuiltinPolyfill>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// countLeadingZeros
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillCountLeadingZeros() {
+  BuiltinPolyfill::Builtins builtins;
+  builtins.count_leading_zeros = true;
+  DataMap data;
+  data.Add<BuiltinPolyfill::Config>(builtins);
+  return data;
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunCountLeadingZeros) {
+  auto* src = R"(
+fn f() {
+  countLeadingZeros(0xf);
+}
+)";
+
+  EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+  EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(src, polyfillCountLeadingZeros()));
+}
+
+TEST_F(BuiltinPolyfillTest, CountLeadingZeros_i32) {
+  auto* src = R"(
+fn f() {
+  let r : i32 = countLeadingZeros(15);
+}
+)";
+
+  auto* expect = R"(
+fn tint_count_leading_zeros(v : i32) -> i32 {
+  var x = u32(v);
+  let b16 = select(0u, 16u, (x <= 65535u));
+  x = (x << b16);
+  let b8 = select(0u, 8u, (x <= 16777215u));
+  x = (x << b8);
+  let b4 = select(0u, 4u, (x <= 268435455u));
+  x = (x << b4);
+  let b2 = select(0u, 2u, (x <= 1073741823u));
+  x = (x << b2);
+  let b1 = select(0u, 1u, (x <= 2147483647u));
+  let is_zero = select(0u, 1u, (x == 0u));
+  return i32((((((b16 | b8) | b4) | b2) | b1) + is_zero));
+}
+
+fn f() {
+  let r : i32 = tint_count_leading_zeros(15);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillCountLeadingZeros());
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, CountLeadingZeros_u32) {
+  auto* src = R"(
+fn f() {
+  let r : u32 = countLeadingZeros(15u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_count_leading_zeros(v : u32) -> u32 {
+  var x = u32(v);
+  let b16 = select(0u, 16u, (x <= 65535u));
+  x = (x << b16);
+  let b8 = select(0u, 8u, (x <= 16777215u));
+  x = (x << b8);
+  let b4 = select(0u, 4u, (x <= 268435455u));
+  x = (x << b4);
+  let b2 = select(0u, 2u, (x <= 1073741823u));
+  x = (x << b2);
+  let b1 = select(0u, 1u, (x <= 2147483647u));
+  let is_zero = select(0u, 1u, (x == 0u));
+  return u32((((((b16 | b8) | b4) | b2) | b1) + is_zero));
+}
+
+fn f() {
+  let r : u32 = tint_count_leading_zeros(15u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillCountLeadingZeros());
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, CountLeadingZeros_vec3_i32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<i32> = countLeadingZeros(vec3<i32>(15));
+}
+)";
+
+  auto* expect = R"(
+fn tint_count_leading_zeros(v : vec3<i32>) -> vec3<i32> {
+  var x = vec3<u32>(v);
+  let b16 = select(vec3<u32>(0u), vec3<u32>(16u), (x <= vec3<u32>(65535u)));
+  x = (x << b16);
+  let b8 = select(vec3<u32>(0u), vec3<u32>(8u), (x <= vec3<u32>(16777215u)));
+  x = (x << b8);
+  let b4 = select(vec3<u32>(0u), vec3<u32>(4u), (x <= vec3<u32>(268435455u)));
+  x = (x << b4);
+  let b2 = select(vec3<u32>(0u), vec3<u32>(2u), (x <= vec3<u32>(1073741823u)));
+  x = (x << b2);
+  let b1 = select(vec3<u32>(0u), vec3<u32>(1u), (x <= vec3<u32>(2147483647u)));
+  let is_zero = select(vec3<u32>(0u), vec3<u32>(1u), (x == vec3<u32>(0u)));
+  return vec3<i32>((((((b16 | b8) | b4) | b2) | b1) + is_zero));
+}
+
+fn f() {
+  let r : vec3<i32> = tint_count_leading_zeros(vec3<i32>(15));
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillCountLeadingZeros());
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, CountLeadingZeros_vec3_u32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<u32> = countLeadingZeros(vec3<u32>(15u));
+}
+)";
+
+  auto* expect = R"(
+fn tint_count_leading_zeros(v : vec3<u32>) -> vec3<u32> {
+  var x = vec3<u32>(v);
+  let b16 = select(vec3<u32>(0u), vec3<u32>(16u), (x <= vec3<u32>(65535u)));
+  x = (x << b16);
+  let b8 = select(vec3<u32>(0u), vec3<u32>(8u), (x <= vec3<u32>(16777215u)));
+  x = (x << b8);
+  let b4 = select(vec3<u32>(0u), vec3<u32>(4u), (x <= vec3<u32>(268435455u)));
+  x = (x << b4);
+  let b2 = select(vec3<u32>(0u), vec3<u32>(2u), (x <= vec3<u32>(1073741823u)));
+  x = (x << b2);
+  let b1 = select(vec3<u32>(0u), vec3<u32>(1u), (x <= vec3<u32>(2147483647u)));
+  let is_zero = select(vec3<u32>(0u), vec3<u32>(1u), (x == vec3<u32>(0u)));
+  return vec3<u32>((((((b16 | b8) | b4) | b2) | b1) + is_zero));
+}
+
+fn f() {
+  let r : vec3<u32> = tint_count_leading_zeros(vec3<u32>(15u));
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillCountLeadingZeros());
+
+  EXPECT_EQ(expect, str(got));
+}
+
+}  // namespace
+}  // namespace transform
+}  // namespace tint

src/tint/transform/glsl.cc

@@ -20,6 +20,7 @@
 #include "src/tint/transform/add_empty_entry_point.h"
 #include "src/tint/transform/add_spirv_block_attribute.h"
 #include "src/tint/transform/binding_remapper.h"
+#include "src/tint/transform/builtin_polyfill.h"
 #include "src/tint/transform/canonicalize_entry_point_io.h"
 #include "src/tint/transform/combine_samplers.h"
 #include "src/tint/transform/decompose_memory_access.h"
@@ -51,6 +52,13 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const {
 
   auto* cfg = inputs.Get<Config>();
 
+  {  // Builtin polyfills
+    BuiltinPolyfill::Builtins polyfills;
+    polyfills.count_leading_zeros = true;
+    data.Add<BuiltinPolyfill::Config>(polyfills);
+    manager.Add<BuiltinPolyfill>();
+  }
+
   if (cfg && !cfg->entry_point.empty()) {
     manager.Add<SingleEntryPoint>();
     data.Add<SingleEntryPoint::Config>(cfg->entry_point);

src/tint/writer/hlsl/generator_impl.cc

@@ -48,6 +48,7 @@
 #include "src/tint/sem/variable.h"
 #include "src/tint/transform/add_empty_entry_point.h"
 #include "src/tint/transform/array_length_from_uniform.h"
+#include "src/tint/transform/builtin_polyfill.h"
 #include "src/tint/transform/calculate_array_length.h"
 #include "src/tint/transform/canonicalize_entry_point_io.h"
 #include "src/tint/transform/decompose_memory_access.h"
@@ -140,6 +141,13 @@ SanitizedResult Sanitize(
   transform::Manager manager;
   transform::DataMap data;
 
+  {  // Builtin polyfills
+    transform::BuiltinPolyfill::Builtins polyfills;
+    polyfills.count_leading_zeros = 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/msl/generator_impl.cc

@@ -1345,6 +1345,9 @@ std::string GeneratorImpl::generate_builtin_name(const sem::Builtin* builtin) {
         out += "abs";
       }
       break;
+    case sem::BuiltinType::kCountLeadingZeros:
+      out += "clz";
+      break;
     case sem::BuiltinType::kCountOneBits:
       out += "popcount";
       break;

src/tint/writer/msl/generator_impl_builtin_test.cc

@@ -124,6 +124,7 @@ const ast::CallExpression* GenerateCall(BuiltinType builtin,
       } else {
         return builder->Call(str.str(), "u2");
       }
+    case BuiltinType::kCountLeadingZeros:
     case BuiltinType::kCountOneBits:
     case BuiltinType::kReverseBits:
       return builder->Call(str.str(), "u2");
@@ -212,6 +213,7 @@ INSTANTIATE_TEST_SUITE_P(
         BuiltinData{BuiltinType::kClamp, ParamType::kU32, "clamp"},
         BuiltinData{BuiltinType::kCos, ParamType::kF32, "cos"},
         BuiltinData{BuiltinType::kCosh, ParamType::kF32, "cosh"},
+        BuiltinData{BuiltinType::kCountLeadingZeros, ParamType::kU32, "clz"},
         BuiltinData{BuiltinType::kCountOneBits, ParamType::kU32, "popcount"},
         BuiltinData{BuiltinType::kCross, ParamType::kF32, "cross"},
         BuiltinData{BuiltinType::kDeterminant, ParamType::kF32, "determinant"},

src/tint/writer/spirv/builder.cc

@@ -43,6 +43,7 @@
 #include "src/tint/sem/vector_type.h"
 #include "src/tint/transform/add_empty_entry_point.h"
 #include "src/tint/transform/add_spirv_block_attribute.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/fold_constants.h"
@@ -258,6 +259,13 @@ SanitizedResult Sanitize(const Program* in,
   transform::Manager manager;
   transform::DataMap data;
 
+  {  // Builtin polyfills
+    transform::BuiltinPolyfill::Builtins polyfills;
+    polyfills.count_leading_zeros = true;
+    data.Add<transform::BuiltinPolyfill::Config>(polyfills);
+    manager.Add<transform::BuiltinPolyfill>();
+  }
+
   manager.Add<transform::Unshadow>();
   if (!disable_workgroup_init) {
     manager.Add<transform::ZeroInitWorkgroupMemory>();

test/builtins/gen/countLeadingZeros/208d46.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(u32) -> u32
+fn countLeadingZeros_208d46() {
+  var res: u32 = countLeadingZeros(1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_208d46();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_208d46();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_208d46();
+}

test/builtins/gen/countLeadingZeros/6d4656.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(i32) -> i32
+fn countLeadingZeros_6d4656() {
+  var res: i32 = countLeadingZeros(1);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_6d4656();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_6d4656();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_6d4656();
+}

test/builtins/gen/countLeadingZeros/70783f.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(vec<2, u32>) -> vec<2, u32>
+fn countLeadingZeros_70783f() {
+  var res: vec2<u32> = countLeadingZeros(vec2<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_70783f();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_70783f();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_70783f();
+}

test/builtins/gen/countLeadingZeros/7c38a6.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(vec<3, i32>) -> vec<3, i32>
+fn countLeadingZeros_7c38a6() {
+  var res: vec3<i32> = countLeadingZeros(vec3<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_7c38a6();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_7c38a6();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_7c38a6();
+}

test/builtins/gen/countLeadingZeros/858d40.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(vec<2, i32>) -> vec<2, i32>
+fn countLeadingZeros_858d40() {
+  var res: vec2<i32> = countLeadingZeros(vec2<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_858d40();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_858d40();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_858d40();
+}

test/builtins/gen/countLeadingZeros/ab6345.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(vec<3, u32>) -> vec<3, u32>
+fn countLeadingZeros_ab6345() {
+  var res: vec3<u32> = countLeadingZeros(vec3<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_ab6345();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_ab6345();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_ab6345();
+}

test/builtins/gen/countLeadingZeros/eab32b.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(vec<4, i32>) -> vec<4, i32>
+fn countLeadingZeros_eab32b() {
+  var res: vec4<i32> = countLeadingZeros(vec4<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_eab32b();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_eab32b();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_eab32b();
+}

test/builtins/gen/countLeadingZeros/f70103.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/builtins/builtins.wgsl.tmpl
+// and the builtin defintion file:
+//   src/builtins.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn countLeadingZeros(vec<4, u32>) -> vec<4, u32>
+fn countLeadingZeros_f70103() {
+  var res: vec4<u32> = countLeadingZeros(vec4<u32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_f70103();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_f70103();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_f70103();
+}

test/tint/BUILD.gn

@@ -310,6 +310,7 @@ tint_unittests_source_set("tint_unittests_transform_src") {
     "../../src/tint/transform/add_empty_entry_point_test.cc",
     "../../src/tint/transform/add_spirv_block_attribute_test.cc",
     "../../src/tint/transform/array_length_from_uniform_test.cc",
+    "../../src/tint/transform/builtin_polyfill_test.cc",
     "../../src/tint/transform/binding_remapper_test.cc",
     "../../src/tint/transform/calculate_array_length_test.cc",
     "../../src/tint/transform/canonicalize_entry_point_io_test.cc",

test/tint/builtins/gen/countLeadingZeros/208d46.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 countLeadingZeros(u32) -> u32
+fn countLeadingZeros_208d46() {
+  var res: u32 = countLeadingZeros(1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  countLeadingZeros_208d46();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  countLeadingZeros_208d46();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  countLeadingZeros_208d46();
+}

test/tint/builtins/gen/countLeadingZeros/208d46.wgsl.expected.glsl

@@ -0,0 +1,93 @@
+#version 310 es
+
+uint tint_count_leading_zeros(uint v) {
+  uint x = uint(v);
+  uint b16 = ((x <= 65535u) ? 16u : 0u);
+  x = (x << b16);
+  uint b8 = ((x <= 16777215u) ? 8u : 0u);
+  x = (x << b8);
+  uint b4 = ((x <= 268435455u) ? 4u : 0u);
+  x = (x << b4);
+  uint b2 = ((x <= 1073741823u) ? 2u : 0u);
+  x = (x << b2);
+  uint b1 = ((x <= 2147483647u) ? 1u : 0u);
+  uint is_zero = ((x == 0u) ? 1u : 0u);
+  return uint((((((b16 | b8) | b4) | b2) | b1) + is_zero));
+}
+
+void countLeadingZeros_208d46() {
+  uint res = tint_count_leading_zeros(1u);
+}
+
+vec4 vertex_main() {
+  countLeadingZeros_208d46();
+  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_count_leading_zeros(uint v) {
+  uint x = uint(v);
+  uint b16 = ((x <= 65535u) ? 16u : 0u);
+  x = (x << b16);
+  uint b8 = ((x <= 16777215u) ? 8u : 0u);
+  x = (x << b8);
+  uint b4 = ((x <= 268435455u) ? 4u : 0u);
+  x = (x << b4);
+  uint b2 = ((x <= 1073741823u) ? 2u : 0u);
+  x = (x << b2);
+  uint b1 = ((x <= 2147483647u) ? 1u : 0u);
+  uint is_zero = ((x == 0u) ? 1u : 0u);
+  return uint((((((b16 | b8) | b4) | b2) | b1) + is_zero));
+}
+
+void countLeadingZeros_208d46() {
+  uint res = tint_count_leading_zeros(1u);
+}
+
+void fragment_main() {
+  countLeadingZeros_208d46();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uint tint_count_leading_zeros(uint v) {
+  uint x = uint(v);
+  uint b16 = ((x <= 65535u) ? 16u : 0u);
+  x = (x << b16);
+  uint b8 = ((x <= 16777215u) ? 8u : 0u);
+  x = (x << b8);
+  uint b4 = ((x <= 268435455u) ? 4u : 0u);
+  x = (x << b4);
+  uint b2 = ((x <= 1073741823u) ? 2u : 0u);
+  x = (x << b2);
+  uint b1 = ((x <= 2147483647u) ? 1u : 0u);
+  uint is_zero = ((x == 0u) ? 1u : 0u);
+  return uint((((((b16 | b8) | b4) | b2) | b1) + is_zero));
+}
+
+void countLeadingZeros_208d46() {
+  uint res = tint_count_leading_zeros(1u);
+}
+
+void compute_main() {
+  countLeadingZeros_208d46();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}