builtins: Add extractBits

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

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

src/tint/builtins.def

@@ -315,6 +315,8 @@ fn exp(f32) -> f32
 fn exp<N: num>(vec<N, f32>) -> vec<N, f32>
 fn exp2(f32) -> f32
 fn exp2<N: num>(vec<N, f32>) -> vec<N, f32>
+fn extractBits<T: iu32>(T, u32, u32) -> T
+fn extractBits<N: num, T: iu32>(vec<N, T>, u32, u32) -> vec<N, T>
 fn faceForward<N: num>(vec<N, f32>, vec<N, f32>, vec<N, f32>) -> vec<N, f32>
 fn firstLeadingBit<T: iu32>(T) -> T
 fn firstLeadingBit<N: num, T: iu32>(vec<N, T>) -> vec<N, T>

src/tint/reader/spirv/function.cc

@@ -463,6 +463,9 @@ 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,9 +901,79 @@ TEST_F(SpvUnaryBitTest, BitReverse_IntVector_IntVector) {
       << body;
 }
 
+TEST_F(SpvUnaryBitTest, ExtractBits_Int) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldSExtract %v2int %int_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("let x_1 : vec2<i32> = extractBits(30, 10u, 20u);"))
+      << body;
+}
+
+TEST_F(SpvUnaryBitTest, ExtractBits_IntVector) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldSExtract %v2int %v2int_30_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> = extractBits(vec2<i32>(30, 40), 10u, 20u);"))
+      << body;
+}
+
+TEST_F(SpvUnaryBitTest, ExtractBits_Uint) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldUExtract %v2uint %uint_20 %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> = extractBits(20u, 10u, 20u);"))
+      << body;
+}
+
+TEST_F(SpvUnaryBitTest, ExtractBits_UintVector) {
+  const auto assembly = BitTestPreamble() + R"(
+     %1 = OpBitFieldUExtract %v2uint %v2uint_10_20 %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> = extractBits(vec2<u32>(10u, 20u), 10u, 20u);"))
+      << body;
+}
+
 // TODO(dneto): OpBitFieldInsert
-// TODO(dneto): OpBitFieldSExtract
-// TODO(dneto): OpBitFieldUExtract
 
 }  // namespace
 }  // namespace spirv

src/tint/sem/builtin_type.cc

@@ -114,6 +114,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
   if (name == "exp2") {
     return BuiltinType::kExp2;
   }
+  if (name == "extractBits") {
+    return BuiltinType::kExtractBits;
+  }
   if (name == "faceForward") {
     return BuiltinType::kFaceForward;
   }
