instrinsics: Implement dot() for integer vector types

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

docs/origin-trial-changes.md

@@ -6,6 +6,10 @@
 
 * Taking the address of a vector component is no longer allowed.
 
+### New Features
+
+* The `dot()` builtin now supports integer vector types.
+
 ## Changes for M97
 
 ### Breaking Changes

src/intrinsic_table.inl

@@ -3780,12 +3780,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [398] */
     /* usage */ ParameterUsage::kNone,
-    /* matcher indices */ &kMatcherIndices[21],
+    /* matcher indices */ &kMatcherIndices[39],
   },
   {
     /* [399] */
     /* usage */ ParameterUsage::kNone,
-    /* matcher indices */ &kMatcherIndices[21],
+    /* matcher indices */ &kMatcherIndices[39],
   },
   {
     /* [400] */
@@ -7791,12 +7791,12 @@ constexpr OverloadInfo kOverloads[] = {
   {
     /* [251] */
     /* num parameters */ 2,
-    /* num open types */ 0,
+    /* num open types */ 1,
     /* num open numbers */ 1,
-    /* open types */ &kOpenTypes[4],
+    /* open types */ &kOpenTypes[1],
     /* open numbers */ &kOpenNumbers[3],
     /* parameters */ &kParameters[398],
-    /* return matcher indices */ &kMatcherIndices[12],
+    /* return matcher indices */ &kMatcherIndices[1],
     /* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute),
     /* is_deprecated */ false,
   },
@@ -8082,7 +8082,7 @@ constexpr IntrinsicInfo kIntrinsics[] = {
   },
   {
     /* [16] */
-    /* fn dot<N : num>(vec<N, f32>, vec<N, f32>) -> f32 */
+    /* fn dot<N : num, T : fiu32>(vec<N, T>, vec<N, T>) -> T */
     /* num overloads */ 1,
     /* overloads */ &kOverloads[251],
   },

src/intrinsics.def

@@ -292,7 +292,7 @@ fn cross(vec3<f32>, vec3<f32>) -> vec3<f32>
 fn determinant<N: num>(mat<N, N, f32>) -> f32
 fn distance(f32, f32) -> f32
 fn distance<N: num>(vec<N, f32>, vec<N, f32>) -> f32
-fn dot<N: num>(vec<N, f32>, vec<N, f32>) -> f32
+fn dot<N: num, T: fiu32>(vec<N, T>, vec<N, T>) -> T
 [[stage("fragment")]] fn dpdx(f32) -> f32
 [[stage("fragment")]] fn dpdx<N: num>(vec<N, f32>) -> vec<N, f32>
 [[stage("fragment")]] fn dpdxCoarse(f32) -> f32

src/resolver/intrinsic_test.cc

@@ -339,7 +339,7 @@ TEST_F(ResolverIntrinsicTest, Dot_Vec2) {
 }
 
 TEST_F(ResolverIntrinsicTest, Dot_Vec3) {
-  Global("my_var", ty.vec3<f32>(), ast::StorageClass::kPrivate);
+  Global("my_var", ty.vec3<i32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call("dot", "my_var", "my_var");
   WrapInFunction(expr);
@@ -347,11 +347,11 @@ TEST_F(ResolverIntrinsicTest, Dot_Vec3) {
   EXPECT_TRUE(r()->Resolve()) << r()->error();
 
   ASSERT_NE(TypeOf(expr), nullptr);
-  EXPECT_TRUE(TypeOf(expr)->Is<sem::F32>());
+  EXPECT_TRUE(TypeOf(expr)->Is<sem::I32>());
 }
 
 TEST_F(ResolverIntrinsicTest, Dot_Vec4) {
-  Global("my_var", ty.vec4<f32>(), ast::StorageClass::kPrivate);
+  Global("my_var", ty.vec4<u32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call("dot", "my_var", "my_var");
   WrapInFunction(expr);
@@ -359,7 +359,7 @@ TEST_F(ResolverIntrinsicTest, Dot_Vec4) {
   EXPECT_TRUE(r()->Resolve()) << r()->error();
 
   ASSERT_NE(TypeOf(expr), nullptr);
-  EXPECT_TRUE(TypeOf(expr)->Is<sem::F32>());
+  EXPECT_TRUE(TypeOf(expr)->Is<sem::U32>());
 }
 
 TEST_F(ResolverIntrinsicTest, Dot_Error_Scalar) {
@@ -372,23 +372,7 @@ TEST_F(ResolverIntrinsicTest, Dot_Error_Scalar) {
             R"(error: no matching call to dot(f32, f32)
 
 1 candidate function:
-  dot(vecN<f32>, vecN<f32>) -> f32
+  dot(vecN<T>, vecN<T>) -> T  where: T is f32, i32 or u32
-)");
-}
-
-TEST_F(ResolverIntrinsicTest, Dot_Error_VectorInt) {
-  Global("my_var", ty.vec4<i32>(), ast::StorageClass::kPrivate);
-
-  auto* expr = Call("dot", "my_var", "my_var");
-  WrapInFunction(expr);
-
-  EXPECT_FALSE(r()->Resolve());
-
-  EXPECT_EQ(r()->error(),
-            R"(error: no matching call to dot(vec4<i32>, vec4<i32>)
-
-1 candidate function:
-  dot(vecN<f32>, vecN<f32>) -> f32
 )");
 }
 

src/writer/glsl/generator_impl.cc

@@ -107,7 +107,11 @@ bool GeneratorImpl::Generate() {
   size_t last_padding_line = 0;
 
   line() << "#version 310 es";
-  line() << "precision mediump float;" << std::endl;
+  line() << "precision mediump float;";
+
+  auto helpers_insertion_point = current_buffer_->lines.size();
+
+  line();
 
   for (auto* decl : builder_.AST().GlobalDeclarations()) {
     if (decl->Is<ast::Alias>()) {
@@ -153,7 +157,8 @@ bool GeneratorImpl::Generate() {
   }
 
   if (!helpers_.lines.empty()) {
-    current_buffer_->Insert(helpers_, 0, 0);
+    current_buffer_->Insert("", helpers_insertion_point++, 0);
+    current_buffer_->Insert(helpers_, helpers_insertion_point++, 0);
   }
 
   return true;
@@ -407,6 +412,8 @@ bool GeneratorImpl::EmitCall(std::ostream& out,
       return EmitTextureCall(out, expr, intrinsic);
     } else if (intrinsic->Type() == sem::IntrinsicType::kSelect) {
       return EmitSelectCall(out, expr);
+    } else if (intrinsic->Type() == sem::IntrinsicType::kDot) {
+      return EmitDotCall(out, expr, intrinsic);
     } else if (intrinsic->Type() == sem::IntrinsicType::kModf) {
       return EmitModfCall(out, expr, intrinsic);
     } else if (intrinsic->Type() == sem::IntrinsicType::kFrexp) {
@@ -671,6 +678,78 @@ bool GeneratorImpl::EmitSelectCall(std::ostream& out,
   return true;
 }
 
+bool GeneratorImpl::EmitDotCall(std::ostream& out,
+                                const ast::CallExpression* expr,
+                                const sem::Intrinsic* intrinsic) {
+  auto* vec_ty = intrinsic->Parameters()[0]->Type()->As<sem::Vector>();
+  std::string fn = "dot";
+  if (vec_ty->type()->is_integer_scalar()) {
+    // GLSL does not have a builtin for dot() with integer vector types.
+    // Generate the helper function if it hasn't been created already
+    fn = utils::GetOrCreate(int_dot_funcs_, vec_ty, [&]() -> std::string {
+      TextBuffer b;
+      TINT_DEFER(helpers_.Append(b));
+
+      auto fn_name = UniqueIdentifier("tint_int_dot");
+
+      std::string v;
+      {
+        std::stringstream s;
+        if (!EmitType(s, vec_ty->type(), ast::StorageClass::kNone,
+                      ast::Access::kRead, "")) {
+          return "";
+        }
+        v = s.str();
+      }
+      {  // (u)int tint_int_dot([i|u]vecN a, [i|u]vecN b) {
+        auto l = line(&b);
+        if (!EmitType(l, vec_ty->type(), ast::StorageClass::kNone,
+                      ast::Access::kRead, "")) {
+          return "";
+        }
+        l << " " << fn_name << "(";
+        if (!EmitType(l, vec_ty, ast::StorageClass::kNone, ast::Access::kRead,
+                      "")) {
+          return "";
+        }
+        l << " a, ";
+        if (!EmitType(l, vec_ty, ast::StorageClass::kNone, ast::Access::kRead,
+                      "")) {
+          return "";
+        }
+        l << " b) {";
+      }
+      {
+        auto l = line(&b);
+        l << "  return ";
+        for (uint32_t i = 0; i < vec_ty->Width(); i++) {
+          if (i > 0) {
+            l << " + ";
+          }
+          l << "a[" << i << "]*b[" << i << "]";
+        }
+        l << ";";
+      }
+      line(&b) << "}";
+      return fn_name;
+    });
+    if (fn.empty()) {
+      return false;
+    }
+  }
+
+  out << fn << "(";
+  if (!EmitExpression(out, expr->args[0])) {
+    return false;
+  }
+  out << ", ";
+  if (!EmitExpression(out, expr->args[1])) {
+    return false;
+  }
+  out << ")";
+  return true;
+}
+
 bool GeneratorImpl::EmitModfCall(std::ostream& out,
                                  const ast::CallExpression* expr,
                                  const sem::Intrinsic* intrinsic) {
@@ -2216,9 +2295,6 @@ bool GeneratorImpl::EmitStatement(const ast::Statement* stmt) {
   }
   if (auto* c = stmt->As<ast::CallStatement>()) {
     auto out = line();
-    if (!TypeOf(c->expr)->Is<sem::Void>()) {
-      out << "(void) ";
-    }
     if (!EmitCall(out, c->expr)) {
       return false;
     }

src/writer/glsl/generator_impl.h

@@ -136,6 +136,14 @@ class GeneratorImpl : public TextGenerator {
   /// @param expr the call expression
   /// @returns true if the call expression is emitted
   bool EmitSelectCall(std::ostream& out, const ast::CallExpression* expr);
+  /// Handles generating a call to the `dot()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitDotCall(std::ostream& out,
+                   const ast::CallExpression* expr,
+                   const sem::Intrinsic* intrinsic);
   /// Handles generating a call to the `modf()` intrinsic
   /// @param out the output of the expression stream
   /// @param expr the call expression
@@ -416,6 +424,7 @@ class GeneratorImpl : public TextGenerator {
   std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
   std::unordered_map<const sem::Struct*, std::string> structure_builders_;
   std::unordered_map<const sem::Vector*, std::string> dynamic_vector_write_;
+  std::unordered_map<const sem::Vector*, std::string> int_dot_funcs_;
 };
 
 }  // namespace glsl

src/writer/glsl/generator_impl_intrinsic_test.cc

@@ -599,6 +599,64 @@ void main() {
 }
 #endif
 
+TEST_F(GlslGeneratorImplTest_Intrinsic, DotI32) {
+  Global("v", ty.vec3<i32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(Call("dot", "v", "v"));
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#version 310 es
+precision mediump float;
+
+int tint_int_dot(ivec3 a, ivec3 b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+
+ivec3 v = ivec3(0, 0, 0);
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void test_function() {
+  tint_int_dot(v, v);
+  return;
+}
+void main() {
+  test_function();
+}
+
+
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, DotU32) {
+  Global("v", ty.vec3<u32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(Call("dot", "v", "v"));
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#version 310 es
+precision mediump float;
+
+uint tint_int_dot(uvec3 a, uvec3 b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+
+uvec3 v = uvec3(0u, 0u, 0u);
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void test_function() {
+  tint_int_dot(v, v);
+  return;
+}
+void main() {
+  test_function();
+}
+
+
+)");
+}
+
 }  // namespace
 }  // namespace glsl
 }  // namespace writer

src/writer/msl/generator_impl.cc

@@ -543,6 +543,8 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
   auto name = generate_builtin_name(intrinsic);
 
   switch (intrinsic->Type()) {
+    case sem::IntrinsicType::kDot:
+      return EmitDotCall(out, expr, intrinsic);
     case sem::IntrinsicType::kModf:
       return EmitModfCall(out, expr, intrinsic);
     case sem::IntrinsicType::kFrexp:
@@ -1005,6 +1007,53 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
   return true;
 }
 
+bool GeneratorImpl::EmitDotCall(std::ostream& out,
+                                const ast::CallExpression* expr,
+                                const sem::Intrinsic* intrinsic) {
+  auto* vec_ty = intrinsic->Parameters()[0]->Type()->As<sem::Vector>();
+  std::string fn = "dot";
+  if (vec_ty->type()->is_integer_scalar()) {
+    // MSL does not have a builtin for dot() with integer vector types.
+    // Generate the helper function if it hasn't been created already
+    fn = utils::GetOrCreate(
+        int_dot_funcs_, vec_ty->Width(), [&]() -> std::string {
+          TextBuffer b;
+          TINT_DEFER(helpers_.Append(b));
+
+          auto fn_name =
+              UniqueIdentifier("tint_dot" + std::to_string(vec_ty->Width()));
+          auto v = "vec<T," + std::to_string(vec_ty->Width()) + ">";
+
+          line(&b) << "template<typename T>";
+          line(&b) << "T " << fn_name << "(" << v << " a, " << v << " b) {";
+          {
+            auto l = line(&b);
+            l << "  return ";
+            for (uint32_t i = 0; i < vec_ty->Width(); i++) {
+              if (i > 0) {
+                l << " + ";
+              }
+              l << "a[" << i << "]*b[" << i << "]";
+            }
+            l << ";";
+          }
+          line(&b) << "}";
+          return fn_name;
+        });
+  }
+
+  out << fn << "(";
+  if (!EmitExpression(out, expr->args[0])) {
+    return false;
+  }
+  out << ", ";
+  if (!EmitExpression(out, expr->args[1])) {
+    return false;
+  }
+  out << ")";
+  return true;
+}
+
 bool GeneratorImpl::EmitModfCall(std::ostream& out,
                                  const ast::CallExpression* expr,
                                  const sem::Intrinsic* intrinsic) {

src/writer/msl/generator_impl.h

@@ -154,6 +154,14 @@ class GeneratorImpl : public TextGenerator {
   bool EmitTextureCall(std::ostream& out,
                        const ast::CallExpression* expr,
                        const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to the `dot()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitDotCall(std::ostream& out,
+                   const ast::CallExpression* expr,
+                   const sem::Intrinsic* intrinsic);
   /// Handles generating a call to the `modf()` intrinsic
   /// @param out the output of the expression stream
   /// @param expr the call expression
@@ -394,6 +402,7 @@ class GeneratorImpl : public TextGenerator {
 
   std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
   std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
+  std::unordered_map<uint32_t, std::string> int_dot_funcs_;
 };
 
 }  // namespace msl

src/writer/msl/generator_impl_intrinsic_test.cc

@@ -343,6 +343,30 @@ TEST_F(MslGeneratorImplTest, Unpack2x16Float) {
   EXPECT_EQ(out.str(), "float2(as_type<half2>(p1))");
 }
 
+TEST_F(MslGeneratorImplTest, DotI32) {
+  Global("v", ty.vec3<i32>(), ast::StorageClass::kPrivate);
+  WrapInFunction(Call("dot", "v", "v"));
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+kernel void test_function() {
+  thread int3 tint_symbol = 0;
+  tint_dot3(tint_symbol, tint_symbol);
+  return;
+}
+
+)");
+}
+
 TEST_F(MslGeneratorImplTest, Ignore) {
   Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
        ty.i32(), {Return(Mul(Add("a", "b"), "c"))});

src/writer/spirv/builder.cc

@@ -2423,9 +2423,48 @@ uint32_t Builder::GenerateIntrinsic(const ast::CallExpression* call,
     case IntrinsicType::kCountOneBits:
       op = spv::Op::OpBitCount;
       break;
-    case IntrinsicType::kDot:
+    case IntrinsicType::kDot: {
       op = spv::Op::OpDot;
+      auto* vec_ty = intrinsic->Parameters()[0]->Type()->As<sem::Vector>();
+      if (vec_ty->type()->is_integer_scalar()) {
+        // TODO(crbug.com/tint/1267): OpDot requires floating-point types, but
+        // WGSL also supports integer types. SPV_KHR_integer_dot_product adds
+        // support for integer vectors. Use it if it is available.
+        auto el_ty = Operand::Int(GenerateTypeIfNeeded(vec_ty->type()));
+        auto vec_a = Operand::Int(get_arg_as_value_id(0));
+        auto vec_b = Operand::Int(get_arg_as_value_id(1));
+        if (vec_a.to_i() == 0 || vec_b.to_i() == 0) {
+          return 0;
+        }
+
+        auto sum = Operand::Int(0);
+        for (uint32_t i = 0; i < vec_ty->Width(); i++) {
+          auto a = result_op();
+          auto b = result_op();
+          auto mul = result_op();
+          if (!push_function_inst(spv::Op::OpCompositeExtract,
+                                  {el_ty, a, vec_a, Operand::Int(i)}) ||
+              !push_function_inst(spv::Op::OpCompositeExtract,
+                                  {el_ty, b, vec_b, Operand::Int(i)}) ||
+              !push_function_inst(spv::Op::OpIMul, {el_ty, mul, a, b})) {
+            return 0;
+          }
+          if (i == 0) {
+            sum = mul;
+          } else {
+            auto prev_sum = sum;
+            auto is_last_el = i == (vec_ty->Width() - 1);
+            sum = is_last_el ? Operand::Int(result_id) : result_op();
+            if (!push_function_inst(spv::Op::OpIAdd,
+                                    {el_ty, sum, prev_sum, mul})) {
+              return 0;
+            }
+          }
+        }
+        return result_id;
+      }
       break;
+    }
     case IntrinsicType::kDpdx:
       op = spv::Op::OpDPdx;
       break;

src/writer/spirv/builder_intrinsic_test.cc

@@ -407,7 +407,7 @@ INSTANTIATE_TEST_SUITE_P(
     testing::Values(IntrinsicData{"countOneBits", "OpBitCount"},
                     IntrinsicData{"reverseBits", "OpBitReverse"}));
 
-TEST_F(IntrinsicBuilderTest, Call_Dot) {
+TEST_F(IntrinsicBuilderTest, Call_Dot_F32) {
   auto* var = Global("v", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call("dot", "v", "v");
@@ -432,6 +432,76 @@ TEST_F(IntrinsicBuilderTest, Call_Dot) {
 )");
 }
 
+TEST_F(IntrinsicBuilderTest, Call_Dot_U32) {
+  auto* var = Global("v", ty.vec3<u32>(), ast::StorageClass::kPrivate);
+
+  auto* expr = Call("dot", "v", "v");
+  WrapInFunction(expr);
+
+  spirv::Builder& b = Build();
+
+  b.push_function(Function{});
+  ASSERT_TRUE(b.GenerateGlobalVariable(var)) << b.error();
+
+  EXPECT_EQ(b.GenerateCallExpression(expr), 6u) << b.error();
+  EXPECT_EQ(DumpInstructions(b.types()), R"(%4 = OpTypeInt 32 0
+%3 = OpTypeVector %4 3
+%2 = OpTypePointer Private %3
+%5 = OpConstantNull %3
+%1 = OpVariable %2 Private %5
+)");
+  EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()),
+            R"(%7 = OpLoad %3 %1
+%8 = OpLoad %3 %1
+%9 = OpCompositeExtract %4 %7 0
+%10 = OpCompositeExtract %4 %8 0
+%11 = OpIMul %4 %9 %10
+%12 = OpCompositeExtract %4 %7 1
+%13 = OpCompositeExtract %4 %8 1
+%14 = OpIMul %4 %12 %13
+%15 = OpIAdd %4 %11 %14
+%16 = OpCompositeExtract %4 %7 2
+%17 = OpCompositeExtract %4 %8 2
+%18 = OpIMul %4 %16 %17
+%6 = OpIAdd %4 %15 %18
+)");
+}
+
+TEST_F(IntrinsicBuilderTest, Call_Dot_I32) {
+  auto* var = Global("v", ty.vec3<i32>(), ast::StorageClass::kPrivate);
+
+  auto* expr = Call("dot", "v", "v");
+  WrapInFunction(expr);
+
+  spirv::Builder& b = Build();
+
+  b.push_function(Function{});
+  ASSERT_TRUE(b.GenerateGlobalVariable(var)) << b.error();
+
+  EXPECT_EQ(b.GenerateCallExpression(expr), 6u) << b.error();
+  EXPECT_EQ(DumpInstructions(b.types()), R"(%4 = OpTypeInt 32 1
+%3 = OpTypeVector %4 3
+%2 = OpTypePointer Private %3
+%5 = OpConstantNull %3
+%1 = OpVariable %2 Private %5
+)");
+  EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()),
+            R"(%7 = OpLoad %3 %1
+%8 = OpLoad %3 %1
+%9 = OpCompositeExtract %4 %7 0
+%10 = OpCompositeExtract %4 %8 0
+%11 = OpIMul %4 %9 %10
+%12 = OpCompositeExtract %4 %7 1
+%13 = OpCompositeExtract %4 %8 1
+%14 = OpIMul %4 %12 %13
+%15 = OpIAdd %4 %11 %14
+%16 = OpCompositeExtract %4 %7 2
+%17 = OpCompositeExtract %4 %8 2
+%18 = OpIMul %4 %16 %17
+%6 = OpIAdd %4 %15 %18
+)");
+}
+
 using IntrinsicDeriveTest = IntrinsicBuilderTestWithParam<IntrinsicData>;
 TEST_P(IntrinsicDeriveTest, Call_Derivative_Scalar) {
   auto param = GetParam();

test/intrinsics/gen/dot/7548a0.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/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<3, u32>, vec<3, u32>) -> u32
+fn dot_7548a0() {
+  var res: u32 = dot(vec3<u32>(), vec3<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  dot_7548a0();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  dot_7548a0();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  dot_7548a0();
+}

test/intrinsics/gen/dot/7548a0.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void dot_7548a0() {
+  uint res = dot(uint3(0u, 0u, 0u), uint3(0u, 0u, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  dot_7548a0();
+  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() {
+  dot_7548a0();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  dot_7548a0();
+  return;
+}

test/intrinsics/gen/dot/7548a0.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void dot_7548a0() {
+  uint res = tint_dot3(uint3(), uint3());
+}
+
+float4 vertex_main_inner() {
+  dot_7548a0();
+  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() {
+  dot_7548a0();
+  return;
+}
+
+kernel void compute_main() {
+  dot_7548a0();
+  return;
+}
+

test/intrinsics/gen/dot/7548a0.wgsl.expected.spvasm

@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %dot_7548a0 "dot_7548a0"
+               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
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+         %16 = OpConstantNull %v3uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %29 = OpConstantNull %uint
+         %30 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %dot_7548a0 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %29
+         %17 = OpCompositeExtract %uint %16 0
+         %18 = OpCompositeExtract %uint %16 0
+         %19 = OpIMul %uint %17 %18
+         %20 = OpCompositeExtract %uint %16 1
+         %21 = OpCompositeExtract %uint %16 1
+         %22 = OpIMul %uint %20 %21
+         %23 = OpIAdd %uint %19 %22
+         %24 = OpCompositeExtract %uint %16 2
+         %25 = OpCompositeExtract %uint %16 2
+         %26 = OpIMul %uint %24 %25
+         %13 = OpIAdd %uint %23 %26
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %30
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %dot_7548a0
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %35 = OpLabel
+         %36 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %36
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %39 = OpLabel
+         %40 = OpFunctionCall %void %dot_7548a0
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %dot_7548a0
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/dot/7548a0.wgsl.expected.wgsl

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

test/intrinsics/gen/dot/97c7ee.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/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<2, u32>, vec<2, u32>) -> u32
+fn dot_97c7ee() {
+  var res: u32 = dot(vec2<u32>(), vec2<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  dot_97c7ee();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  dot_97c7ee();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  dot_97c7ee();
+}

test/intrinsics/gen/dot/97c7ee.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void dot_97c7ee() {
+  uint res = dot(uint2(0u, 0u), uint2(0u, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  dot_97c7ee();
+  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() {
+  dot_97c7ee();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  dot_97c7ee();
+  return;
+}

test/intrinsics/gen/dot/97c7ee.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot2(vec<T,2> a, vec<T,2> b) {
+  return a[0]*b[0] + a[1]*b[1];
+}
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void dot_97c7ee() {
+  uint res = tint_dot2(uint2(), uint2());
+}
+
+float4 vertex_main_inner() {
+  dot_97c7ee();
+  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() {
+  dot_97c7ee();
+  return;
+}
+
+kernel void compute_main() {
+  dot_97c7ee();
+  return;
+}
+

test/intrinsics/gen/dot/97c7ee.wgsl.expected.spvasm

@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %dot_97c7ee "dot_97c7ee"
+               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
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %v2uint = OpTypeVector %uint 2
+         %16 = OpConstantNull %v2uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %25 = OpConstantNull %uint
+         %26 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %dot_97c7ee = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %25
+         %17 = OpCompositeExtract %uint %16 0
+         %18 = OpCompositeExtract %uint %16 0
+         %19 = OpIMul %uint %17 %18
+         %20 = OpCompositeExtract %uint %16 1
+         %21 = OpCompositeExtract %uint %16 1
+         %22 = OpIMul %uint %20 %21
+         %13 = OpIAdd %uint %19 %22
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %26
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %dot_97c7ee
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %32
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %dot_97c7ee
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %38 = OpLabel
+         %39 = OpFunctionCall %void %dot_97c7ee
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/dot/97c7ee.wgsl.expected.wgsl

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

test/intrinsics/gen/dot/e994c7.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/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<4, u32>, vec<4, u32>) -> u32
+fn dot_e994c7() {
+  var res: u32 = dot(vec4<u32>(), vec4<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  dot_e994c7();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  dot_e994c7();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  dot_e994c7();
+}

test/intrinsics/gen/dot/e994c7.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void dot_e994c7() {
+  uint res = dot(uint4(0u, 0u, 0u, 0u), uint4(0u, 0u, 0u, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  dot_e994c7();
+  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() {
+  dot_e994c7();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  dot_e994c7();
+  return;
+}

test/intrinsics/gen/dot/e994c7.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot4(vec<T,4> a, vec<T,4> b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void dot_e994c7() {
+  uint res = tint_dot4(uint4(), uint4());
+}
+
+float4 vertex_main_inner() {
+  dot_e994c7();
+  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() {
+  dot_e994c7();
+  return;
+}
+
+kernel void compute_main() {
+  dot_e994c7();
+  return;
+}
+

test/intrinsics/gen/dot/e994c7.wgsl.expected.spvasm

@@ -0,0 +1,82 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %dot_e994c7 "dot_e994c7"
+               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
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+         %16 = OpConstantNull %v4uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %33 = OpConstantNull %uint
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %dot_e994c7 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %33
+         %17 = OpCompositeExtract %uint %16 0
+         %18 = OpCompositeExtract %uint %16 0
+         %19 = OpIMul %uint %17 %18
+         %20 = OpCompositeExtract %uint %16 1
+         %21 = OpCompositeExtract %uint %16 1
+         %22 = OpIMul %uint %20 %21
+         %23 = OpIAdd %uint %19 %22
+         %24 = OpCompositeExtract %uint %16 2
+         %25 = OpCompositeExtract %uint %16 2
+         %26 = OpIMul %uint %24 %25
+         %27 = OpIAdd %uint %23 %26
+         %28 = OpCompositeExtract %uint %16 3
+         %29 = OpCompositeExtract %uint %16 3
+         %30 = OpIMul %uint %28 %29
+         %13 = OpIAdd %uint %27 %30
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %dot_e994c7
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %dot_e994c7
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %dot_e994c7
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/dot/e994c7.wgsl.expected.wgsl

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

test/intrinsics/gen/dot/ef6b1d.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/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<4, i32>, vec<4, i32>) -> i32
+fn dot_ef6b1d() {
+  var res: i32 = dot(vec4<i32>(), vec4<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  dot_ef6b1d();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  dot_ef6b1d();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  dot_ef6b1d();
+}

test/intrinsics/gen/dot/ef6b1d.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void dot_ef6b1d() {
+  int res = dot(int4(0, 0, 0, 0), int4(0, 0, 0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  dot_ef6b1d();
+  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() {
+  dot_ef6b1d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  dot_ef6b1d();
+  return;
+}

test/intrinsics/gen/dot/ef6b1d.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot4(vec<T,4> a, vec<T,4> b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void dot_ef6b1d() {
+  int res = tint_dot4(int4(), int4());
+}
+
+float4 vertex_main_inner() {
+  dot_ef6b1d();
+  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() {
+  dot_ef6b1d();
+  return;
+}
+
+kernel void compute_main() {
+  dot_ef6b1d();
+  return;
+}
+

test/intrinsics/gen/dot/ef6b1d.wgsl.expected.spvasm

@@ -0,0 +1,82 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %dot_ef6b1d "dot_ef6b1d"
+               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
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
+         %16 = OpConstantNull %v4int
+%_ptr_Function_int = OpTypePointer Function %int
+         %33 = OpConstantNull %int
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %dot_ef6b1d = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %33
+         %17 = OpCompositeExtract %int %16 0
+         %18 = OpCompositeExtract %int %16 0
+         %19 = OpIMul %int %17 %18
+         %20 = OpCompositeExtract %int %16 1
+         %21 = OpCompositeExtract %int %16 1
+         %22 = OpIMul %int %20 %21
+         %23 = OpIAdd %int %19 %22
+         %24 = OpCompositeExtract %int %16 2
+         %25 = OpCompositeExtract %int %16 2
+         %26 = OpIMul %int %24 %25
+         %27 = OpIAdd %int %23 %26
+         %28 = OpCompositeExtract %int %16 3
+         %29 = OpCompositeExtract %int %16 3
+         %30 = OpIMul %int %28 %29
+         %13 = OpIAdd %int %27 %30
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %dot_ef6b1d
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %dot_ef6b1d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %dot_ef6b1d
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/dot/ef6b1d.wgsl.expected.wgsl

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

test/intrinsics/gen/dot/f1312c.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/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<3, i32>, vec<3, i32>) -> i32
+fn dot_f1312c() {
+  var res: i32 = dot(vec3<i32>(), vec3<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  dot_f1312c();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  dot_f1312c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  dot_f1312c();
+}

test/intrinsics/gen/dot/f1312c.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void dot_f1312c() {
+  int res = dot(int3(0, 0, 0), int3(0, 0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  dot_f1312c();
+  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() {
+  dot_f1312c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  dot_f1312c();
+  return;
+}

test/intrinsics/gen/dot/f1312c.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void dot_f1312c() {
+  int res = tint_dot3(int3(), int3());
+}
+
+float4 vertex_main_inner() {
+  dot_f1312c();
+  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() {
+  dot_f1312c();
+  return;
+}
+
+kernel void compute_main() {
+  dot_f1312c();
+  return;
+}
+

test/intrinsics/gen/dot/f1312c.wgsl.expected.spvasm

@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %dot_f1312c "dot_f1312c"
+               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
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+         %16 = OpConstantNull %v3int
+%_ptr_Function_int = OpTypePointer Function %int
+         %29 = OpConstantNull %int
+         %30 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %dot_f1312c = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %29
+         %17 = OpCompositeExtract %int %16 0
+         %18 = OpCompositeExtract %int %16 0
+         %19 = OpIMul %int %17 %18
+         %20 = OpCompositeExtract %int %16 1
+         %21 = OpCompositeExtract %int %16 1
+         %22 = OpIMul %int %20 %21
+         %23 = OpIAdd %int %19 %22
+         %24 = OpCompositeExtract %int %16 2
+         %25 = OpCompositeExtract %int %16 2
+         %26 = OpIMul %int %24 %25
+         %13 = OpIAdd %int %23 %26
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %30
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %dot_f1312c
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %35 = OpLabel
+         %36 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %36
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %39 = OpLabel
+         %40 = OpFunctionCall %void %dot_f1312c
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %dot_f1312c
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/dot/f1312c.wgsl.expected.wgsl

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

test/intrinsics/gen/dot/fc5f7c.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/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<2, i32>, vec<2, i32>) -> i32
+fn dot_fc5f7c() {
+  var res: i32 = dot(vec2<i32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  dot_fc5f7c();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  dot_fc5f7c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  dot_fc5f7c();
+}

test/intrinsics/gen/dot/fc5f7c.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+void dot_fc5f7c() {
+  int res = dot(int2(0, 0), int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  dot_fc5f7c();
+  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() {
+  dot_fc5f7c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  dot_fc5f7c();
+  return;
+}

test/intrinsics/gen/dot/fc5f7c.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot2(vec<T,2> a, vec<T,2> b) {
+  return a[0]*b[0] + a[1]*b[1];
+}
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void dot_fc5f7c() {
+  int res = tint_dot2(int2(), int2());
+}
+
+float4 vertex_main_inner() {
+  dot_fc5f7c();
+  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() {
+  dot_fc5f7c();
+  return;
+}
+
+kernel void compute_main() {
+  dot_fc5f7c();
+  return;
+}
+

test/intrinsics/gen/dot/fc5f7c.wgsl.expected.spvasm

@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %dot_fc5f7c "dot_fc5f7c"
+               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
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %16 = OpConstantNull %v2int
+%_ptr_Function_int = OpTypePointer Function %int
+         %25 = OpConstantNull %int
+         %26 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %dot_fc5f7c = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %25
+         %17 = OpCompositeExtract %int %16 0
+         %18 = OpCompositeExtract %int %16 0
+         %19 = OpIMul %int %17 %18
+         %20 = OpCompositeExtract %int %16 1
+         %21 = OpCompositeExtract %int %16 1
+         %22 = OpIMul %int %20 %21
+         %13 = OpIAdd %int %19 %22
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %26
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %dot_fc5f7c
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %32
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %dot_fc5f7c
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %38 = OpLabel
+         %39 = OpFunctionCall %void %dot_fc5f7c
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/dot/fc5f7c.wgsl.expected.wgsl

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