builtins: Add insertBits

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

Bug: tint:1371
Change-Id: Idd55c0bc9dad1dffb558d0bc57d744f65e9041b5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/81701
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>

src/tint/builtins.def

@@ -336,6 +336,8 @@ fn frexp<N: num>(vec<N, f32>) -> __frexp_result_vec<N>
 [[stage("fragment")]] fn fwidthCoarse<N: num>(vec<N, f32>) -> vec<N, f32>
 [[stage("fragment")]] fn fwidthFine(f32) -> f32
 [[stage("fragment")]] fn fwidthFine<N: num>(vec<N, f32>) -> vec<N, f32>
+fn insertBits<T: iu32>(T, T, u32, u32) -> T
+fn insertBits<N: num, T: iu32>(vec<N, T>, vec<N, T>, u32, u32) -> vec<N, T>
 fn inverseSqrt(f32) -> f32
 fn inverseSqrt<N: num>(vec<N, f32>) -> vec<N, f32>
 [[deprecated]] fn isFinite(f32) -> bool

src/tint/program_builder.h

@@ -1760,6 +1760,16 @@ class ProgramBuilder {
                                          Expr(std::forward<RHS>(rhs)));
   }
 
+  /// @param lhs the left hand argument to the xor operation
+  /// @param rhs the right hand argument to the xor operation
+  /// @returns a `ast::BinaryExpression` bitwise xor-ing `lhs` and `rhs`
+  template <typename LHS, typename RHS>
+  const ast::BinaryExpression* Xor(LHS&& lhs, RHS&& rhs) {
+    return create<ast::BinaryExpression>(ast::BinaryOp::kXor,
+                                         Expr(std::forward<LHS>(lhs)),
+                                         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`

src/tint/reader/spirv/function.cc

@@ -443,6 +443,11 @@ sem::BuiltinType GetBuiltin(SpvOp opcode) {
   switch (opcode) {
     case SpvOpBitCount:
       return sem::BuiltinType::kCountOneBits;
+    case SpvOpBitFieldInsert:
+      return sem::BuiltinType::kInsertBits;
+    case SpvOpBitFieldSExtract:
+    case SpvOpBitFieldUExtract:
+      return sem::BuiltinType::kExtractBits;
     case SpvOpBitReverse:
       return sem::BuiltinType::kReverseBits;
     case SpvOpDot:
@@ -463,9 +468,6 @@ sem::BuiltinType GetBuiltin(SpvOp opcode) {
       return sem::BuiltinType::kDpdxCoarse;
     case SpvOpDPdyCoarse:
       return sem::BuiltinType::kDpdyCoarse;
-    case SpvOpBitFieldSExtract:
-    case SpvOpBitFieldUExtract:
-      return sem::BuiltinType::kExtractBits;
     case SpvOpFwidthCoarse:
       return sem::BuiltinType::kFwidthCoarse;
     default:

src/tint/reader/spirv/function_bit_test.cc

@@ -901,6 +901,78 @@ TEST_F(SpvUnaryBitTest, BitReverse_IntVector_IntVector) {
       << body;
 }
 
+TEST_F(SpvUnaryBitTest, InsertBits_Int) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldInsert %v2int %int_30 %int_40 %uint_10 %uint_20
+     OpReturn
+     OpFunctionEnd
+  )";
+  auto p = parser(test::Assemble(assembly));
+  ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+  auto fe = p->function_emitter(100);
+  EXPECT_TRUE(fe.EmitBody()) << p->error();
+  auto ast_body = fe.ast_body();
+  auto body = test::ToString(p->program(), ast_body);
+  EXPECT_THAT(body,
+              HasSubstr("let x_1 : vec2<i32> = insertBits(30, 40, 10u, 20u);"))
+      << body;
+}
+
+TEST_F(SpvUnaryBitTest, InsertBits_IntVector) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldInsert %v2int %v2int_30_40 %v2int_40_30 %uint_10 %uint_20
+     OpReturn
+     OpFunctionEnd
+  )";
+  auto p = parser(test::Assemble(assembly));
+  ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+  auto fe = p->function_emitter(100);
+  EXPECT_TRUE(fe.EmitBody()) << p->error();
+  auto ast_body = fe.ast_body();
+  auto body = test::ToString(p->program(), ast_body);
+  EXPECT_THAT(
+      body,
+      HasSubstr(
+          R"(let x_1 : vec2<i32> = insertBits(vec2<i32>(30, 40), vec2<i32>(40, 30), 10u, 20u);)"))
+      << body;
+}
+
+TEST_F(SpvUnaryBitTest, InsertBits_Uint) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldInsert %v2uint %uint_20 %uint_10 %uint_10 %uint_20
+     OpReturn
+     OpFunctionEnd
+  )";
+  auto p = parser(test::Assemble(assembly));
+  ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+  auto fe = p->function_emitter(100);
+  EXPECT_TRUE(fe.EmitBody()) << p->error();
+  auto ast_body = fe.ast_body();
+  auto body = test::ToString(p->program(), ast_body);
+  EXPECT_THAT(
+      body, HasSubstr("let x_1 : vec2<u32> = insertBits(20u, 10u, 10u, 20u);"))
+      << body;
+}
+
+TEST_F(SpvUnaryBitTest, InsertBits_UintVector) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldInsert %v2uint %v2uint_10_20 %v2uint_20_10 %uint_10 %uint_20
+     OpReturn
+     OpFunctionEnd
+  )";
+  auto p = parser(test::Assemble(assembly));
+  ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+  auto fe = p->function_emitter(100);
+  EXPECT_TRUE(fe.EmitBody()) << p->error();
+  auto ast_body = fe.ast_body();
+  auto body = test::ToString(p->program(), ast_body);
+  EXPECT_THAT(
+      body,
+      HasSubstr(
+          R"(let x_1 : vec2<u32> = insertBits(vec2<u32>(10u, 20u), vec2<u32>(20u, 10u), 10u, 20u);)"))
+      << body;
+}
+
 TEST_F(SpvUnaryBitTest, ExtractBits_Int) {
   const auto assembly = BitTestPreamble() + R"(
      %1 = OpBitFieldSExtract %v2int %int_30 %uint_10 %uint_20
@@ -973,8 +1045,6 @@ TEST_F(SpvUnaryBitTest, ExtractBits_UintVector) {
       << body;
 }
 
-// TODO(dneto): OpBitFieldInsert
-
 }  // namespace
 }  // namespace spirv
 }  // namespace reader

src/tint/sem/builtin_type.cc

@@ -147,6 +147,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
   if (name == "fwidthFine") {
     return BuiltinType::kFwidthFine;
   }
+  if (name == "insertBits") {
+    return BuiltinType::kInsertBits;
+  }
   if (name == "inverseSqrt") {
     return BuiltinType::kInverseSqrt;
   }
@@ -436,6 +439,8 @@ const char* str(BuiltinType i) {
       return "fwidthCoarse";
     case BuiltinType::kFwidthFine:
       return "fwidthFine";
+    case BuiltinType::kInsertBits:
+      return "insertBits";
     case BuiltinType::kInverseSqrt:
       return "inverseSqrt";
     case BuiltinType::kIsFinite:

src/tint/sem/builtin_type.h

@@ -73,6 +73,7 @@ enum class BuiltinType {
   kFwidth,
   kFwidthCoarse,
   kFwidthFine,
+  kInsertBits,
   kInverseSqrt,
   kIsFinite,
   kIsInf,

src/tint/transform/builtin_polyfill.cc

@@ -55,7 +55,7 @@ struct BuiltinPolyfill::State {
       if (width == 1) {
         return b.ty.u32();
       }
-      return b.ty.vec<ProgramBuilder::u32>(width);
+      return b.ty.vec<u32>(width);
     };
     auto V = [&](uint32_t value) -> const ast::Expression* {
       return ScalarOrVector(width, value);
@@ -117,7 +117,7 @@ struct BuiltinPolyfill::State {
       if (width == 1) {
         return b.ty.u32();
       }
-      return b.ty.vec<ProgramBuilder::u32>(width);
+      return b.ty.vec<u32>(width);
     };
     auto V = [&](uint32_t value) -> const ast::Expression* {
       return ScalarOrVector(width, value);
@@ -187,7 +187,7 @@ struct BuiltinPolyfill::State {
       if (width == 1) {
         return value;
       }
-      return b.Construct(b.ty.vec<ProgramBuilder::u32>(width), value);
+      return b.Construct(b.ty.vec<u32>(width), value);
     };
 
     ast::StatementList body = {
@@ -236,7 +236,7 @@ struct BuiltinPolyfill::State {
       if (width == 1) {
         return b.ty.u32();
       }
-      return b.ty.vec<ProgramBuilder::u32>(width);
+      return b.ty.vec<u32>(width);
     };
     auto V = [&](uint32_t value) -> const ast::Expression* {
       return ScalarOrVector(width, value);
@@ -317,7 +317,7 @@ struct BuiltinPolyfill::State {
       if (width == 1) {
         return b.ty.u32();
       }
-      return b.ty.vec<ProgramBuilder::u32>(width);
+      return b.ty.vec<u32>(width);
     };
     auto V = [&](uint32_t value) -> const ast::Expression* {
       return ScalarOrVector(width, value);
@@ -373,16 +373,90 @@ struct BuiltinPolyfill::State {
     return name;
   }
 
+  /// Builds the polyfill function for the `insertBits` builtin
+  /// @param ty the parameter and return type for the function
+  /// @return the polyfill function name
+  Symbol insertBits(const sem::Type* ty) {
+    auto name = b.Symbols().New("tint_insert_bits");
+    uint32_t width = WidthOf(ty);
+
+    constexpr uint32_t W = 32u;  // 32-bit
+
+    auto V = [&](auto value) -> const ast::Expression* {
+      const ast::Expression* expr = b.Expr(value);
+      if (!ty->is_unsigned_scalar_or_vector()) {
+        expr = b.Construct<i32>(expr);
+      }
+      if (ty->Is<sem::Vector>()) {
+        expr = b.Construct(T(ty), expr);
+      }
+      return expr;
+    };
+    auto U = [&](auto value) -> const ast::Expression* {
+      if (width == 1) {
+        return b.Expr(value);
+      }
+      return b.vec(b.ty.u32(), width, value);
+    };
+
+    ast::StatementList body = {
+        b.Decl(b.Const("s", nullptr, b.Call("min", "offset", W))),
+        b.Decl(b.Const("e", nullptr, b.Call("min", W, b.Add("s", "count")))),
+    };
+
+    switch (polyfill.insert_bits) {
+      case Level::kFull:
+        // let mask = ((1 << s) - 1) ^ ((1 << e) - 1)
+        body.emplace_back(b.Decl(b.Const(
+            "mask", nullptr,
+            b.Xor(b.Sub(b.Shl(1u, "s"), 1u), b.Sub(b.Shl(1u, "e"), 1u)))));
+        // return ((n << s) & mask) | (v & ~mask)
+        body.emplace_back(b.Return(b.Or(b.And(b.Shl("n", U("s")), V("mask")),
+                                        b.And("v", V(b.Complement("mask"))))));
+        break;
+      case Level::kClampParameters:
+        body.emplace_back(
+            b.Return(b.Call("insertBits", "v", "n", "s", b.Sub("e", "s"))));
+        break;
+      default:
+        TINT_ICE(Transform, b.Diagnostics())
+            << "unhandled polyfill level: "
+            << static_cast<int>(polyfill.insert_bits);
+        return {};
+    }
+
+    b.Func(name,
+           {
+               b.Param("v", T(ty)),
+               b.Param("n", T(ty)),
+               b.Param("offset", b.ty.u32()),
+               b.Param("count", b.ty.u32()),
+           },
+           T(ty), body);
+
+    return name;
+  }
+
  private:
+  /// Aliases
+  using u32 = ProgramBuilder::u32;
+  using i32 = ProgramBuilder::i32;
+
+  /// @returns the AST type for the given sem type
   const ast::Type* T(const sem::Type* ty) const {
     return CreateASTTypeFor(ctx, ty);
   }
+
+  /// @returns 1 if `ty` is not a vector, otherwise the vector width
   uint32_t WidthOf(const sem::Type* ty) const {
     if (auto* v = ty->As<sem::Vector>()) {
       return v->Width();
     }
     return 1;
   }
+
+  /// @returns a scalar or vector with the given width, with each element with
+  /// the given value.
   template <typename T>
   const ast::Expression* ScalarOrVector(uint32_t width, T value) const {
     if (width == 1) {
@@ -430,6 +504,11 @@ bool BuiltinPolyfill::ShouldRun(const Program* program,
                 return true;
               }
               break;
+            case sem::BuiltinType::kInsertBits:
+              if (builtins.insert_bits != Level::kNone) {
+                return true;
+              }
+              break;
             default:
               break;
           }
@@ -494,6 +573,13 @@ void BuiltinPolyfill::Run(CloneContext& ctx,
                   });
                 }
                 break;
+              case sem::BuiltinType::kInsertBits:
+                if (builtins.insert_bits != Level::kNone) {
+                  polyfill = utils::GetOrCreate(polyfills, builtin, [&] {
+                    return s.insertBits(builtin->ReturnType());
+                  });
+                }
+                break;
               default:
                 break;
             }

src/tint/transform/builtin_polyfill.h

@@ -50,6 +50,8 @@ class BuiltinPolyfill : public Castable<BuiltinPolyfill, Transform> {
     bool first_leading_bit = false;
     /// Should `firstTrailingBit()` be polyfilled?
     bool first_trailing_bit = false;
+    /// Should `insertBits()` be polyfilled?
+    Level insert_bits = Level::kNone;
   };
 
   /// Config is consumed by the BuiltinPolyfill transform.

src/tint/transform/builtin_polyfill_test.cc

@@ -889,6 +889,233 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// insertBits
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillInsertBits(Level level) {
+  BuiltinPolyfill::Builtins builtins;
+  builtins.insert_bits = level;
+  DataMap data;
+  data.Add<BuiltinPolyfill::Config>(builtins);
+  return data;
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunInsertBits) {
+  auto* src = R"(
+fn f() {
+  insertBits(1234, 5678, 5u, 6u);
+}
+)";
+
+  EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+  EXPECT_FALSE(
+      ShouldRun<BuiltinPolyfill>(src, polyfillInsertBits(Level::kNone)));
+  EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(
+      src, polyfillInsertBits(Level::kClampParameters)));
+  EXPECT_TRUE(
+      ShouldRun<BuiltinPolyfill>(src, polyfillInsertBits(Level::kFull)));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Full_i32) {
+  auto* src = R"(
+fn f() {
+  let r : i32 = insertBits(1234, 5678, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : i32, n : i32, offset : u32, count : u32) -> i32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << s) & i32(mask)) | (v & i32(~(mask))));
+}
+
+fn f() {
+  let r : i32 = tint_insert_bits(1234, 5678, 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Full_u32) {
+  auto* src = R"(
+fn f() {
+  let r : u32 = insertBits(1234u, 5678u, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : u32, n : u32, offset : u32, count : u32) -> u32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << s) & mask) | (v & ~(mask)));
+}
+
+fn f() {
+  let r : u32 = tint_insert_bits(1234u, 5678u, 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Full_vec3_i32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<i32> = insertBits(vec3<i32>(1234), vec3<i32>(5678), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : vec3<i32>, n : vec3<i32>, offset : u32, count : u32) -> vec3<i32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << vec3<u32>(s)) & vec3<i32>(i32(mask))) | (v & vec3<i32>(i32(~(mask)))));
+}
+
+fn f() {
+  let r : vec3<i32> = tint_insert_bits(vec3<i32>(1234), vec3<i32>(5678), 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Full_vec3_u32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<u32> = insertBits(vec3<u32>(1234u), vec3<u32>(5678u), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : vec3<u32>, n : vec3<u32>, offset : u32, count : u32) -> vec3<u32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << vec3<u32>(s)) & vec3<u32>(mask)) | (v & vec3<u32>(~(mask))));
+}
+
+fn f() {
+  let r : vec3<u32> = tint_insert_bits(vec3<u32>(1234u), vec3<u32>(5678u), 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Clamp_i32) {
+  auto* src = R"(
+fn f() {
+  let r : i32 = insertBits(1234, 5678, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : i32, n : i32, offset : u32, count : u32) -> i32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return insertBits(v, n, s, (e - s));
+}
+
+fn f() {
+  let r : i32 = tint_insert_bits(1234, 5678, 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Clamp_u32) {
+  auto* src = R"(
+fn f() {
+  let r : u32 = insertBits(1234u, 5678u, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : u32, n : u32, offset : u32, count : u32) -> u32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return insertBits(v, n, s, (e - s));
+}
+
+fn f() {
+  let r : u32 = tint_insert_bits(1234u, 5678u, 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Clamp_vec3_i32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<i32> = insertBits(vec3<i32>(1234), vec3<i32>(5678), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : vec3<i32>, n : vec3<i32>, offset : u32, count : u32) -> vec3<i32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return insertBits(v, n, s, (e - s));
+}
+
+fn f() {
+  let r : vec3<i32> = tint_insert_bits(vec3<i32>(1234), vec3<i32>(5678), 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, InsertBits_Clamp_vec3_u32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<u32> = insertBits(vec3<u32>(1234u), vec3<u32>(5678u), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_insert_bits(v : vec3<u32>, n : vec3<u32>, offset : u32, count : u32) -> vec3<u32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return insertBits(v, n, s, (e - s));
+}
+
+fn f() {
+  let r : vec3<u32> = tint_insert_bits(vec3<u32>(1234u), vec3<u32>(5678u), 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillInsertBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/tint/transform/glsl.cc

@@ -59,6 +59,7 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const {
     polyfills.extract_bits = BuiltinPolyfill::Level::kClampParameters;
     polyfills.first_leading_bit = true;
     polyfills.first_trailing_bit = true;
+    polyfills.insert_bits = BuiltinPolyfill::Level::kClampParameters;
     data.Add<BuiltinPolyfill::Config>(polyfills);
     manager.Add<BuiltinPolyfill>();
   }

src/tint/writer/glsl/generator_impl.cc

@@ -607,6 +607,9 @@ bool GeneratorImpl::EmitBuiltinCall(std::ostream& out,
   if (builtin->Type() == sem::BuiltinType::kExtractBits) {
     return EmitExtractBits(out, expr);
   }
+  if (builtin->Type() == sem::BuiltinType::kInsertBits) {
+    return EmitInsertBits(out, expr);
+  }
   if (builtin->IsDataPacking()) {
     return EmitDataPackingCall(out, expr, builtin);
   }
@@ -834,6 +837,28 @@ bool GeneratorImpl::EmitExtractBits(std::ostream& out,
   return true;
 }
 
+bool GeneratorImpl::EmitInsertBits(std::ostream& out,
+                                   const ast::CallExpression* expr) {
+  out << "bitfieldInsert(";
+  if (!EmitExpression(out, expr->args[0])) {
+    return false;
+  }
+  out << ", ";
+  if (!EmitExpression(out, expr->args[1])) {
+    return false;
+  }
+  out << ", int(";
+  if (!EmitExpression(out, expr->args[2])) {
+    return false;
+  }
+  out << "), int(";
+  if (!EmitExpression(out, expr->args[3])) {
+    return false;
+  }
+  out << "))";
+  return true;
+}
+
 bool GeneratorImpl::EmitSelectCall(std::ostream& out,
                                    const ast::CallExpression* expr) {
   auto* expr_false = expr->args[0];

src/tint/writer/glsl/generator_impl.h

@@ -176,6 +176,11 @@ class GeneratorImpl : public TextGenerator {
   /// @param expr the call expression
   /// @returns true if the expression is emitted
   bool EmitExtractBits(std::ostream& out, const ast::CallExpression* expr);
+  /// Handles generating a call to `bitfieldInsert`
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @returns true if the expression is emitted
+  bool EmitInsertBits(std::ostream& out, const ast::CallExpression* expr);
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param out the output of the expression stream

src/tint/writer/glsl/generator_impl_builtin_test.cc

@@ -615,6 +615,42 @@ void main() {
 }
 )");
 }
+
+TEST_F(GlslGeneratorImplTest_Builtin, InsertBits) {
+  auto* v = Var("v", ty.vec3<u32>());
+  auto* n = Var("n", ty.vec3<u32>());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("insertBits", v, n, offset, count);
+  WrapInFunction(v, n, offset, count, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#version 310 es
+
+uvec3 tint_insert_bits(uvec3 v, uvec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void test_function() {
+  uvec3 v = uvec3(0u, 0u, 0u);
+  uvec3 n = uvec3(0u, 0u, 0u);
+  uint offset = 0u;
+  uint count = 0u;
+  uvec3 tint_symbol = tint_insert_bits(v, n, offset, count);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  test_function();
+  return;
+}
+)");
+}
+
 #if 0
 TEST_F(GlslGeneratorImplTest_Builtin, Pack4x8Snorm) {
   auto* call = Call("pack4x8snorm", "p1");

src/tint/writer/hlsl/generator_impl.cc

@@ -150,6 +150,7 @@ SanitizedResult Sanitize(
     polyfills.extract_bits = transform::BuiltinPolyfill::Level::kFull;
     polyfills.first_leading_bit = true;
     polyfills.first_trailing_bit = true;
+    polyfills.insert_bits = transform::BuiltinPolyfill::Level::kFull;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);
     manager.Add<transform::BuiltinPolyfill>();
   }

src/tint/writer/msl/generator_impl.cc

@@ -134,6 +134,7 @@ SanitizedResult Sanitize(
         transform::BuiltinPolyfill::Level::kClampParameters;
     polyfills.first_leading_bit = true;
     polyfills.first_trailing_bit = true;
+    polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);
     manager.Add<transform::BuiltinPolyfill>();
   }
@@ -1378,6 +1379,9 @@ std::string GeneratorImpl::generate_builtin_name(const sem::Builtin* builtin) {
     case sem::BuiltinType::kExtractBits:
       out += "extract_bits";
       break;
+    case sem::BuiltinType::kInsertBits:
+      out += "insert_bits";
+      break;
     case sem::BuiltinType::kFwidth:
     case sem::BuiltinType::kFwidthCoarse:
     case sem::BuiltinType::kFwidthFine:

src/tint/writer/msl/generator_impl_builtin_test.cc

@@ -131,6 +131,8 @@ const ast::CallExpression* GenerateCall(BuiltinType builtin,
       return builder->Call(str.str(), "u2");
     case BuiltinType::kExtractBits:
       return builder->Call(str.str(), "u2", "u1", "u1");
+    case BuiltinType::kInsertBits:
+      return builder->Call(str.str(), "u2", "u2", "u1", "u1");
     case BuiltinType::kMax:
     case BuiltinType::kMin:
       if (type == ParamType::kF32) {
@@ -239,6 +241,7 @@ INSTANTIATE_TEST_SUITE_P(
         BuiltinData{BuiltinType::kFwidth, ParamType::kF32, "fwidth"},
         BuiltinData{BuiltinType::kFwidthCoarse, ParamType::kF32, "fwidth"},
         BuiltinData{BuiltinType::kFwidthFine, ParamType::kF32, "fwidth"},
+        BuiltinData{BuiltinType::kInsertBits, ParamType::kU32, "insert_bits"},
         BuiltinData{BuiltinType::kInverseSqrt, ParamType::kF32, "rsqrt"},
         BuiltinData{BuiltinType::kIsFinite, ParamType::kF32, "isfinite"},
         BuiltinData{BuiltinType::kIsInf, ParamType::kF32, "isinf"},

src/tint/writer/spirv/builder.cc

@@ -267,6 +267,7 @@ SanitizedResult Sanitize(const Program* in,
         transform::BuiltinPolyfill::Level::kClampParameters;
     polyfills.first_leading_bit = true;
     polyfills.first_trailing_bit = true;
+    polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);
     manager.Add<transform::BuiltinPolyfill>();
   }
@@ -2520,6 +2521,9 @@ uint32_t Builder::GenerateBuiltinCall(const sem::Call* call,
     case BuiltinType::kFwidthFine:
       op = spv::Op::OpFwidthFine;
       break;
+    case BuiltinType::kInsertBits:
+      op = spv::Op::OpBitFieldInsert;
+      break;
     case BuiltinType::kIsInf:
       op = spv::Op::OpIsInf;
       break;

src/tint/writer/spirv/builder_builtin_test.cc

@@ -2708,6 +2708,176 @@ OpFunctionEnd
 )");
 }
 
+TEST_F(BuiltinBuilderTest, Call_InsertBits_i32) {
+  auto* v = Var("v", ty.i32());
+  auto* n = Var("n", ty.i32());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("insertBits", v, n, offset, count);
+  auto* func = WrapInFunction(v, n, offset, count, call);
+
+  spirv::Builder& b = Build();
+
+  ASSERT_TRUE(b.GenerateFunction(func)) << b.error();
+
+  EXPECT_EQ(DumpBuilder(b), R"(OpEntryPoint GLCompute %3 "test_function"
+OpExecutionMode %3 LocalSize 1 1 1
+OpName %3 "test_function"
+OpName %5 "v"
+OpName %9 "n"
+OpName %10 "offset"
+OpName %14 "count"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%7 = OpTypeInt 32 1
+%6 = OpTypePointer Function %7
+%8 = OpConstantNull %7
+%12 = OpTypeInt 32 0
+%11 = OpTypePointer Function %12
+%13 = OpConstantNull %12
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpVariable %6 Function %8
+%9 = OpVariable %6 Function %8
+%10 = OpVariable %11 Function %13
+%14 = OpVariable %11 Function %13
+%16 = OpLoad %7 %5
+%17 = OpLoad %7 %9
+%18 = OpLoad %12 %10
+%19 = OpLoad %12 %14
+%15 = OpBitFieldInsert %7 %16 %17 %18 %19
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(BuiltinBuilderTest, Call_InsertBits_u32) {
+  auto* v = Var("v", ty.u32());
+  auto* n = Var("n", ty.u32());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("insertBits", v, n, offset, count);
+  auto* func = WrapInFunction(v, n, offset, count, call);
+
+  spirv::Builder& b = Build();
+
+  ASSERT_TRUE(b.GenerateFunction(func)) << b.error();
+
+  EXPECT_EQ(DumpBuilder(b), R"(OpEntryPoint GLCompute %3 "test_function"
+OpExecutionMode %3 LocalSize 1 1 1
+OpName %3 "test_function"
+OpName %5 "v"
+OpName %9 "n"
+OpName %10 "offset"
+OpName %11 "count"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%7 = OpTypeInt 32 0
+%6 = OpTypePointer Function %7
+%8 = OpConstantNull %7
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpVariable %6 Function %8
+%9 = OpVariable %6 Function %8
+%10 = OpVariable %6 Function %8
+%11 = OpVariable %6 Function %8
+%13 = OpLoad %7 %5
+%14 = OpLoad %7 %9
+%15 = OpLoad %7 %10
+%16 = OpLoad %7 %11
+%12 = OpBitFieldInsert %7 %13 %14 %15 %16
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(BuiltinBuilderTest, Call_InsertBits_vec3_i32) {
+  auto* v = Var("v", ty.vec3<i32>());
+  auto* n = Var("n", ty.vec3<i32>());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("insertBits", v, n, offset, count);
+  auto* func = WrapInFunction(v, n, offset, count, call);
+
+  spirv::Builder& b = Build();
+
+  ASSERT_TRUE(b.GenerateFunction(func)) << b.error();
+
+  EXPECT_EQ(DumpBuilder(b), R"(OpEntryPoint GLCompute %3 "test_function"
+OpExecutionMode %3 LocalSize 1 1 1
+OpName %3 "test_function"
+OpName %5 "v"
+OpName %10 "n"
+OpName %11 "offset"
+OpName %15 "count"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeInt 32 1
+%7 = OpTypeVector %8 3
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%13 = OpTypeInt 32 0
+%12 = OpTypePointer Function %13
+%14 = OpConstantNull %13
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpVariable %6 Function %9
+%10 = OpVariable %6 Function %9
+%11 = OpVariable %12 Function %14
+%15 = OpVariable %12 Function %14
+%17 = OpLoad %7 %5
+%18 = OpLoad %7 %10
+%19 = OpLoad %13 %11
+%20 = OpLoad %13 %15
+%16 = OpBitFieldInsert %7 %17 %18 %19 %20
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(BuiltinBuilderTest, Call_InsertBits_vec3_u32) {
+  auto* v = Var("v", ty.vec3<u32>());
+  auto* n = Var("n", ty.vec3<u32>());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("insertBits", v, n, offset, count);
+  auto* func = WrapInFunction(v, n, offset, count, call);
+
+  spirv::Builder& b = Build();
+
+  ASSERT_TRUE(b.GenerateFunction(func)) << b.error();
+
+  EXPECT_EQ(DumpBuilder(b), R"(OpEntryPoint GLCompute %3 "test_function"
+OpExecutionMode %3 LocalSize 1 1 1
+OpName %3 "test_function"
+OpName %5 "v"
+OpName %10 "n"
+OpName %11 "offset"
+OpName %14 "count"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeInt 32 0
+%7 = OpTypeVector %8 3
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%12 = OpTypePointer Function %8
+%13 = OpConstantNull %8
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpVariable %6 Function %9
+%10 = OpVariable %6 Function %9
+%11 = OpVariable %12 Function %13
+%14 = OpVariable %12 Function %13
+%16 = OpLoad %7 %5
+%17 = OpLoad %7 %10
+%18 = OpLoad %8 %11
+%19 = OpLoad %8 %14
+%15 = OpBitFieldInsert %7 %16 %17 %18 %19
+OpReturn
+OpFunctionEnd
+)");
+}
+
 }  // namespace
 }  // namespace spirv
 }  // namespace writer

test/tint/builtins/gen/insertBits/3c7ba5.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 insertBits(vec<2, u32>, vec<2, u32>, u32, u32) -> vec<2, u32>
+fn insertBits_3c7ba5() {
+  var res: vec2<u32> = insertBits(vec2<u32>(), vec2<u32>(), 1u, 1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  insertBits_3c7ba5();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  insertBits_3c7ba5();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  insertBits_3c7ba5();
+}

test/tint/builtins/gen/insertBits/3c7ba5.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+uvec2 tint_insert_bits(uvec2 v, uvec2 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_3c7ba5() {
+  uvec2 res = tint_insert_bits(uvec2(0u, 0u), uvec2(0u, 0u), 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_3c7ba5();
+  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_insert_bits(uvec2 v, uvec2 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_3c7ba5() {
+  uvec2 res = tint_insert_bits(uvec2(0u, 0u), uvec2(0u, 0u), 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_3c7ba5();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uvec2 tint_insert_bits(uvec2 v, uvec2 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_3c7ba5() {
+  uvec2 res = tint_insert_bits(uvec2(0u, 0u), uvec2(0u, 0u), 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_3c7ba5();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/3c7ba5.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+uint2 tint_insert_bits(uint2 v, uint2 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint2((s).xx)) & uint2((mask).xx)) | (v & uint2((~(mask)).xx)));
+}
+
+void insertBits_3c7ba5() {
+  uint2 res = tint_insert_bits(uint2(0u, 0u), uint2(0u, 0u), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_3c7ba5();
+  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() {
+  insertBits_3c7ba5();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_3c7ba5();
+  return;
+}

test/tint/builtins/gen/insertBits/3c7ba5.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint2 tint_insert_bits(uint2 v, uint2 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_3c7ba5() {
+  uint2 res = tint_insert_bits(uint2(), uint2(), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_3c7ba5();
+  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() {
+  insertBits_3c7ba5();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_3c7ba5();
+  return;
+}
+

test/tint/builtins/gen/insertBits/3c7ba5.wgsl.expected.spvasm

@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+         %19 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_3c7ba5 "insertBits_3c7ba5"
+               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 %v2uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+         %30 = OpConstantNull %v2uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %v2uint None %9
+          %v = OpFunctionParameter %v2uint
+          %n = OpFunctionParameter %v2uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %17 = OpLabel
+         %18 = OpExtInst %uint %19 UMin %offset %uint_32
+         %22 = OpIAdd %uint %18 %count
+         %21 = OpExtInst %uint %19 UMin %uint_32 %22
+         %24 = OpISub %uint %21 %18
+         %23 = OpBitFieldInsert %v2uint %v %n %18 %24
+               OpReturnValue %23
+               OpFunctionEnd
+%insertBits_3c7ba5 = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_v2uint Function %30
+         %29 = OpFunctionCall %v2uint %tint_insert_bits %30 %30 %uint_1 %uint_1
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %insertBits_3c7ba5
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %25
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %insertBits_3c7ba5
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %insertBits_3c7ba5
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/3c7ba5.wgsl.expected.wgsl

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

test/tint/builtins/gen/insertBits/428b0b.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 insertBits(vec<3, i32>, vec<3, i32>, u32, u32) -> vec<3, i32>
+fn insertBits_428b0b() {
+  var res: vec3<i32> = insertBits(vec3<i32>(), vec3<i32>(), 1u, 1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  insertBits_428b0b();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  insertBits_428b0b();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  insertBits_428b0b();
+}

test/tint/builtins/gen/insertBits/428b0b.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+ivec3 tint_insert_bits(ivec3 v, ivec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_428b0b() {
+  ivec3 res = tint_insert_bits(ivec3(0, 0, 0), ivec3(0, 0, 0), 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_428b0b();
+  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_insert_bits(ivec3 v, ivec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_428b0b() {
+  ivec3 res = tint_insert_bits(ivec3(0, 0, 0), ivec3(0, 0, 0), 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_428b0b();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+ivec3 tint_insert_bits(ivec3 v, ivec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_428b0b() {
+  ivec3 res = tint_insert_bits(ivec3(0, 0, 0), ivec3(0, 0, 0), 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_428b0b();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/428b0b.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+int3 tint_insert_bits(int3 v, int3 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint3((s).xxx)) & int3((int(mask)).xxx)) | (v & int3((int(~(mask))).xxx)));
+}
+
+void insertBits_428b0b() {
+  int3 res = tint_insert_bits(int3(0, 0, 0), int3(0, 0, 0), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_428b0b();
+  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() {
+  insertBits_428b0b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_428b0b();
+  return;
+}

test/tint/builtins/gen/insertBits/428b0b.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int3 tint_insert_bits(int3 v, int3 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_428b0b() {
+  int3 res = tint_insert_bits(int3(), int3(), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_428b0b();
+  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() {
+  insertBits_428b0b();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_428b0b();
+  return;
+}
+

test/tint/builtins/gen/insertBits/428b0b.wgsl.expected.spvasm

@@ -0,0 +1,90 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; Schema: 0
+               OpCapability Shader
+         %20 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_428b0b "insertBits_428b0b"
+               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
+       %uint = OpTypeInt 32 0
+          %9 = OpTypeFunction %v3int %v3int %v3int %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %26 = OpTypeFunction %void
+         %31 = OpConstantNull %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %35 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %v3int None %9
+          %v = OpFunctionParameter %v3int
+          %n = OpFunctionParameter %v3int
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %18 = OpLabel
+         %19 = OpExtInst %uint %20 UMin %offset %uint_32
+         %23 = OpIAdd %uint %19 %count
+         %22 = OpExtInst %uint %20 UMin %uint_32 %23
+         %25 = OpISub %uint %22 %19
+         %24 = OpBitFieldInsert %v3int %v %n %19 %25
+               OpReturnValue %24
+               OpFunctionEnd
+%insertBits_428b0b = OpFunction %void None %26
+         %29 = OpLabel
+        %res = OpVariable %_ptr_Function_v3int Function %31
+         %30 = OpFunctionCall %v3int %tint_insert_bits %31 %31 %uint_1 %uint_1
+               OpStore %res %30
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %insertBits_428b0b
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %26
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %26
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %insertBits_428b0b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %26
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %insertBits_428b0b
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/428b0b.wgsl.expected.wgsl

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

test/tint/builtins/gen/insertBits/51ede1.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 insertBits(vec<4, u32>, vec<4, u32>, u32, u32) -> vec<4, u32>
+fn insertBits_51ede1() {
+  var res: vec4<u32> = insertBits(vec4<u32>(), vec4<u32>(), 1u, 1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  insertBits_51ede1();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  insertBits_51ede1();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  insertBits_51ede1();
+}

test/tint/builtins/gen/insertBits/51ede1.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+uvec4 tint_insert_bits(uvec4 v, uvec4 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_51ede1() {
+  uvec4 res = tint_insert_bits(uvec4(0u, 0u, 0u, 0u), uvec4(0u, 0u, 0u, 0u), 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_51ede1();
+  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_insert_bits(uvec4 v, uvec4 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_51ede1() {
+  uvec4 res = tint_insert_bits(uvec4(0u, 0u, 0u, 0u), uvec4(0u, 0u, 0u, 0u), 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_51ede1();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uvec4 tint_insert_bits(uvec4 v, uvec4 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_51ede1() {
+  uvec4 res = tint_insert_bits(uvec4(0u, 0u, 0u, 0u), uvec4(0u, 0u, 0u, 0u), 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_51ede1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/51ede1.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+uint4 tint_insert_bits(uint4 v, uint4 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint4((s).xxxx)) & uint4((mask).xxxx)) | (v & uint4((~(mask)).xxxx)));
+}
+
+void insertBits_51ede1() {
+  uint4 res = tint_insert_bits(uint4(0u, 0u, 0u, 0u), uint4(0u, 0u, 0u, 0u), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_51ede1();
+  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() {
+  insertBits_51ede1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_51ede1();
+  return;
+}

test/tint/builtins/gen/insertBits/51ede1.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint4 tint_insert_bits(uint4 v, uint4 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_51ede1() {
+  uint4 res = tint_insert_bits(uint4(), uint4(), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_51ede1();
+  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() {
+  insertBits_51ede1();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_51ede1();
+  return;
+}
+

test/tint/builtins/gen/insertBits/51ede1.wgsl.expected.spvasm

@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+         %19 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_51ede1 "insertBits_51ede1"
+               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 %v4uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+         %30 = OpConstantNull %v4uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %v4uint None %9
+          %v = OpFunctionParameter %v4uint
+          %n = OpFunctionParameter %v4uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %17 = OpLabel
+         %18 = OpExtInst %uint %19 UMin %offset %uint_32
+         %22 = OpIAdd %uint %18 %count
+         %21 = OpExtInst %uint %19 UMin %uint_32 %22
+         %24 = OpISub %uint %21 %18
+         %23 = OpBitFieldInsert %v4uint %v %n %18 %24
+               OpReturnValue %23
+               OpFunctionEnd
+%insertBits_51ede1 = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_v4uint Function %30
+         %29 = OpFunctionCall %v4uint %tint_insert_bits %30 %30 %uint_1 %uint_1
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %insertBits_51ede1
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %25
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %insertBits_51ede1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %insertBits_51ede1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/51ede1.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/insertBits/65468b.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+int tint_insert_bits(int v, int n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_65468b() {
+  int res = tint_insert_bits(1, 1, 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_65468b();
+  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_insert_bits(int v, int n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_65468b() {
+  int res = tint_insert_bits(1, 1, 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_65468b();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+int tint_insert_bits(int v, int n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_65468b() {
+  int res = tint_insert_bits(1, 1, 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_65468b();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/65468b.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+int tint_insert_bits(int v, int n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << s) & int(mask)) | (v & int(~(mask))));
+}
+
+void insertBits_65468b() {
+  int res = tint_insert_bits(1, 1, 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_65468b();
+  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() {
+  insertBits_65468b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_65468b();
+  return;
+}

test/tint/builtins/gen/insertBits/65468b.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_insert_bits(int v, int n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_65468b() {
+  int res = tint_insert_bits(1, 1, 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_65468b();
+  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() {
+  insertBits_65468b();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_65468b();
+  return;
+}
+

test/tint/builtins/gen/insertBits/65468b.wgsl.expected.spvasm

@@ -0,0 +1,90 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; Schema: 0
+               OpCapability Shader
+         %19 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_65468b "insertBits_65468b"
+               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
+       %uint = OpTypeInt 32 0
+          %9 = OpTypeFunction %int %int %int %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+      %int_1 = OpConstant %int 1
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %34 = OpConstantNull %int
+         %35 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %int None %9
+          %v = OpFunctionParameter %int
+          %n = OpFunctionParameter %int
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %17 = OpLabel
+         %18 = OpExtInst %uint %19 UMin %offset %uint_32
+         %22 = OpIAdd %uint %18 %count
+         %21 = OpExtInst %uint %19 UMin %uint_32 %22
+         %24 = OpISub %uint %21 %18
+         %23 = OpBitFieldInsert %int %v %n %18 %24
+               OpReturnValue %23
+               OpFunctionEnd
+%insertBits_65468b = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %34
+         %29 = OpFunctionCall %int %tint_insert_bits %int_1 %int_1 %uint_1 %uint_1
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %insertBits_65468b
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %25
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %insertBits_65468b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %insertBits_65468b
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/65468b.wgsl.expected.wgsl

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

test/tint/builtins/gen/insertBits/87826b.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 insertBits(vec<3, u32>, vec<3, u32>, u32, u32) -> vec<3, u32>
+fn insertBits_87826b() {
+  var res: vec3<u32> = insertBits(vec3<u32>(), vec3<u32>(), 1u, 1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  insertBits_87826b();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  insertBits_87826b();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  insertBits_87826b();
+}

test/tint/builtins/gen/insertBits/87826b.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+uvec3 tint_insert_bits(uvec3 v, uvec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_87826b() {
+  uvec3 res = tint_insert_bits(uvec3(0u, 0u, 0u), uvec3(0u, 0u, 0u), 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_87826b();
+  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_insert_bits(uvec3 v, uvec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_87826b() {
+  uvec3 res = tint_insert_bits(uvec3(0u, 0u, 0u), uvec3(0u, 0u, 0u), 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_87826b();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uvec3 tint_insert_bits(uvec3 v, uvec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_87826b() {
+  uvec3 res = tint_insert_bits(uvec3(0u, 0u, 0u), uvec3(0u, 0u, 0u), 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_87826b();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/87826b.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+uint3 tint_insert_bits(uint3 v, uint3 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint3((s).xxx)) & uint3((mask).xxx)) | (v & uint3((~(mask)).xxx)));
+}
+
+void insertBits_87826b() {
+  uint3 res = tint_insert_bits(uint3(0u, 0u, 0u), uint3(0u, 0u, 0u), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_87826b();
+  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() {
+  insertBits_87826b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_87826b();
+  return;
+}

test/tint/builtins/gen/insertBits/87826b.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint3 tint_insert_bits(uint3 v, uint3 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_87826b() {
+  uint3 res = tint_insert_bits(uint3(), uint3(), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_87826b();
+  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() {
+  insertBits_87826b();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_87826b();
+  return;
+}
+

test/tint/builtins/gen/insertBits/87826b.wgsl.expected.spvasm

@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+         %19 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_87826b "insertBits_87826b"
+               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 %v3uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+         %30 = OpConstantNull %v3uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %v3uint None %9
+          %v = OpFunctionParameter %v3uint
+          %n = OpFunctionParameter %v3uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %17 = OpLabel
+         %18 = OpExtInst %uint %19 UMin %offset %uint_32
+         %22 = OpIAdd %uint %18 %count
+         %21 = OpExtInst %uint %19 UMin %uint_32 %22
+         %24 = OpISub %uint %21 %18
+         %23 = OpBitFieldInsert %v3uint %v %n %18 %24
+               OpReturnValue %23
+               OpFunctionEnd
+%insertBits_87826b = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_v3uint Function %30
+         %29 = OpFunctionCall %v3uint %tint_insert_bits %30 %30 %uint_1 %uint_1
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %insertBits_87826b
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %25
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %insertBits_87826b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %insertBits_87826b
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/87826b.wgsl.expected.wgsl

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

test/tint/builtins/gen/insertBits/d86978.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 insertBits(vec<4, i32>, vec<4, i32>, u32, u32) -> vec<4, i32>
+fn insertBits_d86978() {
+  var res: vec4<i32> = insertBits(vec4<i32>(), vec4<i32>(), 1u, 1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  insertBits_d86978();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  insertBits_d86978();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  insertBits_d86978();
+}

test/tint/builtins/gen/insertBits/d86978.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+ivec4 tint_insert_bits(ivec4 v, ivec4 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_d86978() {
+  ivec4 res = tint_insert_bits(ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_d86978();
+  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_insert_bits(ivec4 v, ivec4 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_d86978() {
+  ivec4 res = tint_insert_bits(ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_d86978();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+ivec4 tint_insert_bits(ivec4 v, ivec4 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_d86978() {
+  ivec4 res = tint_insert_bits(ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_d86978();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/d86978.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+int4 tint_insert_bits(int4 v, int4 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint4((s).xxxx)) & int4((int(mask)).xxxx)) | (v & int4((int(~(mask))).xxxx)));
+}
+
+void insertBits_d86978() {
+  int4 res = tint_insert_bits(int4(0, 0, 0, 0), int4(0, 0, 0, 0), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_d86978();
+  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() {
+  insertBits_d86978();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_d86978();
+  return;
+}

test/tint/builtins/gen/insertBits/d86978.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int4 tint_insert_bits(int4 v, int4 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_d86978() {
+  int4 res = tint_insert_bits(int4(), int4(), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_d86978();
+  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() {
+  insertBits_d86978();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_d86978();
+  return;
+}
+

test/tint/builtins/gen/insertBits/d86978.wgsl.expected.spvasm

@@ -0,0 +1,90 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; Schema: 0
+               OpCapability Shader
+         %20 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_d86978 "insertBits_d86978"
+               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
+       %uint = OpTypeInt 32 0
+          %9 = OpTypeFunction %v4int %v4int %v4int %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %26 = OpTypeFunction %void
+         %31 = OpConstantNull %v4int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %35 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %v4int None %9
+          %v = OpFunctionParameter %v4int
+          %n = OpFunctionParameter %v4int
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %18 = OpLabel
+         %19 = OpExtInst %uint %20 UMin %offset %uint_32
+         %23 = OpIAdd %uint %19 %count
+         %22 = OpExtInst %uint %20 UMin %uint_32 %23
+         %25 = OpISub %uint %22 %19
+         %24 = OpBitFieldInsert %v4int %v %n %19 %25
+               OpReturnValue %24
+               OpFunctionEnd
+%insertBits_d86978 = OpFunction %void None %26
+         %29 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %31
+         %30 = OpFunctionCall %v4int %tint_insert_bits %31 %31 %uint_1 %uint_1
+               OpStore %res %30
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %insertBits_d86978
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %26
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %26
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %insertBits_d86978
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %26
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %insertBits_d86978
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/d86978.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/insertBits/e3e3a2.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+uint tint_insert_bits(uint v, uint n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_e3e3a2() {
+  uint res = tint_insert_bits(1u, 1u, 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_e3e3a2();
+  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_insert_bits(uint v, uint n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_e3e3a2() {
+  uint res = tint_insert_bits(1u, 1u, 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_e3e3a2();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+uint tint_insert_bits(uint v, uint n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_e3e3a2() {
+  uint res = tint_insert_bits(1u, 1u, 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_e3e3a2();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/e3e3a2.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+uint tint_insert_bits(uint v, uint n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << s) & mask) | (v & ~(mask)));
+}
+
+void insertBits_e3e3a2() {
+  uint res = tint_insert_bits(1u, 1u, 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_e3e3a2();
+  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() {
+  insertBits_e3e3a2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_e3e3a2();
+  return;
+}

test/tint/builtins/gen/insertBits/e3e3a2.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint tint_insert_bits(uint v, uint n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_e3e3a2() {
+  uint res = tint_insert_bits(1u, 1u, 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_e3e3a2();
+  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() {
+  insertBits_e3e3a2();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_e3e3a2();
+  return;
+}
+

test/tint/builtins/gen/insertBits/e3e3a2.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 47
+; Schema: 0
+               OpCapability Shader
+         %18 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_e3e3a2 "insertBits_e3e3a2"
+               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 %uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %24 = OpTypeFunction %void
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %32 = OpConstantNull %uint
+         %33 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %uint None %9
+          %v = OpFunctionParameter %uint
+          %n = OpFunctionParameter %uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %16 = OpLabel
+         %17 = OpExtInst %uint %18 UMin %offset %uint_32
+         %21 = OpIAdd %uint %17 %count
+         %20 = OpExtInst %uint %18 UMin %uint_32 %21
+         %23 = OpISub %uint %20 %17
+         %22 = OpBitFieldInsert %uint %v %n %17 %23
+               OpReturnValue %22
+               OpFunctionEnd
+%insertBits_e3e3a2 = OpFunction %void None %24
+         %27 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %32
+         %28 = OpFunctionCall %uint %tint_insert_bits %uint_1 %uint_1 %uint_1 %uint_1
+               OpStore %res %28
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %33
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %insertBits_e3e3a2
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %24
+         %38 = OpLabel
+         %39 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %39
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %24
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %insertBits_e3e3a2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %24
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %insertBits_e3e3a2
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/e3e3a2.wgsl.expected.wgsl

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

test/tint/builtins/gen/insertBits/fe6ba6.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 insertBits(vec<2, i32>, vec<2, i32>, u32, u32) -> vec<2, i32>
+fn insertBits_fe6ba6() {
+  var res: vec2<i32> = insertBits(vec2<i32>(), vec2<i32>(), 1u, 1u);
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  insertBits_fe6ba6();
+  return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+  insertBits_fe6ba6();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  insertBits_fe6ba6();
+}

test/tint/builtins/gen/insertBits/fe6ba6.wgsl.expected.glsl

@@ -0,0 +1,66 @@
+#version 310 es
+
+ivec2 tint_insert_bits(ivec2 v, ivec2 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_fe6ba6() {
+  ivec2 res = tint_insert_bits(ivec2(0, 0), ivec2(0, 0), 1u, 1u);
+}
+
+vec4 vertex_main() {
+  insertBits_fe6ba6();
+  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_insert_bits(ivec2 v, ivec2 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_fe6ba6() {
+  ivec2 res = tint_insert_bits(ivec2(0, 0), ivec2(0, 0), 1u, 1u);
+}
+
+void fragment_main() {
+  insertBits_fe6ba6();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+ivec2 tint_insert_bits(ivec2 v, ivec2 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void insertBits_fe6ba6() {
+  ivec2 res = tint_insert_bits(ivec2(0, 0), ivec2(0, 0), 1u, 1u);
+}
+
+void compute_main() {
+  insertBits_fe6ba6();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/insertBits/fe6ba6.wgsl.expected.hlsl

@@ -0,0 +1,37 @@
+int2 tint_insert_bits(int2 v, int2 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint2((s).xx)) & int2((int(mask)).xx)) | (v & int2((int(~(mask))).xx)));
+}
+
+void insertBits_fe6ba6() {
+  int2 res = tint_insert_bits(int2(0, 0), int2(0, 0), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  insertBits_fe6ba6();
+  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() {
+  insertBits_fe6ba6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  insertBits_fe6ba6();
+  return;
+}

test/tint/builtins/gen/insertBits/fe6ba6.wgsl.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int2 tint_insert_bits(int2 v, int2 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void insertBits_fe6ba6() {
+  int2 res = tint_insert_bits(int2(), int2(), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  insertBits_fe6ba6();
+  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() {
+  insertBits_fe6ba6();
+  return;
+}
+
+kernel void compute_main() {
+  insertBits_fe6ba6();
+  return;
+}
+

test/tint/builtins/gen/insertBits/fe6ba6.wgsl.expected.spvasm

@@ -0,0 +1,90 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; Schema: 0
+               OpCapability Shader
+         %20 = OpExtInstImport "GLSL.std.450"
+               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_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %insertBits_fe6ba6 "insertBits_fe6ba6"
+               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
+       %uint = OpTypeInt 32 0
+          %9 = OpTypeFunction %v2int %v2int %v2int %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %26 = OpTypeFunction %void
+         %31 = OpConstantNull %v2int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %35 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_insert_bits = OpFunction %v2int None %9
+          %v = OpFunctionParameter %v2int
+          %n = OpFunctionParameter %v2int
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %18 = OpLabel
+         %19 = OpExtInst %uint %20 UMin %offset %uint_32
+         %23 = OpIAdd %uint %19 %count
+         %22 = OpExtInst %uint %20 UMin %uint_32 %23
+         %25 = OpISub %uint %22 %19
+         %24 = OpBitFieldInsert %v2int %v %n %19 %25
+               OpReturnValue %24
+               OpFunctionEnd
+%insertBits_fe6ba6 = OpFunction %void None %26
+         %29 = OpLabel
+        %res = OpVariable %_ptr_Function_v2int Function %31
+         %30 = OpFunctionCall %v2int %tint_insert_bits %31 %31 %uint_1 %uint_1
+               OpStore %res %30
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %insertBits_fe6ba6
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %26
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %26
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %insertBits_fe6ba6
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %26
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %insertBits_fe6ba6
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/insertBits/fe6ba6.wgsl.expected.wgsl

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

test/tint/builtins/insertBits/scalar/i32.spvasm

@@ -0,0 +1,33 @@
+                      OpCapability Shader
+                %16 = OpExtInstImport "GLSL.std.450"
+                      OpMemoryModel Logical GLSL450
+                      OpEntryPoint GLCompute %f "f"
+                      OpExecutionMode %f LocalSize 1 1 1
+                      OpName %f "f"
+                      OpName %v "v"
+                      OpName %n "n"
+                      OpName %offset "offset"
+                      OpName %count "count"
+              %void = OpTypeVoid
+                 %1 = OpTypeFunction %void
+               %int = OpTypeInt 32 1
+ %_ptr_Function_int = OpTypePointer Function %int
+                 %8 = OpConstantNull %int
+              %uint = OpTypeInt 32 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+                %13 = OpConstantNull %uint
+           %uint_31 = OpConstant %uint 31
+           %uint_32 = OpConstant %uint 32
+                 %f = OpFunction %void None %1
+                 %4 = OpLabel
+                 %v = OpVariable %_ptr_Function_int Function %8
+                 %n = OpVariable %_ptr_Function_int Function %8
+            %offset = OpVariable %_ptr_Function_uint Function %13
+             %count = OpVariable %_ptr_Function_uint Function %13
+                %17 = OpLoad %int %v
+                %18 = OpLoad %int %n
+                %19 = OpLoad %uint %offset
+                %20 = OpLoad %uint %count
+                %15 = OpBitFieldInsert %int %17 %18 %19 %20
+                      OpReturn
+                      OpFunctionEnd

test/tint/builtins/insertBits/scalar/i32.spvasm.expected.glsl

@@ -0,0 +1,26 @@
+#version 310 es
+
+int tint_insert_bits(int v, int n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void f_1() {
+  int v = 0;
+  int n = 0;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  int x_15 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+void f() {
+  f_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f();
+  return;
+}

test/tint/builtins/insertBits/scalar/i32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+int tint_insert_bits(int v, int n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << s) & int(mask)) | (v & int(~(mask))));
+}
+
+void f_1() {
+  int v = 0;
+  int n = 0;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const int x_15 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

test/tint/builtins/insertBits/scalar/i32.spvasm.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_insert_bits(int v, int n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void f_1() {
+  int v = 0;
+  int n = 0;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  int const x_17 = v;
+  int const x_18 = n;
+  uint const x_19 = offset_1;
+  uint const x_20 = count;
+  int const x_15 = tint_insert_bits(x_17, x_18, x_19, x_20);
+  return;
+}
+
+kernel void f() {
+  f_1();
+  return;
+}
+

test/tint/builtins/insertBits/scalar/i32.spvasm.expected.spvasm

@@ -0,0 +1,68 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 39
+; Schema: 0
+               OpCapability Shader
+         %11 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %tint_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               OpName %n_0 "n"
+               OpName %offset_1 "offset_1"
+               OpName %count_0 "count"
+               OpName %f "f"
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+          %1 = OpTypeFunction %int %int %int %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %17 = OpTypeFunction %void
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %24 = OpConstantNull %int
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %29 = OpConstantNull %uint
+%tint_insert_bits = OpFunction %int None %1
+          %v = OpFunctionParameter %int
+          %n = OpFunctionParameter %int
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+          %9 = OpLabel
+         %10 = OpExtInst %uint %11 UMin %offset %uint_32
+         %14 = OpIAdd %uint %10 %count
+         %13 = OpExtInst %uint %11 UMin %uint_32 %14
+         %16 = OpISub %uint %13 %10
+         %15 = OpBitFieldInsert %int %v %n %10 %16
+               OpReturnValue %15
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %17
+         %20 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_int Function %24
+        %n_0 = OpVariable %_ptr_Function_int Function %24
+   %offset_1 = OpVariable %_ptr_Function_uint Function %29
+    %count_0 = OpVariable %_ptr_Function_uint Function %29
+               OpStore %v_0 %int_0
+               OpStore %n_0 %int_0
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %31 = OpLoad %int %v_0
+         %32 = OpLoad %int %n_0
+         %33 = OpLoad %uint %offset_1
+         %34 = OpLoad %uint %count_0
+         %35 = OpFunctionCall %int %tint_insert_bits %31 %32 %33 %34
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %17
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/insertBits/scalar/i32.spvasm.expected.wgsl

@@ -0,0 +1,17 @@
+fn f_1() {
+  var v : i32 = 0;
+  var n : i32 = 0;
+  var offset_1 : u32 = 0u;
+  var count : u32 = 0u;
+  let x_17 : i32 = v;
+  let x_18 : i32 = n;
+  let x_19 : u32 = offset_1;
+  let x_20 : u32 = count;
+  let x_15 : i32 = insertBits(x_17, x_18, x_19, x_20);
+  return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+  f_1();
+}

test/tint/builtins/insertBits/scalar/u32.spvasm

@@ -0,0 +1,30 @@
+                      OpCapability Shader
+                %13 = OpExtInstImport "GLSL.std.450"
+                      OpMemoryModel Logical GLSL450
+                      OpEntryPoint GLCompute %f "f"
+                      OpExecutionMode %f LocalSize 1 1 1
+                      OpName %f "f"
+                      OpName %v "v"
+                      OpName %n "n"
+                      OpName %offset "offset"
+                      OpName %count "count"
+              %void = OpTypeVoid
+                 %1 = OpTypeFunction %void
+              %uint = OpTypeInt 32 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+                 %8 = OpConstantNull %uint
+           %uint_31 = OpConstant %uint 31
+           %uint_32 = OpConstant %uint 32
+                 %f = OpFunction %void None %1
+                 %4 = OpLabel
+                 %v = OpVariable %_ptr_Function_uint Function %8
+                 %n = OpVariable %_ptr_Function_uint Function %8
+            %offset = OpVariable %_ptr_Function_uint Function %8
+             %count = OpVariable %_ptr_Function_uint Function %8
+                %14 = OpLoad %uint %v
+                %15 = OpLoad %uint %n
+                %16 = OpLoad %uint %offset
+                %17 = OpLoad %uint %count
+                %12 = OpBitFieldInsert %uint %14 %15 %16 %17
+                      OpReturn
+                      OpFunctionEnd

test/tint/builtins/insertBits/scalar/u32.spvasm.expected.glsl

@@ -0,0 +1,26 @@
+#version 310 es
+
+uint tint_insert_bits(uint v, uint n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void f_1() {
+  uint v = 0u;
+  uint n = 0u;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uint x_12 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+void f() {
+  f_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f();
+  return;
+}

test/tint/builtins/insertBits/scalar/u32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+uint tint_insert_bits(uint v, uint n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << s) & mask) | (v & ~(mask)));
+}
+
+void f_1() {
+  uint v = 0u;
+  uint n = 0u;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const uint x_12 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

test/tint/builtins/insertBits/scalar/u32.spvasm.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint tint_insert_bits(uint v, uint n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void f_1() {
+  uint v = 0u;
+  uint n = 0u;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uint const x_14 = v;
+  uint const x_15 = n;
+  uint const x_16 = offset_1;
+  uint const x_17 = count;
+  uint const x_12 = tint_insert_bits(x_14, x_15, x_16, x_17);
+  return;
+}
+
+kernel void f() {
+  f_1();
+  return;
+}
+

test/tint/builtins/insertBits/scalar/u32.spvasm.expected.spvasm

@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 35
+; Schema: 0
+               OpCapability Shader
+         %10 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %tint_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               OpName %n_0 "n"
+               OpName %offset_1 "offset_1"
+               OpName %count_0 "count"
+               OpName %f "f"
+       %uint = OpTypeInt 32 0
+          %1 = OpTypeFunction %uint %uint %uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %23 = OpConstantNull %uint
+%tint_insert_bits = OpFunction %uint None %1
+          %v = OpFunctionParameter %uint
+          %n = OpFunctionParameter %uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+          %8 = OpLabel
+          %9 = OpExtInst %uint %10 UMin %offset %uint_32
+         %13 = OpIAdd %uint %9 %count
+         %12 = OpExtInst %uint %10 UMin %uint_32 %13
+         %15 = OpISub %uint %12 %9
+         %14 = OpBitFieldInsert %uint %v %n %9 %15
+               OpReturnValue %14
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %16
+         %19 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_uint Function %23
+        %n_0 = OpVariable %_ptr_Function_uint Function %23
+   %offset_1 = OpVariable %_ptr_Function_uint Function %23
+    %count_0 = OpVariable %_ptr_Function_uint Function %23
+               OpStore %v_0 %uint_0
+               OpStore %n_0 %uint_0
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %27 = OpLoad %uint %v_0
+         %28 = OpLoad %uint %n_0
+         %29 = OpLoad %uint %offset_1
+         %30 = OpLoad %uint %count_0
+         %31 = OpFunctionCall %uint %tint_insert_bits %27 %28 %29 %30
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %16
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/insertBits/scalar/u32.spvasm.expected.wgsl

@@ -0,0 +1,17 @@
+fn f_1() {
+  var v : u32 = 0u;
+  var n : u32 = 0u;
+  var offset_1 : u32 = 0u;
+  var count : u32 = 0u;
+  let x_14 : u32 = v;
+  let x_15 : u32 = n;
+  let x_16 : u32 = offset_1;
+  let x_17 : u32 = count;
+  let x_12 : u32 = insertBits(x_14, x_15, x_16, x_17);
+  return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+  f_1();
+}

test/tint/builtins/insertBits/vec3/i32.spvasm

@@ -0,0 +1,34 @@
+                       OpCapability Shader
+                 %17 = OpExtInstImport "GLSL.std.450"
+                       OpMemoryModel Logical GLSL450
+                       OpEntryPoint GLCompute %f "f"
+                       OpExecutionMode %f LocalSize 1 1 1
+                       OpName %f "f"
+                       OpName %v "v"
+                       OpName %n "n"
+                       OpName %offset "offset"
+                       OpName %count "count"
+               %void = OpTypeVoid
+                  %1 = OpTypeFunction %void
+                %int = OpTypeInt 32 1
+              %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+                  %9 = OpConstantNull %v3int
+               %uint = OpTypeInt 32 0
+ %_ptr_Function_uint = OpTypePointer Function %uint
+                 %14 = OpConstantNull %uint
+            %uint_31 = OpConstant %uint 31
+            %uint_32 = OpConstant %uint 32
+                  %f = OpFunction %void None %1
+                  %4 = OpLabel
+                  %v = OpVariable %_ptr_Function_v3int Function %9
+                  %n = OpVariable %_ptr_Function_v3int Function %9
+             %offset = OpVariable %_ptr_Function_uint Function %14
+              %count = OpVariable %_ptr_Function_uint Function %14
+                 %18 = OpLoad %v3int %v
+                 %19 = OpLoad %v3int %n
+                 %20 = OpLoad %uint %offset
+                 %21 = OpLoad %uint %count
+                 %16 = OpBitFieldInsert %v3int %18 %19 %20 %21
+                       OpReturn
+                       OpFunctionEnd

test/tint/builtins/insertBits/vec3/i32.spvasm.expected.glsl

@@ -0,0 +1,26 @@
+#version 310 es
+
+ivec3 tint_insert_bits(ivec3 v, ivec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void f_1() {
+  ivec3 v = ivec3(0, 0, 0);
+  ivec3 n = ivec3(0, 0, 0);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  ivec3 x_16 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+void f() {
+  f_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f();
+  return;
+}

test/tint/builtins/insertBits/vec3/i32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+int3 tint_insert_bits(int3 v, int3 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint3((s).xxx)) & int3((int(mask)).xxx)) | (v & int3((int(~(mask))).xxx)));
+}
+
+void f_1() {
+  int3 v = int3(0, 0, 0);
+  int3 n = int3(0, 0, 0);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const int3 x_16 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

test/tint/builtins/insertBits/vec3/i32.spvasm.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int3 tint_insert_bits(int3 v, int3 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void f_1() {
+  int3 v = int3();
+  int3 n = int3();
+  uint offset_1 = 0u;
+  uint count = 0u;
+  int3 const x_18 = v;
+  int3 const x_19 = n;
+  uint const x_20 = offset_1;
+  uint const x_21 = count;
+  int3 const x_16 = tint_insert_bits(x_18, x_19, x_20, x_21);
+  return;
+}
+
+kernel void f() {
+  f_1();
+  return;
+}
+

test/tint/builtins/insertBits/vec3/i32.spvasm.expected.spvasm

@@ -0,0 +1,68 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 39
+; Schema: 0
+               OpCapability Shader
+         %12 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %tint_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               OpName %n_0 "n"
+               OpName %offset_1 "offset_1"
+               OpName %count_0 "count"
+               OpName %f "f"
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+          %1 = OpTypeFunction %v3int %v3int %v3int %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %18 = OpTypeFunction %void
+         %22 = OpConstantNull %v3int
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %29 = OpConstantNull %uint
+%tint_insert_bits = OpFunction %v3int None %1
+          %v = OpFunctionParameter %v3int
+          %n = OpFunctionParameter %v3int
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %10 = OpLabel
+         %11 = OpExtInst %uint %12 UMin %offset %uint_32
+         %15 = OpIAdd %uint %11 %count
+         %14 = OpExtInst %uint %12 UMin %uint_32 %15
+         %17 = OpISub %uint %14 %11
+         %16 = OpBitFieldInsert %v3int %v %n %11 %17
+               OpReturnValue %16
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %18
+         %21 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_v3int Function %22
+        %n_0 = OpVariable %_ptr_Function_v3int Function %22
+   %offset_1 = OpVariable %_ptr_Function_uint Function %29
+    %count_0 = OpVariable %_ptr_Function_uint Function %29
+               OpStore %v_0 %22
+               OpStore %n_0 %22
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %31 = OpLoad %v3int %v_0
+         %32 = OpLoad %v3int %n_0
+         %33 = OpLoad %uint %offset_1
+         %34 = OpLoad %uint %count_0
+         %35 = OpFunctionCall %v3int %tint_insert_bits %31 %32 %33 %34
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %18
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/insertBits/vec3/i32.spvasm.expected.wgsl

@@ -0,0 +1,17 @@
+fn f_1() {
+  var v : vec3<i32> = vec3<i32>();
+  var n : vec3<i32> = vec3<i32>();
+  var offset_1 : u32 = 0u;
+  var count : u32 = 0u;
+  let x_18 : vec3<i32> = v;
+  let x_19 : vec3<i32> = n;
+  let x_20 : u32 = offset_1;
+  let x_21 : u32 = count;
+  let x_16 : vec3<i32> = insertBits(x_18, x_19, x_20, x_21);
+  return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+  f_1();
+}

test/tint/builtins/insertBits/vec3/u32.spvasm

@@ -0,0 +1,33 @@
+                        OpCapability Shader
+                  %16 = OpExtInstImport "GLSL.std.450"
+                        OpMemoryModel Logical GLSL450
+                        OpEntryPoint GLCompute %f "f"
+                        OpExecutionMode %f LocalSize 1 1 1
+                        OpName %f "f"
+                        OpName %v "v"
+                        OpName %n "n"
+                        OpName %offset "offset"
+                        OpName %count "count"
+                %void = OpTypeVoid
+                   %1 = OpTypeFunction %void
+                %uint = OpTypeInt 32 0
+              %v3uint = OpTypeVector %uint 3
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+                   %9 = OpConstantNull %v3uint
+  %_ptr_Function_uint = OpTypePointer Function %uint
+                  %13 = OpConstantNull %uint
+             %uint_31 = OpConstant %uint 31
+             %uint_32 = OpConstant %uint 32
+                   %f = OpFunction %void None %1
+                   %4 = OpLabel
+                   %v = OpVariable %_ptr_Function_v3uint Function %9
+                   %n = OpVariable %_ptr_Function_v3uint Function %9
+              %offset = OpVariable %_ptr_Function_uint Function %13
+               %count = OpVariable %_ptr_Function_uint Function %13
+                  %17 = OpLoad %v3uint %v
+                  %18 = OpLoad %v3uint %n
+                  %19 = OpLoad %uint %offset
+                  %20 = OpLoad %uint %count
+                  %15 = OpBitFieldInsert %v3uint %17 %18 %19 %20
+                        OpReturn
+                        OpFunctionEnd

test/tint/builtins/insertBits/vec3/u32.spvasm.expected.glsl

@@ -0,0 +1,26 @@
+#version 310 es
+
+uvec3 tint_insert_bits(uvec3 v, uvec3 n, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldInsert(v, n, int(s), int((e - s)));
+}
+
+void f_1() {
+  uvec3 v = uvec3(0u, 0u, 0u);
+  uvec3 n = uvec3(0u, 0u, 0u);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uvec3 x_15 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+void f() {
+  f_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f();
+  return;
+}

test/tint/builtins/insertBits/vec3/u32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+uint3 tint_insert_bits(uint3 v, uint3 n, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint mask = (((1u << s) - 1u) ^ ((1u << e) - 1u));
+  return (((n << uint3((s).xxx)) & uint3((mask).xxx)) | (v & uint3((~(mask)).xxx)));
+}
+
+void f_1() {
+  uint3 v = uint3(0u, 0u, 0u);
+  uint3 n = uint3(0u, 0u, 0u);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const uint3 x_15 = tint_insert_bits(v, n, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

test/tint/builtins/insertBits/vec3/u32.spvasm.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint3 tint_insert_bits(uint3 v, uint3 n, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return insert_bits(v, n, s, (e - s));
+}
+
+void f_1() {
+  uint3 v = uint3();
+  uint3 n = uint3();
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uint3 const x_17 = v;
+  uint3 const x_18 = n;
+  uint const x_19 = offset_1;
+  uint const x_20 = count;
+  uint3 const x_15 = tint_insert_bits(x_17, x_18, x_19, x_20);
+  return;
+}
+
+kernel void f() {
+  f_1();
+  return;
+}
+

test/tint/builtins/insertBits/vec3/u32.spvasm.expected.spvasm

@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+         %11 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %tint_insert_bits "tint_insert_bits"
+               OpName %v "v"
+               OpName %n "n"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               OpName %n_0 "n"
+               OpName %offset_1 "offset_1"
+               OpName %count_0 "count"
+               OpName %f "f"
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+          %1 = OpTypeFunction %v3uint %v3uint %v3uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %17 = OpTypeFunction %void
+         %21 = OpConstantNull %v3uint
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %28 = OpConstantNull %uint
+%tint_insert_bits = OpFunction %v3uint None %1
+          %v = OpFunctionParameter %v3uint
+          %n = OpFunctionParameter %v3uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+          %9 = OpLabel
+         %10 = OpExtInst %uint %11 UMin %offset %uint_32
+         %14 = OpIAdd %uint %10 %count
+         %13 = OpExtInst %uint %11 UMin %uint_32 %14
+         %16 = OpISub %uint %13 %10
+         %15 = OpBitFieldInsert %v3uint %v %n %10 %16
+               OpReturnValue %15
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %17
+         %20 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_v3uint Function %21
+        %n_0 = OpVariable %_ptr_Function_v3uint Function %21
+   %offset_1 = OpVariable %_ptr_Function_uint Function %28
+    %count_0 = OpVariable %_ptr_Function_uint Function %28
+               OpStore %v_0 %21
+               OpStore %n_0 %21
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %30 = OpLoad %v3uint %v_0
+         %31 = OpLoad %v3uint %n_0
+         %32 = OpLoad %uint %offset_1
+         %33 = OpLoad %uint %count_0
+         %34 = OpFunctionCall %v3uint %tint_insert_bits %30 %31 %32 %33
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %17
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/insertBits/vec3/u32.spvasm.expected.wgsl

@@ -0,0 +1,17 @@
+fn f_1() {
+  var v : vec3<u32> = vec3<u32>();
+  var n : vec3<u32> = vec3<u32>();
+  var offset_1 : u32 = 0u;
+  var count : u32 = 0u;
+  let x_17 : vec3<u32> = v;
+  let x_18 : vec3<u32> = n;
+  let x_19 : u32 = offset_1;
+  let x_20 : u32 = count;
+  let x_15 : vec3<u32> = insertBits(x_17, x_18, x_19, x_20);
+  return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+  f_1();
+}