@@ -411,6 +414,8 @@ const char* str(BuiltinType i) {
       return "exp";
     case BuiltinType::kExp2:
       return "exp2";
+    case BuiltinType::kExtractBits:
+      return "extractBits";
     case BuiltinType::kFaceForward:
       return "faceForward";
     case BuiltinType::kFirstLeadingBit:

src/tint/sem/builtin_type.h

@@ -62,6 +62,7 @@ enum class BuiltinType {
   kDpdyFine,
   kExp,
   kExp2,
+  kExtractBits,
   kFaceForward,
   kFirstLeadingBit,
   kFirstTrailingBit,

src/tint/transform/builtin_polyfill.cc

@@ -31,10 +31,13 @@ namespace transform {
 struct BuiltinPolyfill::State {
   /// Constructor
   /// @param c the CloneContext
-  explicit State(CloneContext& c) : ctx(c) {}
+  /// @param p the builtins to polyfill
+  State(CloneContext& c, Builtins p) : ctx(c), polyfill(p) {}
 
   /// The clone context
   CloneContext& ctx;
+  /// The builtins to polyfill
+  Builtins polyfill;
   /// The destination program builder
   ProgramBuilder& b = *ctx.dst;
   /// The source clone context
@@ -170,6 +173,57 @@ struct BuiltinPolyfill::State {
     return name;
   }
 
+  /// Builds the polyfill function for the `extractBits` builtin
+  /// @param ty the parameter and return type for the function
+  /// @return the polyfill function name
+  Symbol extractBits(const sem::Type* ty) {
+    auto name = b.Symbols().New("tint_extract_bits");
+    uint32_t width = WidthOf(ty);
+
+    constexpr uint32_t W = 32u;  // 32-bit
+
+    auto vecN_u32 =
+        [&](const ast::Expression* value) -> const ast::Expression* {
+      if (width == 1) {
+        return value;
+      }
+      return b.Construct(b.ty.vec<ProgramBuilder::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.extract_bits) {
+      case Level::kFull:
+        body.emplace_back(b.Decl(b.Const("shl", nullptr, b.Sub(W, "e"))));
+        body.emplace_back(b.Decl(b.Const("shr", nullptr, b.Add("shl", "s"))));
+        body.emplace_back(b.Return(b.Shr(b.Shl("v", vecN_u32(b.Expr("shl"))),
+                                         vecN_u32(b.Expr("shr")))));
+        break;
+      case Level::kClampParameters:
+        body.emplace_back(
+            b.Return(b.Call("extractBits", "v", "s", b.Sub("e", "s"))));
+        break;
+      default:
+        TINT_ICE(Transform, b.Diagnostics())
+            << "unhandled polyfill level: "
+            << static_cast<int>(polyfill.extract_bits);
+        return {};
+    }
+
+    b.Func(name,
+           {
+               b.Param("v", T(ty)),
+               b.Param("offset", b.ty.u32()),
+               b.Param("count", b.ty.u32()),
+           },
+           T(ty), body);
+
+    return name;
+  }
+
   /// Builds the polyfill function for the `firstLeadingBit` builtin
   /// @param ty the parameter and return type for the function
   /// @return the polyfill function name
@@ -361,6 +415,11 @@ bool BuiltinPolyfill::ShouldRun(const Program* program,
                 return true;
               }
               break;
+            case sem::BuiltinType::kExtractBits:
+              if (builtins.extract_bits != Level::kNone) {
+                return true;
+              }
+              break;
             case sem::BuiltinType::kFirstLeadingBit:
               if (builtins.first_leading_bit) {
                 return true;
@@ -395,7 +454,7 @@ void BuiltinPolyfill::Run(CloneContext& ctx,
   ctx.ReplaceAll(
       [&](const ast::CallExpression* expr) -> const ast::CallExpression* {
         auto builtins = cfg->builtins;
-        State s{ctx};
+        State s{ctx, builtins};
         if (auto* call = s.sem.Get<sem::Call>(expr)) {
           if (auto* builtin = call->Target()->As<sem::Builtin>()) {
             Symbol polyfill;
@@ -414,6 +473,13 @@ void BuiltinPolyfill::Run(CloneContext& ctx,
                   });
                 }
                 break;
+              case sem::BuiltinType::kExtractBits:
+                if (builtins.extract_bits != Level::kNone) {
+                  polyfill = utils::GetOrCreate(polyfills, builtin, [&] {
+                    return s.extractBits(builtin->ReturnType());
+                  });
+                }
+                break;
               case sem::BuiltinType::kFirstLeadingBit:
                 if (builtins.first_leading_bit) {
                   polyfill = utils::GetOrCreate(polyfills, builtin, [&] {

src/tint/transform/builtin_polyfill.h

@@ -28,12 +28,24 @@ class BuiltinPolyfill : public Castable<BuiltinPolyfill, Transform> {
   /// Destructor
   ~BuiltinPolyfill() override;
 
+  /// Enumerator of polyfill levels
+  enum class Level {
+    /// No polyfill needed, supported by the backend.
+    kNone,
+    /// Clamp the parameters to the inner implementation.
+    kClampParameters,
+    /// Polyfill the entire function
+    kFull,
+  };
+
   /// Specifies the builtins that should be polyfilled by the transform.
   struct Builtins {
     /// Should `countLeadingZeros()` be polyfilled?
     bool count_leading_zeros = false;
     /// Should `countTrailingZeros()` be polyfilled?
     bool count_trailing_zeros = false;
+    /// What level should `extractBits()` be polyfilled?
+    Level extract_bits = Level::kNone;
     /// Should `firstLeadingBit()` be polyfilled?
     bool first_leading_bit = false;
     /// Should `firstTrailingBit()` be polyfilled?

src/tint/transform/builtin_polyfill_test.cc

@@ -22,6 +22,8 @@ namespace tint {
 namespace transform {
 namespace {
 
+using Level = BuiltinPolyfill::Level;
+
 using BuiltinPolyfillTest = TransformTest;
 
 TEST_F(BuiltinPolyfillTest, ShouldRunEmptyModule) {
@@ -348,6 +350,237 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// extractBits
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillExtractBits(Level level) {
+  BuiltinPolyfill::Builtins builtins;
+  builtins.extract_bits = level;
+  DataMap data;
+  data.Add<BuiltinPolyfill::Config>(builtins);
+  return data;
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunExtractBits) {
+  auto* src = R"(
+fn f() {
+  extractBits(1234, 5u, 6u);
+}
+)";
+
+  EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+  EXPECT_FALSE(
+      ShouldRun<BuiltinPolyfill>(src, polyfillExtractBits(Level::kNone)));
+  EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(
+      src, polyfillExtractBits(Level::kClampParameters)));
+  EXPECT_TRUE(
+      ShouldRun<BuiltinPolyfill>(src, polyfillExtractBits(Level::kFull)));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Full_i32) {
+  auto* src = R"(
+fn f() {
+  let r : i32 = extractBits(1234, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : i32, offset : u32, count : u32) -> i32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let shl = (32u - e);
+  let shr = (shl + s);
+  return ((v << shl) >> shr);
+}
+
+fn f() {
+  let r : i32 = tint_extract_bits(1234, 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Full_u32) {
+  auto* src = R"(
+fn f() {
+  let r : u32 = extractBits(1234u, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : u32, offset : u32, count : u32) -> u32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let shl = (32u - e);
+  let shr = (shl + s);
+  return ((v << shl) >> shr);
+}
+
+fn f() {
+  let r : u32 = tint_extract_bits(1234u, 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Full_vec3_i32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<i32> = extractBits(vec3<i32>(1234), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : vec3<i32>, offset : u32, count : u32) -> vec3<i32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let shl = (32u - e);
+  let shr = (shl + s);
+  return ((v << vec3<u32>(shl)) >> vec3<u32>(shr));
+}
+
+fn f() {
+  let r : vec3<i32> = tint_extract_bits(vec3<i32>(1234), 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Full_vec3_u32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<u32> = extractBits(vec3<u32>(1234u), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : vec3<u32>, offset : u32, count : u32) -> vec3<u32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  let shl = (32u - e);
+  let shr = (shl + s);
+  return ((v << vec3<u32>(shl)) >> vec3<u32>(shr));
+}
+
+fn f() {
+  let r : vec3<u32> = tint_extract_bits(vec3<u32>(1234u), 5u, 6u);
+}
+)";
+
+  auto got = Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kFull));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Clamp_i32) {
+  auto* src = R"(
+fn f() {
+  let r : i32 = extractBits(1234, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : i32, offset : u32, count : u32) -> i32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return extractBits(v, s, (e - s));
+}
+
+fn f() {
+  let r : i32 = tint_extract_bits(1234, 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Clamp_u32) {
+  auto* src = R"(
+fn f() {
+  let r : u32 = extractBits(1234u, 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : u32, offset : u32, count : u32) -> u32 {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return extractBits(v, s, (e - s));
+}
+
+fn f() {
+  let r : u32 = tint_extract_bits(1234u, 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Clamp_vec3_i32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<i32> = extractBits(vec3<i32>(1234), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : vec3<i32>, offset : u32, count : u32) -> vec3<i32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return extractBits(v, s, (e - s));
+}
+
+fn f() {
+  let r : vec3<i32> = tint_extract_bits(vec3<i32>(1234), 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, ExtractBits_Clamp_vec3_u32) {
+  auto* src = R"(
+fn f() {
+  let r : vec3<u32> = extractBits(vec3<u32>(1234u), 5u, 6u);
+}
+)";
+
+  auto* expect = R"(
+fn tint_extract_bits(v : vec3<u32>, offset : u32, count : u32) -> vec3<u32> {
+  let s = min(offset, 32u);
+  let e = min(32u, (s + count));
+  return extractBits(v, s, (e - s));
+}
+
+fn f() {
+  let r : vec3<u32> = tint_extract_bits(vec3<u32>(1234u), 5u, 6u);
+}
+)";
+
+  auto got =
+      Run<BuiltinPolyfill>(src, polyfillExtractBits(Level::kClampParameters));
+
+  EXPECT_EQ(expect, str(got));
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // firstLeadingBit
 ////////////////////////////////////////////////////////////////////////////////

src/tint/transform/glsl.cc

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

src/tint/writer/glsl/generator_impl.cc

@@ -604,6 +604,9 @@ bool GeneratorImpl::EmitBuiltinCall(std::ostream& out,
   if (builtin->Type() == sem::BuiltinType::kArrayLength) {
     return EmitArrayLength(out, expr);
   }
+  if (builtin->Type() == sem::BuiltinType::kExtractBits) {
+    return EmitExtractBits(out, expr);
+  }
   if (builtin->IsDataPacking()) {
     return EmitDataPackingCall(out, expr, builtin);
   }
@@ -813,6 +816,24 @@ bool GeneratorImpl::EmitArrayLength(std::ostream& out,
   return true;
 }
 
+bool GeneratorImpl::EmitExtractBits(std::ostream& out,
+                                    const ast::CallExpression* expr) {
+  out << "bitfieldExtract(";
+  if (!EmitExpression(out, expr->args[0])) {
+    return false;
+  }
+  out << ", int(";
+  if (!EmitExpression(out, expr->args[1])) {
+    return false;
+  }
+  out << "), int(";
+  if (!EmitExpression(out, expr->args[2])) {
+    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

@@ -171,6 +171,11 @@ class GeneratorImpl : public TextGenerator {
   /// @param expr the call expression
   /// @returns true if the array length expression is emitted
   bool EmitArrayLength(std::ostream& out, const ast::CallExpression* expr);
+  /// Handles generating a call to `bitfieldExtract`
+  /// @param out the output of the expression stream
+  /// @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 a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param out the output of the expression stream

src/tint/writer/glsl/generator_impl_builtin_test.cc

@@ -583,6 +583,38 @@ void main() {
 )");
 }
 
+TEST_F(GlslGeneratorImplTest_Builtin, ExtractBits) {
+  auto* v = Var("v", ty.vec3<u32>());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("extractBits", v, offset, count);
+  WrapInFunction(v, offset, count, call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#version 310 es
+
+uvec3 tint_extract_bits(uvec3 v, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldExtract(v, int(s), int((e - s)));
+}
+
+void test_function() {
+  uvec3 v = uvec3(0u, 0u, 0u);
+  uint offset = 0u;
+  uint count = 0u;
+  uvec3 tint_symbol = tint_extract_bits(v, 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

@@ -147,6 +147,7 @@ SanitizedResult Sanitize(
     // and `firstbithigh`.
     polyfills.count_leading_zeros = true;
     polyfills.count_trailing_zeros = true;
+    polyfills.extract_bits = transform::BuiltinPolyfill::Level::kFull;
     polyfills.first_leading_bit = true;
     polyfills.first_trailing_bit = true;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);

src/tint/writer/msl/generator_impl.cc

@@ -130,6 +130,8 @@ SanitizedResult Sanitize(
 
   {  // Builtin polyfills
     transform::BuiltinPolyfill::Builtins polyfills;
+    polyfills.extract_bits =
+        transform::BuiltinPolyfill::Level::kClampParameters;
     polyfills.first_leading_bit = true;
     polyfills.first_trailing_bit = true;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);
@@ -1373,6 +1375,9 @@ std::string GeneratorImpl::generate_builtin_name(const sem::Builtin* builtin) {
     case sem::BuiltinType::kDpdyFine:
       out += "dfdy";
       break;
+    case sem::BuiltinType::kExtractBits:
+      out += "extract_bits";
+      break;
     case sem::BuiltinType::kFwidth:
     case sem::BuiltinType::kFwidthCoarse:
     case sem::BuiltinType::kFwidthFine:

src/tint/writer/msl/generator_impl_builtin_test.cc

@@ -129,6 +129,8 @@ const ast::CallExpression* GenerateCall(BuiltinType builtin,
     case BuiltinType::kCountTrailingZeros:
     case BuiltinType::kReverseBits:
       return builder->Call(str.str(), "u2");
+    case BuiltinType::kExtractBits:
+      return builder->Call(str.str(), "u2", "u1", "u1");
     case BuiltinType::kMax:
     case BuiltinType::kMin:
       if (type == ParamType::kF32) {
@@ -229,6 +231,7 @@ INSTANTIATE_TEST_SUITE_P(
         BuiltinData{BuiltinType::kDpdyFine, ParamType::kF32, "dfdy"},
         BuiltinData{BuiltinType::kExp, ParamType::kF32, "exp"},
         BuiltinData{BuiltinType::kExp2, ParamType::kF32, "exp2"},
+        BuiltinData{BuiltinType::kExtractBits, ParamType::kU32, "extract_bits"},
         BuiltinData{BuiltinType::kFaceForward, ParamType::kF32, "faceforward"},
         BuiltinData{BuiltinType::kFloor, ParamType::kF32, "floor"},
         BuiltinData{BuiltinType::kFma, ParamType::kF32, "fma"},

src/tint/writer/spirv/builder.cc

@@ -263,6 +263,8 @@ SanitizedResult Sanitize(const Program* in,
     transform::BuiltinPolyfill::Builtins polyfills;
     polyfills.count_leading_zeros = true;
     polyfills.count_trailing_zeros = true;
+    polyfills.extract_bits =
+        transform::BuiltinPolyfill::Level::kClampParameters;
     polyfills.first_leading_bit = true;
     polyfills.first_trailing_bit = true;
     data.Add<transform::BuiltinPolyfill::Config>(polyfills);
@@ -2504,6 +2506,11 @@ uint32_t Builder::GenerateBuiltinCall(const sem::Call* call,
     case BuiltinType::kDpdyFine:
       op = spv::Op::OpDPdyFine;
       break;
+    case BuiltinType::kExtractBits:
+      op = builtin->Parameters()[0]->Type()->is_unsigned_scalar_or_vector()
+               ? spv::Op::OpBitFieldUExtract
+               : spv::Op::OpBitFieldSExtract;
+      break;
     case BuiltinType::kFwidth:
       op = spv::Op::OpFwidth;
       break;

src/tint/writer/spirv/builder_builtin_test.cc

@@ -2554,6 +2554,160 @@ OpReturn
   Validate(b);
 }
 
+TEST_F(BuiltinBuilderTest, Call_ExtractBits_i32) {
+  auto* v = Var("v", ty.i32());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("extractBits", v, offset, count);
+  auto* func = WrapInFunction(v, 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 "offset"
+OpName %13 "count"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%7 = OpTypeInt 32 1
+%6 = OpTypePointer Function %7
+%8 = OpConstantNull %7
+%11 = OpTypeInt 32 0
+%10 = OpTypePointer Function %11
+%12 = OpConstantNull %11
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpVariable %6 Function %8
+%9 = OpVariable %10 Function %12
+%13 = OpVariable %10 Function %12
+%15 = OpLoad %7 %5
+%16 = OpLoad %11 %9
+%17 = OpLoad %11 %13
+%14 = OpBitFieldSExtract %7 %15 %16 %17
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(BuiltinBuilderTest, Call_ExtractBits_u32) {
+  auto* v = Var("v", ty.u32());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("extractBits", v, offset, count);
+  auto* func = WrapInFunction(v, 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 "offset"
+OpName %10 "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
+%12 = OpLoad %7 %5
+%13 = OpLoad %7 %9
+%14 = OpLoad %7 %10
+%11 = OpBitFieldUExtract %7 %12 %13 %14
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(BuiltinBuilderTest, Call_ExtractBits_vec3_i32) {
+  auto* v = Var("v", ty.vec3<i32>());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("extractBits", v, offset, count);
+  auto* func = WrapInFunction(v, 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 "offset"
+OpName %14 "count"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeInt 32 1
+%7 = OpTypeVector %8 3
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%12 = OpTypeInt 32 0
+%11 = OpTypePointer Function %12
+%13 = OpConstantNull %12
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpVariable %6 Function %9
+%10 = OpVariable %11 Function %13
+%14 = OpVariable %11 Function %13
+%16 = OpLoad %7 %5
+%17 = OpLoad %12 %10
+%18 = OpLoad %12 %14
+%15 = OpBitFieldSExtract %7 %16 %17 %18
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(BuiltinBuilderTest, Call_ExtractBits_vec3_u32) {
+  auto* v = Var("v", ty.vec3<u32>());
+  auto* offset = Var("offset", ty.u32());
+  auto* count = Var("count", ty.u32());
+  auto* call = Call("extractBits", v, offset, count);
+  auto* func = WrapInFunction(v, 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 "offset"
+OpName %13 "count"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeInt 32 0
+%7 = OpTypeVector %8 3
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%11 = OpTypePointer Function %8
+%12 = OpConstantNull %8
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpVariable %6 Function %9
+%10 = OpVariable %11 Function %12
+%13 = OpVariable %11 Function %12
+%15 = OpLoad %7 %5
+%16 = OpLoad %8 %10
+%17 = OpLoad %8 %13
+%14 = OpBitFieldUExtract %7 %15 %16 %17
+OpReturn
+OpFunctionEnd
+)");
+}
+
 }  // namespace
 }  // namespace spirv
 }  // namespace writer

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

@@ -0,0 +1,30 @@
+                      OpCapability Shader
+                %15 = OpExtInstImport "GLSL.std.450"
+                      OpMemoryModel Logical GLSL450
+                      OpEntryPoint GLCompute %f "f"
+                      OpExecutionMode %f LocalSize 1 1 1
+                      OpName %f "f"
+                      OpName %v "v"
+                      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
+                %12 = 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
+            %offset = OpVariable %_ptr_Function_uint Function %12
+             %count = OpVariable %_ptr_Function_uint Function %12
+                %16 = OpLoad %int %v
+                %17 = OpLoad %uint %offset
+                %18 = OpLoad %uint %count
+                %14 = OpBitFieldSExtract %int %16 %17 %18
+                      OpReturn
+                      OpFunctionEnd

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

@@ -0,0 +1,25 @@
+#version 310 es
+
+int tint_extract_bits(int v, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldExtract(v, int(s), int((e - s)));
+}
+
+void f_1() {
+  int v = 0;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  int x_14 = tint_extract_bits(v, 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/extractBits/scalar/i32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+int tint_extract_bits(int v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << shl) >> shr);
+}
+
+void f_1() {
+  int v = 0;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const int x_14 = tint_extract_bits(v, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

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

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int tint_extract_bits(int v, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return extract_bits(v, s, (e - s));
+}
+
+void f_1() {
+  int v = 0;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  int const x_16 = v;
+  uint const x_17 = offset_1;
+  uint const x_18 = count;
+  int const x_14 = tint_extract_bits(x_16, x_17, x_18);
+  return;
+}
+
+kernel void f() {
+  f_1();
+  return;
+}
+

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

@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+         %10 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %tint_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               OpName %offset_1 "offset_1"
+               OpName %count_0 "count"
+               OpName %f "f"
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+          %1 = OpTypeFunction %int %int %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %23 = OpConstantNull %int
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %27 = OpConstantNull %uint
+%tint_extract_bits = OpFunction %int None %1
+          %v = OpFunctionParameter %int
+     %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 = OpBitFieldSExtract %int %v %9 %15
+               OpReturnValue %14
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %16
+         %19 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_int Function %23
+   %offset_1 = OpVariable %_ptr_Function_uint Function %27
+    %count_0 = OpVariable %_ptr_Function_uint Function %27
+               OpStore %v_0 %int_0
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %29 = OpLoad %int %v_0
+         %30 = OpLoad %uint %offset_1
+         %31 = OpLoad %uint %count_0
+         %32 = OpFunctionCall %int %tint_extract_bits %29 %30 %31
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %16
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,15 @@
+fn f_1() {
+  var v : i32 = 0;
+  var offset_1 : u32 = 0u;
+  var count : u32 = 0u;
+  let x_16 : i32 = v;
+  let x_17 : u32 = offset_1;
+  let x_18 : u32 = count;
+  let x_14 : i32 = extractBits(x_16, x_17, x_18);
+  return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+  f_1();
+}

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

@@ -0,0 +1,27 @@
+                      OpCapability Shader
+                %12 = OpExtInstImport "GLSL.std.450"
+                      OpMemoryModel Logical GLSL450
+                      OpEntryPoint GLCompute %f "f"
+                      OpExecutionMode %f LocalSize 1 1 1
+                      OpName %f "f"
+                      OpName %v "v"
+                      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
+            %offset = OpVariable %_ptr_Function_uint Function %8
+             %count = OpVariable %_ptr_Function_uint Function %8
+                %13 = OpLoad %uint %v
+                %14 = OpLoad %uint %offset
+                %15 = OpLoad %uint %count
+                %11 = OpBitFieldUExtract %uint %13 %14 %15
+                      OpReturn
+                      OpFunctionEnd

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

@@ -0,0 +1,25 @@
+#version 310 es
+
+uint tint_extract_bits(uint v, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldExtract(v, int(s), int((e - s)));
+}
+
+void f_1() {
+  uint v = 0u;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uint x_11 = tint_extract_bits(v, 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/extractBits/scalar/u32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+uint tint_extract_bits(uint v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << shl) >> shr);
+}
+
+void f_1() {
+  uint v = 0u;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const uint x_11 = tint_extract_bits(v, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

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

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint tint_extract_bits(uint v, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return extract_bits(v, s, (e - s));
+}
+
+void f_1() {
+  uint v = 0u;
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uint const x_13 = v;
+  uint const x_14 = offset_1;
+  uint const x_15 = count;
+  uint const x_11 = tint_extract_bits(x_13, x_14, x_15);
+  return;
+}
+
+kernel void f() {
+  f_1();
+  return;
+}
+

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

@@ -0,0 +1,58 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+          %9 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %tint_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               OpName %offset_1 "offset_1"
+               OpName %count_0 "count"
+               OpName %f "f"
+       %uint = OpTypeInt 32 0
+          %1 = OpTypeFunction %uint %uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %22 = OpConstantNull %uint
+%tint_extract_bits = OpFunction %uint None %1
+          %v = OpFunctionParameter %uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+          %7 = OpLabel
+          %8 = OpExtInst %uint %9 UMin %offset %uint_32
+         %12 = OpIAdd %uint %8 %count
+         %11 = OpExtInst %uint %9 UMin %uint_32 %12
+         %14 = OpISub %uint %11 %8
+         %13 = OpBitFieldUExtract %uint %v %8 %14
+               OpReturnValue %13
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %15
+         %18 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_uint Function %22
+   %offset_1 = OpVariable %_ptr_Function_uint Function %22
+    %count_0 = OpVariable %_ptr_Function_uint Function %22
+               OpStore %v_0 %uint_0
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %25 = OpLoad %uint %v_0
+         %26 = OpLoad %uint %offset_1
+         %27 = OpLoad %uint %count_0
+         %28 = OpFunctionCall %uint %tint_extract_bits %25 %26 %27
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %15
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,15 @@
+fn f_1() {
+  var v : u32 = 0u;
+  var offset_1 : u32 = 0u;
+  var count : u32 = 0u;
+  let x_13 : u32 = v;
+  let x_14 : u32 = offset_1;
+  let x_15 : u32 = count;
+  let x_11 : u32 = extractBits(x_13, x_14, x_15);
+  return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+  f_1();
+}

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

@@ -0,0 +1,31 @@
+                       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 %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
+                 %13 = 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
+             %offset = OpVariable %_ptr_Function_uint Function %13
+              %count = OpVariable %_ptr_Function_uint Function %13
+                 %17 = OpLoad %v3int %v
+                 %18 = OpLoad %uint %offset
+                 %19 = OpLoad %uint %count
+                 %15 = OpBitFieldSExtract %v3int %17 %18 %19
+                       OpReturn
+                       OpFunctionEnd

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

@@ -0,0 +1,25 @@
+#version 310 es
+
+ivec3 tint_extract_bits(ivec3 v, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldExtract(v, int(s), int((e - s)));
+}
+
+void f_1() {
+  ivec3 v = ivec3(0, 0, 0);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  ivec3 x_15 = tint_extract_bits(v, 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/extractBits/vec3/i32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+int3 tint_extract_bits(int3 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint3((shl).xxx)) >> uint3((shr).xxx));
+}
+
+void f_1() {
+  int3 v = int3(0, 0, 0);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const int3 x_15 = tint_extract_bits(v, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

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

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

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

@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+         %11 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %tint_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %17 = OpTypeFunction %void
+         %21 = OpConstantNull %v3int
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %27 = OpConstantNull %uint
+%tint_extract_bits = OpFunction %v3int None %1
+          %v = OpFunctionParameter %v3int
+     %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 = OpBitFieldSExtract %v3int %v %10 %16
+               OpReturnValue %15
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %17
+         %20 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_v3int Function %21
+   %offset_1 = OpVariable %_ptr_Function_uint Function %27
+    %count_0 = OpVariable %_ptr_Function_uint Function %27
+               OpStore %v_0 %21
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %29 = OpLoad %v3int %v_0
+         %30 = OpLoad %uint %offset_1
+         %31 = OpLoad %uint %count_0
+         %32 = OpFunctionCall %v3int %tint_extract_bits %29 %30 %31
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %17
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

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

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

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

@@ -0,0 +1,28 @@
+                        OpCapability Shader
+                  %15 = OpExtInstImport "GLSL.std.450"
+                        OpMemoryModel Logical GLSL450
+                        OpEntryPoint GLCompute %f "f"
+                        OpExecutionMode %f LocalSize 1 1 1
+                        OpName %f "f"
+                        OpName %v "v"
+                        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
+                  %12 = OpConstantNull %uint
+                   %f = OpFunction %void None %1
+                   %4 = OpLabel
+                   %v = OpVariable %_ptr_Function_v3uint Function %9
+              %offset = OpVariable %_ptr_Function_uint Function %12
+               %count = OpVariable %_ptr_Function_uint Function %12
+                  %16 = OpLoad %v3uint %v
+                  %17 = OpLoad %uint %offset
+                  %18 = OpLoad %uint %count
+                  %14 = OpBitFieldUExtract %v3uint %16 %17 %18
+                        OpReturn
+                        OpFunctionEnd

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

@@ -0,0 +1,25 @@
+#version 310 es
+
+uvec3 tint_extract_bits(uvec3 v, uint offset, uint count) {
+  uint s = min(offset, 32u);
+  uint e = min(32u, (s + count));
+  return bitfieldExtract(v, int(s), int((e - s)));
+}
+
+void f_1() {
+  uvec3 v = uvec3(0u, 0u, 0u);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uvec3 x_14 = tint_extract_bits(v, 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/extractBits/vec3/u32.spvasm.expected.hlsl

@@ -0,0 +1,21 @@
+uint3 tint_extract_bits(uint3 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint3((shl).xxx)) >> uint3((shr).xxx));
+}
+
+void f_1() {
+  uint3 v = uint3(0u, 0u, 0u);
+  uint offset_1 = 0u;
+  uint count = 0u;
+  const uint3 x_14 = tint_extract_bits(v, offset_1, count);
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  f_1();
+  return;
+}

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

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint3 tint_extract_bits(uint3 v, uint offset, uint count) {
+  uint const s = min(offset, 32u);
+  uint const e = min(32u, (s + count));
+  return extract_bits(v, s, (e - s));
+}
+
+void f_1() {
+  uint3 v = uint3();
+  uint offset_1 = 0u;
+  uint count = 0u;
+  uint3 const x_16 = v;
+  uint const x_17 = offset_1;
+  uint const x_18 = count;
+  uint3 const x_14 = tint_extract_bits(x_16, x_17, x_18);
+  return;
+}
+
+kernel void f() {
+  f_1();
+  return;
+}
+

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

@@ -0,0 +1,61 @@
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %f_1 "f_1"
+               OpName %v_0 "v"
+               OpName %offset_1 "offset_1"
+               OpName %count_0 "count"
+               OpName %f "f"
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+          %1 = OpTypeFunction %v3uint %v3uint %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+         %20 = OpConstantNull %v3uint
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %26 = OpConstantNull %uint
+%tint_extract_bits = OpFunction %v3uint None %1
+          %v = OpFunctionParameter %v3uint
+     %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 = OpBitFieldUExtract %v3uint %v %9 %15
+               OpReturnValue %14
+               OpFunctionEnd
+        %f_1 = OpFunction %void None %16
+         %19 = OpLabel
+        %v_0 = OpVariable %_ptr_Function_v3uint Function %20
+   %offset_1 = OpVariable %_ptr_Function_uint Function %26
+    %count_0 = OpVariable %_ptr_Function_uint Function %26
+               OpStore %v_0 %20
+               OpStore %offset_1 %uint_0
+               OpStore %count_0 %uint_0
+         %28 = OpLoad %v3uint %v_0
+         %29 = OpLoad %uint %offset_1
+         %30 = OpLoad %uint %count_0
+         %31 = OpFunctionCall %v3uint %tint_extract_bits %28 %29 %30
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %16
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %f_1
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,15 @@
+fn f_1() {
+  var v : vec3<u32> = vec3<u32>();
+  var offset_1 : u32 = 0u;
+  var count : u32 = 0u;
+  let x_16 : vec3<u32> = v;
+  let x_17 : u32 = offset_1;
+  let x_18 : u32 = count;
+  let x_14 : vec3<u32> = extractBits(x_16, x_17, x_18);
+  return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+  f_1();
+}

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

test/tint/builtins/gen/extractBits/12b197.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/12b197.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+uint3 tint_extract_bits(uint3 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint3((shl).xxx)) >> uint3((shr).xxx));
+}
+
+void extractBits_12b197() {
+  uint3 res = tint_extract_bits(uint3(0u, 0u, 0u), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_12b197();
+  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() {
+  extractBits_12b197();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_12b197();
+  return;
+}

test/tint/builtins/gen/extractBits/12b197.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/12b197.wgsl.expected.spvasm

@@ -0,0 +1,87 @@
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_12b197 "extractBits_12b197"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %24 = OpTypeFunction %void
+         %29 = OpConstantNull %v3uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %33 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %v3uint None %9
+          %v = OpFunctionParameter %v3uint
+     %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 = OpBitFieldUExtract %v3uint %v %17 %23
+               OpReturnValue %22
+               OpFunctionEnd
+%extractBits_12b197 = OpFunction %void None %24
+         %27 = OpLabel
+        %res = OpVariable %_ptr_Function_v3uint Function %29
+         %28 = OpFunctionCall %v3uint %tint_extract_bits %29 %uint_1 %uint_1
+               OpStore %res %28
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %33
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %extractBits_12b197
+               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 %extractBits_12b197
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %24
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %extractBits_12b197
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/12b197.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/extractBits/249874.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/249874.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+int tint_extract_bits(int v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << shl) >> shr);
+}
+
+void extractBits_249874() {
+  int res = tint_extract_bits(1, 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_249874();
+  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() {
+  extractBits_249874();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_249874();
+  return;
+}

test/tint/builtins/gen/extractBits/249874.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/249874.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_249874 "extractBits_249874"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %24 = OpTypeFunction %void
+      %int_1 = OpConstant %int 1
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %33 = OpConstantNull %int
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %int None %9
+          %v = OpFunctionParameter %int
+     %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 = OpBitFieldSExtract %int %v %17 %23
+               OpReturnValue %22
+               OpFunctionEnd
+%extractBits_249874 = OpFunction %void None %24
+         %27 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %33
+         %28 = OpFunctionCall %int %tint_extract_bits %int_1 %uint_1 %uint_1
+               OpStore %res %28
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %extractBits_249874
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %24
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %24
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %extractBits_249874
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %24
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %extractBits_249874
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/249874.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/extractBits/631377.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/631377.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+uint4 tint_extract_bits(uint4 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint4((shl).xxxx)) >> uint4((shr).xxxx));
+}
+
+void extractBits_631377() {
+  uint4 res = tint_extract_bits(uint4(0u, 0u, 0u, 0u), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_631377();
+  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() {
+  extractBits_631377();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_631377();
+  return;
+}

test/tint/builtins/gen/extractBits/631377.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/631377.wgsl.expected.spvasm

@@ -0,0 +1,87 @@
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_631377 "extractBits_631377"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %24 = OpTypeFunction %void
+         %29 = OpConstantNull %v4uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %33 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %v4uint None %9
+          %v = OpFunctionParameter %v4uint
+     %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 = OpBitFieldUExtract %v4uint %v %17 %23
+               OpReturnValue %22
+               OpFunctionEnd
+%extractBits_631377 = OpFunction %void None %24
+         %27 = OpLabel
+        %res = OpVariable %_ptr_Function_v4uint Function %29
+         %28 = OpFunctionCall %v4uint %tint_extract_bits %29 %uint_1 %uint_1
+               OpStore %res %28
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %33
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %extractBits_631377
+               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 %extractBits_631377
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %24
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %extractBits_631377
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/631377.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/extractBits/a99a8d.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/a99a8d.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+int2 tint_extract_bits(int2 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint2((shl).xx)) >> uint2((shr).xx));
+}
+
+void extractBits_a99a8d() {
+  int2 res = tint_extract_bits(int2(0, 0), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_a99a8d();
+  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() {
+  extractBits_a99a8d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_a99a8d();
+  return;
+}

test/tint/builtins/gen/extractBits/a99a8d.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/a99a8d.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_a99a8d "extractBits_a99a8d"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+         %30 = OpConstantNull %v2int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %v2int None %9
+          %v = OpFunctionParameter %v2int
+     %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 = OpBitFieldSExtract %v2int %v %18 %24
+               OpReturnValue %23
+               OpFunctionEnd
+%extractBits_a99a8d = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_v2int Function %30
+         %29 = OpFunctionCall %v2int %tint_extract_bits %30 %uint_1 %uint_1
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %extractBits_a99a8d
+               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 %extractBits_a99a8d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %extractBits_a99a8d
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/a99a8d.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/extractBits/ce81f8.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/ce81f8.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+uint tint_extract_bits(uint v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << shl) >> shr);
+}
+
+void extractBits_ce81f8() {
+  uint res = tint_extract_bits(1u, 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_ce81f8();
+  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() {
+  extractBits_ce81f8();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_ce81f8();
+  return;
+}

test/tint/builtins/gen/extractBits/ce81f8.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/ce81f8.wgsl.expected.spvasm

@@ -0,0 +1,86 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 46
+; Schema: 0
+               OpCapability Shader
+         %17 = 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_ce81f8 "extractBits_ce81f8"
+               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_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %23 = OpTypeFunction %void
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %31 = OpConstantNull %uint
+         %32 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %uint None %9
+          %v = OpFunctionParameter %uint
+     %offset = OpFunctionParameter %uint
+      %count = OpFunctionParameter %uint
+         %15 = OpLabel
+         %16 = OpExtInst %uint %17 UMin %offset %uint_32
+         %20 = OpIAdd %uint %16 %count
+         %19 = OpExtInst %uint %17 UMin %uint_32 %20
+         %22 = OpISub %uint %19 %16
+         %21 = OpBitFieldUExtract %uint %v %16 %22
+               OpReturnValue %21
+               OpFunctionEnd
+%extractBits_ce81f8 = OpFunction %void None %23
+         %26 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %31
+         %27 = OpFunctionCall %uint %tint_extract_bits %uint_1 %uint_1 %uint_1
+               OpStore %res %27
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %32
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %extractBits_ce81f8
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %23
+         %37 = OpLabel
+         %38 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %38
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %23
+         %41 = OpLabel
+         %42 = OpFunctionCall %void %extractBits_ce81f8
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %23
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %extractBits_ce81f8
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/ce81f8.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/extractBits/e04f5d.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/e04f5d.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+int3 tint_extract_bits(int3 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint3((shl).xxx)) >> uint3((shr).xxx));
+}
+
+void extractBits_e04f5d() {
+  int3 res = tint_extract_bits(int3(0, 0, 0), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_e04f5d();
+  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() {
+  extractBits_e04f5d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_e04f5d();
+  return;
+}

test/tint/builtins/gen/extractBits/e04f5d.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/e04f5d.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_e04f5d "extractBits_e04f5d"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+         %30 = OpConstantNull %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %v3int None %9
+          %v = OpFunctionParameter %v3int
+     %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 = OpBitFieldSExtract %v3int %v %18 %24
+               OpReturnValue %23
+               OpFunctionEnd
+%extractBits_e04f5d = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_v3int Function %30
+         %29 = OpFunctionCall %v3int %tint_extract_bits %30 %uint_1 %uint_1
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %extractBits_e04f5d
+               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 %extractBits_e04f5d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %extractBits_e04f5d
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/e04f5d.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/extractBits/f28f69.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/f28f69.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+uint2 tint_extract_bits(uint2 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint2((shl).xx)) >> uint2((shr).xx));
+}
+
+void extractBits_f28f69() {
+  uint2 res = tint_extract_bits(uint2(0u, 0u), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_f28f69();
+  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() {
+  extractBits_f28f69();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_f28f69();
+  return;
+}

test/tint/builtins/gen/extractBits/f28f69.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/f28f69.wgsl.expected.spvasm

@@ -0,0 +1,87 @@
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_f28f69 "extractBits_f28f69"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %24 = OpTypeFunction %void
+         %29 = OpConstantNull %v2uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %33 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %v2uint None %9
+          %v = OpFunctionParameter %v2uint
+     %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 = OpBitFieldUExtract %v2uint %v %17 %23
+               OpReturnValue %22
+               OpFunctionEnd
+%extractBits_f28f69 = OpFunction %void None %24
+         %27 = OpLabel
+        %res = OpVariable %_ptr_Function_v2uint Function %29
+         %28 = OpFunctionCall %v2uint %tint_extract_bits %29 %uint_1 %uint_1
+               OpStore %res %28
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %33
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %extractBits_f28f69
+               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 %extractBits_f28f69
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %24
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %extractBits_f28f69
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/f28f69.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/extractBits/fb850f.wgsl.expected.glsl

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

test/tint/builtins/gen/extractBits/fb850f.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+int4 tint_extract_bits(int4 v, uint offset, uint count) {
+  const uint s = min(offset, 32u);
+  const uint e = min(32u, (s + count));
+  const uint shl = (32u - e);
+  const uint shr = (shl + s);
+  return ((v << uint4((shl).xxxx)) >> uint4((shr).xxxx));
+}
+
+void extractBits_fb850f() {
+  int4 res = tint_extract_bits(int4(0, 0, 0, 0), 1u, 1u);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  extractBits_fb850f();
+  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() {
+  extractBits_fb850f();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  extractBits_fb850f();
+  return;
+}

test/tint/builtins/gen/extractBits/fb850f.wgsl.expected.msl

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

test/tint/builtins/gen/extractBits/fb850f.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; 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_extract_bits "tint_extract_bits"
+               OpName %v "v"
+               OpName %offset "offset"
+               OpName %count "count"
+               OpName %extractBits_fb850f "extractBits_fb850f"
+               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 %uint %uint
+    %uint_32 = OpConstant %uint 32
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+         %30 = OpConstantNull %v4int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%tint_extract_bits = OpFunction %v4int None %9
+          %v = OpFunctionParameter %v4int
+     %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 = OpBitFieldSExtract %v4int %v %18 %24
+               OpReturnValue %23
+               OpFunctionEnd
+%extractBits_fb850f = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %30
+         %29 = OpFunctionCall %v4int %tint_extract_bits %30 %uint_1 %uint_1
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %extractBits_fb850f
+               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 %extractBits_fb850f
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %extractBits_fb850f
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/extractBits/fb850f.wgsl.expected.wgsl

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